OPENCL - 转移环球存储工作组+边境本地内存

Question

下面的代码草稿我公司生产的：

void __kernel myKernel(__global const short* input, 
         __global short* output, 
         const int width, 
         const int height){       


    // Always square. (and 16x16 in our example) 
    const uint local_size = get_local_size(0); 

    // Get the work-item col/row index 
    const uint wi_c = get_local_id(0); 
    const uint wi_r = get_local_id(1); 

    // Get the global col/row index 
    const uint g_c = get_global_id(0); 
    const uint g_r = get_global_id(1); 

    // Declare a local array NxN 
    const uint arr_size = local_size *local_size ; 
    __local short local_in[arr_size]; 

    // Transfer the global memory for into a local one. 
    local_in[wi_c + wi_r*local_size ] = input[g_c + g_r*width]; 

    // Wait that all the work-item are sync 
    barrier(CLK_LOCAL_MEM_FENCE); 

    // Now add code to process on the local array (local_in). 

据我了解的OpenCL工作组/工作项目，这正是我需要的可以将全球范围内的全局16x16 ROI复制到本地内存。 （如果我错了，请纠正我，因为我是从这开始的）。

因此，在屏障之后，local_in中的每个元素都可以通过wi_c + wi_r*local_size访问。

但是，现在让我们做一些棘手。如果我想让我的工作组中的每个工作项目在3x3邻域上工作，我需要一个18x18的local_in数组。

但如何建立呢？由于我只有16x16 = 256个工作项目（线程），但我需要18x18 = 324（缺少68个线程来完成）。

我的基本思路应该是这样做：

if(wi_c == 0 && wi_r == 0){ 
    // Code that copy the border into the new array that should be 
    // local_in[(local_size+2)*(local_size+2)];   
} 

但是，这是可怕的，因为第一个工作项（第1线）将不得不处理所有的边界和工作项的休息这个小组将只是等待这第一个工作项目完成。 （同样，这是我对OpenCL的理解，可能是错误的）。

因此，这里是我的真正的问题：

1. 是有这种问题的另一个简单的解决方案？像改变NDRange本地大小重叠或什么的？
2. 我开始阅读有关凝聚的内存访问，是我的代码看起来像它初稿？我不这么认为，因为我正在使用“跨步”方法来加载全局内存。但我不明白我如何才能将代码的第一部分更改为高效。
3. 一旦达到屏障，每个工作项的处理都会继续，以获得需要存储回全局输出数组的最终值。我应该在这个“写作”还是所有的好东西之前再次设置一个障碍，让所有的工作项目完成他们的自我？

Answer 1

我尝试不同的方法和我带着最后的版本，这是少的“如果”，并使用线程尽可能多地（在第二阶段，可能无法完全有效的，因为一些线程是空闲的，但它的最好的我能够得到）。

原理是在左上角设置原点（start pos），并使用循环索引从此位置创建读/写索引。循环从2D中的本地ID位置开始。所以，所有256个工作项目写自己的第一个元素，并在第二阶段只在256 68工作项目将完成2个底部行+ 2个右列。我不是OpenCL专业版，所以这仍然可以有更多的改进（也许循环展开，我不知道）。

__local float wrkSrc[324]; 
    const int lpitch = 18; 

    // Add halfROI to handle the corner 
    const int lcol = get_local_id(0); 
    const int lrow = get_local_id(1); 

    const int2 gid = { col, row }; 
    const int2 lid = { lcol, lrow }; 

    // Always get the most Top-left corner of that ROI to extract. 
    const int2 startPos = gid - lid - halfROI; 

    // Loop on each thread to get their right ID. 
    // Thread with id < 2 * halfROI will process more then others, but not that much an issue. 
    for (int x = lid.x; x < lpitch; x += 16) { 
     for (int y = lid.y; y < lpitch; y += 16) { 

      // Get the position to write into the local array. 
      const int lidx = x + y * lpitch; 

      // Get the position to read into the global memory (src) 
      const int2 readPos = startPos + (int2)(x, y); 

      // Is inside ? 
      if (readPos.x >= 0 && readPos.x < width && readPos.y >= 0 && readPos.y < height) 
       wrkSrc[lidx] = src[readPos.x + readPos.y * lab_new_pitch]; 
      else 
       wrkSrc[lidx] = 0.0f; 
     } 
    }