nVidia 上的 OpenCL NDRange 尺寸顺序错误？
OpenCL NDRange dimensions order bug on nVidia?
我知道这些天 OpenCL 相当空闲——尤其是 NVidia 的 CUDA 实现。也就是说，我认为我在 Nvidia 中发现了一个重大错误，我想看看其他人是否注意到了同样的错误。使用 Linux 平台版本 OpenCL 1.2 CUDA 10.1.0 和 C++ 绑定 我一直遇到 NDRange 顺序的各种问题，我终于有了一个可以明确重现的简单内核问题：
void kernel test()
{
printf("G0:%d G1:%d G2:%d L0:%d L1:%d L2:%d\n",
get_global_id(0),
get_global_id(1),
get_global_id(2),
get_local_id(0),
get_local_id(1),
get_local_id(2));
}
如果我用 3 个维度对这个内核进行排队：全局 (4,3,2) 和局部 (1,1,1)：
queue.enqueueNDRangeKernel(kernel, cl::NullRange,
cl::NDRange(4, 3, 2),
cl::NDRange(1, 1, 1),
NULL, events);
它在 AMD/Intel 上随机正确输出以下内容（为清楚起见，随机输出排序）：
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0
G0:0 G1:1 G2:0 L0:0 L1:0 L2:0
G0:0 G1:1 G2:1 L0:0 L1:0 L2:0
G0:0 G1:2 G2:0 L0:0 L1:0 L2:0
G0:0 G1:2 G2:1 L0:0 L1:0 L2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0
G0:1 G1:1 G2:0 L0:0 L1:0 L2:0
G0:1 G1:1 G2:1 L0:0 L1:0 L2:0
G0:1 G1:2 G2:0 L0:0 L1:0 L2:0
G0:1 G1:2 G2:1 L0:0 L1:0 L2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0
G0:2 G1:1 G2:0 L0:0 L1:0 L2:0
G0:2 G1:1 G2:1 L0:0 L1:0 L2:0
G0:2 G1:2 G2:0 L0:0 L1:0 L2:0
G0:2 G1:2 G2:1 L0:0 L1:0 L2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0
G0:3 G1:1 G2:0 L0:0 L1:0 L2:0
G0:3 G1:1 G2:1 L0:0 L1:0 L2:0
G0:3 G1:2 G2:0 L0:0 L1:0 L2:0
G0:3 G1:2 G2:1 L0:0 L1:0 L2:0
这符合规范。但是，如果我使用 NVidia 安排具有相同尺寸的完全相同的内核，我将得到以下输出：
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0
G0:0 G1:0 G2:0 L0:0 L1:1 L2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0
G0:0 G1:0 G2:1 L0:0 L1:1 L2:0
G0:0 G1:0 G2:2 L0:0 L1:0 L2:0
G0:0 G1:0 G2:2 L0:0 L1:1 L2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0
G0:1 G1:0 G2:0 L0:0 L1:1 L2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0
G0:1 G1:0 G2:1 L0:0 L1:1 L2:0
G0:1 G1:0 G2:2 L0:0 L1:0 L2:0
G0:1 G1:0 G2:2 L0:0 L1:1 L2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0
G0:2 G1:0 G2:0 L0:0 L1:1 L2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0
G0:2 G1:0 G2:1 L0:0 L1:1 L2:0
G0:2 G1:0 G2:2 L0:0 L1:0 L2:0
G0:2 G1:0 G2:2 L0:0 L1:1 L2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0
G0:3 G1:0 G2:0 L0:0 L1:1 L2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0
G0:3 G1:0 G2:1 L0:0 L1:1 L2:0
G0:3 G1:0 G2:2 L0:0 L1:0 L2:0
G0:3 G1:0 G2:2 L0:0 L1:1 L2:0
NVidia 对 global/local 维度的解释似乎是交错的，这与规范不符。这似乎也不涉及 C++ 绑定。本地 ID 只能为零，并且 get_global_id(1) 始终为零。
我知道 NVidia 不太关心 OpenCL，但这似乎是一个相当大的问题。还有其他人遇到过这样的事情吗？这不是 printf 的同步问题。我在实际数据用例中注意到了它，构建这个内核只是为了演示它。
虽然很难详细验证，但我会post它作为答案，因为根据我的观察，似乎说明问题：
tl;dr: 原因几乎可以肯定是由于 printf.
中缺少同步
首先，我观察到与您相同的行为：在 AMD 上，输出似乎是正确的。在 NVIDIA 上，这似乎是令人恼火的错误。所以我很好奇，扩展了内核，也打印了 get_local_size:
void kernel test()
{
printf("G0:%d G1:%d G2:%d L0:%d L1:%d L2:%d S0:%d S1:%d S2:%d\n",
get_global_id(0),
get_global_id(1),
get_global_id(2),
get_local_id(0),
get_local_id(1),
get_local_id(2),
get_local_size(0),
get_local_size(1),
get_local_size(2));
}
现在，get_local_id当然必须小于这个大小，否则大多数内核都会崩溃。在 AMD 上，输出非常干净：
platform AMD Accelerated Parallel Processing
device Spectre
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
在 NVIDIA 上，输出为
platform NVIDIA CUDA
device GeForce GTX 970
G0:3 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
现在，不可能是正确的：本地工作大小始终为 0！
经过一些进一步的测试（例如使用 2D 内核和不同的数字），输出 generally 似乎根本没有任何意义。所以我尝试了这个内核：
void kernel test()
{
printf("G0:%d\n", get_global_id(0));
printf("G1:%d\n", get_global_id(1));
printf("G2:%d\n", get_global_id(2));
printf("L0:%d\n", get_local_id(0));
printf("L1:%d\n", get_local_id(1));
printf("L2:%d\n", get_local_id(2));
printf("S0:%d\n", get_local_size(0));
printf("S1:%d\n", get_local_size(1));
printf("S2:%d\n", get_local_size(2));
}
在NVIDIA上，则输出为
platform NVIDIA CUDA
device GeForce GTX 970
G0:1
G0:1
G0:1
G0:2
G0:2
G0:2
G0:2
G0:2
G0:3
G0:2
G0:3
G0:3
G0:0
G0:3
G0:3
G0:0
G0:0
G0:3
G0:0
G0:0
G0:0
G0:1
G0:1
G0:1
G1:2
G1:2
G1:0
G1:0
G1:1
G1:2
G1:2
G1:1
G1:1
G1:1
G1:0
G1:0
G1:2
G1:1
G1:0
G1:0
G1:2
G1:1
G1:1
G1:0
G1:2
G1:2
G1:0
G1:1
G2:0
G2:0
G2:1
G2:1
G2:0
G2:0
G2:1
G2:0
G2:0
G2:0
G2:0
G2:0
G2:1
G2:1
G2:0
G2:1
G2:1
G2:1
G2:1
G2:0
G2:1
G2:0
G2:1
G2:1
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L2:0
L1:0
L1:0
L1:0
L1:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S2:1
S2:1
S1:1
S1:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
重点是：每个人输出正确!。问题似乎是将所有内容都放入一个 printf 会弄乱一些内部缓冲区。
这当然很可惜。它基本上不可能将 printf 用于内核内部合理使用的唯一目的，即 调试 ...
旁白：规范在这一点上仍然有点难以解释——至少在决定观察到的行为是"right"还是"wrong"时是这样。来自 Khronos documentation of printf :
In the case that printf is executed from multiple work-items concurrently, there is no guarantee of ordering with respect to written data. For example, it is valid for the output of a work-item with a global id (0,0,1) to appear intermixed with the output of a work-item with a global id (0,0,4) and so on.
NVIDIA documentation of the CUDA printf implementation 还包含一些免责声明并讨论了一些可能被覆盖的缓冲区，但是将其（在规范的技术层面上）映射到 OpenCL 行为是困难的...
我知道这些天 OpenCL 相当空闲——尤其是 NVidia 的 CUDA 实现。也就是说，我认为我在 Nvidia 中发现了一个重大错误，我想看看其他人是否注意到了同样的错误。使用 Linux 平台版本 OpenCL 1.2 CUDA 10.1.0 和 C++ 绑定 我一直遇到 NDRange 顺序的各种问题，我终于有了一个可以明确重现的简单内核问题：
void kernel test()
{
printf("G0:%d G1:%d G2:%d L0:%d L1:%d L2:%d\n",
get_global_id(0),
get_global_id(1),
get_global_id(2),
get_local_id(0),
get_local_id(1),
get_local_id(2));
}
如果我用 3 个维度对这个内核进行排队：全局 (4,3,2) 和局部 (1,1,1)：
queue.enqueueNDRangeKernel(kernel, cl::NullRange,
cl::NDRange(4, 3, 2),
cl::NDRange(1, 1, 1),
NULL, events);
它在 AMD/Intel 上随机正确输出以下内容（为清楚起见，随机输出排序）：
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0
G0:0 G1:1 G2:0 L0:0 L1:0 L2:0
G0:0 G1:1 G2:1 L0:0 L1:0 L2:0
G0:0 G1:2 G2:0 L0:0 L1:0 L2:0
G0:0 G1:2 G2:1 L0:0 L1:0 L2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0
G0:1 G1:1 G2:0 L0:0 L1:0 L2:0
G0:1 G1:1 G2:1 L0:0 L1:0 L2:0
G0:1 G1:2 G2:0 L0:0 L1:0 L2:0
G0:1 G1:2 G2:1 L0:0 L1:0 L2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0
G0:2 G1:1 G2:0 L0:0 L1:0 L2:0
G0:2 G1:1 G2:1 L0:0 L1:0 L2:0
G0:2 G1:2 G2:0 L0:0 L1:0 L2:0
G0:2 G1:2 G2:1 L0:0 L1:0 L2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0
G0:3 G1:1 G2:0 L0:0 L1:0 L2:0
G0:3 G1:1 G2:1 L0:0 L1:0 L2:0
G0:3 G1:2 G2:0 L0:0 L1:0 L2:0
G0:3 G1:2 G2:1 L0:0 L1:0 L2:0
这符合规范。但是，如果我使用 NVidia 安排具有相同尺寸的完全相同的内核，我将得到以下输出：
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0
G0:0 G1:0 G2:0 L0:0 L1:1 L2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0
G0:0 G1:0 G2:1 L0:0 L1:1 L2:0
G0:0 G1:0 G2:2 L0:0 L1:0 L2:0
G0:0 G1:0 G2:2 L0:0 L1:1 L2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0
G0:1 G1:0 G2:0 L0:0 L1:1 L2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0
G0:1 G1:0 G2:1 L0:0 L1:1 L2:0
G0:1 G1:0 G2:2 L0:0 L1:0 L2:0
G0:1 G1:0 G2:2 L0:0 L1:1 L2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0
G0:2 G1:0 G2:0 L0:0 L1:1 L2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0
G0:2 G1:0 G2:1 L0:0 L1:1 L2:0
G0:2 G1:0 G2:2 L0:0 L1:0 L2:0
G0:2 G1:0 G2:2 L0:0 L1:1 L2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0
G0:3 G1:0 G2:0 L0:0 L1:1 L2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0
G0:3 G1:0 G2:1 L0:0 L1:1 L2:0
G0:3 G1:0 G2:2 L0:0 L1:0 L2:0
G0:3 G1:0 G2:2 L0:0 L1:1 L2:0
NVidia 对 global/local 维度的解释似乎是交错的，这与规范不符。这似乎也不涉及 C++ 绑定。本地 ID 只能为零，并且 get_global_id(1) 始终为零。
我知道 NVidia 不太关心 OpenCL，但这似乎是一个相当大的问题。还有其他人遇到过这样的事情吗？这不是 printf 的同步问题。我在实际数据用例中注意到了它，构建这个内核只是为了演示它。
虽然很难详细验证，但我会post它作为答案，因为根据我的观察，似乎说明问题：
tl;dr: 原因几乎可以肯定是由于 printf.
首先，我观察到与您相同的行为：在 AMD 上，输出似乎是正确的。在 NVIDIA 上，这似乎是令人恼火的错误。所以我很好奇，扩展了内核，也打印了 get_local_size:
void kernel test()
{
printf("G0:%d G1:%d G2:%d L0:%d L1:%d L2:%d S0:%d S1:%d S2:%d\n",
get_global_id(0),
get_global_id(1),
get_global_id(2),
get_local_id(0),
get_local_id(1),
get_local_id(2),
get_local_size(0),
get_local_size(1),
get_local_size(2));
}
现在，get_local_id当然必须小于这个大小，否则大多数内核都会崩溃。在 AMD 上，输出非常干净：
platform AMD Accelerated Parallel Processing
device Spectre
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:1 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:2 G2:0 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:1 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:0 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:1 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:2 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
G0:3 G1:2 G2:1 L0:0 L1:0 L2:0 S0:1 S1:1 S2:1
在 NVIDIA 上，输出为
platform NVIDIA CUDA
device GeForce GTX 970
G0:3 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:2 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:3 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:1 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:0 G1:0 G2:0 L0:0 L1:0 L2:0 S0:0 S1:0 S2:0
G0:2 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:2 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:1 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
G0:1 G1:0 G2:0 L0:0 L1:1 L2:0 S0:0 S1:0 S2:0
现在，不可能是正确的：本地工作大小始终为 0！
经过一些进一步的测试（例如使用 2D 内核和不同的数字），输出 generally 似乎根本没有任何意义。所以我尝试了这个内核：
void kernel test()
{
printf("G0:%d\n", get_global_id(0));
printf("G1:%d\n", get_global_id(1));
printf("G2:%d\n", get_global_id(2));
printf("L0:%d\n", get_local_id(0));
printf("L1:%d\n", get_local_id(1));
printf("L2:%d\n", get_local_id(2));
printf("S0:%d\n", get_local_size(0));
printf("S1:%d\n", get_local_size(1));
printf("S2:%d\n", get_local_size(2));
}
在NVIDIA上，则输出为
platform NVIDIA CUDA
device GeForce GTX 970
G0:1
G0:1
G0:1
G0:2
G0:2
G0:2
G0:2
G0:2
G0:3
G0:2
G0:3
G0:3
G0:0
G0:3
G0:3
G0:0
G0:0
G0:3
G0:0
G0:0
G0:0
G0:1
G0:1
G0:1
G1:2
G1:2
G1:0
G1:0
G1:1
G1:2
G1:2
G1:1
G1:1
G1:1
G1:0
G1:0
G1:2
G1:1
G1:0
G1:0
G1:2
G1:1
G1:1
G1:0
G1:2
G1:2
G1:0
G1:1
G2:0
G2:0
G2:1
G2:1
G2:0
G2:0
G2:1
G2:0
G2:0
G2:0
G2:0
G2:0
G2:1
G2:1
G2:0
G2:1
G2:1
G2:1
G2:1
G2:0
G2:1
G2:0
G2:1
G2:1
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L2:0
L1:0
L1:0
L1:0
L1:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S2:1
S2:1
S1:1
S1:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
重点是：每个人输出正确!。问题似乎是将所有内容都放入一个 printf 会弄乱一些内部缓冲区。
这当然很可惜。它基本上不可能将 printf 用于内核内部合理使用的唯一目的，即 调试 ...
旁白：规范在这一点上仍然有点难以解释——至少在决定观察到的行为是"right"还是"wrong"时是这样。来自 Khronos documentation of printf :
In the case that printf is executed from multiple work-items concurrently, there is no guarantee of ordering with respect to written data. For example, it is valid for the output of a work-item with a global id (0,0,1) to appear intermixed with the output of a work-item with a global id (0,0,4) and so on.
NVIDIA documentation of the CUDA printf implementation 还包含一些免责声明并讨论了一些可能被覆盖的缓冲区，但是将其（在规范的技术层面上）映射到 OpenCL 行为是困难的...