clover: Switch command queues to the new model.

Tested-by: Tom Stellard <[email protected]>

4 files changed, 194 insertions, 184 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<cl_command_queue>() = ev.queue();
+      buf.as_scalar<cl_command_queue>() = 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<hard_event> hev = transfer(new hard_event(*d_q, 0, { }));
+   ref_ptr<hard_event> 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<void (event &)>
-   kernel_op(cl_command_queue q, cl_kernel kern,
+   kernel_op(cl_command_queue d_q, cl_kernel kern,
              const std::vector<size_t> &grid_offset,
              const std::vector<size_t> &grid_size,
              const std::vector<size_t> &block_size) {
+      auto &q = obj(d_q);
       const std::vector<size_t> 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<cl_context>() = &q->ctx;
+      buf.as_scalar<cl_context>() = desc(q.ctx);
       break;
 
    case CL_QUEUE_DEVICE:
-      buf.as_scalar<cl_device_id>() = &q->dev;
+      buf.as_scalar<cl_device_id>() = desc(q.dev);
       break;
 
    case CL_QUEUE_REFERENCE_COUNT:
-      buf.as_scalar<cl_uint>() = q->ref_count();
+      buf.as_scalar<cl_uint>() = q.ref_count();
       break;
 
    case CL_QUEUE_PROPERTIES:
-      buf.as_scalar<cl_command_queue_properties>() = q->props();
+      buf.as_scalar<cl_command_queue_properties>() = 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<memory_obj *> {
       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<typename T, typename S>
    std::function<void (event &)>
-   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 &region) {
-      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<wait_list_tag>(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<wait_list_tag>(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<wait_list_tag>(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<wait_list_tag>(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<wait_list_tag>(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<wait_list_tag>(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<wait_list_tag>(d_deps, num_deps);
-   image *img = dynamic_cast<image *>(obj);
+   image *img = dynamic_cast<image *>(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<wait_list_tag>(d_deps, num_deps);
-   image *img = dynamic_cast<image *>(obj);
+   image *img = dynamic_cast<image *>(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<wait_list_tag>(d_deps, num_deps);
-   image *src_img = dynamic_cast<image *>(src_obj);
-   image *dst_img = dynamic_cast<image *>(dst_obj);
+   image *src_img = dynamic_cast<image *>(src_mem);
+   image *dst_img = dynamic_cast<image *>(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<wait_list_tag>(d_deps, num_deps);
-   image *src_img = dynamic_cast<image *>(src_obj);
+   image *src_img = dynamic_cast<image *>(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<wait_list_tag>(d_deps, num_deps);
-   image *dst_img = dynamic_cast<image *>(dst_obj);
+   image *dst_img = dynamic_cast<image *>(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<wait_list_tag>(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<wait_list_tag>(d_deps, num_deps);
-   image *img = dynamic_cast<image *>(obj);
+   image *img = dynamic_cast<image *>(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<wait_list_tag>(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) {