cuda线栅栏

Question

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

我正在与__threadfence()和第一个维度(N<192)进行同步，这是可以的，但是如果我在这段代码中插入其他__threadfence()，结果对于更多维度是正确的。

一个threadfence()不足以进行同步吗？此外，数据结果在相同的块中使用。

在编程指南中，信息表明threadfence等待所有内存空间都已就绪(共享和全局)

Answer 1

没有很好的方法来执行块之间的同步。你可以使用一种老套的方法--旋转等待和消耗你的GPU内存带宽，或者你可以终止你的内核并开始一个新的内核。

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

您可以检查以下问题：

• cuda block synchronization
• CUDA __threadfence()

Answer 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