opencv - OpenCL :Access proper index by using globalid(.)

Question

你好，
我在 OpenCL 中编码。
我正在转换具有从 i=1 和 j=1 .PFB 开始的二维数组的“C 函数”。

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 

这是我想在 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;
    }
}

...
主机代码:

//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 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;                                

    } 

在这里，我无法按预期获得输出。
我有一些问题

我的 for 循环是从 i=1 而不是零开始，那么我如何通过在 x 和 y 方向使用 global_id() 来获得正确的索引

我上面的内核代码出了什么问题:(

我怀疑缓冲步幅存在问题，但由于一整天都已经折断了，因此无法进一步折断我的头:(
我观察到，在某些 7/8 帧序列之后，以下逻辑输出会跳过一帧或两帧。
我添加了与引用输出进行比较的输出的屏幕截图。
我的上述逻辑是对我的输入进行部分sobelling。我将宽度更改为-

const uint width = get_global_size(0)+1;

PFB

非常欢迎您的建议!!!

Answer 1

看起来您可能在您的 opencl 版本中以 (y,x) 格式获取值。此外，您需要将 1 添加到全局 id 以复制从 1 而不是 0 开始的 for 循环。

我不知道为什么有一个未使用的 iOffset 变量。也许您的错误与此有关？我在我的版本中删除了它。

这个内核对你来说更好吗？

__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 ;                                     
}