CUDA 动态并行和全局内存同步答案

【问题标题】：CUDA dynamic parallelism and global memory synchronizationCUDA 动态并行和全局内存同步
【发布时间】：2013-11-13 11:48:24
【问题描述】：

我无法弄清楚以下内容。

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

我的猜测是没有。事实上，在 CUDA C 编程指南的 synchronization functions 部分中，声明：

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

但是，当谈到动态并行中的 global memory consistency 时，CUDA C 编程指南指出：

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

那么__syncthreads() 是否也会在涉及动态并行时跨块提供更改？

谢谢

【问题讨论】：

关于您的第一个问题，您的结论是正确的，因为 CUDA 不允许以稳健的方式启用跨块同步。 NVIDIA 论坛上有一个名为 Synchronize all blocks in CUDA 的讨论，您可能感兴趣。
关于内存一致性，参考你说的具体例子，我认为指南只是说你需要第一个__synchthreads()来保证父子内核的全局内存一致，因为所有设备端内核启动都是异步的（这在 CUDA Dynamic Parallelism Programming Guide 中提到）。需要第二个__synchthreads() 调用来确保同一线程块内的内核之间的全局内存一致性，因为启动的子内核可能需要不同的处理时间才能完成。
查看Memory Fence Functions 中的__threadfence_system()，了解一种使一个线程编写的内容在整个设备上可见的方法。

标签： cuda

【解决方案1】：

__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);
}

这里，parent_launch 内核的所有256 线程都在data 中写入了一些内容。之后，线程0 调用child_launch。第一个__syncthreads() 需要确保在子内核调用之前所有内存写入都已完成。在这一点上引用指南：

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

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

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

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

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

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

请注意，在具体示例中，host_launch 函数只启动了一个线程块。这可能会误导您。

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

Synchronize all blocks in CUDA

【讨论】：