OpenCL 中的全局内存是不是连续

Posted 2023-02-16

【中文标题】OpenCL 中的全局内存是不是连续

【英文标题】：Is global memory contiguous in OpenCLOpenCL 中的全局内存是否连续

【发布时间】：2016-01-11 23:26:37

【问题描述】：

我是 OpenCL 新手，正在尝试使用 OpenCL 中的全局内存进行 2D 扫描。

我在输出数组中有一些无效值，这让人怀疑 如果全局内存是连续的。因为通过运行以下内核，我在输出数组中发现了一些垃圾值。

这是我的内核。输入和输出都是 8 x 8 二维数组。

#define SWAP(a,b) __global uint *tmp=a;a=b;b=tmp;
__kernel void 2dScan(
    const __global  uint * const input,
    __global  uint * const output,
    __global uint *lb,
    __global uint *lc
    )

    const uint x = get_global_id(0);
    const uint y = get_global_id(1);
    const uint xm = get_global_size(0);
    const uint ym = get_global_size(1);
    uint gs = get_global_size(0) * get_global_size(1);
    uint index = y * xm + x;

    lb[index] = lc[index] = input[index];
    barrier(CLK_GLOBAL_MEM_FENCE);

    for(uint s = 1; s < gs; s <<= 1) 
        if(index > (s-1)) 
            lc[index] = lb[index]+lb[index-s];
         else 
            lc[index] = lb[index];
        
        barrier(CLK_GLOBAL_MEM_FENCE);
        SWAP(lb,lc);
    
    output[index]= lb[index];

如果您能提供一些建议，非常感谢。

【问题讨论】：

我认为你的意思是“连续”。

【参考方案1】：

我有一个疑问，什么是垃圾价值？一个未初始化的还是不正确的？

在我看来，你的 for 循环有问题。您的访问不是确定性的，也就是说，您无法确定输出[2] 的结果。

例如，工作项 #2 在循环内执行 (s = 1)：lc[2] = lb[2] + lb[1]; 好的，问题是：lb[1] 的值是否已被工作项 #1 修改？您的屏障同步了属于同一工作组的工作项。我认为lc[index] = lb[index]+lb[index-s]; 中的访问超出了工作组的“限制”，因为s 值导致了非确定性问题。

【讨论】：

非常感谢您的回答。我所说的垃圾值是指未初始化的值。现在我明白了问题所在。谢谢！

【参考方案2】：

您正在覆盖每个工作项中的 lb 和 lc。因此结果是不确定的，因为 OpenCL 在全局项之间没有同步。 barrier(CLK_GLOBAL_MEM_FENCE); 在全局内存范围内同步本地工作组。

您将获得的结果基于计划模式，并且无法以任何方式预测。但是，它只会发生在本地组边界中。

你需要改变你的算法。

【讨论】：

感谢您的回答！现在我明白barrier(CLK_GLOCAL_MEM_FENCE); 不会同步全局项目。非常感谢！