c - 我的 OpenCL 代码基于一个看似 noop 更改输出

Question

我在 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。我检查了(使用一些 printf，请注意)并且 left 始终为真，并且永远不会执行循环内 else 分支上的代码。尽管如此，如果我删除或注释掉那里的 by = i+ty 行，代码就可以工作。为什么？我还不知道，但我认为这可能与 by 没有分配预期值有关。

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

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 循环内的障碍没有被击中正确的次数。特别是，有一个错误涉及每个本地组的第一个元素，这使得它运行的迭代比其余的少，从而引发未定义的行为。事后看来，这是显而易见的痛苦。

感谢大家的回答和时间。

Answer 1

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

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

您的 EnqueueNDRange 调用是什么？非常感谢获得示例输出所需的调用示例(主要对全局和局部大小参数感兴趣)。

(我会在评论中提出这个问题，但我是新用户)。

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

通过使用它，我得到了一个提示，即不一致的控制流可能会达到障碍。

这完全取决于 dim、nota 和 upper get 的值。你能提供一些例子吗？

我做了一些测试。假设 left = 1.nota != upper 和 dim = 32，行为 16 或 32 或诸如此类，仍然有效并得到以下结果:

...
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
...

因此，如果我对变量值的假设甚至接近于正确，那么你就会遇到障碍分歧问题。一些线程遇到了另一个线程永远不会遇到的障碍。我很惊讶它没有死锁。