CODEWITHSUNDEEP
crashkernelopencl
shunyo
shunyo

Reputation: 1307

OpenCL - Kernel crashes on the second run

I am trying to run a code which works the first time but crashes the second time that it runs. The function which causes the crash is part of the class Octree_GPU and is this:

int Octree_GPU::runCreateNodeKernel(int length)
{
      cl_uint nodeLength;
      if(nodeNumsArray[length-1] == 0)
            nodeLength = nodeAddArray[length-1];
      else
            nodeLength = nodeAddArray[length-1]+8;
      nodeArray = (cl_uint*)malloc(sizeof(cl_uint)*nodeLength);
      nodePointsArray = (cl_int*)malloc(sizeof(cl_uint)*nodeLength);
      startIndexArray = (cl_int*)malloc(sizeof(cl_int)*nodeLength);

      d_nodeAdd = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_uint)*length, NULL, &err);
      d_nodeArray = clCreateBuffer(context,CL_MEM_READ_WRITE, sizeof(cl_uint)*temp_length, NULL, &err);
      d_numPoints = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_uint)*length, NULL, &err);
      d_pointIndex = clCreateBuffer(context, CL_MEM_READ_WRITE,sizeof(cl_uint)*length,NULL, &err);
      d_nodePointsArray = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int)*temp_length, NULL, &err);
      d_nodeIndexArray = clCreateBuffer(context,CL_MEM_READ_WRITE, sizeof(cl_int)*temp_length, NULL, &err);

      err |= clEnqueueWriteBuffer(commands, d_nodeAdd, CL_TRUE, 0, sizeof(cl_uint)*length, nodeAddArray, 0, NULL,NULL);
      err |= clEnqueueWriteBuffer(commands, d_numPoints,CL_TRUE, 0, sizeof(cl_uint)*length,numPointsArray,0,NULL,NULL);
      err |= clEnqueueWriteBuffer(commands, d_pointIndex, CL_TRUE, 0, sizeof(cl_uint)*length,pointStartIndexArray,0, NULL, NULL);
      clFinish(commands);

      err  = clSetKernelArg(createNodeKernel, 0, sizeof(cl_mem), &d_odata);
      err |= clSetKernelArg(createNodeKernel, 1, sizeof(cl_mem), &d_nodeNums);
      err |= clSetKernelArg(createNodeKernel, 2, sizeof(cl_mem), &d_nodeAdd);
      err |= clSetKernelArg(createNodeKernel, 3, sizeof(cl_mem), &d_numPoints);
      err |= clSetKernelArg(createNodeKernel, 4, sizeof(cl_mem), &d_pointIndex);
      err |= clSetKernelArg(createNodeKernel, 5, sizeof(cl_mem), &d_nodeArray);
      err |= clSetKernelArg(createNodeKernel, 6, sizeof(cl_mem), &d_nodePointsArray);
      err |= clSetKernelArg(createNodeKernel, 7, sizeof(cl_mem), &d_nodeIndexArray);
      clFinish(commands);

      if(err != CL_SUCCESS) {
            printf("Cannot set Kernel Arg \n");
            exit(1);
      }
      size_t global_size[1] = {limit-1};
      err = clEnqueueNDRangeKernel(commands, createNodeKernel, 1, NULL, global_size, NULL, 0, NULL, NULL);
      if(err != CL_SUCCESS) {
            printf(" Kernel does not work \n");
            exit(1);
      }
      clFinish(commands);

      err = clEnqueueReadBuffer(commands, d_nodeArray, CL_TRUE, 0, sizeof(cl_uint)*temp_length, nodeArray, 0, NULL, NULL);
      err|= clEnqueueReadBuffer(commands, d_nodePointsArray, CL_TRUE, 0, sizeof(cl_int)*nodeLength, nodePointsArray, 0, NULL, NULL);
      err|= clEnqueueReadBuffer(commands, d_nodeIndexArray, CL_TRUE, 0, sizeof(cl_int)*nodeLength, startIndexArray, 0, NULL, NULL);
      clFinish(commands);

      clReleaseMemObject(d_nodeAdd);
      clReleaseMemObject(d_numPoints);
      clReleaseMemObject(d_nodeArray);
      clReleaseMemObject(d_nodePointsArray);
      clFinish(commands);
      return 0;
}

Please note that d_odata and d_nodeNums have been declared in the previous functions. The kernel code is given below for the same:

__kernel void createNode(__global int* uniqueCode, __global int* nodeNums,__global int* nodeAdd, __global int* numPoints, __global int* pointIndex,__global int* nodeArray, __global int* nodePoints,__global int* nodeIndex)
{
    int ig = get_global_id(0);
    int add;
    int num = uniqueCode[ig];
    int pt = numPoints[ig];
    int ind = pointIndex[ig];
    int temp,j;
    if(nodeNums[ig] == 8) 
    {
           for(int i=0;i<8;i++)
           {
                    temp = ((int)num/10)*10+i;
                    add = nodeAdd[ig] + i;
                    nodeArray[add] = temp;
                    nodePoints[add] = select(0, pt, temp==num);
                    nodeIndex[add] = select(-1, ind, temp==num);
                    barrier(CLK_LOCAL_MEM_FENCE);
           }
    }
    else
    {
           j = num % 10;
           nodeAdd[ig] = nodeAdd[ig-1];
           add = nodeAdd[ig]+j;
           nodePoints[add] = pt;
           nodeIndex[add] = ind;
           barrier(CLK_LOCAL_MEM_FENCE);
    }

}

I have tried to find out why but have not succeeded. I might be overlooking something really simple. Thank you for your help.

Upvotes: 0

Views: 537

Answers (1)

James B
James B

Reputation: 333

I'm not 100% sure this is causing the crash, but where you've written

if(nodeNums[ig] == 8) 
{
       for(int i=0;i<8;i++)
       {
                barrier(CLK_LOCAL_MEM_FENCE);
       }
}
else
{
       barrier(CLK_LOCAL_MEM_FENCE);
}

This means that different threads in a work group will be executing different numbers of barriers, which may cause a hang/crash. A barrier (with CLK_LOCAL_MEM_FENCE) is for synchronising accesses to local memory, so all work items in a group must execute this before continuing

On a non crash note, it looks like you're using CLK_LOCAL_MEM_FENCE (ensure that local memory accesses are visible across threads) when you mean CLK_GLOBAL_MEM_FENCE (ensure that global memory accesses are visible across threads)

Also

nodeAdd[ig] = nodeAdd[ig-1];

Is not correct for ig == 0. This may not be causing the actual crash (because I've found that OpenCL can be unfortunately quite forgiving), but its worth fixing

Upvotes: 2

Related Questions