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

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

2 files changed, 171 insertions, 202 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: