void cleanup()
{
if(handle) {
CHECK_CL(clfftDestroyPlan(&handle));
handle = 0;
}
}
/* H2D */
void copyPitched(cl_mem dst,
const void* src,
cl_command_queue queue,
size_t w,
size_t h,
size_t pitch) const
{
const size_t offset[3] = {0, 0, 0};
size_t region[3] = {w, h, 1};
CHECK_CL(clEnqueueWriteBufferRect( queue,
dst,
Blocking<T_Async>::value,
offset, // buffer origin
offset, // host origin
region,
pitch, // buffer row pitch
0, // buffer slice pitch
0, // host row pitch
0, // host slice pitch
src,
0, // num_events_in_wait_list
nullptr, // event_wait_list
nullptr )); // event
}
void CLKernel::Init()
{
Deinit();
cl_context & ctx = CLManager::Get()->GetContext();
const cl_device_id *device = &(CLManager::Get()->GetDeviceId());
cl_int ciErrNum; // Error code var
printf("CL Compiling file '%s', kernel '%s'....\n", m_filename.c_str(), m_kernelName.c_str() );
while (true)
{
int numBytes; // Byte size of kernel code
const char * cSourceCL = ::TextFileRead( m_filename.c_str(), &numBytes );
// Create the program
size_t programSize = numBytes;
m_clProgram = clCreateProgramWithSource(ctx, 1, (const char **)&cSourceCL, &programSize, &ciErrNum);
CHECK_CL(ciErrNum);
delete[] cSourceCL;
// build program
ciErrNum = clBuildProgram(m_clProgram, 1, device, NULL, NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
size_t len;
char buffer[4096];
clGetProgramBuildInfo(m_clProgram, CLManager::Get()->GetDeviceId(), CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
MESSAGE_BOX("CL Compile Error", "CL '%s'\nCompile Error: %s", m_filename.c_str(), buffer );
}
else
{
printf("CL Compiled file '%s', kernel '%s' successfully!\n", m_filename.c_str(), m_kernelName.c_str() );
// compiled successfully, quit!
break;
}
}
// create kernel
m_clKernel = clCreateKernel(m_clProgram, m_kernelName.c_str(), &ciErrNum);
CHECK_CL(ciErrNum);
}
void CLKernel::Dispatch1D( int cmdQueueIndex, size_t globalWorkSize, size_t localWorkSize, size_t globalWorkOffset )
{
// handle local work size
size_t * localWorkSizePtr = NULL;
if ( localWorkSize > 0 ) localWorkSizePtr = &localWorkSize;
// run it!
cl_command_queue cmdQueue = CLManager::Get()->GetCmdQueue(cmdQueueIndex);
cl_int ciErrNum = clEnqueueNDRangeKernel( cmdQueue, m_clKernel, 1, &globalWorkOffset, &globalWorkSize, localWorkSizePtr, 0, 0, 0 );
CHECK_CL(ciErrNum);
}
/* H2D */
void copy(cl_mem dst, const void* src, size_t size, cl_command_queue queue) const
{
CHECK_CL( clEnqueueWriteBuffer(queue, dst,
Blocking<T_Async>::value,
0, size, src, 0, NULL, NULL) );
}
/* D2D */
void copy(cl_mem dst, cl_mem src, size_t size, cl_command_queue queue) const
{
CHECK_CL( clEnqueueCopyBuffer(queue, src, dst, 0, 0, size, 0, NULL, NULL) );
}
/* D2H */
void copy(void* dst, cl_mem src, size_t size, cl_command_queue queue) const
{
CHECK_CL( clEnqueueReadBuffer(queue, src,
Blocking<T_Async>::value,
0, size, dst, 0, NULL, NULL) );
}
void CLKernel::DispatchND( int cmdQueueIndex, size_t dimens, const size_t *globalWorkSize, const size_t *localWorkSize, const size_t * globalWorkOffset )
{
cl_int ciErrNum;
ciErrNum = clEnqueueNDRangeKernel( CLManager::Get()->GetCmdQueue(cmdQueueIndex), m_clKernel, dimens, globalWorkOffset, globalWorkSize, localWorkSize, 0, 0, 0 );
CHECK_CL(ciErrNum);
}
