如何将映像从全局内存复制到openCL中的本地内存

Question

我有两个输入映像，我从主机传递给内核。 我的图像尺寸是370x427。

我想让自己熟悉本地内存，所以我将本地图像传递给内核，并尝试将全局图像复制到本地。

我送我的形象1D array.when我试图显示它不work.My全球worksize结果是{width*height}和我通过null本地大小clEnqueueNDRangeKernel假设opencl会选择本地存储适当的大小。

下面是我的内核代码。

请如果有人能给我一个提示。

 __kernel void computeSSD(__global unsigned char *imgL,__global unsigned char *imgR,__global unsigned char *result,__global unsigned char *diff,int width,int MAX_DISP,int WIN_SIZE,__local unsigned char *localLeft,__local unsigned char *localRight) 
     { 

      int xCord=get_global_id(0); 
      int yCord=get_local_id(0); 

      // copy both images to local memory 

      localRight[yCord]= imgR[yCord]; 
      localLeft[yCord] = imgL[yCord]; 

      barrier(CLK_LOCAL_MEM_FENCE); 
// do operation on local images 
        result[xCord]=localRight[yCord]; 
// 

     } 

Answer 1

如果您使用3x3过滤器过滤图像。每个工作组都需要工作组像素+1边距。

所以你的内核可以是这样的：

__kernel filter(...){ 
    int x_start = get_group_id(0)*get_local_size(0)-1; 
    int y_start = get_group_id(1)*get_local_size(1)-1; 
    int x_end = (get_group_id(0)+1)*get_local_size(0)+1; 
    int y_end = (get_group_id(1)+1)*get_local_size(1)+1; 

    __local mytype l[18][18]; //just an example for work sizes 16x16! 

    //Fetch 
    //Normally a direct operation per work item is preferred, since it is simpler and the driver will pack all the memory accesses together. 
    //So just follow coalesced access 
    //Using CL async workgroup copy has a complex sintax 
    for(int j=y_start+get_local_id(1); j<y_end; j+=get_local_size(1)){ 
     for(int i=x_start+get_local_id(0); i<x_end; i+=get_local_size(0)){ 
      l[j-y_start][i-x_start] = global[j][i]; 
     } 
    } 
    barrier(CLK_GLOBAL_MEM); 

    //Use the memory for your filtering! (remember to substract y_start & x_start from your original indexes) 
    //.... 
} 

本地内存是不是一个“耶让所有复制到本地，然后用它，所以它会更快。”这是关于复制该组需要的区域，并在本地重新使用这些读取，因此避免从全局（即3x3过滤器，从全局8次冗余读取，将内存需求减少到1/9）进行大量读取。

此外，如果您没有真正重复使用工作项之间的读取，本地内存可能会变慢。如果每个项目只读写一个内存位置，本地内存将使内核变得更慢，而不是更快。