Cuda 并行执行答案

【问题标题】：Cuda Parallel executionCuda 并行执行
【发布时间】：2012-08-02 19:33:42
【问题描述】：

有人可以告诉我：块在 CUDA 中是并行/并发执行的吗？换句话说，如果两个不同的块尝试写入同一个全局地址，即 globalPtr[12]，是否存在丢失更新问题？

（我在问这个问题，因为我读到 CUDA 中的并行执行单元是 warp=32 线程。）

【问题讨论】：

标签： cuda

【解决方案1】：

是的，多个块并行执行，因此如果多个线程需要访问同一地址，则对全局内存的访问需要是原子的。这适用于同一块中的两个线程或不同块中的两个线程。

【讨论】：

【解决方案2】：

是的，如果 CUDA 设备有多个 warp 调度程序，您可以在多个块之间并行执行。

具有计算能力 2.1 的 CUDA 设备有两个 warp 调度程序，因此来自两个不同 warp（来自同一个块或来自不同块，无关紧要）的指令可以同时执行。

具有计算能力 3.0 的 CUDA 设备具有四个 warp 调度程序，并且可以为每个准备执行的 warp 发出两条独立的指令。

请注意，即使 warp 之间没有并发执行，调度程序可以使用多个块是有利的，这样如果一个 warp 被阻塞以等待内存操作完成，调度程序可以切换到另一个 warp 执行，以便核心不要闲着。

可以驻留在内核上供调度程序切换到的 warp 数量因计算能力而异。

如果您只定义与调度程序一样多的块，您将无法充分发挥设备的计算潜力。如果您的代码具有大量内存 I/O，则尤其如此——“隐藏”内存延迟的一种方法是确保有足够的块/warp 可用，以便调度程序始终有一个准备好的 warp 切换到何时其中一个经线空闲等待内存 I/O。

当您有多个 warp 读取和写入相同的内存地址时，您应该使用原子 I/O 或锁定，无论您当前的硬件是否可以同时执行多个 warp。写后写工件（“丢失的更新”）甚至可以在任务切换的单核执行中表现出来。

【讨论】：

非常清楚，谢谢。这让我去阅读相关的 Cuda 文档以获得更好的理解。我正在编写一个使用嵌套循环来更新 32*32 共享内存的内核。我知道如果内核是用 32 个线程发布的，那么我可以保证在我的共享内存中没有竞争条件。但是我用 256 个线程发布它。除了使用诸如 if(threadIdx.x
好吧，你可以把它拉回到每个块只运行 32 个线程，但在更大的 CUDA 铁上，我怀疑这会比使用原子写入产生更多的性能问题。