我在 Intel CPU 和 NVIDIA GPU 上运行相同的 OpenCL 内核代码，结果在第一个上是错误的，但在后者上是正确的；奇怪的是，如果我做了一些看似无关的更改，那么两种情况下的输出都会按预期工作。
该函数的目标是计算 A（三角形）和 B（正则）之间的矩阵乘法，其中 A 在运算中的位置由变量 的值决定left。该错误仅在left为 true 并且 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;
}
现在变得很奇怪。
如您所见，如果为真，则by始终等于。我检查了（用一些s，请注意）并且总是正确的，并且循环内 else 分支上的代码永远不会执行。不过，如果我删除或注释掉那里的行，代码就可以工作。为什么？我还不知道，但我认为这可能与没有分配预期值有关。yleftprintfleftby = i+tyby
我的思路带我检查 和 之间是否存在差异by，y因为它们应该始终具有相同的值；by != y正如预期的那样，我添加了一条检查 if但该比较始终返回 false 的行。所以我继续改变了byfor yso 的外观
if(bx >= m || by >= n)
转化成
if(bx >= m || y >= n)
它再次起作用，即使我仍然在by下面的三行中正确使用变量。
以开放的心态，我尝试了一些其他的事情，我发现如果我在循环中添加以下行，代码就可以工作，只要它位于初始 if/else 之后和 if 条件之前的任何点我刚才提到了。
if(y >= n) left = 1;
( ) 中的代码left = 1可以替换任何东西（a 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 内部发生了什么以及为什么。
只是为了确保我没有忘记任何东西，这里是完整的功能代码和带有示例输入和输出的要点。
非常感谢。
解决方案
用户 DarkZeros 和Sharpneli 都对他们的假设是正确的：for 循环内的障碍没有被击中正确的次数。特别是，存在一个涉及每个本地组的第一个元素的错误，使其运行的迭代次数少于其余部分，从而引发了未定义的行为。事后看来，这很痛苦。
谢谢大家的回答和时间。