内核中的opencl局部数组

Question

我正在尝试使用hough变换来检测二值图像中的圆。并在这里发布我的问题Hough Transform: improving algorithm efficiency over OpenCL

内核代码在执行时非常慢。原子函数中的等待时间和我尝试使用本地内存的速度更快，但是输出累加器in[]内核代码中的问题

   kernel void hough_circle(read_only image2d_t imageIn, global int* in,const int w_hough,__global int * circle)
  {
         sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
     int gid0 = get_global_id(0);
     int gid1 = get_global_id(1);
         const int nWidth = get_local_size(0);

int i = get_group_id(0);
int j = get_group_id(1); 
int idX = get_local_id(0);
int idY = get_local_id(1);
    int ii = i*nWidth + idX; 
int jj = j*nWidth + idY; 
     uint4 pixel;
     int x0=0,y0=0,r;
     int maxval=0;
        local int localImage[25];
        localImage[idY*5+idX] = 0;   
        barrier(CLK_LOCAL_MEM_FENCE);
     pixel=read_imageui(imageIn,sampler,(int2)(gid0,gid1));
     if(pixel.x==255)
     {
    //  #pragma unroll 20
        for(int r=90;r<110;r+=1)
     {
         // int r=190;

            //    #pragma unroll 360
               for(int theta=0; theta<360;theta++)
              {

                          //     x0=(int) round(gid0-r*cos( (float) radians( (float) theta) ));
                          //    y0=(int) round(gid1-r*sin( (float) radians( (float) theta) ));

                              x0=(int) round(gid0-r*sin_parameter[theta] );
                            y0=(int) round(gid1-r*cos_parameter[theta] );
                                  if((x0>0) && (x0<get_global_size(0)) && (y0>0)&&(y0<get_global_size(1)))
                             //  in[w_hough*y0+x0]++;
                                   localImage[5*y0+x0]++;
                            //  atom_inc(&in[w_hough*y0+x0]);
               }
              }

     }


 in[jj*w_hough+ii]=  localImage[jj*5+ii];

  }  

输出累加器in[]中的局部问题显示不正确的值随机谁能帮我做些什么改变来解决这个问题

Answer 1

本地内存是跨工作组共享的，所以我猜第一个问题是如何将线程划分到工作组中？

如果您只是使用clEnqueueNDRangeKernel中的默认值，您将让OpenCL决定哪个线程可以访问本地内存，这样看起来很奇怪的事情就会发生。

也许你的意思是使用private内存？如果是这样，请在声明localImage之前删除local。本地内存的声明类似于全局内存：

__kernel void function(__global float *Input, __local float *WorkgroupBuffer){ }