cuda 线程栅栏

Posted 2023-03-23

【中文标题】cuda 线程栅栏

【英文标题】：cuda threadfence

【发布时间】：2012-07-19 05:23:52

【问题描述】：

我正在编写一个必须执行块间同步（N 维和其他内存传输操作的总和）的代码。当我增加问题的维度时，结果是错误的。

我正在与__threadfence() 和第一个维度（N__threadfence()，则结果对于更多维度是正确的。

一个threadfence() 不够同步？另外，数据结果在同一个block中使用。

在编程指南中的信息表明threadfence 等待所有内存空间都准备好（共享和全局）

【问题讨论】：

欢迎来到 Stack Overflow。我认为您的问题中没有足够的信息让某人能够给出一个好的答复。您是否检查过 CUDA C 编程指南中关于 __threadfence() 的部分，看看它是否按您期望的方式工作？

请为您的问题提供示例代码。

【参考方案1】：

没有很好的方法来执行块之间的同步。您可以采用一种 hacky 方法来等待自旋并占用您的 GPU 内存带宽，也可以终止您的内核并启动一个新内核。

__threadfence() 不是用于块之间的同步。 __threadfence() 用于暂停当前线程，直到其他线程可以看到所有先前对共享和全局内存的写入。它不会停止也不会影响其他线程的位置！

您可以检查以下问题：

cuda block synchronization CUDA __threadfence()

【讨论】：

感谢您的回答。但在我的问题中，只有一个块进程生成全局结果，在这一点上不了解函数 threadfence。如果结果仅受一个块影响，则无需等待结果。这种方法的原因是产生最好的性能，因为内核的调用是迭代的，应该分配动态内存。

@CygnusX1：假设内核执行中的所有线程都在对全局内存执行一些写操作，并且我们不想继续，直到所有线程都可以看到所有写操作。写后所有线程都不会调用__threadfence() 来确保同步吗？

没有。 __threadfence() 只保证 current 线程不会继续，直到当前线程对全局内存的所有先前更改都被其他线程可见。然而，其他线程可能会尝试在此之前从相同的内存单元读取，读取旧数据并继续执行超出__threadfence()。

【参考方案2】：

合作组将允许同一内核中不同块之间的同步。现在也很容易使用。

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

// and then in your code
cg::grid_group grid = cg::this_grid();
grid.sync(); // All threads in all blocks must run this line

整个内核中每个块中的每个线程都必须运行grid.sync()。 只有在所有线程都运行完下一行之后，才会继续执行下一行。

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cooperative-groups