代码之家 › 专栏 › 技术社区 › Vuwox

Opencl-将全局内存工作组+边界转移到本地内存

opencl memory c++

Vuwox · 技术社区 · 7 年前

这里是我编写的代码草案:

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).

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

但是现在让我们做一些棘手的事情。如果我想让我的工作组中的每个工作项在3x3邻域上工作,我需要一个18x18 local\u阵列。

但如何创造这个呢?因为我只有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)];         
}

这是我真正的问题:

这类问题还有更简单的解决方案吗?比如将NDRange局部大小更改为重叠或其他什么?
我开始阅读关于合并内存访问的文章,我的第一个代码草稿是这样的吗?我不这么认为,因为我正在使用“跨步”方法来加载全局内存。但我不明白我怎么能把代码的第一部分也改成高效的。
一旦到达障碍,将继续处理每个工作项,以获得需要存储回全局输出数组的最终值。我应该在这个“写”之前再设置一个障碍,还是让所有好的工作项目都自己完成?

1 回复 | 直到 7 年前

Vuwox 7 年前

我尝试了不同的方法,最终得到了一个版本,它不那么“如果”,并且尽可能多地使用线程(在第二阶段,可能没有完全有效,因为很少有线程空闲,但这是我能得到的最好的版本)。

其原理是在左上角设置原点(起始位置),并使用循环索引从此位置创建读/写索引。循环从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;
        }
    }