From 9968d9daf264b726ee50bbc97937daac4e9c1811 Mon Sep 17 00:00:00 2001 From: Francisco Jerez Date: Tue, 1 Oct 2013 11:57:32 -0700 Subject: [PATCH] clover: Switch command queues to the new model. Tested-by: Tom Stellard --- .../state_trackers/clover/api/event.cpp | 25 +- .../state_trackers/clover/api/kernel.cpp | 28 +- .../state_trackers/clover/api/queue.cpp | 52 ++-- .../state_trackers/clover/api/transfer.cpp | 273 +++++++++--------- .../state_trackers/clover/core/device.hpp | 2 +- .../state_trackers/clover/core/error.hpp | 2 +- .../state_trackers/clover/core/event.hpp | 2 +- .../state_trackers/clover/core/kernel.cpp | 14 +- .../state_trackers/clover/core/memory.cpp | 32 +- .../state_trackers/clover/core/memory.hpp | 8 +- .../state_trackers/clover/core/object.hpp | 3 + .../state_trackers/clover/core/queue.cpp | 14 +- .../state_trackers/clover/core/queue.hpp | 59 ++-- .../state_trackers/clover/core/timestamp.hpp | 2 +- 14 files changed, 264 insertions(+), 252 deletions(-) diff --git a/src/gallium/state_trackers/clover/api/event.cpp b/src/gallium/state_trackers/clover/api/event.cpp index 4f689419c99..1395c54db08 100644 --- a/src/gallium/state_trackers/clover/api/event.cpp +++ b/src/gallium/state_trackers/clover/api/event.cpp @@ -92,7 +92,7 @@ clGetEventInfo(cl_event d_ev, cl_event_info param, switch (param) { case CL_EVENT_COMMAND_QUEUE: - buf.as_scalar() = ev.queue(); + buf.as_scalar() = desc(ev.queue()); break; case CL_EVENT_CONTEXT: @@ -167,13 +167,12 @@ clReleaseEvent(cl_event d_ev) try { PUBLIC cl_int clEnqueueMarker(cl_command_queue d_q, cl_event *rd_ev) try { - if (!d_q) - throw error(CL_INVALID_COMMAND_QUEUE); + auto &q = obj(d_q); if (!rd_ev) throw error(CL_INVALID_VALUE); - *rd_ev = desc(new hard_event(*d_q, CL_COMMAND_MARKER, {})); + *rd_ev = desc(new hard_event(q, CL_COMMAND_MARKER, {})); return CL_SUCCESS; @@ -182,22 +181,21 @@ clEnqueueMarker(cl_command_queue d_q, cl_event *rd_ev) try { } PUBLIC cl_int -clEnqueueBarrier(cl_command_queue d_q) { - if (!d_q) - return CL_INVALID_COMMAND_QUEUE; +clEnqueueBarrier(cl_command_queue d_q) try { + obj(d_q); // No need to do anything, q preserves data ordering strictly. return CL_SUCCESS; + +} catch (error &e) { + return e.get(); } PUBLIC cl_int clEnqueueWaitForEvents(cl_command_queue d_q, cl_uint num_evs, const cl_event *d_evs) try { - if (!d_q) - throw error(CL_INVALID_COMMAND_QUEUE); - - auto &q = *d_q; + auto &q = obj(d_q); auto evs = objs(d_evs, num_evs); for (auto &ev : evs) { @@ -260,12 +258,11 @@ clGetEventProfilingInfo(cl_event d_ev, cl_profiling_info param, PUBLIC cl_int clFinish(cl_command_queue d_q) try { - if (!d_q) - throw error(CL_INVALID_COMMAND_QUEUE); + auto &q = obj(d_q); // Create a temporary hard event -- it implicitly depends on all // the previously queued hard events. - ref_ptr hev = transfer(new hard_event(*d_q, 0, { })); + ref_ptr hev = transfer(new hard_event(q, 0, { })); // And wait on it. hev->wait(); diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp index 3335ee6a713..99e090b857d 100644 --- a/src/gallium/state_trackers/clover/api/kernel.cpp +++ b/src/gallium/state_trackers/clover/api/kernel.cpp @@ -204,20 +204,19 @@ namespace { /// Common argument checking shared by kernel invocation commands. /// void - kernel_validate(cl_command_queue q, cl_kernel kern, + kernel_validate(cl_command_queue d_q, cl_kernel kern, cl_uint dims, const size_t *grid_offset, const size_t *grid_size, const size_t *block_size, cl_uint num_deps, const cl_event *deps, cl_event *ev) { - if (!q) - throw error(CL_INVALID_COMMAND_QUEUE); + auto &q = obj(d_q); if (!kern) throw error(CL_INVALID_KERNEL); - if (&kern->prog.ctx != &q->ctx || + if (&kern->prog.ctx != &q.ctx || any_of([&](const cl_event ev) { - return &obj(ev).ctx != &q->ctx; + return &obj(ev).ctx != &q.ctx; }, range(deps, num_deps))) throw error(CL_INVALID_CONTEXT); @@ -230,10 +229,10 @@ namespace { }, kern->args)) throw error(CL_INVALID_KERNEL_ARGS); - if (!kern->prog.binaries().count(&q->dev)) + if (!kern->prog.binaries().count(&q.dev)) throw error(CL_INVALID_PROGRAM_EXECUTABLE); - if (dims < 1 || dims > q->dev.max_block_size().size()) + if (dims < 1 || dims > q.dev.max_block_size().size()) throw error(CL_INVALID_WORK_DIMENSION); if (!grid_size || any_of(is_zero(), range(grid_size, dims))) @@ -243,7 +242,7 @@ namespace { if (any_of([](size_t b, size_t max) { return b == 0 || b > max; }, range(block_size, dims), - q->dev.max_block_size())) + q.dev.max_block_size())) throw error(CL_INVALID_WORK_ITEM_SIZE); if (any_of(modulus(), range(grid_size, dims), @@ -251,7 +250,7 @@ namespace { throw error(CL_INVALID_WORK_GROUP_SIZE); if (fold(multiplies(), 1u, range(block_size, dims)) > - q->dev.max_threads_per_block()) + q.dev.max_threads_per_block()) throw error(CL_INVALID_WORK_GROUP_SIZE); } } @@ -260,15 +259,16 @@ namespace { /// Common event action shared by kernel invocation commands. /// std::function - kernel_op(cl_command_queue q, cl_kernel kern, + kernel_op(cl_command_queue d_q, cl_kernel kern, const std::vector &grid_offset, const std::vector &grid_size, const std::vector &block_size) { + auto &q = obj(d_q); const std::vector reduced_grid_size = map(divides(), grid_size, block_size); - return [=](event &) { - kern->launch(*q, grid_offset, reduced_grid_size, block_size); + return [=, &q](event &) { + kern->launch(q, grid_offset, reduced_grid_size, block_size); }; } @@ -296,7 +296,7 @@ clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern, num_deps, d_deps, ev); hard_event *hev = new hard_event( - *q, CL_COMMAND_NDRANGE_KERNEL, deps, + obj(q), CL_COMMAND_NDRANGE_KERNEL, deps, kernel_op(q, kern, grid_offset, grid_size, block_size)); ret_object(ev, hev); @@ -319,7 +319,7 @@ clEnqueueTask(cl_command_queue q, cl_kernel kern, block_size.data(), num_deps, d_deps, ev); hard_event *hev = new hard_event( - *q, CL_COMMAND_TASK, deps, + obj(q), CL_COMMAND_TASK, deps, kernel_op(q, kern, grid_offset, grid_size, block_size)); ret_object(ev, hev); diff --git a/src/gallium/state_trackers/clover/api/queue.cpp b/src/gallium/state_trackers/clover/api/queue.cpp index ba459250f94..b68dfa1f62a 100644 --- a/src/gallium/state_trackers/clover/api/queue.cpp +++ b/src/gallium/state_trackers/clover/api/queue.cpp @@ -28,7 +28,7 @@ using namespace clover; PUBLIC cl_command_queue clCreateCommandQueue(cl_context d_ctx, cl_device_id d_dev, cl_command_queue_properties props, - cl_int *errcode_ret) try { + cl_int *r_errcode) try { auto &ctx = obj(d_ctx); auto &dev = obj(d_dev); @@ -39,57 +39,55 @@ clCreateCommandQueue(cl_context d_ctx, cl_device_id d_dev, CL_QUEUE_PROFILING_ENABLE)) throw error(CL_INVALID_VALUE); - ret_error(errcode_ret, CL_SUCCESS); + ret_error(r_errcode, CL_SUCCESS); return new command_queue(ctx, dev, props); } catch (error &e) { - ret_error(errcode_ret, e); + ret_error(r_errcode, e); return NULL; } PUBLIC cl_int -clRetainCommandQueue(cl_command_queue q) { - if (!q) - return CL_INVALID_COMMAND_QUEUE; - - q->retain(); +clRetainCommandQueue(cl_command_queue d_q) try { + obj(d_q).retain(); return CL_SUCCESS; + +} catch (error &e) { + return e.get(); } PUBLIC cl_int -clReleaseCommandQueue(cl_command_queue q) { - if (!q) - return CL_INVALID_COMMAND_QUEUE; - - if (q->release()) - delete q; +clReleaseCommandQueue(cl_command_queue d_q) try { + if (obj(d_q).release()) + delete pobj(d_q); return CL_SUCCESS; + +} catch (error &e) { + return e.get(); } PUBLIC cl_int -clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param, +clGetCommandQueueInfo(cl_command_queue d_q, cl_command_queue_info param, size_t size, void *r_buf, size_t *r_size) try { property_buffer buf { r_buf, size, r_size }; - - if (!q) - return CL_INVALID_COMMAND_QUEUE; + auto &q = obj(d_q); switch (param) { case CL_QUEUE_CONTEXT: - buf.as_scalar() = &q->ctx; + buf.as_scalar() = desc(q.ctx); break; case CL_QUEUE_DEVICE: - buf.as_scalar() = &q->dev; + buf.as_scalar() = desc(q.dev); break; case CL_QUEUE_REFERENCE_COUNT: - buf.as_scalar() = q->ref_count(); + buf.as_scalar() = q.ref_count(); break; case CL_QUEUE_PROPERTIES: - buf.as_scalar() = q->props(); + buf.as_scalar() = q.props(); break; default: @@ -103,10 +101,10 @@ clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param, } PUBLIC cl_int -clFlush(cl_command_queue q) { - if (!q) - return CL_INVALID_COMMAND_QUEUE; - - q->flush(); +clFlush(cl_command_queue d_q) try { + obj(d_q).flush(); return CL_SUCCESS; + +} catch (error &e) { + return e.get(); } diff --git a/src/gallium/state_trackers/clover/api/transfer.cpp b/src/gallium/state_trackers/clover/api/transfer.cpp index 62f9d326ddf..f91da617b68 100644 --- a/src/gallium/state_trackers/clover/api/transfer.cpp +++ b/src/gallium/state_trackers/clover/api/transfer.cpp @@ -40,16 +40,13 @@ namespace { /// Common argument checking shared by memory transfer commands. /// void - validate_base(cl_command_queue q, cl_uint num_deps, const cl_event *deps) { - if (!q) - throw error(CL_INVALID_COMMAND_QUEUE); - + validate_base(command_queue &q, cl_uint num_deps, const cl_event *deps) { if (bool(num_deps) != bool(deps) || any_of(is_zero(), range(deps, num_deps))) throw error(CL_INVALID_EVENT_WAIT_LIST); if (any_of([&](const cl_event ev) { - return &obj(ev).ctx != &q->ctx; + return &obj(ev).ctx != &q.ctx; }, range(deps, num_deps))) throw error(CL_INVALID_CONTEXT); } @@ -59,11 +56,11 @@ namespace { /// transfer commands. /// void - validate_obj(cl_command_queue q, cl_mem obj) { - if (!obj) + validate_obj(command_queue &q, cl_mem mem) { + if (!mem) throw error(CL_INVALID_MEM_OBJECT); - if (&obj->ctx != &q->ctx) + if (&mem->ctx != &q.ctx) throw error(CL_INVALID_CONTEXT); } @@ -92,9 +89,9 @@ namespace { template<> struct _map { static mapping - get(cl_command_queue q, memory_obj *obj, cl_map_flags flags, - size_t offset, size_t size) { - return { *q, obj->resource(q), flags, true, + get(cl_command_queue q, memory_obj *mem, cl_map_flags flags, + size_t offset, size_t size) {< + return { obj(q), mem->resource(obj(q)), flags, true, {{ offset }}, {{ size, 1, 1 }}}; } }; @@ -134,36 +131,37 @@ namespace { /// template std::function - hard_copy_op(cl_command_queue q, T dst_obj, const vector_t &dst_orig, + hard_copy_op(command_queue &q, T dst_obj, const vector_t &dst_orig, S src_obj, const vector_t &src_orig, const vector_t ®ion) { - return [=](event &) { - dst_obj->resource(q).copy(*q, dst_orig, region, + return [=, &q](event &) { + dst_obj->resource(q).copy(q, dst_orig, region, src_obj->resource(q), src_orig); }; } } PUBLIC cl_int -clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueReadBuffer(cl_command_queue d_q, cl_mem mem, cl_bool blocking, size_t offset, size_t size, void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); - if (!ptr || offset > obj->size() || offset + size > obj->size()) + if (!ptr || offset > mem->size() || offset + size > mem->size()) throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_READ_BUFFER, deps, - soft_copy_op(q, + q, CL_COMMAND_READ_BUFFER, deps, + soft_copy_op(d_q, ptr, {{ 0 }}, {{ 1 }}, - obj, {{ offset }}, {{ 1 }}, + mem, {{ offset }}, {{ 1 }}, {{ size, 1, 1 }})); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -171,26 +169,27 @@ clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueWriteBuffer(cl_command_queue d_q, cl_mem mem, cl_bool blocking, size_t offset, size_t size, const void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); - if (!ptr || offset > obj->size() || offset + size > obj->size()) + if (!ptr || offset > mem->size() || offset + size > mem->size()) throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_WRITE_BUFFER, deps, - soft_copy_op(q, - obj, {{ offset }}, {{ 1 }}, + q, CL_COMMAND_WRITE_BUFFER, deps, + soft_copy_op(d_q, + mem, {{ offset }}, {{ 1 }}, ptr, {{ 0 }}, {{ 1 }}, {{ size, 1, 1 }})); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -198,32 +197,34 @@ clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, - const size_t *obj_origin, const size_t *host_origin, +clEnqueueReadBufferRect(cl_command_queue d_q, cl_mem mem, cl_bool blocking, + const size_t *obj_origin, + const size_t *host_origin, const size_t *region, size_t obj_row_pitch, size_t obj_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); if (!ptr) throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_READ_BUFFER_RECT, deps, - soft_copy_op(q, + q, CL_COMMAND_READ_BUFFER_RECT, deps, + soft_copy_op(d_q, ptr, vector(host_origin), {{ 1, host_row_pitch, host_slice_pitch }}, - obj, vector(obj_origin), + mem, vector(obj_origin), {{ 1, obj_row_pitch, obj_slice_pitch }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -231,32 +232,34 @@ clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, - const size_t *obj_origin, const size_t *host_origin, +clEnqueueWriteBufferRect(cl_command_queue d_q, cl_mem mem, cl_bool blocking, + const size_t *obj_origin, + const size_t *host_origin, const size_t *region, size_t obj_row_pitch, size_t obj_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, const void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); if (!ptr) throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_WRITE_BUFFER_RECT, deps, - soft_copy_op(q, - obj, vector(obj_origin), + q, CL_COMMAND_WRITE_BUFFER_RECT, deps, + soft_copy_op(d_q, + mem, vector(obj_origin), {{ 1, obj_row_pitch, obj_slice_pitch }}, ptr, vector(host_origin), {{ 1, host_row_pitch, host_slice_pitch }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -264,23 +267,24 @@ clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, +clEnqueueCopyBuffer(cl_command_queue d_q, cl_mem src_mem, cl_mem dst_mem, size_t src_offset, size_t dst_offset, size_t size, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, src_obj); - validate_obj(q, dst_obj); + validate_obj(q, src_mem); + validate_obj(q, dst_mem); hard_event *hev = new hard_event( - *q, CL_COMMAND_COPY_BUFFER, deps, - hard_copy_op(q, dst_obj, {{ dst_offset }}, - src_obj, {{ src_offset }}, + q, CL_COMMAND_COPY_BUFFER, deps, + hard_copy_op(q, dst_mem, {{ dst_offset }}, + src_mem, {{ src_offset }}, {{ size, 1, 1 }})); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -288,29 +292,31 @@ clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, } PUBLIC cl_int -clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, +clEnqueueCopyBufferRect(cl_command_queue d_q, cl_mem src_mem, + cl_mem dst_mem, const size_t *src_origin, const size_t *dst_origin, const size_t *region, size_t src_row_pitch, size_t src_slice_pitch, size_t dst_row_pitch, size_t dst_slice_pitch, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, src_obj); - validate_obj(q, dst_obj); + validate_obj(q, src_mem); + validate_obj(q, dst_mem); hard_event *hev = new hard_event( - *q, CL_COMMAND_COPY_BUFFER_RECT, deps, - soft_copy_op(q, - dst_obj, vector(dst_origin), + q, CL_COMMAND_COPY_BUFFER_RECT, deps, + soft_copy_op(d_q, + dst_mem, vector(dst_origin), {{ 1, dst_row_pitch, dst_slice_pitch }}, - src_obj, vector(src_origin), + src_mem, vector(src_origin), {{ 1, src_row_pitch, src_slice_pitch }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -318,13 +324,14 @@ clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, } PUBLIC cl_int -clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueReadImage(cl_command_queue d_q, cl_mem mem, cl_bool blocking, const size_t *origin, const size_t *region, size_t row_pitch, size_t slice_pitch, void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *img = dynamic_cast(obj); + image *img = dynamic_cast(mem); validate_base(q, num_deps, d_deps); validate_obj(q, img); @@ -333,15 +340,15 @@ clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking, throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_READ_IMAGE, deps, - soft_copy_op(q, + q, CL_COMMAND_READ_IMAGE, deps, + soft_copy_op(d_q, ptr, {}, {{ 1, row_pitch, slice_pitch }}, - obj, vector(origin), + mem, vector(origin), {{ 1, img->row_pitch(), img->slice_pitch() }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -349,13 +356,14 @@ clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueWriteImage(cl_command_queue d_q, cl_mem mem, cl_bool blocking, const size_t *origin, const size_t *region, size_t row_pitch, size_t slice_pitch, const void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *img = dynamic_cast(obj); + image *img = dynamic_cast(mem); validate_base(q, num_deps, d_deps); validate_obj(q, img); @@ -364,15 +372,15 @@ clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking, throw error(CL_INVALID_VALUE); hard_event *hev = new hard_event( - *q, CL_COMMAND_WRITE_IMAGE, deps, - soft_copy_op(q, - obj, vector(origin), + q, CL_COMMAND_WRITE_IMAGE, deps, + soft_copy_op(d_q, + mem, vector(origin), {{ 1, img->row_pitch(), img->slice_pitch() }}, ptr, {}, {{ 1, row_pitch, slice_pitch }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -380,27 +388,28 @@ clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking, } PUBLIC cl_int -clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, +clEnqueueCopyImage(cl_command_queue d_q, cl_mem src_mem, cl_mem dst_mem, const size_t *src_origin, const size_t *dst_origin, const size_t *region, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *src_img = dynamic_cast(src_obj); - image *dst_img = dynamic_cast(dst_obj); + image *src_img = dynamic_cast(src_mem); + image *dst_img = dynamic_cast(dst_mem); validate_base(q, num_deps, d_deps); validate_obj(q, src_img); validate_obj(q, dst_img); hard_event *hev = new hard_event( - *q, CL_COMMAND_COPY_IMAGE, deps, + q, CL_COMMAND_COPY_IMAGE, deps, hard_copy_op(q, - dst_obj, vector(dst_origin), - src_obj, vector(src_origin), + dst_img, vector(dst_origin), + src_img, vector(src_origin), vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -408,28 +417,30 @@ clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, } PUBLIC cl_int -clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, +clEnqueueCopyImageToBuffer(cl_command_queue d_q, + cl_mem src_mem, cl_mem dst_mem, const size_t *src_origin, const size_t *region, size_t dst_offset, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *src_img = dynamic_cast(src_obj); + image *src_img = dynamic_cast(src_mem); validate_base(q, num_deps, d_deps); validate_obj(q, src_img); - validate_obj(q, dst_obj); + validate_obj(q, dst_mem); hard_event *hev = new hard_event( - *q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, deps, - soft_copy_op(q, - dst_obj, {{ dst_offset }}, + q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, deps, + soft_copy_op(d_q, + dst_mem, {{ dst_offset }}, {{ 0, 0, 0 }}, - src_obj, vector(src_origin), + src_mem, vector(src_origin), {{ 1, src_img->row_pitch(), src_img->slice_pitch() }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -437,28 +448,30 @@ clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, } PUBLIC cl_int -clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, +clEnqueueCopyBufferToImage(cl_command_queue d_q, + cl_mem src_mem, cl_mem dst_mem, size_t src_offset, const size_t *dst_origin, const size_t *region, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *dst_img = dynamic_cast(dst_obj); + image *dst_img = dynamic_cast(dst_mem); validate_base(q, num_deps, d_deps); - validate_obj(q, src_obj); + validate_obj(q, src_mem); validate_obj(q, dst_img); hard_event *hev = new hard_event( - *q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, deps, - soft_copy_op(q, - dst_obj, vector(dst_origin), + q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, deps, + soft_copy_op(d_q, + dst_mem, vector(dst_origin), {{ 1, dst_img->row_pitch(), dst_img->slice_pitch() }}, - src_obj, {{ src_offset }}, + src_mem, {{ src_offset }}, {{ 0, 0, 0 }}, vector(region))); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { @@ -466,72 +479,74 @@ clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, } PUBLIC void * -clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueMapBuffer(cl_command_queue d_q, cl_mem mem, cl_bool blocking, cl_map_flags flags, size_t offset, size_t size, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev, cl_int *errcode_ret) try { + cl_event *rd_ev, cl_int *r_errcode) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); + validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); - if (offset > obj->size() || offset + size > obj->size()) + if (offset > mem->size() || offset + size > mem->size()) throw error(CL_INVALID_VALUE); - void *map = obj->resource(q).add_map( - *q, flags, blocking, {{ offset }}, {{ size }}); + void *map = mem->resource(q).add_map( + q, flags, blocking, {{ offset }}, {{ size }}); - ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER, - deps)); - ret_error(errcode_ret, CL_SUCCESS); + ret_object(rd_ev, new hard_event(q, CL_COMMAND_MAP_BUFFER, deps)); + ret_error(r_errcode, CL_SUCCESS); return map; } catch (error &e) { - ret_error(errcode_ret, e); + ret_error(r_errcode, e); return NULL; } PUBLIC void * -clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking, +clEnqueueMapImage(cl_command_queue d_q, cl_mem mem, cl_bool blocking, cl_map_flags flags, const size_t *origin, const size_t *region, size_t *row_pitch, size_t *slice_pitch, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev, cl_int *errcode_ret) try { + cl_event *rd_ev, cl_int *r_errcode) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); - image *img = dynamic_cast(obj); + image *img = dynamic_cast(mem); validate_base(q, num_deps, d_deps); validate_obj(q, img); - void *map = obj->resource(q).add_map( - *q, flags, blocking, vector(origin), vector(region)); + void *map = img->resource(q).add_map( + q, flags, blocking, vector(origin), vector(region)); - ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE, - deps)); - ret_error(errcode_ret, CL_SUCCESS); + ret_object(rd_ev, new hard_event(q, CL_COMMAND_MAP_IMAGE, deps)); + ret_error(r_errcode, CL_SUCCESS); return map; } catch (error &e) { - ret_error(errcode_ret, e); + ret_error(r_errcode, e); return NULL; } PUBLIC cl_int -clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr, +clEnqueueUnmapMemObject(cl_command_queue d_q, cl_mem mem, void *ptr, cl_uint num_deps, const cl_event *d_deps, - cl_event *ev) try { + cl_event *rd_ev) try { + auto &q = obj(d_q); auto deps = objs(d_deps, num_deps); validate_base(q, num_deps, d_deps); - validate_obj(q, obj); + validate_obj(q, mem); hard_event *hev = new hard_event( - *q, CL_COMMAND_UNMAP_MEM_OBJECT, deps, - [=](event &) { - obj->resource(q).del_map(ptr); + q, CL_COMMAND_UNMAP_MEM_OBJECT, deps, + [=, &q, &mem](event &) { + mem->resource(q).del_map(ptr); }); - ret_object(ev, hev); + ret_object(rd_ev, hev); return CL_SUCCESS; } catch (error &e) { diff --git a/src/gallium/state_trackers/clover/core/device.hpp b/src/gallium/state_trackers/clover/core/device.hpp index c490b20edc2..95669bdeb64 100644 --- a/src/gallium/state_trackers/clover/core/device.hpp +++ b/src/gallium/state_trackers/clover/core/device.hpp @@ -66,7 +66,7 @@ namespace clover { std::string ir_target() const; enum pipe_endian endianness() const; - friend struct ::_cl_command_queue; + friend class command_queue; friend class root_resource; friend class hard_event; friend std::set diff --git a/src/gallium/state_trackers/clover/core/error.hpp b/src/gallium/state_trackers/clover/core/error.hpp index bc9f161c4db..fa43c1a5eed 100644 --- a/src/gallium/state_trackers/clover/core/error.hpp +++ b/src/gallium/state_trackers/clover/core/error.hpp @@ -28,7 +28,7 @@ #include "util/compat.hpp" namespace clover { - typedef struct _cl_command_queue command_queue; + class command_queue; class context; class device; class event; diff --git a/src/gallium/state_trackers/clover/core/event.hpp b/src/gallium/state_trackers/clover/core/event.hpp index c4e1bb71de0..ee25d19c1e5 100644 --- a/src/gallium/state_trackers/clover/core/event.hpp +++ b/src/gallium/state_trackers/clover/core/event.hpp @@ -110,7 +110,7 @@ namespace clover { const lazy &time_start() const; const lazy &time_end() const; - friend struct ::_cl_command_queue; + friend class command_queue; private: virtual void fence(pipe_fence_handle *fence); diff --git a/src/gallium/state_trackers/clover/core/kernel.cpp b/src/gallium/state_trackers/clover/core/kernel.cpp index 1a75cd48a00..5663f1f8b2e 100644 --- a/src/gallium/state_trackers/clover/core/kernel.cpp +++ b/src/gallium/state_trackers/clover/core/kernel.cpp @@ -318,7 +318,7 @@ _cl_kernel::global_argument::bind(exec_context &ctx, const clover::module::argument &marg) { align(ctx.input, marg.target_align); ctx.g_handles.push_back(allocate(ctx.input, marg.target_size)); - ctx.g_buffers.push_back(obj->resource(ctx.q).pipe); + ctx.g_buffers.push_back(obj->resource(*ctx.q).pipe); } void @@ -378,13 +378,13 @@ _cl_kernel::constant_argument::bind(exec_context &ctx, align(ctx.input, marg.target_align); insert(ctx.input, v); - st = obj->resource(ctx.q).bind_surface(*ctx.q, false); + st = obj->resource(*ctx.q).bind_surface(*ctx.q, false); ctx.resources.push_back(st); } void _cl_kernel::constant_argument::unbind(exec_context &ctx) { - obj->resource(ctx.q).unbind_surface(*ctx.q, st); + obj->resource(*ctx.q).unbind_surface(*ctx.q, st); } void @@ -409,13 +409,13 @@ _cl_kernel::image_rd_argument::bind(exec_context &ctx, align(ctx.input, marg.target_align); insert(ctx.input, v); - st = obj->resource(ctx.q).bind_sampler_view(*ctx.q); + st = obj->resource(*ctx.q).bind_sampler_view(*ctx.q); ctx.sviews.push_back(st); } void _cl_kernel::image_rd_argument::unbind(exec_context &ctx) { - obj->resource(ctx.q).unbind_sampler_view(*ctx.q, st); + obj->resource(*ctx.q).unbind_sampler_view(*ctx.q, st); } void @@ -440,13 +440,13 @@ _cl_kernel::image_wr_argument::bind(exec_context &ctx, align(ctx.input, marg.target_align); insert(ctx.input, v); - st = obj->resource(ctx.q).bind_surface(*ctx.q, true); + st = obj->resource(*ctx.q).bind_surface(*ctx.q, true); ctx.resources.push_back(st); } void _cl_kernel::image_wr_argument::unbind(exec_context &ctx) { - obj->resource(ctx.q).unbind_surface(*ctx.q, st); + obj->resource(*ctx.q).unbind_surface(*ctx.q, st); } void diff --git a/src/gallium/state_trackers/clover/core/memory.cpp b/src/gallium/state_trackers/clover/core/memory.cpp index 742e3d3106a..38da67b8fef 100644 --- a/src/gallium/state_trackers/clover/core/memory.cpp +++ b/src/gallium/state_trackers/clover/core/memory.cpp @@ -74,19 +74,19 @@ root_buffer::root_buffer(clover::context &ctx, cl_mem_flags flags, } clover::resource & -root_buffer::resource(cl_command_queue q) { +root_buffer::resource(command_queue &q) { // Create a new resource if there's none for this device yet. - if (!resources.count(&q->dev)) { + if (!resources.count(&q.dev)) { auto r = (!resources.empty() ? - new root_resource(q->dev, *this, *resources.begin()->second) : - new root_resource(q->dev, *this, *q, data)); + new root_resource(q.dev, *this, *resources.begin()->second) : + new root_resource(q.dev, *this, q, data)); - resources.insert(std::make_pair(&q->dev, + resources.insert(std::make_pair(&q.dev, std::unique_ptr(r))); data.clear(); } - return *resources.find(&q->dev)->second; + return *resources.find(&q.dev)->second; } sub_buffer::sub_buffer(clover::root_buffer &parent, cl_mem_flags flags, @@ -97,16 +97,16 @@ sub_buffer::sub_buffer(clover::root_buffer &parent, cl_mem_flags flags, } clover::resource & -sub_buffer::resource(cl_command_queue q) { +sub_buffer::resource(command_queue &q) { // Create a new resource if there's none for this device yet. - if (!resources.count(&q->dev)) { + if (!resources.count(&q.dev)) { auto r = new sub_resource(parent.resource(q), {{ offset() }}); - resources.insert(std::make_pair(&q->dev, + resources.insert(std::make_pair(&q.dev, std::unique_ptr(r))); } - return *resources.find(&q->dev)->second; + return *resources.find(&q.dev)->second; } size_t @@ -125,19 +125,19 @@ image::image(clover::context &ctx, cl_mem_flags flags, } clover::resource & -image::resource(cl_command_queue q) { +image::resource(command_queue &q) { // Create a new resource if there's none for this device yet. - if (!resources.count(&q->dev)) { + if (!resources.count(&q.dev)) { auto r = (!resources.empty() ? - new root_resource(q->dev, *this, *resources.begin()->second) : - new root_resource(q->dev, *this, *q, data)); + new root_resource(q.dev, *this, *resources.begin()->second) : + new root_resource(q.dev, *this, q, data)); - resources.insert(std::make_pair(&q->dev, + resources.insert(std::make_pair(&q.dev, std::unique_ptr(r))); data.clear(); } - return *resources.find(&q->dev)->second; + return *resources.find(&q.dev)->second; } cl_image_format diff --git a/src/gallium/state_trackers/clover/core/memory.hpp b/src/gallium/state_trackers/clover/core/memory.hpp index f495c45ed33..de0b531a9d3 100644 --- a/src/gallium/state_trackers/clover/core/memory.hpp +++ b/src/gallium/state_trackers/clover/core/memory.hpp @@ -47,7 +47,7 @@ public: virtual ~_cl_mem(); virtual cl_mem_object_type type() const = 0; - virtual clover::resource &resource(cl_command_queue q) = 0; + virtual clover::resource &resource(clover::command_queue &q) = 0; void destroy_notify(std::function f); cl_mem_flags flags() const; @@ -81,7 +81,7 @@ namespace clover { root_buffer(clover::context &ctx, cl_mem_flags flags, size_t size, void *host_ptr); - virtual clover::resource &resource(cl_command_queue q); + virtual clover::resource &resource(clover::command_queue &q); private: std::map {}; +struct _cl_command_queue : + public clover::descriptor {}; + #endif diff --git a/src/gallium/state_trackers/clover/core/queue.cpp b/src/gallium/state_trackers/clover/core/queue.cpp index 62a59f8350b..084e3c3f734 100644 --- a/src/gallium/state_trackers/clover/core/queue.cpp +++ b/src/gallium/state_trackers/clover/core/queue.cpp @@ -29,20 +29,20 @@ using namespace clover; -_cl_command_queue::_cl_command_queue(context &ctx, device &dev, - cl_command_queue_properties props) : +command_queue::command_queue(context &ctx, device &dev, + cl_command_queue_properties props) : ctx(ctx), dev(dev), _props(props) { pipe = dev.pipe->context_create(dev.pipe, NULL); if (!pipe) throw error(CL_INVALID_DEVICE); } -_cl_command_queue::~_cl_command_queue() { +command_queue::~command_queue() { pipe->destroy(pipe); } void -_cl_command_queue::flush() { +command_queue::flush() { pipe_screen *screen = dev.pipe; pipe_fence_handle *fence = NULL; @@ -61,17 +61,17 @@ _cl_command_queue::flush() { } cl_command_queue_properties -_cl_command_queue::props() const { +command_queue::props() const { return _props; } bool -_cl_command_queue::profiling_enabled() const { +command_queue::profiling_enabled() const { return _props & CL_QUEUE_PROFILING_ENABLE; } void -_cl_command_queue::sequence(clover::hard_event *ev) { +command_queue::sequence(hard_event *ev) { if (!queued_events.empty()) queued_events.back()->chain(ev); diff --git a/src/gallium/state_trackers/clover/core/queue.hpp b/src/gallium/state_trackers/clover/core/queue.hpp index a02de95cf8f..4a2d02251b1 100644 --- a/src/gallium/state_trackers/clover/core/queue.hpp +++ b/src/gallium/state_trackers/clover/core/queue.hpp @@ -29,46 +29,45 @@ #include "pipe/p_context.h" namespace clover { - typedef struct _cl_command_queue command_queue; class resource; class mapping; class hard_event; -} -struct _cl_command_queue : public clover::ref_counter { -public: - _cl_command_queue(clover::context &ctx, clover::device &dev, - cl_command_queue_properties props); - _cl_command_queue(const _cl_command_queue &q) = delete; - ~_cl_command_queue(); + class command_queue : public ref_counter, public _cl_command_queue { + public: + command_queue(context &ctx, device &dev, + cl_command_queue_properties props); + command_queue(const command_queue &q) = delete; + ~command_queue(); - void flush(); + void flush(); - cl_command_queue_properties props() const; - bool profiling_enabled() const; + cl_command_queue_properties props() const; + bool profiling_enabled() const; - clover::context &ctx; - clover::device &dev; + context &ctx; + device &dev; - friend class clover::resource; - friend class clover::root_resource; - friend class clover::mapping; - friend class clover::hard_event; - friend struct _cl_sampler; - friend struct _cl_kernel; - friend class clover::timestamp::query; - friend class clover::timestamp::current; + friend class resource; + friend class root_resource; + friend class mapping; + friend class hard_event; + friend struct ::_cl_sampler; + friend struct ::_cl_kernel; + friend class clover::timestamp::query; + friend class clover::timestamp::current; -private: - /// Serialize a hardware event with respect to the previous ones, - /// and push it to the pending list. - void sequence(clover::hard_event *ev); + private: + /// Serialize a hardware event with respect to the previous ones, + /// and push it to the pending list. + void sequence(hard_event *ev); - cl_command_queue_properties _props; - pipe_context *pipe; + cl_command_queue_properties _props; + pipe_context *pipe; - typedef clover::ref_ptr event_ptr; - std::vector queued_events; -}; + typedef ref_ptr event_ptr; + std::vector queued_events; + }; +} #endif diff --git a/src/gallium/state_trackers/clover/core/timestamp.hpp b/src/gallium/state_trackers/clover/core/timestamp.hpp index bf9e204108b..11c8ef0195d 100644 --- a/src/gallium/state_trackers/clover/core/timestamp.hpp +++ b/src/gallium/state_trackers/clover/core/timestamp.hpp @@ -28,7 +28,7 @@ struct pipe_query; namespace clover { - typedef struct _cl_command_queue command_queue; + class command_queue; namespace timestamp { /// -- 2.30.2