2

我无法弄清楚以下内容。

例如，如果我启动内核并考虑0块中的线程0，在调用之后，所有其他块中的所有其他线程是否会看到块__syncthreads()中线程对全局内存所做的更改？00

我的猜测是否定的。实际上，在 CUDA C 编程指南的同步函数部分中，声明：

void __syncthreads(); 等待直到线程块中的所有线程都达到这一点，并且这些线程之前进行的所有全局和共享内存访问对块中的所有线程__syncthreads()都是可见的。

但是，当谈到动态并行中的全局内存一致性时，CUDA C Programming Guide 指出：

只有在第二次调用之后，这些修改才可用于父网格的其他线程。__syncthreads()

那么__syncthreads()当涉及动态并行时，是否也可以跨块进行更改？

谢谢

4

1 回答 1

4

唯一执行的操作__syncthreads()是您在 CUDA C 编程指南中描述的引用。除了在多个内核启动中划分内核执行的幼稚方法之外，CUDA 没有办法跨块同步，在性能方面存在所有缺点。因此，您的第一个问题的答案（您也猜到了）是否定的。

在您帖子的第二部分中，您指的是 CUDA C 编程指南的一个具体示例，即

__global__ void child_launch(int *data) {
    data[threadIdx.x] = data[threadIdx.x]+1;
}

__global__ void parent_launch(int *data) { 
    data[threadIdx.x] = threadIdx.x;

    __syncthreads();

    if (threadIdx.x == 0) {
        child_launch<<< 1, 256 >>>(data);
        cudaDeviceSynchronize();
    }

    __syncthreads();
}

void host_launch(int *data) {
    parent_launch<<< 1, 256 >>>(data);
}

在这里，内核的所有256线程都在. 之后，线程调用. 第一个是需要确保在子内核调用之前所有内存写入都已完成。在这一点上引用指南：parent_launchdata0child_launch__syncthreads()

由于第一次__syncthreads()调用，孩子会看到data[0]=0, data[1]=1, ..., data[255]=255（没有__syncthreads()调用，只能data[0]保证被孩子看到）。

关于第二个__syncthreads()，指南解释说

当子网格返回时，线程0保证可以看到其子网格中的线程所做的修改。__syncthreads()只有在第二次调用之后，这些修改才可用于父网格的其他线程。

在该特定示例中，第二个__syncthreads()是多余的，因为由于内核终止存在隐式同步，但是__syncthreads()当子内核启动后必须执行其他操作时，第二个变得需要。

最后，关于您在帖子中引用的句子：

__syncthreads()只有在第二次调用之后，这些修改才可用于父网格的其他线程

请注意，在具体示例中，该host_launch函数只启动了一个线程块。这或许有些误导了你。

在 NVIDIA 论坛上有一个有趣的讨论（可能不止一个）关于跨块的线程同步，标题为

同步 CUDA 中的所有块

于 2013-11-03T22:24:57.173 回答