我的 OpenCL 代码根据看似 noop 更改输出

Posted 2023-04-15

【中文标题】我的 OpenCL 代码根据看似 noop 更改输出

【英文标题】：My OpenCL code changes the output based on a seemingly noop

【发布时间】：2013-11-15 00:35:10

【问题描述】：

我在 Intel CPU 和 NVIDIA GPU 上运行相同的 OpenCL 内核代码，结果第一个错误但后者正确；奇怪的是，如果我做了一些看似不相关的更改，那么两种情况下的输出都会按预期工作。

函数的目标是计算A（三角形）和B（正则）之间的矩阵相乘，其中A在运算中的位置由变量left的值决定。只有当left 为真并且for 循环至少迭代两次时才会出现该错误。

这里是一段代码，为了清楚起见，省略了一些不应该影响的部分。

__kernel void blas_strmm(int left, int upper, int nota, int unit, int row, int dim, int m, int n,
                         float alpha, __global const float *a, __global const float *b, __global float *c) 

  /* [...] */
  int ty = get_local_id(1);
  int y = ty + BLOCK_SIZE * get_group_id(1);
  int by = y;
  __local float Bs[BLOCK_SIZE][BLOCK_SIZE];
  /* [...] */

  for(int i=start; i<end; i+=BLOCK_SIZE) 
    if(left) 
      ay = i+ty;
      bx = i+tx;
       
    else 
      ax = i+tx;
      by = i+ty;
       

    barrier(CLK_LOCAL_MEM_FENCE);
    /* [...] (Load As) */
    if(bx >= m || by >= n)
      Bs[tx][ty] = 0;
    else
      Bs[tx][ty] = b[bx*n+by];
    barrier(CLK_LOCAL_MEM_FENCE);

    /* [...] (Calculate Csub) */
  

  if(y < n && x < (left ? row : m)) // In bounds
    c[x*n+y] = alpha*Csub;

现在变得奇怪了。

如您所见，如果left 为真，by 始终等于y。我检查了（有一些printfs，请注意）并且left 总是正确的，并且循环内else 分支上的代码永远不会执行。不过，如果我删除或注释掉那里的by = i+ty 行，代码就可以工作。为什么？我还不知道，但我认为这可能与 by 没有分配预期值有关。

我的思路是检查by 和y 之间是否存在差异，因为它们应该始终具有相同的值；我添加了一行检查by != y 是否像预期的那样比较总是返回假。所以我继续将by 的外观更改为y 所以这条线

if(bx >= m || by >= n)

变成了

if(bx >= m || y >= n)

它再次起作用，即使我仍在正确使用变量 by 下面三行。

怀着开放的心态，我尝试了其他一些事情，我发现如果我在循环中添加以下行，代码就可以工作，只要它位于初始 if/else 之后和之前的任何点我刚才提到的 if 条件。

if(y >= n) left = 1;

里面的代码 (left = 1) 可以替换任何东西（printf、另一个无用 赋值等），但条件限制性更强。以下是一些使代码输出正确值的示例：

if(y >= n) left = 1;
if(y < n) left = 1;
if(y+1 < n+1) left = 1;
if(n > y) left = 1;

还有一些不起作用，请注意我正在测试的特定示例中的m = n：

if(y >= n+1) left = 1;
if(y > n) left = 1;
if(y >= m) left = 1;
/* etc. */

这就是我现在所处的位置。我添加了一条根本不应该影响程序但它使它工作的行。这个神奇的解决方案让我不满意，我想知道我的 CPU 内部发生了什么以及为什么。

为了确保我没有忘记任何事情，这里是 full function code 和 gist with example inputs and outputs。

非常感谢。

---

解决方案

用户 DarkZeros 和Sharpneli 的假设都是正确的：for 循环内的障碍没有被正确地击中。特别是，存在一个涉及每个本地组的第一个元素的错误，使其运行的迭代次数少于其余部分，从而引发了未定义的行为。事后看来，这一点非常明显。

感谢大家的回答和时间。

【问题讨论】：

