opencl - 如何在 OpenCL 中验证波前/扭曲大小？

Question

我使用的是 AMD Radeon HD 7700 GPU。我想使用以下内核来验证波前大小为 64。

__kernel
void kernel__test_warpsize(
        __global T* dataSet,
        uint size
        )
{   
    size_t idx = get_global_id(0);

    T value = dataSet[idx];
    if (idx<size-1)
        dataSet[idx+1] = value;
}

在主程序中，我传递了一个包含 128 个元素的数组。初始值为 dataSet[i]=i。在内核之后，我期望以下值: 数据集[0]=0 数据集[1]=0 数据集[2]=1 ... 数据集[63]=62 数据集[64]=63 数据集[65]=63 数据集[66]=65 ... 数据集[127]=126

但是，我发现dataSet[65]是64，而不是63，这不是我的预期。

我的理解是，第一个 wavefront(64 个线程)应该将 dataSet[64] 更改为 63。因此，当执行第二个 wavefront 时，线程 #64 应该得到 63 并将其写入 dataSet[65]。但是我看到 dataSet[65] 仍然是 64。为什么？

Answer 1

您正在调用未定义的行为。如果您希望访问工作组中另一个线程正在写入的内存，您必须使用屏障。

另外假设 GPU 同时运行 2 个波前。那么 dataSet[65] 确实包含了正确的值，只是第一个波阵面还没有完成。

此外，根据规范，所有项目的输出均为 0 也是有效结果。这是因为一切也可以完全串行执行。这就是您需要障碍的原因。

根据您的意见我编辑了这部分:

安装 http://developer.amd.com/tools-and-sdks/heterogeneous-computing/codexl/阅读:http://developer.amd.com/download/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide.pdf

在一定数量的线程内优化分支只是优化的一小部分。您应该了解 AMD HW 如何在工作组内安排波前，以及它如何通过交错执行波前(在工作组内)来隐藏内存延迟。分支也会影响整个工作组的执行，因为运行它的有效时间与执行单个最长运行波前的时间基本相同(它不能释放本地内存等，直到组中的所有内容都完成，所以它不能安排另一个工作组)。但这也取决于您的本地内存和寄存器使用情况等。要查看实际发生的情况，只需获取 CodeXL 并运行 GPU 分析运行。这将准确显示设备上发生的情况。

即使这仅适用于当前这一代的硬件。这就是为什么这个概念不在 OpenCL 规范本身中的原因。这些属性变化很大，并且在很大程度上取决于硬件。

但如果您真的想知道什么是 AMD 波前大小，答案几乎都是 64(参见 http://devgurus.amd.com/thread/159153 以引用他们的 OpenCL 编程指南)。构成其整个当前阵容的所有 GCN 设备均为 64。也许一些旧设备有 16 或 32，但现在一切都只有 64(对于 nvidia，它通常是 32)。