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)];         
}

但这很糟糕，因为第一个工作项（第一个线程）必须处理所有边界，而该组中的其余工作项将只是等待第一个工作项完成。 （再次声明，这是我对 OpenCL 的理解，可能有误）。

所以这是我真正的问题：

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

Answer 1

我尝试了不同的方法，最终得到了最终版本，它减少了“如果”并尽可能多地使用线程（在第二阶段，由于很少有线程空闲，因此可能效率不高，但这是我最好的能够得到）。

原理是在左上角设置一个原点（起始位置），并使用循环索引从这个位置创建读/写索引。循环从 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;
        }
    }