OpenCL 内核的非确定性行为 我注意到调试的事实：问题：

如何解决OpenCL 内核的非确定性行为 我注意到调试的事实：问题：

考虑以下 OpenCL 内核。这是一个从长期的简化过程中产生的 MWE。当然拥有这样的内核是没有意义的，但这不是重点。保证MAX_ORDER >= num_fields，特别是MAX_ORDER=15（它通过clBuildProgram作为编译时常量传递给-D）。

__kernel void MWE_kernel(__global const double* const x,__global const double* alpha,const int num_fields,__global double* const restrict result,__local double* const restrict result_local)
{
    const int id          = get_global_id(0);
    const int global_size = get_global_size(0);

    double sum[MAX_ORDER];
    for (unsigned int i = 0; i < num_fields; i++)
        sum[i] = 0.0;
    
    double sum_sc=0.0;

    for (unsigned int id_mem = id; id_mem < 128 * num_fields; id_mem += global_size) {
        const int alpha_idx = id_mem / 128;
        //sum[0] += 3.0;
        sum[alpha_idx] += 3.0;
        sum_sc += 3.0;
        /*
        if (id == 8){
            printf("sum[%d]=%f   sum_sc=%f\n",alpha_idx,sum[alpha_idx],sum_sc);
        }
        */
    }
    
    if(id == 8){
        result[0] = sum[0];
        result[1] = sum_sc;
    }
}

我用 Intel(R) Gen9 HD Graphics NEO 和 local_size=64 以及 global_size=64 在 GPU 设备（设备名称：num_fields=1）上排队。

我遇到的奇怪行为（我无法解释）是，如果我尝试使用 sum 作为索引填充 alpha_idx 私有数组，那么这对某些线程不起作用，在sum 条目没有增加的感觉。显示这种行为的第一个线程是数字 8，最后一个 if-clause 设置 result[0] = 0 和 result[1]=6，因为我可以在主机上检查。

我注意到调试的事实：

• 设置 MAX_ODER=1（在本例中可能，因为 num_fields=1），使奇怪的行为消失（result[0] 变为 6）。但是，对于 1<MAX_ORDER<21 没有任何变化。令人惊讶的是，MAX_ORDER>=22 更改了内核行为，我得到了 result[0]=6。
• 改变 local_size 和/或 global_size 不会改变任何东西，除非 local_size 减少到 8 或更少（见下一个项目符号）。
• 将用于将内核排入队列的 local_size 减少到 8 会使奇怪的行为消失（在上面的示例中很简单，因为不再有线程号 8，但是在更复杂的完整版本中，我得到了预期的结果）。
• 使用 this one 等在线工具没有发现任何问题。

问题：

1. 您有什么想法可以解释我正在经历的行为吗？你能重现这种行为吗？
2. 私有内存使用安全吗？或者 sum 数组声明在某种程度上很重要？
3. 如果我没有做错任何事，这种行为是否是 (GPU) 编译器错误！？

解决方法

在这个问题上做更多的工作后，我可能对正在发生的事情有一种解释，尽管这不是一个真正的证据，而是更多的事实。

起点是让内核按预期运行的奇怪行为，要么减少它运行时使用的 local_size，要么扩大私有内存阵列 sum 的大小。 OpenCL 编译器如何处理私有内存？阅读AMD OpenCL User Guide，我在第 25-26 页找到了以下句子（重点是我的）。

私有内存中的数据首先放在寄存器中。如果使用的私有内存多于可以放置在寄存器中的数量，或者在私有数组上使用动态索引，则溢出数据将放置（溢出）到暂存内存中。暂存内存是全局内存的私有子集，因此如果发生溢出，性能会显着降低。

奇怪的是，完整指南中没有定义动态索引的含义。然而，我发现 an interesting article 由 nVIDIA 人提供，这几乎解释了这个想法。 动态索引发生在编译器无法将数组索引解析为常量时。在这种情况下，它将被迫不再使用寄存器。

在当前内核中，通过在设置 alpha_idx 中使用 sum 来强制动态索引。这在某种程度上是一个突破点，虽然在那里放置一个 sum[0] 会使内核按预期工作，但它也会从动态索引变为静态索引！

为了进一步检查这个猜测，我尝试实现内核在本地内存中存储私有 sum。这可以通过本地内存中的缓冲区来实现，每个工作项仅访问其中的给定部分。有趣的是，这样做内核按预期工作，并且不再出现奇怪的行为。

在这一点上，为什么通过将 local_size 减少到 8 或更小或将 MAX_ORDER 的值设置为大于 21 来修复代码仍不清楚。在 this interesting SO answer 中作者写道

如果从本地转到私有效果不好，您应该将本地线程组大小从 256 减少到 64，例如。因此每个线程有更多的私有寄存器可用。

这基本上是说减少 local_size 使编译器能够以不同的方式处理私有内存中的变量。

总而言之，我感觉所用 GPU（Intel(R) Gen9 HD Graphics NEO，实际上不是真正的 GPGPU）上的编译器正在做一些奇怪的事情来处理内核中的动态索引。究竟是什么，我真的不知道，但在我看来，这确实是一个编译器错误。一方面，较小的 local_size 值隐藏了问题，因为 - 只是猜测 - 每个工作组有更多内存，编译器错误不会被击中。另一方面，增加 MAX_ORDER 也使问题消失，因为 - 再次猜测 - 编译器做了一些不同的事情（例如，使用不同的策略，内核需要更多内存），并且错误没有被击中。

为了进一步支持所有这一系列的想法，我在真实的 GPGPU 上测试了原始代码。在 AMD Radeon Instinct MI50 上运行代码没有发生任何奇怪的事情。