OpenCL :Access proper index by using globalid(.)

Question

Hi,
I am coding in OpenCL.
I am converting a "C function" having 2D array starting from i=1 and j=1 .PFB .

cv::Mat input; //Input :having some data in it ..
//Image input size is :input.rows=288 ,input.cols =640
cv::Mat output(input.rows-2,input.cols-2,CV_32F); //Output buffer
//Image output size is :output.rows=286 ,output.cols =638 

This is a code Which I want to modify in OpenCL:

for(int i=1;i<output.rows-1;i++)
{
  for(int j=1;j<output.cols-1;j++)
    {
        float xVal = input.at<uchar>(i-1,j-1)-input.at<uchar>(i-1,j+1)+ 2*(input.at<uchar>(i,j-1)-input.at<uchar>(i,j+1))+input.at<uchar>(i+1,j-1) - input.at<uchar>(i+1,j+1);
        float yVal = input.at<uchar>(i-1,j-1) - input.at<uchar>(i+1,j-1)+ 2*(input.at<uchar>(i-1,j)   - input.at<uchar>(i+1,j))+input.at<uchar>(i-1,j+1)-input.at<uchar>(i+1,j+1);
        output.at<float>(i-1,j-1) = xVal*xVal+yVal*yVal;
    }
}

... Host code :

//Input Image size is :input.rows=288 ,input.cols =640 
//Output Image size is :output.rows=286 ,output.cols =638 
 OclStr->global_work_size[0] =(input.cols);
 OclStr->global_work_size[1] =(input.rows);

 size_t outBufSize = (output.rows) * (output.cols) * 4;//4 as I am copying all 4 uchar values into one float variable space

    cl_mem cl_input_buffer = clCreateBuffer(
        OclStr->context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR ,
        (input.rows) * (input.cols),
        static_cast<void *>(input.data), &OclStr->returnstatus);

    cl_mem cl_output_buffer = clCreateBuffer(
        OclStr->context, CL_MEM_WRITE_ONLY| CL_MEM_USE_HOST_PTR , 
        (output.rows) * (output.cols) * sizeof(float), 
        static_cast<void *>(output.data), &OclStr->returnstatus);

OclStr->returnstatus = clSetKernelArg(OclStr->objkernel, 0, sizeof(cl_mem), (void *)&cl_input_buffer);
OclStr->returnstatus = clSetKernelArg(OclStr->objkernel, 1, sizeof(cl_mem), (void *)&cl_output_buffer); 

    OclStr->returnstatus = clEnqueueNDRangeKernel(
        OclStr->command_queue, 
        OclStr->objkernel, 
        2, 
        NULL, 
        OclStr->global_work_size, 
        NULL, 
        0, 
        NULL, 
        NULL
        );
clEnqueueMapBuffer(OclStr->command_queue, cl_output_buffer, true, CL_MAP_READ, 0, outBufSize, 0, NULL, NULL, &OclStr->returnstatus);    

kernel Code :

__kernel void Sobel_uchar (__global uchar *pSrc, __global float *pDstImage)              
{                                                                                      
const uint cols = get_global_id(0)+1;                                              
const uint rows = get_global_id(1)+1;                                              
const uint width= get_global_size(0);                                              
uchar Opsoble[8];                                                                  
Opsoble[0] = pSrc[(cols-1)+((rows-1)*width)];                           
Opsoble[1] = pSrc[(cols+1)+((rows-1)*width)];                           
Opsoble[2] = pSrc[(cols-1)+((rows+0)*width)];                           
Opsoble[3] = pSrc[(cols+1)+((rows+0)*width)];                           
Opsoble[4] = pSrc[(cols-1)+((rows+1)*width)];                           
Opsoble[5] = pSrc[(cols+1)+((rows+1)*width)];                           
Opsoble[6] = pSrc[(cols+0)+((rows-1)*width)];                           
Opsoble[7] = pSrc[(cols+0)+((rows+1)*width)];                           
float gx =   Opsoble[0]-Opsoble[1]+2*(Opsoble[2]-Opsoble[3])+Opsoble[4]-Opsoble[5];
float gy =   Opsoble[0]-Opsoble[4]+2*(Opsoble[6]-Opsoble[7])+Opsoble[1]-Opsoble[5];
pDstImage[(cols-1)+(rows-1)*width] = gx*gx + gy*gy;                                

    } 

Here I am not able to get the output as expected. I am having some questions that

