clover: Switch kernel and program objects to the new model.

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

11 files changed, 458 insertions, 492 deletions

diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp
index 99e090b857d..15b4c14e1d9 100644
--- a/src/gallium/state_trackers/clover/api/kernel.cpp
+++ b/src/gallium/state_trackers/clover/api/kernel.cpp
@@ -27,122 +27,119 @@
 using namespace clover;
 
 PUBLIC cl_kernel
-clCreateKernel(cl_program prog, const char *name,
-               cl_int *errcode_ret) try {
-   if (!prog)
-      throw error(CL_INVALID_PROGRAM);
+clCreateKernel(cl_program d_prog, const char *name, cl_int *r_errcode) try {
+   auto &prog = obj(d_prog);
 
    if (!name)
       throw error(CL_INVALID_VALUE);
 
-   if (prog->binaries().empty())
+   if (prog.binaries().empty())
       throw error(CL_INVALID_PROGRAM_EXECUTABLE);
 
-   auto sym = prog->binaries().begin()->second.sym(name);
+   auto sym = prog.binaries().begin()->second.sym(name);
 
-   ret_error(errcode_ret, CL_SUCCESS);
-   return new kernel(*prog, name, { sym.args.begin(), sym.args.end() });
+   ret_error(r_errcode, CL_SUCCESS);
+   return new kernel(prog, name, range(sym.args));
 
 } catch (module::noent_error &e) {
-   ret_error(errcode_ret, CL_INVALID_KERNEL_NAME);
+   ret_error(r_errcode, CL_INVALID_KERNEL_NAME);
    return NULL;
 
-} catch(error &e) {
-   ret_error(errcode_ret, e);
+} catch (error &e) {
+   ret_error(r_errcode, e);
    return NULL;
 }
 
 PUBLIC cl_int
-clCreateKernelsInProgram(cl_program prog, cl_uint count,
-                         cl_kernel *kerns, cl_uint *count_ret) {
-   if (!prog)
-      throw error(CL_INVALID_PROGRAM);
+clCreateKernelsInProgram(cl_program d_prog, cl_uint count,
+                         cl_kernel *rd_kerns, cl_uint *r_count) try {
+   auto &prog = obj(d_prog);
 
-   if (prog->binaries().empty())
+   if (prog.binaries().empty())
       throw error(CL_INVALID_PROGRAM_EXECUTABLE);
 
-   auto &syms = prog->binaries().begin()->second.syms;
+   auto &syms = prog.binaries().begin()->second.syms;
 
-   if (kerns && count < syms.size())
+   if (rd_kerns && count < syms.size())
       throw error(CL_INVALID_VALUE);
 
-   if (kerns)
-      std::transform(syms.begin(), syms.end(), kerns,
-                     [=](const module::symbol &sym) {
-                        return new kernel(*prog, compat::string(sym.name),
-                                          { sym.args.begin(), sym.args.end() });
-                     });
+   if (rd_kerns)
+      copy(map([&](const module::symbol &sym) {
+               return desc(new kernel(prog, compat::string(sym.name),
+                                      range(sym.args)));
+            }, syms),
+         rd_kerns);
 
-   if (count_ret)
-      *count_ret = syms.size();
+   if (r_count)
+      *r_count = syms.size();
 
    return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
 }
 
 PUBLIC cl_int
-clRetainKernel(cl_kernel kern) {
-   if (!kern)
-      return CL_INVALID_KERNEL;
-
-   kern->retain();
+clRetainKernel(cl_kernel d_kern) try {
+   obj(d_kern).retain();
    return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
 }
 
 PUBLIC cl_int
-clReleaseKernel(cl_kernel kern) {
-   if (!kern)
-      return CL_INVALID_KERNEL;
-
-   if (kern->release())
-      delete kern;
+clReleaseKernel(cl_kernel d_kern) try {
+   if (obj(d_kern).release())
+      delete pobj(d_kern);
 
    return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
 }
 
 PUBLIC cl_int
-clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size,
+clSetKernelArg(cl_kernel d_kern, cl_uint idx, size_t size,
                const void *value) try {
-   if (!kern)
-      throw error(CL_INVALID_KERNEL);
+   auto &kern = obj(d_kern);
 
-   if (idx >= kern->args.size())
+   if (idx >= kern.args.size())
       throw error(CL_INVALID_ARG_INDEX);
 
-   kern->args[idx]->set(size, value);
+   kern.args[idx]->set(size, value);
 
    return CL_SUCCESS;
 
-} catch(error &e) {
+} catch (error &e) {
    return e.get();
 }
 
 PUBLIC cl_int
-clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
+clGetKernelInfo(cl_kernel d_kern, cl_kernel_info param,
                 size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
-
-   if (!kern)
-      return CL_INVALID_KERNEL;
+   auto &kern = obj(d_kern);
 
    switch (param) {
    case CL_KERNEL_FUNCTION_NAME:
-      buf.as_string() = kern->name();
+      buf.as_string() = kern.name();
       break;
 
    case CL_KERNEL_NUM_ARGS:
-      buf.as_scalar<cl_uint>() = kern->args.size();
+      buf.as_scalar<cl_uint>() = kern.args.size();
       break;
 
    case CL_KERNEL_REFERENCE_COUNT:
-      buf.as_scalar<cl_uint>() = kern->ref_count();
+      buf.as_scalar<cl_uint>() = kern.ref_count();
       break;
 
    case CL_KERNEL_CONTEXT:
-      buf.as_scalar<cl_context>() = &kern->prog.ctx;
+      buf.as_scalar<cl_context>() = desc(kern.prog.ctx);
       break;
 
    case CL_KERNEL_PROGRAM:
-      buf.as_scalar<cl_program>() = &kern->prog;
+      buf.as_scalar<cl_program>() = desc(kern.prog);
       break;
 
    default:
@@ -156,29 +153,28 @@ clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
 }
 
 PUBLIC cl_int
-clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev,
+clGetKernelWorkGroupInfo(cl_kernel d_kern, cl_device_id d_dev,
                          cl_kernel_work_group_info param,
                          size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
+   auto &kern = obj(d_kern);
+   auto pdev = pobj(d_dev);
 
-   if (!kern)
-      return CL_INVALID_KERNEL;
-
-   if ((!dev && kern->prog.binaries().size() != 1) ||
-       (dev && !kern->prog.binaries().count(pobj(dev))))
-      return CL_INVALID_DEVICE;
+   if ((!pdev && kern.prog.binaries().size() != 1) ||
+       (pdev && !kern.prog.binaries().count(pdev)))
+      throw error(CL_INVALID_DEVICE);
 
    switch (param) {
    case CL_KERNEL_WORK_GROUP_SIZE:
-      buf.as_scalar<size_t>() = kern->max_block_size();
+      buf.as_scalar<size_t>() = kern.max_block_size();
       break;
 
    case CL_KERNEL_COMPILE_WORK_GROUP_SIZE:
-      buf.as_vector<size_t>() = kern->block_size();
+      buf.as_vector<size_t>() = kern.block_size();
       break;
 
    case CL_KERNEL_LOCAL_MEM_SIZE:
-      buf.as_scalar<cl_ulong>() = kern->mem_local();
+      buf.as_scalar<cl_ulong>() = kern.mem_local();
       break;
 
    case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
@@ -186,7 +182,7 @@ clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev,
       break;
 
    case CL_KERNEL_PRIVATE_MEM_SIZE:
-      buf.as_scalar<cl_ulong>() = kern->mem_private();
+      buf.as_scalar<cl_ulong>() = kern.mem_private();
       break;
 
    default:
@@ -204,76 +200,52 @@ namespace {
    /// Common argument checking shared by kernel invocation commands.
    ///
    void
-   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) {
-      auto &q = obj(d_q);
-
-      if (!kern)
-         throw error(CL_INVALID_KERNEL);
-
-      if (&kern->prog.ctx != &q.ctx ||
-          any_of([&](const cl_event ev) {
-                return &obj(ev).ctx != &q.ctx;
-             }, range(deps, num_deps)))
+   validate_common(command_queue &q, kernel &kern,
+                   const ref_vector<event> &deps) {
+      if (&kern.prog.ctx != &q.ctx ||
+          any_of([&](const event &ev) {
+                return &ev.ctx != &q.ctx;
+             }, deps))
          throw error(CL_INVALID_CONTEXT);
 
-      if (bool(num_deps) != bool(deps) ||
-          any_of(is_zero(), range(deps, num_deps)))
-         throw error(CL_INVALID_EVENT_WAIT_LIST);
-
-      if (any_of([](std::unique_ptr<kernel::argument> &arg) {
-               return !arg->set();
-            }, kern->args))
+      if (any_of([](kernel::argument &arg) {
+               return !arg.set();
+            }, map(derefs(), 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);
+   }
+
+   void
+   validate_grid(command_queue &q, cl_uint dims,
+                 const size_t *d_grid_size, const size_t *d_block_size) {
+      auto grid_size = range(d_grid_size, dims);
 
       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)))
+      if (!d_grid_size || any_of(is_zero(), grid_size))
          throw error(CL_INVALID_GLOBAL_WORK_SIZE);
 
-      if (block_size) {
-         if (any_of([](size_t b, size_t max) {
-                  return b == 0 || b > max;
-               }, range(block_size, dims),
-               q.dev.max_block_size()))
+      if (d_block_size) {
+         auto block_size = range(d_block_size, dims);
+
+         if (any_of(is_zero(), block_size) ||
+             any_of(greater(), block_size, q.dev.max_block_size()))
             throw error(CL_INVALID_WORK_ITEM_SIZE);
 
-         if (any_of(modulus(), range(grid_size, dims),
-                    range(block_size, dims)))
+         if (any_of(modulus(), grid_size, block_size))
             throw error(CL_INVALID_WORK_GROUP_SIZE);
 
-         if (fold(multiplies(), 1u, range(block_size, dims)) >
+         if (fold(multiplies(), 1u, block_size) >
              q.dev.max_threads_per_block())
             throw error(CL_INVALID_WORK_GROUP_SIZE);
       }
    }
 
-   ///
-   /// Common event action shared by kernel invocation commands.
-   ///
-   std::function<void (event &)>
-   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 [=, &q](event &) {
-         kern->launch(q, grid_offset, reduced_grid_size, block_size);
-      };
-   }
-
    std::vector<size_t>
-   opt_vector(const size_t *p, unsigned n, size_t x) {
+   pad_vector(const size_t *p, unsigned n, size_t x) {
       if (p)
          return { p, p + n };
       else
@@ -282,58 +254,62 @@ namespace {
 }
 
 PUBLIC cl_int
-clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern,
-                       cl_uint dims, const size_t *pgrid_offset,
-                       const size_t *pgrid_size, const size_t *pblock_size,
+clEnqueueNDRangeKernel(cl_command_queue d_q, cl_kernel d_kern,
+                       cl_uint dims, const size_t *d_grid_offset,
+                       const size_t *d_grid_size, const size_t *d_block_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 &kern = obj(d_kern);
    auto deps = objs<wait_list_tag>(d_deps, num_deps);
-   auto grid_offset = opt_vector(pgrid_offset, dims, 0);
-   auto grid_size = opt_vector(pgrid_size, dims, 1);
-   auto block_size = opt_vector(pblock_size, dims, 1);
 
-   kernel_validate(q, kern, dims, pgrid_offset, pgrid_size, pblock_size,
-                   num_deps, d_deps, ev);
+   validate_common(q, kern, deps);
+   validate_grid(q, dims, d_grid_size, d_block_size);
 
+   auto grid_offset = pad_vector(d_grid_offset, dims, 0);
+   auto grid_size = pad_vector(d_grid_size, dims, 1);
+   auto block_size = pad_vector(d_block_size, dims, 1);
    hard_event *hev = new hard_event(
-      obj(q), CL_COMMAND_NDRANGE_KERNEL, deps,
-      kernel_op(q, kern, grid_offset, grid_size, block_size));
+      q, CL_COMMAND_NDRANGE_KERNEL, deps,
+      [=, &kern, &q](event &) {
+         kern.launch(q, grid_offset, grid_size, block_size);
+      });
 
-   ret_object(ev, hev);
+   ret_object(rd_ev, hev);
    return CL_SUCCESS;
 
-} catch(error &e) {
+} catch (error &e) {
    return e.get();
 }
 
 PUBLIC cl_int
-clEnqueueTask(cl_command_queue q, cl_kernel kern,
+clEnqueueTask(cl_command_queue d_q, cl_kernel d_kern,
               cl_uint num_deps, const cl_event *d_deps,
-              cl_event *ev) try {
+              cl_event *rd_ev) try {
+   auto &q = obj(d_q);
+   auto &kern = obj(d_kern);
    auto deps = objs<wait_list_tag>(d_deps, num_deps);
-   const std::vector<size_t> grid_offset = { 0 };
-   const std::vector<size_t> grid_size = { 1 };
-   const std::vector<size_t> block_size = { 1 };
 
-   kernel_validate(q, kern, 1, grid_offset.data(), grid_size.data(),
-                   block_size.data(), num_deps, d_deps, ev);
+   validate_common(q, kern, deps);
 
    hard_event *hev = new hard_event(
-      obj(q), CL_COMMAND_TASK, deps,
-      kernel_op(q, kern, grid_offset, grid_size, block_size));
+      q, CL_COMMAND_TASK, deps,
+      [=, &kern, &q](event &) {
+         kern.launch(q, { 0 }, { 1 }, { 1 });
+      });
 
-   ret_object(ev, hev);
+   ret_object(rd_ev, hev);
    return CL_SUCCESS;
 
-} catch(error &e) {
+} catch (error &e) {
    return e.get();
 }
 
 PUBLIC cl_int
-clEnqueueNativeKernel(cl_command_queue q, void (*func)(void *),
+clEnqueueNativeKernel(cl_command_queue d_q, void (*func)(void *),
                       void *args, size_t args_size,
-                      cl_uint obj_count, const cl_mem *obj_list,
-                      const void **obj_args, cl_uint num_deps,
-                      const cl_event *deps, cl_event *ev) {
+                      cl_uint num_mems, const cl_mem *d_mems,
+                      const void **mem_handles, cl_uint num_deps,
+                      const cl_event *d_deps, cl_event *rd_ev) {
    return CL_INVALID_OPERATION;
 }
diff --git a/src/gallium/state_trackers/clover/api/program.cpp b/src/gallium/state_trackers/clover/api/program.cpp
index f6c12f40367..84260472953 100644
--- a/src/gallium/state_trackers/clover/api/program.cpp
+++ b/src/gallium/state_trackers/clover/api/program.cpp
@@ -28,7 +28,7 @@ using namespace clover;
 PUBLIC cl_program
 clCreateProgramWithSource(cl_context d_ctx, cl_uint count,
                           const char **strings, const size_t *lengths,
-                          cl_int *errcode_ret) try {
+                          cl_int *r_errcode) try {
    auto &ctx = obj(d_ctx);
    std::string source;
 
@@ -43,19 +43,20 @@ clCreateProgramWithSource(cl_context d_ctx, cl_uint count,
                     std::string(strings[i]));
 
    // ...and create a program object for them.
-   ret_error(errcode_ret, CL_SUCCESS);
+   ret_error(r_errcode, CL_SUCCESS);
    return new program(ctx, source);
 
 } catch (error &e) {
-   ret_error(errcode_ret, e);
+   ret_error(r_errcode, e);
    return NULL;
 }
 
 PUBLIC cl_program
 clCreateProgramWithBinary(cl_context d_ctx, cl_uint n,
-                          const cl_device_id *d_devs, const size_t *lengths,
-                          const unsigned char **binaries, cl_int *status_ret,
-                          cl_int *errcode_ret) try {
+                          const cl_device_id *d_devs,
+                          const size_t *lengths,
+                          const unsigned char **binaries,
+                          cl_int *r_status, cl_int *r_errcode) try {
    auto &ctx = obj(d_ctx);
    auto devs = objs(d_devs, n);
 
@@ -68,7 +69,7 @@ clCreateProgramWithBinary(cl_context d_ctx, cl_uint n,
       throw error(CL_INVALID_DEVICE);
 
    // Deserialize the provided binaries,
-   auto modules = map(
+   auto result = map(
       [](const unsigned char *p, size_t l) -> std::pair<cl_int, module> {
          if (!p || !l)
             return { CL_INVALID_VALUE, {} };
@@ -87,69 +88,64 @@ clCreateProgramWithBinary(cl_context d_ctx, cl_uint n,
       range(lengths, n));
 
    // update the status array,
-   if (status_ret)
-      copy(map(keys(), modules), status_ret);
+   if (r_status)
+      copy(map(keys(), result), r_status);
 
-   if (any_of(key_equals(CL_INVALID_VALUE), modules))
+   if (any_of(key_equals(CL_INVALID_VALUE), result))
       throw error(CL_INVALID_VALUE);
 
-   if (any_of(key_equals(CL_INVALID_BINARY), modules))
+   if (any_of(key_equals(CL_INVALID_BINARY), result))
       throw error(CL_INVALID_BINARY);
 
    // initialize a program object with them.
-   ret_error(errcode_ret, CL_SUCCESS);
-   return new program(ctx, map(addresses(), devs), map(values(), modules));
+   ret_error(r_errcode, CL_SUCCESS);
+   return new program(ctx, devs, map(values(), result));
 
 } catch (error &e) {
-   ret_error(errcode_ret, e);
+   ret_error(r_errcode, e);
    return NULL;
 }
 
 PUBLIC cl_int
-clRetainProgram(cl_program prog) {
-   if (!prog)
-      return CL_INVALID_PROGRAM;
-
-   prog->retain();
+clRetainProgram(cl_program d_prog) try {
+   obj(d_prog).retain();
    return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
 }
 
 PUBLIC cl_int
-clReleaseProgram(cl_program prog) {
-   if (!prog)
-      return CL_INVALID_PROGRAM;
-
-   if (prog->release())
-      delete prog;
+clReleaseProgram(cl_program d_prog) try {
+   if (obj(d_prog).release())
+      delete pobj(d_prog);
 
    return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
 }
 
 PUBLIC cl_int
-clBuildProgram(cl_program prog, cl_uint count, const cl_device_id *devs,
-               const char *opts, void (*pfn_notify)(cl_program, void *),
+clBuildProgram(cl_program d_prog, cl_uint num_devs,
+               const cl_device_id *d_devs, const char *p_opts,
+               void (*pfn_notify)(cl_program, void *),
                void *user_data) try {
-   if (!prog)
-      throw error(CL_INVALID_PROGRAM);
+   auto &prog = obj(d_prog);
+   auto devs = (d_devs ? objs(d_devs, num_devs) :
+                ref_vector<device>(map(derefs(), prog.ctx.devs)));
+   auto opts = (p_opts ? p_opts : "");
 
-   if (bool(count) != bool(devs) ||
+   if (bool(num_devs) != bool(d_devs) ||
        (!pfn_notify && user_data))
       throw error(CL_INVALID_VALUE);
 
-   if (!opts)
-      opts = "";
-
-   if (devs) {
-      if (any_of([&](const cl_device_id dev) {
-               return !prog->ctx.has_device(obj(dev));
-            }, range(devs, count)))
-         throw error(CL_INVALID_DEVICE);
-
-      prog->build(map(addresses(), objs(devs, count)), opts);
-   } else {
-      prog->build(prog->ctx.devs, opts);
-   }
+   if (any_of([&](device &dev) {
+            return !prog.ctx.has_device(dev);
+         }, devs))
+      throw error(CL_INVALID_DEVICE);
 
+   prog.build(devs, opts);
    return CL_SUCCESS;
 
 } catch (error &e) {
@@ -162,32 +158,30 @@ clUnloadCompiler() {
 }
 
 PUBLIC cl_int
-clGetProgramInfo(cl_program prog, cl_program_info param,
+clGetProgramInfo(cl_program d_prog, cl_program_info param,
                  size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
-
-   if (!prog)
-      return CL_INVALID_PROGRAM;
+   auto &prog = obj(d_prog);
 
    switch (param) {
    case CL_PROGRAM_REFERENCE_COUNT:
-      buf.as_scalar<cl_uint>() = prog->ref_count();
+      buf.as_scalar<cl_uint>() = prog.ref_count();
       break;
 
    case CL_PROGRAM_CONTEXT:
-      buf.as_scalar<cl_context>() = &prog->ctx;
+      buf.as_scalar<cl_context>() = desc(prog.ctx);
       break;
 
    case CL_PROGRAM_NUM_DEVICES:
-      buf.as_scalar<cl_uint>() = prog->binaries().size();
+      buf.as_scalar<cl_uint>() = prog.binaries().size();
       break;
 
    case CL_PROGRAM_DEVICES:
-      buf.as_vector<cl_device_id>() = map(keys(), prog->binaries());
+      buf.as_vector<cl_device_id>() = map(keys(), prog.binaries());
       break;
 
    case CL_PROGRAM_SOURCE:
-      buf.as_string() = prog->source();
+      buf.as_string() = prog.source();
       break;
 
    case CL_PROGRAM_BINARY_SIZES:
@@ -198,7 +192,7 @@ clGetProgramInfo(cl_program prog, cl_program_info param,
                ent.second.serialize(s);
                return bin.size();
             },
-            prog->binaries());
+            prog.binaries());
       break;
 
    case CL_PROGRAM_BINARIES:
@@ -209,7 +203,7 @@ clGetProgramInfo(cl_program prog, cl_program_info param,
                ent.second.serialize(s);
                return bin;
             },
-            prog->binaries());
+            prog.binaries());
       break;
 
    default:
@@ -223,28 +217,27 @@ clGetProgramInfo(cl_program prog, cl_program_info param,
 }
 
 PUBLIC cl_int
-clGetProgramBuildInfo(cl_program prog, cl_device_id dev,
+clGetProgramBuildInfo(cl_program d_prog, cl_device_id d_dev,
                       cl_program_build_info param,
                       size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
+   auto &prog = obj(d_prog);
+   auto &dev = obj(d_dev);
 
-   if (!prog)
-      return CL_INVALID_PROGRAM;
-
-   if (!prog->ctx.has_device(obj(dev)))
+   if (!prog.ctx.has_device(dev))
       return CL_INVALID_DEVICE;
 
    switch (param) {
    case CL_PROGRAM_BUILD_STATUS:
-      buf.as_scalar<cl_build_status>() = prog->build_status(pobj(dev));
+      buf.as_scalar<cl_build_status>() = prog.build_status(dev);
       break;
 
    case CL_PROGRAM_BUILD_OPTIONS:
-      buf.as_string() = prog->build_opts(pobj(dev));
+      buf.as_string() = prog.build_opts(dev);
       break;
 
    case CL_PROGRAM_BUILD_LOG:
-      buf.as_string() = prog->build_log(pobj(dev));
+      buf.as_string() = prog.build_log(dev);
       break;
 
    default:
diff --git a/src/gallium/state_trackers/clover/core/error.hpp b/src/gallium/state_trackers/clover/core/error.hpp
index fa43c1a5eed..088bdac3ef3 100644
--- a/src/gallium/state_trackers/clover/core/error.hpp
+++ b/src/gallium/state_trackers/clover/core/error.hpp
@@ -34,7 +34,7 @@ namespace clover {
    class event;
    class hard_event;
    class soft_event;
-   typedef struct _cl_kernel kernel;
+   class kernel;
    typedef struct _cl_mem memory_obj;
    class buffer;
    class root_buffer;
@@ -43,7 +43,7 @@ namespace clover {
    class image2d;
    class image3d;
    class platform;
-   typedef struct _cl_program program;
+   class program;
    typedef struct _cl_sampler sampler;
 
    ///
diff --git a/src/gallium/state_trackers/clover/core/kernel.cpp b/src/gallium/state_trackers/clover/core/kernel.cpp
index 5663f1f8b2e..9f9577b1921 100644
--- a/src/gallium/state_trackers/clover/core/kernel.cpp
+++ b/src/gallium/state_trackers/clover/core/kernel.cpp
@@ -22,15 +22,14 @@
 
 #include "core/kernel.hpp"
 #include "core/resource.hpp"
-#include "util/algorithm.hpp"
 #include "util/u_math.h"
 #include "pipe/p_context.h"
 
 using namespace clover;
 
-_cl_kernel::_cl_kernel(clover::program &prog,
-                       const std::string &name,
-                       const std::vector<clover::module::argument> &margs) :
+kernel::kernel(program &prog,
+               const std::string &name,
+               const std::vector<module::argument> &margs) :
    prog(prog), _name(name), exec(*this) {
    for (auto marg : margs) {
       if (marg.type == module::argument::scalar)
@@ -56,17 +55,17 @@ _cl_kernel::_cl_kernel(clover::program &prog,
 
 template<typename T, typename V>
 static inline std::vector<T>
-pad_vector(clover::command_queue &q, const V &v, T x) {
+pad_vector(command_queue &q, const V &v, T x) {
    std::vector<T> w { v.begin(), v.end() };
    w.resize(q.dev.max_block_size().size(), x);
    return w;
 }
 
 void
-_cl_kernel::launch(clover::command_queue &q,
-                   const std::vector<size_t> &grid_offset,
-                   const std::vector<size_t> &grid_size,
-                   const std::vector<size_t> &block_size) {
+kernel::launch(command_queue &q,
+               const std::vector<size_t> &grid_offset,
+               const std::vector<size_t> &grid_size,
+               const std::vector<size_t> &block_size) {
    void *st = exec.bind(&q);
    std::vector<uint32_t *> g_handles = map([&](size_t h) {
          return (uint32_t *)&exec.input[h];
@@ -80,7 +79,7 @@ _cl_kernel::launch(clover::command_queue &q,
    q.pipe->set_compute_sampler_views(q.pipe, 0, exec.sviews.size(),
                                      exec.sviews.data());
    q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(),
-                                     exec.resources.data());
+                                 exec.resources.data());
    q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(),
                               exec.g_buffers.data(), g_handles.data());
 
@@ -99,7 +98,7 @@ _cl_kernel::launch(clover::command_queue &q,
 }
 
 size_t
-_cl_kernel::mem_local() const {
+kernel::mem_local() const {
    size_t sz = 0;
 
    for (auto &arg : args) {
@@ -111,49 +110,49 @@ _cl_kernel::mem_local() const {
 }
 
 size_t
-_cl_kernel::mem_private() const {
+kernel::mem_private() const {
    return 0;
 }
 
 size_t
-_cl_kernel::max_block_size() const {
+kernel::max_block_size() const {
    return std::numeric_limits<std::size_t>::max();
 }
 
 const std::string &
-_cl_kernel::name() const {
+kernel::name() const {
    return _name;
 }
 
 std::vector<size_t>
-_cl_kernel::block_size() const {
+kernel::block_size() const {
    return { 0, 0, 0 };
 }
 
-const clover::module &
-_cl_kernel::module(const clover::command_queue &q) const {
+const module &
+kernel::module(const command_queue &q) const {
    return prog.binaries().find(&q.dev)->second;
 }
 
-_cl_kernel::exec_context::exec_context(clover::kernel &kern) :
+kernel::exec_context::exec_context(kernel &kern) :
    kern(kern), q(NULL), mem_local(0), st(NULL) {
 }
 
-_cl_kernel::exec_context::~exec_context() {
+kernel::exec_context::~exec_context() {
    if (st)
       q->pipe->delete_compute_state(q->pipe, st);
 }
 
 void *
-_cl_kernel::exec_context::bind(clover::command_queue *_q) {
+kernel::exec_context::bind(command_queue *_q) {
    std::swap(q, _q);
 
    // Bind kernel arguments.
    auto margs = kern.module(*q).sym(kern.name()).args;
    for_each([=](std::unique_ptr<kernel::argument> &karg,
                 const module::argument &marg) {
-         karg->bind(*this, marg);
-      }, kern.args, margs);
+               karg->bind(*this, marg);
+            }, kern.args, margs);
 
    // Create a new compute state if anything changed.
    if (!st || q != _q ||
@@ -172,7 +171,7 @@ _cl_kernel::exec_context::bind(clover::command_queue *_q) {
 }
 
 void
-_cl_kernel::exec_context::unbind() {
+kernel::exec_context::unbind() {
    for (auto &arg : kern.args)
       arg->unbind(*this);
 
@@ -226,7 +225,7 @@ namespace {
    ///
    template<typename T>
    void
-   extend(T &v, enum clover::module::argument::ext_type ext, size_t n) {
+   extend(T &v, enum module::argument::ext_type ext, size_t n) {
       const size_t m = std::min(v.size(), n);
       const bool sign_ext = (ext == module::argument::sign_ext);
       const uint8_t fill = (sign_ext && msb(v) ? ~0 : 0);
@@ -261,24 +260,24 @@ namespace {
    }
 }
 
-_cl_kernel::argument::argument() : _set(false) {
+kernel::argument::argument() : _set(false) {
 }
 
 bool
-_cl_kernel::argument::set() const {
+kernel::argument::set() const {
    return _set;
 }
 
 size_t
-_cl_kernel::argument::storage() const {
+kernel::argument::storage() const {
    return 0;
 }
 
-_cl_kernel::scalar_argument::scalar_argument(size_t size) : size(size) {
+kernel::scalar_argument::scalar_argument(size_t size) : size(size) {
 }
 
 void
-_cl_kernel::scalar_argument::set(size_t size, const void *value) {
+kernel::scalar_argument::set(size_t size, const void *value) {
    if (size != this->size)
       throw error(CL_INVALID_ARG_SIZE);
 
@@ -287,8 +286,8 @@ _cl_kernel::scalar_argument::set(size_t size, const void *value) {
 }
 
 void
-_cl_kernel::scalar_argument::bind(exec_context &ctx,
-                                  const clover::module::argument &marg) {
+kernel::scalar_argument::bind(exec_context &ctx,
+                              const module::argument &marg) {
    auto w = v;
 
    extend(w, marg.ext_type, marg.target_size);
@@ -298,40 +297,40 @@ _cl_kernel::scalar_argument::bind(exec_context &ctx,
 }
 
 void
-_cl_kernel::scalar_argument::unbind(exec_context &ctx) {
+kernel::scalar_argument::unbind(exec_context &ctx) {
 }
 
 void
-_cl_kernel::global_argument::set(size_t size, const void *value) {
+kernel::global_argument::set(size_t size, const void *value) {
    if (size != sizeof(cl_mem))
       throw error(CL_INVALID_ARG_SIZE);
 
-   obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
-   if (!obj)
+   buf = dynamic_cast<buffer *>(*(cl_mem *)value);
+   if (!buf)
       throw error(CL_INVALID_MEM_OBJECT);
 
    _set = true;
 }
 
 void
-_cl_kernel::global_argument::bind(exec_context &ctx,
-                                  const clover::module::argument &marg) {
+kernel::global_argument::bind(exec_context &ctx,
+                              const 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(buf->resource(*ctx.q).pipe);
 }
 
 void
-_cl_kernel::global_argument::unbind(exec_context &ctx) {
+kernel::global_argument::unbind(exec_context &ctx) {
 }
 
 size_t
-_cl_kernel::local_argument::storage() const {
+kernel::local_argument::storage() const {
    return _storage;
 }
 
 void
-_cl_kernel::local_argument::set(size_t size, const void *value) {
+kernel::local_argument::set(size_t size, const void *value) {
    if (value)
       throw error(CL_INVALID_ARG_VALUE);
 
@@ -340,8 +339,8 @@ _cl_kernel::local_argument::set(size_t size, const void *value) {
 }
 
 void
-_cl_kernel::local_argument::bind(exec_context &ctx,
-                                 const clover::module::argument &marg) {
+kernel::local_argument::bind(exec_context &ctx,
+                             const module::argument &marg) {
    auto v = bytes(ctx.mem_local);
 
    extend(v, module::argument::zero_ext, marg.target_size);
@@ -353,24 +352,24 @@ _cl_kernel::local_argument::bind(exec_context &ctx,
 }
 
 void
-_cl_kernel::local_argument::unbind(exec_context &ctx) {
+kernel::local_argument::unbind(exec_context &ctx) {
 }
 
 void
-_cl_kernel::constant_argument::set(size_t size, const void *value) {
+kernel::constant_argument::set(size_t size, const void *value) {
    if (size != sizeof(cl_mem))
       throw error(CL_INVALID_ARG_SIZE);
 
-   obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
-   if (!obj)
+   buf = dynamic_cast<buffer *>(*(cl_mem *)value);
+   if (!buf)
       throw error(CL_INVALID_MEM_OBJECT);
 
    _set = true;
 }
 
 void
-_cl_kernel::constant_argument::bind(exec_context &ctx,
-                                    const clover::module::argument &marg) {
+kernel::constant_argument::bind(exec_context &ctx,
+                                const module::argument &marg) {
    auto v = bytes(ctx.resources.size() << 24);
 
    extend(v, module::argument::zero_ext, marg.target_size);
@@ -378,30 +377,30 @@ _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 = buf->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);
+kernel::constant_argument::unbind(exec_context &ctx) {
+   buf->resource(*ctx.q).unbind_surface(*ctx.q, st);
 }
 
 void
-_cl_kernel::image_rd_argument::set(size_t size, const void *value) {
+kernel::image_rd_argument::set(size_t size, const void *value) {
    if (size != sizeof(cl_mem))
       throw error(CL_INVALID_ARG_SIZE);
 
-   obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
-   if (!obj)
+   img = dynamic_cast<image *>(*(cl_mem *)value);
+   if (!img)
       throw error(CL_INVALID_MEM_OBJECT);
 
    _set = true;
 }
 
 void
-_cl_kernel::image_rd_argument::bind(exec_context &ctx,
-                                    const clover::module::argument &marg) {
+kernel::image_rd_argument::bind(exec_context &ctx,
+                                const module::argument &marg) {
    auto v = bytes(ctx.sviews.size());
 
    extend(v, module::argument::zero_ext, marg.target_size);
@@ -409,30 +408,30 @@ _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 = img->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);
+kernel::image_rd_argument::unbind(exec_context &ctx) {
+   img->resource(*ctx.q).unbind_sampler_view(*ctx.q, st);
 }
 
 void
-_cl_kernel::image_wr_argument::set(size_t size, const void *value) {
+kernel::image_wr_argument::set(size_t size, const void *value) {
    if (size != sizeof(cl_mem))
       throw error(CL_INVALID_ARG_SIZE);
 
-   obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
-   if (!obj)
+   img = dynamic_cast<image *>(*(cl_mem *)value);
+   if (!img)
       throw error(CL_INVALID_MEM_OBJECT);
 
    _set = true;
 }
 
 void
-_cl_kernel::image_wr_argument::bind(exec_context &ctx,
-                                    const clover::module::argument &marg) {
+kernel::image_wr_argument::bind(exec_context &ctx,
+                                const module::argument &marg) {
    auto v = bytes(ctx.resources.size());
 
    extend(v, module::argument::zero_ext, marg.target_size);
@@ -440,32 +439,32 @@ _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 = img->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);
+kernel::image_wr_argument::unbind(exec_context &ctx) {
+   img->resource(*ctx.q).unbind_surface(*ctx.q, st);
 }
 
 void
-_cl_kernel::sampler_argument::set(size_t size, const void *value) {
+kernel::sampler_argument::set(size_t size, const void *value) {
    if (size != sizeof(cl_sampler))
       throw error(CL_INVALID_ARG_SIZE);
 
-   obj = *(cl_sampler *)value;
+   s = *(cl_sampler *)value;
    _set = true;
 }
 
 void
-_cl_kernel::sampler_argument::bind(exec_context &ctx,
-                                   const clover::module::argument &marg) {
-   st = obj->bind(*ctx.q);
+kernel::sampler_argument::bind(exec_context &ctx,
+                               const module::argument &marg) {
+   st = s->bind(*ctx.q);
    ctx.samplers.push_back(st);
 }
 
 void
-_cl_kernel::sampler_argument::unbind(exec_context &ctx) {
-   obj->unbind(*ctx.q, st);
+kernel::sampler_argument::unbind(exec_context &ctx) {
+   s->unbind(*ctx.q, st);
 }
diff --git a/src/gallium/state_trackers/clover/core/kernel.hpp b/src/gallium/state_trackers/clover/core/kernel.hpp
index 984e2139c7b..e469108d4be 100644
--- a/src/gallium/state_trackers/clover/core/kernel.hpp
+++ b/src/gallium/state_trackers/clover/core/kernel.hpp
@@ -32,176 +32,173 @@
 #include "pipe/p_state.h"
 
 namespace clover {
-   typedef struct _cl_kernel kernel;
-   class argument;
-}
-
-struct _cl_kernel : public clover::ref_counter {
-private:
-   ///
-   /// Class containing all the state required to execute a compute
-   /// kernel.
-   ///
-   struct exec_context {
-      exec_context(clover::kernel &kern);
-      ~exec_context();
-
-      void *bind(clover::command_queue *q);
-      void unbind();
-
-      clover::kernel &kern;
-      clover::command_queue *q;
-
-      std::vector<uint8_t> input;
-      std::vector<void *> samplers;
-      std::vector<pipe_sampler_view *> sviews;
-      std::vector<pipe_surface *> resources;
-      std::vector<pipe_resource *> g_buffers;
-      std::vector<size_t> g_handles;
-      size_t mem_local;
-
-   private:
-      void *st;
-      pipe_compute_state cs;
-   };
-
-public:
-   class argument {
-   public:
-      argument();
-
-      /// \a true if the argument has been set.
-      bool set() const;
-
-      /// Storage space required for the referenced object.
-      virtual size_t storage() const;
-
-      /// Set this argument to some object.
-      virtual void set(size_t size, const void *value) = 0;
-
-      /// Allocate the necessary resources to bind the specified
-      /// object to this argument, and update \a ctx accordingly.
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg) = 0;
-
-      /// Free any resources that were allocated in bind().
-      virtual void unbind(exec_context &ctx) = 0;
-
-   protected:
-      bool _set;
-   };
-
-   _cl_kernel(clover::program &prog,
-              const std::string &name,
-              const std::vector<clover::module::argument> &margs);
-
-   void launch(clover::command_queue &q,
-               const std::vector<size_t> &grid_offset,
-               const std::vector<size_t> &grid_size,
-               const std::vector<size_t> &block_size);
-
-   size_t mem_local() const;
-   size_t mem_private() const;
-   size_t max_block_size() const;
-
-   const std::string &name() const;
-   std::vector<size_t> block_size() const;
-
-   clover::program &prog;
-   std::vector<std::unique_ptr<argument>> args;
-
-private:
-   const clover::module &
-   module(const clover::command_queue &q) const;
-
-   class scalar_argument : public argument {
-   public:
-      scalar_argument(size_t size);
-
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
-
+   class kernel : public ref_counter, public _cl_kernel {
    private:
-      size_t size;
-      std::vector<uint8_t> v;
-   };
+      ///
+      /// Class containing all the state required to execute a compute
+      /// kernel.
+      ///
+      struct exec_context {
+         exec_context(kernel &kern);
+         ~exec_context();
+
+         void *bind(command_queue *q);
+         void unbind();
+
+         kernel &kern;
+         command_queue *q;
+
+         std::vector<uint8_t> input;
+         std::vector<void *> samplers;
+         std::vector<pipe_sampler_view *> sviews;
+         std::vector<pipe_surface *> resources;
+         std::vector<pipe_resource *> g_buffers;
+         std::vector<size_t> g_handles;
+         size_t mem_local;
+
+      private:
+         void *st;
+         pipe_compute_state cs;
+      };
 
-   class global_argument : public argument {
    public:
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+      class argument {
+      public:
+         argument();
 
-   private:
-      clover::buffer *obj;
-   };
+         /// \a true if the argument has been set.
+         bool set() const;
 
-   class local_argument : public argument {
-   public:
-      virtual size_t storage() const;
+         /// Storage space required for the referenced object.
+         virtual size_t storage() const;
 
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+         /// Set this argument to some object.
+         virtual void set(size_t size, const void *value) = 0;
 
-   private:
-      size_t _storage;
-   };
+         /// Allocate the necessary resources to bind the specified
+         /// object to this argument, and update \a ctx accordingly.
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg) = 0;
 
-   class constant_argument : public argument {
-   public:
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+         /// Free any resources that were allocated in bind().
+         virtual void unbind(exec_context &ctx) = 0;
 
-   private:
-      clover::buffer *obj;
-      pipe_surface *st;
-   };
+      protected:
+         bool _set;
+      };
 
-   class image_rd_argument : public argument {
-   public:
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+      kernel(program &prog,
+             const std::string &name,
+             const std::vector<module::argument> &margs);
 
-   private:
-      clover::image *obj;
-      pipe_sampler_view *st;
-   };
+      void launch(command_queue &q,
+                  const std::vector<size_t> &grid_offset,
+                  const std::vector<size_t> &grid_size,
+                  const std::vector<size_t> &block_size);
 
-   class image_wr_argument : public argument {
-   public:
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+      size_t mem_local() const;
+      size_t mem_private() const;
+      size_t max_block_size() const;
 
-   private:
-      clover::image *obj;
-      pipe_surface *st;
-   };
+      const std::string &name() const;
+      std::vector<size_t> block_size() const;
 
-   class sampler_argument : public argument {
-   public:
-      virtual void set(size_t size, const void *value);
-      virtual void bind(exec_context &ctx,
-                        const clover::module::argument &marg);
-      virtual void unbind(exec_context &ctx);
+      program &prog;
+      std::vector<std::unique_ptr<argument>> args;
 
    private:
-      clover::sampler *obj;
-      void *st;
+      const clover::module &
+      module(const command_queue &q) const;
+
+      class scalar_argument : public argument {
+      public:
+         scalar_argument(size_t size);
+
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         size_t size;
+         std::vector<uint8_t> v;
+      };
+
+      class global_argument : public argument {
+      public:
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         buffer *buf;
+      };
+
+      class local_argument : public argument {
+      public:
+         virtual size_t storage() const;
+
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         size_t _storage;
+      };
+
+      class constant_argument : public argument {
+      public:
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         buffer *buf;
+         pipe_surface *st;
+      };
+
+      class image_rd_argument : public argument {
+      public:
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         image *img;
+         pipe_sampler_view *st;
+      };
+
+      class image_wr_argument : public argument {
+      public:
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         image *img;
+         pipe_surface *st;
+      };
+
+      class sampler_argument : public argument {
+      public:
+         virtual void set(size_t size, const void *value);
+         virtual void bind(exec_context &ctx,
+                           const module::argument &marg);
+         virtual void unbind(exec_context &ctx);
+
+      private:
+         sampler *s;
+         void *st;
+      };
+
+      std::string _name;
+      exec_context exec;
    };
-
-   std::string _name;
-   exec_context exec;
-};
+}
 
 #endif
diff --git a/src/gallium/state_trackers/clover/core/object.hpp b/src/gallium/state_trackers/clover/core/object.hpp
index 6a99f19bd1e..9c2180f3b29 100644
--- a/src/gallium/state_trackers/clover/core/object.hpp
+++ b/src/gallium/state_trackers/clover/core/object.hpp
@@ -188,9 +188,15 @@ struct _cl_device_id :
 struct _cl_event :
    public clover::descriptor<clover::event, _cl_event> {};
 
+struct _cl_kernel :
+   public clover::descriptor<clover::kernel, _cl_kernel> {};
+
 struct _cl_platform_id :
    public clover::descriptor<clover::platform, _cl_platform_id> {};
 
+struct _cl_program :
+   public clover::descriptor<clover::program, _cl_program> {};
+
 struct _cl_command_queue :
    public clover::descriptor<clover::command_queue, _cl_command_queue> {};
 
diff --git a/src/gallium/state_trackers/clover/core/program.cpp b/src/gallium/state_trackers/clover/core/program.cpp
index 42b301497b5..8082cf0f6f6 100644
--- a/src/gallium/state_trackers/clover/core/program.cpp
+++ b/src/gallium/state_trackers/clover/core/program.cpp
@@ -22,70 +22,67 @@
 
 #include "core/program.hpp"
 #include "core/compiler.hpp"
-#include "util/algorithm.hpp"
 
 using namespace clover;
 
-_cl_program::_cl_program(clover::context &ctx,
-                         const std::string &source) :
+program::program(context &ctx, const std::string &source) :
    ctx(ctx), _source(source) {
 }
 
-_cl_program::_cl_program(clover::context &ctx,
-                         const std::vector<clover::device *> &devs,
-                         const std::vector<clover::module> &binaries) :
+program::program(context &ctx,
+                 const ref_vector<device> &devs,
+                 const std::vector<module> &binaries) :
    ctx(ctx) {
-   for_each([&](clover::device *dev, const clover::module &bin) {
-         _binaries.insert({ dev, bin });
+   for_each([&](device &dev, const module &bin) {
+         _binaries.insert({ &dev, bin });
       },
       devs, binaries);
 }
 
 void
-_cl_program::build(const std::vector<clover::device *> &devs,
-                   const char *opts) {
+program::build(const ref_vector<device> &devs, const char *opts) {
+   for (auto &dev : devs) {
+      _binaries.erase(&dev);
+      _logs.erase(&dev);
+      _opts.erase(&dev);
 
-   for (auto dev : devs) {
-      _binaries.erase(dev);
-      _logs.erase(dev);
-      _opts.erase(dev);
+      _opts.insert({ &dev, opts });
 
-      _opts.insert({ dev, opts });
       try {
-         auto module = (dev->ir_format() == PIPE_SHADER_IR_TGSI ?
+         auto module = (dev.ir_format() == PIPE_SHADER_IR_TGSI ?
                         compile_program_tgsi(_source) :
-                        compile_program_llvm(_source, dev->ir_format(),
-                        dev->ir_target(), build_opts(dev)));
-         _binaries.insert({ dev, module });
+                        compile_program_llvm(_source, dev.ir_format(),
+                                             dev.ir_target(), build_opts(dev)));
+         _binaries.insert({ &dev, module });
 
       } catch (build_error &e) {
-         _logs.insert({ dev, e.what() });
+         _logs.insert({ &dev, e.what() });
          throw;
       }
    }
 }
 
 const std::string &
-_cl_program::source() const {
+program::source() const {
    return _source;
 }
 
-const std::map<clover::device *, clover::module> &
-_cl_program::binaries() const {
+const std::map<device *, module> &
+program::binaries() const {
    return _binaries;
 }
 
 cl_build_status
-_cl_program::build_status(clover::device *dev) const {
-   return _binaries.count(dev) ? CL_BUILD_SUCCESS : CL_BUILD_NONE;
+program::build_status(device &dev) const {
+   return _binaries.count(&dev) ? CL_BUILD_SUCCESS : CL_BUILD_NONE;
 }
 
 std::string
-_cl_program::build_opts(clover::device *dev) const {
-   return _opts.count(dev) ? _opts.find(dev)->second : "";
+program::build_opts(device &dev) const {
+   return _opts.count(&dev) ? _opts.find(&dev)->second : "";
 }
 
 std::string
-_cl_program::build_log(clover::device *dev) const {
-   return _logs.count(dev) ? _logs.find(dev)->second : "";
+program::build_log(device &dev) const {
+   return _logs.count(&dev) ? _logs.find(&dev)->second : "";
 }
diff --git a/src/gallium/state_trackers/clover/core/program.hpp b/src/gallium/state_trackers/clover/core/program.hpp
index 0d7bf372f6d..fa1afa7c66e 100644
--- a/src/gallium/state_trackers/clover/core/program.hpp
+++ b/src/gallium/state_trackers/clover/core/program.hpp
@@ -30,33 +30,31 @@
 #include "core/module.hpp"
 
 namespace clover {
-   typedef struct _cl_program program;
+   class program : public ref_counter, public _cl_program {
+   public:
+      program(context &ctx,
+              const std::string &source);
+      program(context &ctx,
+              const ref_vector<device> &devs,
+              const std::vector<module> &binaries);
+
+      void build(const ref_vector<device> &devs, const char *opts);
+
+      const std::string &source() const;
+      const std::map<device *, module> &binaries() const;
+
+      cl_build_status build_status(device &dev) const;
+      std::string build_opts(device &dev) const;
+      std::string build_log(device &dev) const;
+
+      context &ctx;
+
+   private:
+      std::map<device *, module> _binaries;
+      std::map<device *, std::string> _logs;
+      std::map<device *, std::string> _opts;
+      std::string _source;
+   };
 }
 
-struct _cl_program : public clover::ref_counter {
-public:
-   _cl_program(clover::context &ctx,
-               const std::string &source);
-   _cl_program(clover::context &ctx,
-               const std::vector<clover::device *> &devs,
-               const std::vector<clover::module> &binaries);
-
-   void build(const std::vector<clover::device *> &devs, const char *opts);
-
-   const std::string &source() const;
-   const std::map<clover::device *, clover::module> &binaries() const;
-
-   cl_build_status build_status(clover::device *dev) const;
-   std::string build_opts(clover::device *dev) const;
-   std::string build_log(clover::device *dev) const;
-
-   clover::context &ctx;
-
-private:
-   std::map<clover::device *, clover::module> _binaries;
-   std::map<clover::device *, std::string> _logs;
-   std::map<clover::device *, std::string> _opts;
-   std::string _source;
-};
-
 #endif
diff --git a/src/gallium/state_trackers/clover/core/queue.hpp b/src/gallium/state_trackers/clover/core/queue.hpp
index 4a2d02251b1..65f2d63ffc8 100644
--- a/src/gallium/state_trackers/clover/core/queue.hpp
+++ b/src/gallium/state_trackers/clover/core/queue.hpp
@@ -53,7 +53,7 @@ namespace clover {
       friend class mapping;
       friend class hard_event;
       friend struct ::_cl_sampler;
-      friend struct ::_cl_kernel;
+      friend class kernel;
       friend class clover::timestamp::query;
       friend class clover::timestamp::current;
 
diff --git a/src/gallium/state_trackers/clover/core/resource.hpp b/src/gallium/state_trackers/clover/core/resource.hpp
index b2eddc08fa6..8fcfb496a33 100644
--- a/src/gallium/state_trackers/clover/core/resource.hpp
+++ b/src/gallium/state_trackers/clover/core/resource.hpp
@@ -57,7 +57,7 @@ namespace clover {
 
       friend class sub_resource;
       friend class mapping;
-      friend struct ::_cl_kernel;
+      friend class kernel;
 
    protected:
       resource(clover::device &dev, clover::memory_obj &obj);
diff --git a/src/gallium/state_trackers/clover/core/sampler.hpp b/src/gallium/state_trackers/clover/core/sampler.hpp
index 9716aabd22b..ad1531914b6 100644
--- a/src/gallium/state_trackers/clover/core/sampler.hpp
+++ b/src/gallium/state_trackers/clover/core/sampler.hpp
@@ -41,7 +41,7 @@ public:
 
    clover::context &ctx;
 
-   friend class _cl_kernel;
+   friend class clover::kernel;
 
 private:
    void *bind(clover::command_queue &q);