我有一个带有"while"循环的内核，它使用邻居的信息迭代更新数组的元素（下面的示例代码中只有一个邻居）。当当前迭代中没有元素更改时，此循环停止。
不幸的是，在某些情况下，部分线程会提前退出这个循环（比如忽略同步屏障）。有些输入每次都被正确处理，而其他输入（其中许多）每次都被错误处理（即不存在随机因素）。奇怪的是，这个错误只发生在发布版本，而调试版本总是工作良好。更准确地说，CUDA编译器选项"-G（生成GPU调试信息）"确定处理是正确的。大小为32x32或更小的数组始终得到正确处理。
这是一个示例代码：
__global__ void kernel(int *source, int size, unsigned char *result, unsigned char *alpha)
{
int x = threadIdx.x, y0 = threadIdx.y * 4;
int i, y;
__shared__ bool alpha_changed;
// Zero intermediate array using margins for safe access to neighbors
const int stride = MAX_SIZE + 2;
for (i = threadIdx.x + threadIdx.y * blockDim.x; i < stride * (stride + 3); i += blockDim.x * blockDim.y)
{
alpha[i] = 0;
}
__syncthreads();
for (int bit = MAX_BITS - 1; bit >= 0; bit--)
{
__syncthreads();
// Fill intermediate array with bit values from input array
alpha_changed = true;
alpha[(x + 1) + (y0 + 1) * stride] = (source[x + (y0 + 0) * size] & (1 << bit)) != 0;
alpha[(x + 1) + (y0 + 2) * stride] = (source[x + (y0 + 1) * size] & (1 << bit)) != 0;
alpha[(x + 1) + (y0 + 3) * stride] = (source[x + (y0 + 2) * size] & (1 << bit)) != 0;
alpha[(x + 1) + (y0 + 4) * stride] = (source[x + (y0 + 3) * size] & (1 << bit)) != 0;
__syncthreads();
// The loop in question
while (alpha_changed)
{
alpha_changed = false;
__syncthreads();
if (alpha[(x + 0) + (y0 + 1) * stride] != 0 && alpha[(x + 1) + (y0 + 1) * stride] == 0)
{
alpha_changed = true;
alpha[(x + 1) + (y0 + 1) * stride] = 1;
}
__syncthreads();
if (alpha[(x + 0) + (y0 + 2) * stride] != 0 && alpha[(x + 1) + (y0 + 2) * stride] == 0)
{
alpha_changed = true;
alpha[(x + 1) + (y0 + 2) * stride] = 1;
}
__syncthreads();
if (alpha[(x + 0) + (y0 + 3) * stride] != 0 && alpha[(x + 1) + (y0 + 3) * stride] == 0)
{
alpha_changed = true;
alpha[(x + 1) + (y0 + 3) * stride] = 1;
}
__syncthreads();
if (alpha[(x + 0) + (y0 + 4) * stride] != 0 && alpha[(x + 1) + (y0 + 4) * stride] == 0)
{
alpha_changed = true;
alpha[(x + 1) + (y0 + 4) * stride] = 1;
}
__syncthreads();
}
__syncthreads();
// Save result
result[x + (y0 + 0) * size + bit * size * size] = alpha[(x + 1) + (y0 + 1) * stride];
result[x + (y0 + 1) * size + bit * size * size] = alpha[(x + 1) + (y0 + 2) * stride];
result[x + (y0 + 2) * size + bit * size * size] = alpha[(x + 1) + (y0 + 3) * stride];
result[x + (y0 + 3) * size + bit * size * size] = alpha[(x + 1) + (y0 + 4) * stride];
__syncthreads();
}
}
// Run only 1 thread block, where size equals 64.
kernel <<< 1, dim3(size, size / 4) >>> (source_gpu, size, result_gpu, alpha_gpu);
这个示例内核的预期结果是数组，其中每一行只能包含连续的间隔的值。但取而代之的是，我得到了一些行，其中"0"one_answers"1"以某种方式交替出现。
此错误在我的移动GPU GeForce 740M（开普勒）、Windows 7 x64 SP1、CUDA 6.0或6.5上重现，使用Visual C++2012或2013。我还可以提供一个带有示例输入数组的示例VisualStudio项目（即处理不正确）。
我已经尝试过syncthreads（）、fences和"volatile"限定符的不同配置，但这个错误保留。
感谢您的帮助。
我认为问题出在您对alpha_changed的访问上。请记住，这只是一个块中所有线程的一个值。在一个扭曲重置此变量和另一个扭曲检查循环条件之间存在竞争条件：
// The loop in question
while (alpha_changed)
{
alpha_changed = false;
// ...
// alpha_changed may be set to true here
// ...
__syncthreads();
// race condition window here. Another warp may already execute
// the alpha_changed = false; line before this warp continues.
}
关键是在将共享变量设置为false之前先执行__syncthreads()。
您可以在循环中使用局部变量来判断该线程是否进行了任何更改。这避免了必须在所有地方使用__syncthreads()。然后在循环结束时进行缩减：
// The loop in question
while (alpha_changed)
{
bool alpha_changed_here = false;
// ...
// alpha_changed_here may be set to true here
// ...
__syncthreads();
alpha_changed = false;
__syncthreads();
// I think you can get away with a simple if-statement here
// instead of a proper reduction
if (alpha_changed_here) alpha_changed = true;
__syncthreads();
}
据我所知，这种在共享内存中只使用一个变量的方法目前是有效的。如果您想确定，请使用适当的归约算法。可以使用__any()将一条指令中的32个值减少一个扭曲。要使用的算法取决于块的大小（我不知道确切的行为是大小不是32的倍数）。