在概括将二维数组的值向右移动一个空间（环绕行边界）的内核时，我遇到了扭曲同步问题。完整的代码附在下面。
该代码适用于任意数组宽度、数组高度、线程块数和每个块的线程数。选择33的线程大小（即比完整扭曲更长的线程），第33个线程不会__syncthreads()调用。这会导致输出数据出现问题。该问题仅在存在多个经纱且阵列的宽度大于线程数（例如，宽度 = 35 和 34 线程）时出现。
下面是一个缩小的例子，说明发生了什么（实际上，数组需要有更多的元素才能让内核产生错误）。
初始数组：
0 1 2 3 4
5 6 7 8 9
预期结果：
4 0 1 2 3
9 5 6 7 8
内核产生：
4 0 1 2 3
8 5 6 7 8
第一行正确完成（如果有多个块，则为每个块），所有后续行都具有重复的倒数第二个值。我已经测试了这两种不同的卡（8600GT 和 GTX280）并得到相同的结果。我想知道这是否只是我的内核的错误，还是无法通过调整我的代码来修复的问题？
完整的源文件包含在下面。
谢谢你。
#include <cstdio>
#include <cstdlib>
// A method to ensure all reads use the same logical layout.
inline __device__ __host__ int loc(int x, int y, int width)
{
return y*width + x;
} …