如何在 OpenCL 中验证波前/翘曲尺寸？答案

【问题标题】：How to verify wavefront/warp size in OpenCL?如何在 OpenCL 中验证波前/翘曲尺寸？
【发布时间】：2013-11-21 04:46:00
【问题描述】：

我正在使用 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，这和我的预期不太一样。

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

【问题讨论】：

您不应该尝试验证经线或波前尺寸。如果您编写的代码测试经线大小为 32 和 64，那么当您使用的设备的经线大小为 8、16 或 48 时会发生什么？如果 Nvidia 或 AMD 改变它们的经线/波前尺寸会发生什么。如果您试图找到最佳的工作组规模，那么最好的解决方案是编写一个迷你基准测试来测试所有合理的配置（或至少一个足够的子集）。

标签： opencl

【解决方案1】：

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

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

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

根据您的 cmets 我编辑了这部分：

安装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）。

【讨论】：

首先感谢您的回复！我知道使用屏障会强制同步。我想要的只是通过利用属性“同一波前的所有线程执行相同的指令”来避免障碍（出于性能原因）。但是，结果并不如预期。
关于“get_local_size(0)”：我认为它应该返回工作组中的线程数（由我设置），而不是波前的大小。正确的？例如，如果我将工作组大小设置为 512，get_local_size(0) 将返回 512，而不是 64。
如果您将其设置为 512，它几乎肯定会失败，规范不需要实现来支持任意本地大小。在 AMD HW 中，局部尺寸正好是波前尺寸。同样适用于英伟达。一般来说，你并不需要关心实现将如何处理它。
是的，512 会失败，因为 CL_KERNEL_WORK_GROUP_SIZE 返回 256。但是，我可以将工作组大小设置为 256。您的意思是这意味着波前大小为 256？您是否有任何官方链接来支持您的主张“本地尺寸正是波前尺寸”？
我现在编辑了主要答案，因为您想知道波前大小，所以我有了更好的视图。从锁步中执行指令的线程数量来看，它确实是 64，但它不是单个处理器一次执行的线程数量。例如，如果您有 if(get_global_id(0) == 0) do_something_massive;否则返回；您实际上将停止 256 个线程（或构成该工作组的线程数量，然后由波前组成）。

【解决方案2】：

CUDA model - what is warp size? 我认为这是一个很好的答案，它简要地解释了扭曲。

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

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

【讨论】：