void Kernel::execute (NDRange global_work_size) {
cl_int err = clEnqueueNDRangeKernel(_controller.cmdQueue(), _id,
global_work_size.dims(), NULL, global_work_size.sizes(),
NULL, 0, NULL, NULL);
clFinish(_controller.cmdQueue());
if (err != CL_SUCCESS)
throw Exception("Could not execute kernel.");
}
bool Kernel::NDRange::divisible (NDRange& range) const {
if (range.dims() != dims())
return false;
for (int i = 0; i < dims(); i++)
if (sizes()[i] % range.sizes()[i] != 0)
return false;
return true;
}
Event CommandQueue::enqueueNDRangeKernel(const Kernel& kernel, const NDRange& globalWorkOffset,
const NDRange& globalWorkSize, const NDRange& localWorkSize,
const Epic::Core::Array<Event>& eventWaitList) const
{
Event ret;
cl_int err = 0;
Array<cl_event> events = eventArrayToCLEventArray(eventWaitList);
err = clEnqueueNDRangeKernel(queueHandle(), kernel.kernelHandle(), globalWorkSize.dimension(),
globalWorkOffset.sizes(), globalWorkSize.sizes(), localWorkSize.sizes(),
events.count(), events.data(), ret.eventPointer());
if(err != CL_SUCCESS) {
throw OpenCLException(err);
}
return ret;
}
void Kernel::execute (NDRange global_work_size, NDRange local_work_size) {
if (global_work_size.dims() != local_work_size.dims())
throw Exception("Number of specified dimensions of global and local work range is not equal!");
if (!global_work_size.divisible(local_work_size))
throw Exception("Global group size not divisible with local group size.");
cl_int err = clEnqueueNDRangeKernel(_controller.cmdQueue(), _id,
global_work_size.dims(), // working dimensions
NULL, // offset
global_work_size.sizes(), // global work size
local_work_size.sizes(), // local work size
0, NULL, NULL);
clFinish(_controller.cmdQueue());
switch (err) {
case CL_SUCCESS:
return;
break;
case CL_INVALID_KERNEL_ARGS:
throw Exception("Invalid kernel arguments.");
case CL_INVALID_WORK_DIMENSION:
throw Exception("Invalid work dimension.");
case CL_INVALID_WORK_GROUP_SIZE:
throw Exception("Invalid work group size.");
case CL_INVALID_WORK_ITEM_SIZE:
throw Exception("Invalid local work group size.");
default:
throw Exception("Could not execute kernel.");
}
}
void CommandQueue::enqueue(Kernel const & kernel, NDRange global, driver::NDRange local, std::vector<Event> const *, Event* event)
{
switch(backend_)
{
case CUDA:
if(event)
check(dispatch::cuEventRecord(event->h_.cu().first, h_.cu()));
check(dispatch::cuLaunchKernel(kernel.h_.cu(), global[0]/local[0], global[1]/local[1], global[2]/local[2],
local[0], local[1], local[2], 0, h_.cu(),(void**)&kernel.cu_params_[0], NULL));
if(event)
check(dispatch::cuEventRecord(event->h_.cu().second, h_.cu()));
break;
case OPENCL:
check(dispatch::clEnqueueNDRangeKernel(h_.cl(), kernel.h_.cl(), global.dimension(), NULL, (const size_t *)global, (const size_t *) local, 0, NULL, event?&event->h_.cl():NULL));
break;
default: throw;
}
}
Event CommandQueue::enqueueNDRangeKernel(const Kernel& kernel, const NDRange& offset, const NDRange& global, const NDRange& local, const vector<Event> *events) {
Event ev;
ClCheck(::clEnqueueNDRangeKernel(Handle(), kernel(), global.dimensions(), offset, global, local, NumEvents(events), ToEventWaitList(events), &ev.m_h));
return ev;
}
