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:
}
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:
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:
/// 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
}
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;
}
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;
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);
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, {} };
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) {
}
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:
ent.second.serialize(s);
return bin.size();
},
- prog->binaries());
+ prog.binaries());
break;
case CL_PROGRAM_BINARIES:
ent.second.serialize(s);
return bin;
},
- prog->binaries());
+ prog.binaries());
break;
default:
}
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:
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;
class image2d;
class image3d;
class platform;
- typedef struct _cl_program program;
+ class program;
typedef struct _cl_sampler sampler;
///
#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)
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];
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());
}
size_t
-_cl_kernel::mem_local() const {
+kernel::mem_local() const {
size_t sz = 0;
for (auto &arg : args) {
}
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 ||
}
void
-_cl_kernel::exec_context::unbind() {
+kernel::exec_context::unbind() {
for (auto &arg : kern.args)
arg->unbind(*this);
///
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);
}
}
-_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);
}
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);
}
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);
}
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);
}
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);
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);
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);
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);
}
#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
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> {};
#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 : "";
}
#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
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;
friend class sub_resource;
friend class mapping;
- friend struct ::_cl_kernel;
+ friend class kernel;
protected:
resource(clover::device &dev, clover::memory_obj &obj);
clover::context &ctx;
- friend class _cl_kernel;
+ friend class clover::kernel;
private:
void *bind(clover::command_queue &q);
projects / mesa.git / commitdiff