nVidia 上的 OpenCL NDRange 尺寸顺序错误？

Question

我知道 OpenCL 这些天相当空闲——尤其是 NVidia 的 CUDA 实现。也就是说，我认为我在 Nvidia 中发现了一个重大错误，我想看看其他人是否注意到了同样的问题。使用带有 C++ 绑定的 Linux Platform Version OpenCL 1.2 CUDA 10.1.0 我一直遇到 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 对全局/局部维度的解释似乎是交错的，与规范不符。这似乎也不涉及 C++ 绑定。本地 ID 不应为零，而 get_global_id(1) 始终为零。

我知道 NVidia 不太关心 OpenCL，但这似乎是一个相当大的问题。其他人遇到这样的事情吗？这不是 printf 的同步问题。我在实际的数据用例中注意到了这一点，并且构建这个内核只是为了演示它。

Answer 1

^{虽然很难详细验证这一点，但我将其作为答案发布，因为根据我的观察，它似乎可以解释这个问题：}

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 内核和不同的数字），输出通常似乎根本没有任何意义。所以我尝试了这个内核：

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 仅用于内核内部的唯一目的，即调试...

---

顺便说一句：在这一点上，规范仍然有点难以解释 - 至少在决定观察到的行为是“正确”还是“错误”时是如此。来自Khronos documentation of printf：

如果同时从多个工作项执行 printf，则无法保证写入数据的顺序。例如，具有全局 id (0,0,1) 的工作项的输出与具有全局 id (0,0,4) 的工作项的输出混合出现是有效的，依此类推.

NVIDIA documentation of the CUDA printf implementation 还包含一些免责声明，并讨论了一些可能被覆盖的缓冲区，但将其（在规范的技术层面上）映射到 OpenCL 行为是困难的......