3

我正在使用 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。在内核之后，我期望以下值： dataSet[0]=0 dataSet[1]=0 dataSet[2]=1 ... dataSet[63]=62 dataSet[64]=63 dataSet[65]=63 dataSet [66]=65 ...数据集[127]=126

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

我的理解是第一个波前（64 个线程）应该将 dataSet[64] 更改为 63。所以当执行第二个波前时，线程 #64 应该得到 63 并将其写入 dataSet[65]。但我看到 dataSet[65] 仍然是 64。为什么？

4

2 回答 2

2

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

另外假设 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）。

于 2013-11-09T10:38:59.363 回答
0

CUDA 模型 - 什么是经纱尺寸？ 我认为这是一个很好的答案，它简要解释了经线。

但是我对Sharpneli所说的话有点困惑，例如“ [如果将其设置为512，它几乎肯定会失败，规范不需要实现来支持任意局部大小。在AMD HW中，局部大小正是波前大小.同样适用于Nvidia。一般来说，你并不需要关心实现将如何处理它。]"。

我认为本地大小意味着组大小是由程序员设置的。但是当执行器发生时，细分组由诸如warp之类的硬件设置。

于 2016-12-29T08:26:28.420 回答