Kernel randomly crashes with errors 36 and 5_问答_开发者

I have the problem, that the kernel I am developing is crashing randomly. Meaning that it crashes every 10th run. I think, 开发者_如何转开发my kernel could be too complex, but reducing the complexity did not really help.

When it crashes clEnqueueNDRangeKernel returns no error, but the following clFinish returns a -36 (CL_INVALID_COMMAND_QUEUE) and the following clEnqueueReadBuffer return a -5 error (CL_OUT_OF_RESOURCES).

So my questions are:

What could be the reason for the error messages and the crashes?
Is my kernel already 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 are information missing, that you think would be helpfull, please leave a comment.

I tried to reduce my kernel to a minimal example, that stills shows the 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 use a Quadro FX 580 with driver version 260.19.21 under Ubuntu 10.10 64Bit.

Thanks for your time reading this!!!

[update]

The SDK examples like oclBandwidthTest work and I check errors after every cl command, compare my command queue creation and launch:

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;      // OpenCL platform
    cl_device_id cdDevice;          // OpenCL device

    //Get an OpenCL platform
    _err = clGetPlatformIDs(1, &cpPlatform, NULL);
    if(_err || VERBOSE)printf("clGetPlatformIDs:%i\n",_err);

    //Get the devices
    _err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
    if(_err || VERBOSE)printf("clGetDeviceIDs:%i\n",_err);

    // Create a context to run OpenCL on our CUDA-enabled NVIDIA GPU
    *GPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &_err);
    if(_err || VERBOSE)printf("clCreateContextFromType:%i\n",_err);

    // Get the list of GPU devices associated with this context
    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);

    // Create a command-queue on the first GPU device
    *GPUCommandQueue = clCreateCommandQueue(*GPUContext, GPUDevices[0], 0, &_err);
    if(_err || VERBOSE)printf("clCreateCommandQueue:%i\n",_err);

    // Create OpenCL program with source code
    *OpenCLProgram = clCreateProgramWithSource(*GPUContext, sizeof(OpenCLSource)/sizeof(char *), OpenCLSource, NULL, &_err);
    if(_err || VERBOSE)printf("CreateProgramWithSource:%i\n",_err);

    //build OpenCl program
    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);
    }

    //create Kernel
    *cl_myKernel1 = clCreateKernel(*(OpenCLProgram), "myKernel1", &_err);
    if(_err || VERBOSE)printf("clCreateKernel:%i\n",_err);

    //output system info
    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);

CL_OUT_OF_RESOURCES is used by NVIDIA as a "generic error", which means something went wrong, as reading or writing a buffer out of bounds. You should check that.

I know this is a rather vague answer, but that's what this error means (or what is used for).