如何暂停所有 CUDA 线程，直到线程 0 完成一些代码？

Question

我在下面有一个简单的 CUDA 代码。我的问题是如何在线程 0 初始化数组数据时使所有线程暂停。所以在那之后，所有线程都可以访问数据中的元素。

__device__ int *data;

__global__ void test() {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx == 0) {
        data = (int *) malloc(10 * sizeof (int));
    }

    data[idx] = idx;
    printf("%d", data[idx]);
}

int main() {
    test << <1, 10 >> >();
    return 0;
}

Answer 1

CUDA 提供的唯一方法是通过cooperative groups grid synchronization 来执行此操作（在一般情况下，内核启动有超过 1 个块）。 This answer 提供了一个使用协作组网格范围同步的示例。

由于使用协作组非常复杂，您可能需要考虑在启动内核之前仅在主机代码中初始化此指针。 This answer 展示了如何做到这一点。

在您的特定情况下，您只启动一个块，您可以使用__syncthreads()：

if (idx == 0) {
    data = (int *) malloc(10 * sizeof (int));
}
__syncthreads();
data[idx] = idx;
printf("%d", data[idx]);