clover: Switch command queues to the new model.

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

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<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) {
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<cl_image_format>
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<cl_ulong> &time_start() const;
       const lazy<cl_ulong> &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<root_resource>(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<sub_resource>(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<root_resource>(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<void ()> 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<clover::device *,
@@ -93,7 +93,7 @@ namespace clover {
       sub_buffer(clover::root_buffer &parent, cl_mem_flags flags,
                  size_t offset, size_t size);
 
-      virtual clover::resource &resource(cl_command_queue q);
+      virtual clover::resource &resource(clover::command_queue &q);
       size_t offset() const;
 
       clover::root_buffer &parent;
@@ -113,7 +113,7 @@ namespace clover {
             void *host_ptr);
 
    public:
-      virtual clover::resource &resource(cl_command_queue q);
+      virtual clover::resource &resource(clover::command_queue &q);
       cl_image_format format() const;
       size_t width() const;
       size_t height() const;
diff --git a/src/gallium/state_trackers/clover/core/object.hpp b/src/gallium/state_trackers/clover/core/object.hpp
index 6a916b25bef..6a99f19bd1e 100644
--- a/src/gallium/state_trackers/clover/core/object.hpp
+++ b/src/gallium/state_trackers/clover/core/object.hpp
@@ -191,4 +191,7 @@ struct _cl_event :
 struct _cl_platform_id :
    public clover::descriptor<clover::platform, _cl_platform_id> {};
 
+struct _cl_command_queue :
+   public clover::descriptor<clover::command_queue, _cl_command_queue> {};
+
 #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<clover::hard_event> event_ptr;
-   std::vector<event_ptr> queued_events;
-};
+      typedef ref_ptr<hard_event> event_ptr;
+      std::vector<event_ptr> 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 {
       ///