I have a problem that the kernel I'm developing accidentally crashes. This means that it drops every 10th run. I think my kernel may be too complex, but reducing complexity didn’t really help.
On failure clEnqueueNDRangeKernel
, an error is not returned, but the next clFinish
returns -36
( CL_INVALID_COMMAND_QUEUE
), and the next clEnqueueReadBuffer
returns an error -5
( CL_OUT_OF_RESOURCES
).
So my questions are:
- What is the cause of error messages and crashes?
- Is my core too complex? What is your experience?
- Is there a way to find out how complex a kernel can be without waiting for a crash?
If there is no information that you find useful, leave a comment.
I tried to reduce my kernel to a minimal example that shows an error. It looks like this:
" __kernel void myKernel1(",
" __local float* x,",
" __local float* y,",
" __global float* z,",
" __global float* b1,",
" __global float* data,",
" __global float* classes,",
" int nsamples,",
" int nfeatures,",
" int sizeb,",
" int nclasses,",
" int groupsize,",
" )",
" {",
" int local_id = get_local_id(0); ",
" int j = get_global_id(0); ",
" int cls = classes[j];",
" for(int k = 0; k<nfeatures; k++) x[k] = data[j*nfeatures+k];",
" float target[20];",
" for(int k = 0; k<nclasses; k++) target[k] = 0;",
" z[1]=z[1]+(float)get_local_id(0);",
" b1[2]=b1[2]+(float)local_id+get_group_id(0)*groupsize;",
" target[cls] =1;",
" int l1 = 0;",
" for(int l1 = 0; l1<sizeb ; l1++) {",
" y[l1+local_id*groupsize]=b1[l1];",
" for(int l2 = 0; l2<nfeatures; l2++){",
" }",
" }",
Buffers:
cl_mem z_cl = clCreateBuffer(GPUContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, niters * nclasses * sizeof(*z), z, &_err);
cl_mem b1_cl = clCreateBuffer(GPUContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeb*sizeof(*b1), b1, &_err);
cl_mem gpu_data_cl = clCreateBuffer(GPUContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, nsamples * nfeatures * sizeof(*gpu_data), gpu_data, &_err);
cl_mem gpu_classes_cl = clCreateBuffer(GPUContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, nsamples * sizeof(*gpu_classes), gpu_classes, &_err);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(float)*nfeatures, NULL);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(float)*nhidden*groupsize,NULL);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(cl_mem),(void*)&z_cl);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(cl_mem),(void*)&b1_cl);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(cl_mem),(void*)&gpu_data_cl);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(cl_mem),(void*)&gpu_classes_cl);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(int), (void *) &nsamples);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(int), (void *) &nfeatures);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(int), (void *) &sizeb);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(int), (void *) &nclasses);
_err = clSetKernelArg(myKernel1, ArgCounter++, sizeof(int), (void *) &groupsize);
WITH
- niters ~ = 10000
- nclasses ~ = 10
- sizeb ~ = 80
- nsamples ~ = 50000
- nfeatures ~ = 10
I am using Quadro FX 580 with driver version 260.19.21 under Ubuntu 10.10 64Bit.
Thanks for your time reading this !!!
[update]
Examples of the SDK, such as oclBandwidthTest, and I check for errors after each cl command, compare the creation and launch of the command queue:
cl_device_id* init_opencl(cl_context *GPUContext,cl_command_queue *GPUCommandQueue, cl_kernel* cl_myKernel1,cl_program *OpenCLProgram){
cl_int _err=0;
cl_platform_id cpPlatform;
cl_device_id cdDevice;
_err = clGetPlatformIDs(1, &cpPlatform, NULL);
if(_err || VERBOSE)printf("clGetPlatformIDs:%i\n",_err);
_err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
if(_err || VERBOSE)printf("clGetDeviceIDs:%i\n",_err);
*GPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &_err);
if(_err || VERBOSE)printf("clCreateContextFromType:%i\n",_err);
size_t ParmDataBytes;
_err = clGetContextInfo(*GPUContext, CL_CONTEXT_DEVICES, 0, NULL, &ParmDataBytes);
if(_err || VERBOSE)printf("clGetContextInfo:%i\n",_err);
cl_device_id* GPUDevices;
GPUDevices = (cl_device_id*)malloc(ParmDataBytes);
_err = clGetContextInfo(*GPUContext, CL_CONTEXT_DEVICES, ParmDataBytes, GPUDevices, NULL);
if(_err || VERBOSE)printf("clGetContextInfo:%i\n",_err);
*GPUCommandQueue = clCreateCommandQueue(*GPUContext, GPUDevices[0], 0, &_err);
if(_err || VERBOSE)printf("clCreateCommandQueue:%i\n",_err);
*OpenCLProgram = clCreateProgramWithSource(*GPUContext, sizeof(OpenCLSource)/sizeof(char *), OpenCLSource, NULL, &_err);
if(_err || VERBOSE)printf("CreateProgramWithSource:%i\n",_err);
char * buildoptions= "-Werror";
_err= clBuildProgram(*(OpenCLProgram), 0, NULL, buildoptions, NULL, NULL);
if(_err != CL_SUCCESS){
if(_err || VERBOSE)printf("clBuildProgram:%i\n",_err);
cl_build_status build_status;
_err = clGetProgramBuildInfo(*(OpenCLProgram), GPUDevices[0], CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &build_status, NULL);
char *build_log;
size_t ret_val_size;
_err = clGetProgramBuildInfo(*(OpenCLProgram), GPUDevices[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
build_log = (char*)malloc(ret_val_size+1);
_err = clGetProgramBuildInfo(*(OpenCLProgram), GPUDevices[0], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
build_log[ret_val_size] = '\0';
printf("BUILD LOG: \n %s", build_log);
}
*cl_myKernel1 = clCreateKernel(*(OpenCLProgram), "myKernel1", &_err);
if(_err || VERBOSE)printf("clCreateKernel:%i\n",_err);
if (VERBOSE){
size_t workgroupsize;
cl_uint devicedata;
size_t maxitems[3];
clGetKernelWorkGroupInfo(*cl_myKernel1,GPUDevices[0], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroupsize, NULL);
printf("CL_KERNEL_WORK_GROUP_SIZE:%i (recommended workgroupsize for the used kernel)\n",workgroupsize);
clGetDeviceInfo(GPUDevices[0], CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint), &devicedata, NULL);
printf("CL_DEVICE_ADDRESS_BITS:%i\n",devicedata);
clGetDeviceInfo(GPUDevices[0], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &devicedata, NULL);
printf("CL_DEVICE_MAX_COMPUTE_UNITS:%i\n",devicedata);
_err= clGetDeviceInfo(GPUDevices[0], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof( maxitems), &maxitems, NULL);
printf("CL_DEVICE_MAX_WORK_ITEM_SIZES:%i,%i,%i error=%i\n",maxitems[0],maxitems[1],maxitems[2],_err);
_err= clGetDeviceInfo(GPUDevices[0], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof( maxitems), &maxitems, NULL);
printf("CL_DEVICE_MAX_WORK_GROUP_SIZE:%i,%i,%i error=%i\n",maxitems[0],maxitems[1],maxitems[2],_err);
printf("Lines of CL code: %i\n",sizeof(OpenCLSource)/sizeof(char*));
getchar();
}
return GPUDevices;
}
Launch:
clEnqueueNDRangeKernel(GPUCommandQueue, cl_myKernel1, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
if(_err!=CL_SUCCESS)printf("\nclEnqueueNDRangeKernel:%i\n",_err);
_err = clFinish(GPUCommandQueue);
if(_err!=CL_SUCCESS)printf("\nclFinish GPUCommandQueue:%i\n",_err);