c++ - cuda __syncthreads() 在我的代码中不起作用

Question

情况是这样的。

我有一个运行 while 循环的线程 block ，当且仅当这些线程中的任何一个满足某些条件时，我才需要循环继续。为此，我使用一个共享变量作为继续标志，该标志在每次迭代开始时由线程 #0 清除，后跟 __syncthreads()，并且可以由任何线程在迭代期间设置如果满足继续条件，则迭代。然后在下一次迭代的检查点之前再次调用 __syncthreads() 以确保线程同步。内核基本上是这样的:

__global__ void foo(void* data) {
    __shared__ int blockContinueFlag;
    do {
        if (threadIdx.x || threadIdx.y || threadIdx.z) {
            blockContinueFlag = 0;
        }
        __syncthreads(); //synch1
        //some data manipulations...
        if(some predicate) {
            blockContinueFlag = true;
        }
        //some data manipulations...
        __syncthreads(); //synch2
    } while (blockContinueFlag);
}

问题是屏障 synch2 似乎在我的代码中不起作用，有时即使某些线程满足继续条件，内核也会终止(我通过检查主机端返回的数据知道这一点)。为了进一步检查这一点，我在 do-while 循环之后设置了一个断点，如下面的代码，其中有时 blockContinueFlag 表示为 true(我只能假设 block 退出在某些线程可以设置 blockContinueFlag 之前循环。

__global__ void foo(void* data) {
    __shared__ int blockContinueFlag;
    do {
        if (threadIdx.x || threadIdx.y || threadIdx.z) {
            blockContinueFlag = 0;
        }
        __syncthreads(); //synch1
        //some data manipulations...
        if(some predicate) {
            blockContinueFlag = true;
        }
        //some data manipulations...
        __syncthreads(); //synch2
    } while (blockContinueFlag);
    //a break point is set here
}

我记得从 cuda 手册中读到，如果谓词对所有线程的评估都相同，则条件子句中允许使用 __syncthreads()，在这种情况下应该如此。

我有另一个简化版本的代码作为说明。

__global__ void foo(int* data, int kernelSize, int threshold) {
    __shared__ int blockContinueFlag;
    do {
        if (threadIdx.x == 0) {
            blockContinueFlag = 0;
        }
        __syncthreads();
        if (threadIdx.x < kernelSize)  {
            data[threadIdx.x]--;
            for (int i = 0; i < threadIdx.x; i++);
            if (data[threadIdx.x] > threshold)
                blockContinueFlag = true;
        }
        __syncthreads();
    } while (blockContinueFlag);
}

int main()
{
    int hostData[1024], *deviceData;
    for (int i = 0; i < 1024; i++)
        hostData[i] = i;
    cudaMalloc(&deviceData, 1024 * sizeof(int));
    cudaMemcpy(deviceData, hostData, 1024 * sizeof(int), cudaMemcpyHostToDevice);
    foo << <1, 1024 >> >(deviceData, 512, 0);
    cudaDeviceSynchronize();
    cudaMemcpy(hostData, deviceData, 1024 * sizeof(int), cudaMemcpyDeviceToHost);
    fprintf(stderr, cudaGetErrorString(cudaGetLastError()));
    return 0;

}

hostData[] 的预期值为 {-511, -510, -509, ..., 0, 512, 513, 514,..., 1023} 在 main() 的末尾，这有时是实际情况。但在某些情况下，它会在 VS 2013 Debug模式下产生以下值

[0]: -95
[1]: -94
...
[29]: -66
[30]: -65
[31]: -64
[32]: 31
[33]: 32
[34]: 33
...
[61]: 60
[62]: 61
[63]: 62
[64]: -31
[65]: -30
[66]: -29
...
[92]: -3
[93]: -2
[94]: -1
[95]: 0
[96]: 95
[97]: 96
[98]: 97
...

，这表明 warp 实际上并未同步。

那么有没有人知道这个的原因和/或是否有办法让线程屏障正常工作？

如有任何帮助，我们将不胜感激。提前致谢。

Answer 1

所以这是我的解决方案，使用一个 __syncthreads_or() 而不是请求的三个 __syncthreads()。

__global__ void foo(void* data) {
    int blockContinueFlag;
    do {
        blockContinueFlag = 0;
        //some data manipulations...
        if(some predicate) {
            blockContinueFlag = true;
        }
        //some data manipulations...
    } while (__syncthreads_or(blockContinueFlag));
}

在实践中，这比三个同步线程的速度稍快。

再次感谢您的帖子。