1. My for loop is starting from i=1 instead of zero, then How can I get proper index by using the global_id() in x and y direction
2. What is going wrong in my above kernel code :(

I am suspecting there is a problem in buffer stride but not able to further break my head as already broke it throughout a day :( I have observed that with below logic output is skipping one or two frames after some 7/8 frames sequence. I have added the screen shot of my output which is compared with the reference output. My above logic is doing partial sobelling on my input .I changed the width as -

const uint width = get_global_size(0)+1;

PFB

Your suggestions are most welcome !!!

Answer 1

It looks like you may be fetching values in (y,x) format in your opencl version. Also, you need to add 1 to the global id to replicate your for loops starting from 1 rather than 0.

I don't know why there is an unused iOffset variable. Maybe your bug is related to this? I removed it in my version.

Does this kernel work better for you?

__kernel void simple(__global uchar *pSrc, __global float *pDstImage)              
{                                                                                      
   const uint i = get_global_id(0) +1;                                                  
   const uint j = get_global_id(1) +1;                                                  
   const uint width = get_global_size(0) +2;                                              

   uchar Opsoble[8];                                                                   
   Opsoble[0] = pSrc[(i-1) + (j - 1)*width];                                           
   Opsoble[1] = pSrc[(i-1) + (j + 1)*width];                                           
   Opsoble[2] = pSrc[i + (j-1)*width];                                                 
   Opsoble[3] = pSrc[i + (j+1)*width];                                                 
   Opsoble[4] = pSrc[(i+1) + (j - 1)*width];                                           
   Opsoble[5] = pSrc[(i+1) + (j + 1)*width];                                           
   Opsoble[6] = pSrc[(i-1) + (j)*width];                                               
   Opsoble[7] = pSrc[(i+1) + (j)*width];
   float gx =   Opsoble[0]-Opsoble[1]+2*(Opsoble[2]-Opsoble[3])+Opsoble[4]-Opsoble[5]; 
   float gy =   Opsoble[0]-Opsoble[4]+2*(Opsoble[6]-Opsoble[7])+Opsoble[1]-Opsoble[5]; 
   pDstImage[(i-1) + (j-1)*width] = gx*gx + gy*gy ;                                     
}

Answer 2

I am a bit apprehensive about posting an answer suggesting optimizations to your kernel, seeing as the original output has not been reproduced exactly as of yet. There is a major improvement available to be made for problems related to image processing/filtering.

Using local memory will help you out by reducing the number of global reads by a factor of eight, as well as grouping the global writes together for potential gains with the single write-per-pixel output.

The kernel below reads a block of up to 34x34 from pSrc, and outputs a 32x32(max) area of the pDstImage. I hope the comments in the code are enough to guide you in using the kernel. I have not been able to give this a complete test, so there could be changes required. Any comments are appreciated as well.

__kernel void sobel_uchar_wlocal (__global uchar *pSrc, __global float *pDstImage, __global uint2 dimDstImage)
{
    //call this kernel 1-dimensional work group size: 32x1
    //calculates 32x32 region of output with 32 work items

    const uint wid = get_local_id(0);
    const uint wid_1 = wid+1; // corrected for the calculation step
    const uint2 gid = (uint2)(get_group_id(0),get_group_id(1));
    const uint localDim = get_local_size(0);

    const uint2 globalTopLeft = (uint2)(localDim * gid.x, localDim * gid.y); //position in pSrc to copy from/to

    //dimLocalBuff is used for the right and bottom edges of the image, where the work group may run over the border
    const uint2 dimLocalBuff = (uint2)(localDim,localDim);
    if(dimDstImage.x - globalTopLeft.x < dimLocalBuff.x){
        dimLocalBuff.x = dimDstImage.x - globalTopLeft.x;
    }
    if(dimDstImage.y - globalTopLeft.y < dimLocalBuff.y){
        dimLocalBuff.y = dimDstImage.y - globalTopLeft.y;
    }

    int i,j;

    //save region of data into local memory
    __local uchar srcBuff[34][34]; //34^2 uchar = 1156 bytes
    for(j=-1;j<dimLocalBuff.y+1;j++){
        for(i=x-1;i<dimLocalBuff.x+1;i+=localDim){
            srcBuff[i+1][j+1] = pSrc[globalTopLeft.x+i][globalTopLeft.y+j];
        }
    }
    mem_fence(CLK_LOCAL_MEM_FENCE);

    //compute output and store locally
    __local float dstBuff[32][32]; //32^2 float = 4096 bytes
    if(wid_1 < dimLocalBuff.x){
        for(i=0;i<dimLocalBuff.y;i++){
            float gx = srcBuff[(wid_1-1)+ (i - 1)]-srcBuff[(wid_1-1)+ (i + 1)]+2*(srcBuff[wid_1+ (i-1)]-srcBuff[wid_1+ (i+1)])+srcBuff[(wid_1+1)+ (i - 1)]-srcBuff[(wid_1+1)+ (i + 1)]; 
            float gy = srcBuff[(wid_1-1)+ (i - 1)]-srcBuff[(wid_1+1)+ (i - 1)]+2*(srcBuff[(wid_1-1)+ (i)]-srcBuff[(wid_1+1)+ (i)])+srcBuff[(wid_1-1)+ (i + 1)]-srcBuff[(wid_1+1)+ (i + 1)]; 
            dstBuff[wid][i] = gx*gx + gy*gy;
        }
    }
    mem_fence(CLK_LOCAL_MEM_FENCE);

    //copy results to output
    for(j=0;j<dimLocalBuff.y;j++){
        for(i=0;i<dimLocalBuff.x;i+=localDim){
            srcBuff[i][j] = pSrc[globalTopLeft.x+i][globalTopLeft.y+j];
        }
    }
}