CPU 版本您使用的是哪个 OpenCL 工具链？您是否尝试过查看发出的编译器输出？这听起来像是一个微妙的编译器错误或需要非常仔细的分析才能理解的东西

@talonmies 我在 Linux (1.2-3.0.67279) 上使用英特尔 OpenCL 1.2 SDK。我保存了编译后的二进制文件，但我不太确定如何对它们进行分析，我将寻找可以帮助我解决此问题的工具。

您可以使用clGetProgramInfo 检索JIT 编译发出的代码。请查看我的回答 here，了解如何转储编译器输出的示例。尝试比较好/坏情况的输出，看看有什么不同。我不确定您将从英特尔 SDK 中得到什么，它可能是 LLVM 类型的 SSA 或实际的 EMT64 汇编器，但无论如何，它应该很容易理解

【参考方案1】：

您是否检查过 get_local_size 始终返回正确的值？

您说“简而言之，矩阵的全长被划分为 BLOCK_SIZE 的局部块并并行运行；”。请记住，OpenCL 只允许工作组内的任何并发。因此，如果您使用全局大小为 [32,32] 和本地大小为 [16,16] 调用 enqueueNDrange，则第一个线程块可能从头到尾运行，然后是第二个，然后是第三个等等。您无法在工作组。

您的 EnqueueNDRange 调用是什么？获取示例输出所需的调用示例将不胜感激（主要对全局和本地大小参数感兴趣）。

（我会在评论中问这个问题，但我是新用户）。

E（有答案，经验证没有，还需要更多信息）： http://multicore.doc.ic.ac.uk/tools/GPUVerify/

通过使用它，我收到了一个投诉，即非统一控制流可能会到达障碍。

这完全取决于 dim、nota 和 upper 得到的值。可以举一些例子吗？

我做了一些测试。假设 left = 1. nota != upper 和 dim = 32, row as 16 or 32 or whatnot, 仍然有效并得到以下结果：

...
gid0: 2 gid1: 0 lid0: 14 lid1: 13 start:  0  end: 32
gid0: 2 gid1: 0 lid0: 14 lid1: 14 start:  0  end: 32
gid0: 2 gid1: 0 lid0: 14 lid1: 15 start:  0  end: 32
gid0: 2 gid1: 0 lid0: 15 lid1:  0 start:  0  end: 48
gid0: 2 gid1: 0 lid0: 15 lid1:  1 start:  0  end: 48
gid0: 2 gid1: 0 lid0: 15 lid1:  2 start:  0  end: 48
...

因此，如果我对变量值的假设甚至接近正确，那么您就存在障碍分歧问题。一些线程遇到另一个线程永远不会遇到的障碍。我很惊讶它没有陷入僵局。

【讨论】：

get_local_size() 似乎按预期工作；它为 16 人的本地工作组返回 0 到 15 之间的数字。我的代码中的确切调用是 clEnqueueNDRangeKernel(command_queues[i], kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL)。在这种特殊情况下，只有一个命令队列，global_work_size 是 [32,32]，local_work_size 等于 [16,16]。我只尝试在同一组的元素之间进行同步，以尽量减少本地 Bs 和 As 数组上的内容。

如果将 local_work_size 设置为 [16,16]，

get_local_size 应始终返回 16。应该没有例外。也许你把它误认为是 get_local_id？

没错，我正在检查get_local_id。 get_local_size 为两个维度返回 16。

编辑了主帖。我现在有一个更大的答案。

你明白了！我会在几分钟内用答案更新主帖。非常感谢。

【参考方案2】：

我看到的第一件事可能会非常失败，那就是您在 for 循环中使用了屏障。

如果所有线程没有进入相同数量的 for 循环。然后结果是完全未定义的。并且您明确指出，只有在 for 循环多次运行时才会出现问题。

你保证这个条件吗？

【讨论】：

是的，确保每个线程执行循环的次数完全相同。简而言之，将矩阵的全长划分为BLOCK_SIZE的局部块，并行运行；我什至创建了额外的本地线程，所以总是有BLOCK_SIZE 元素同时迭代。