From 35307f540fedf9680ce8b05d0784c5b0d5b0f6a7 Mon Sep 17 00:00:00 2001 From: Francisco Jerez Date: Tue, 17 Sep 2013 23:20:11 -0700 Subject: [PATCH] clover: Switch kernel and program objects to the new model. Tested-by: Tom Stellard --- .../state_trackers/clover/api/kernel.cpp | 256 +++++++-------- .../state_trackers/clover/api/program.cpp | 117 ++++--- .../state_trackers/clover/core/error.hpp | 4 +- .../state_trackers/clover/core/kernel.cpp | 149 +++++---- .../state_trackers/clover/core/kernel.hpp | 305 +++++++++--------- .../state_trackers/clover/core/object.hpp | 6 + .../state_trackers/clover/core/program.cpp | 55 ++-- .../state_trackers/clover/core/program.hpp | 52 ++- .../state_trackers/clover/core/queue.hpp | 2 +- .../state_trackers/clover/core/resource.hpp | 2 +- .../state_trackers/clover/core/sampler.hpp | 2 +- 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() = kern->args.size(); + buf.as_scalar() = kern.args.size(); break; case CL_KERNEL_REFERENCE_COUNT: - buf.as_scalar() = kern->ref_count(); + buf.as_scalar() = kern.ref_count(); break; case CL_KERNEL_CONTEXT: - buf.as_scalar() = &kern->prog.ctx; + buf.as_scalar() = desc(kern.prog.ctx); break; case CL_KERNEL_PROGRAM: - buf.as_scalar() = &kern->prog; + buf.as_scalar() = 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() = kern->max_block_size(); + buf.as_scalar() = kern.max_block_size(); break; case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: - buf.as_vector() = kern->block_size(); + buf.as_vector() = kern.block_size(); break; case CL_KERNEL_LOCAL_MEM_SIZE: - buf.as_scalar() = kern->mem_local(); + buf.as_scalar() = 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() = kern->mem_private(); + buf.as_scalar() = 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 &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 &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 - kernel_op(cl_command_queue d_q, cl_kernel kern, - const std::vector &grid_offset, - const std::vector &grid_size, - const std::vector &block_size) { - auto &q = obj(d_q); - const std::vector reduced_grid_size = - map(divides(), grid_size, block_size); - - return [=, &q](event &) { - kern->launch(q, grid_offset, reduced_grid_size, block_size); - }; - } - std::vector - 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(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(d_deps, num_deps); - const std::vector grid_offset = { 0 }; - const std::vector grid_size = { 1 }; - const std::vector 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 { 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(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() = prog->ref_count(); + buf.as_scalar() = prog.ref_count(); break; case CL_PROGRAM_CONTEXT: - buf.as_scalar() = &prog->ctx; + buf.as_scalar() = desc(prog.ctx); break; case CL_PROGRAM_NUM_DEVICES: - buf.as_scalar() = prog->binaries().size(); + buf.as_scalar() = prog.binaries().size(); break; case CL_PROGRAM_DEVICES: - buf.as_vector() = map(keys(), prog->binaries()); + buf.as_vector() = 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() = prog->build_status(pobj(dev)); + buf.as_scalar() = 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 &margs) : +kernel::kernel(program &prog, + const std::string &name, + const std::vector &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 static inline std::vector -pad_vector(clover::command_queue &q, const V &v, T x) { +pad_vector(command_queue &q, const V &v, T x) { std::vector 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 &grid_offset, - const std::vector &grid_size, - const std::vector &block_size) { +kernel::launch(command_queue &q, + const std::vector &grid_offset, + const std::vector &grid_size, + const std::vector &block_size) { void *st = exec.bind(&q); std::vector 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::max(); } const std::string & -_cl_kernel::name() const { +kernel::name() const { return _name; } std::vector -_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 &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 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(*(cl_mem *)value); - if (!obj) + buf = dynamic_cast(*(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(*(cl_mem *)value); - if (!obj) + buf = dynamic_cast(*(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(*(cl_mem *)value); - if (!obj) + img = dynamic_cast(*(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(*(cl_mem *)value); - if (!obj) + img = dynamic_cast(*(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 input; - std::vector samplers; - std::vector sviews; - std::vector resources; - std::vector g_buffers; - std::vector 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 &margs); - - void launch(clover::command_queue &q, - const std::vector &grid_offset, - const std::vector &grid_size, - const std::vector &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 block_size() const; - - clover::program &prog; - std::vector> 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 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 input; + std::vector samplers; + std::vector sviews; + std::vector resources; + std::vector g_buffers; + std::vector 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 &margs); - private: - clover::image *obj; - pipe_sampler_view *st; - }; + void launch(command_queue &q, + const std::vector &grid_offset, + const std::vector &grid_size, + const std::vector &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 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> 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 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 {}; +struct _cl_kernel : + public clover::descriptor {}; + struct _cl_platform_id : public clover::descriptor {}; +struct _cl_program : + public clover::descriptor {}; + struct _cl_command_queue : public clover::descriptor {}; 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 &devs, - const std::vector &binaries) : +program::program(context &ctx, + const ref_vector &devs, + const std::vector &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 &devs, - const char *opts) { +program::build(const ref_vector &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 & -_cl_program::binaries() const { +const std::map & +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 &devs, + const std::vector &binaries); + + void build(const ref_vector &devs, const char *opts); + + const std::string &source() const; + const std::map &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 _binaries; + std::map _logs; + std::map _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 &devs, - const std::vector &binaries); - - void build(const std::vector &devs, const char *opts); - - const std::string &source() const; - const std::map &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 _binaries; - std::map _logs; - std::map _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); -- 2.30.2