与 cuda 合作组的僵局

Question

在 CUDA 编程指南中关于合作组的部分中，有一个网格本地同步的示例：

grid_group grid = this_grid();
grid.sync();

不幸的是，我没有找到grid.sync() 行为的精确定义。将__syncthreads 的以下定义扩展到网格级别是否正确？

void __syncthreads();

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

所以，我的问题是正确的：

this_grid().sync();

等到 grid 中的所有线程都有 达到这一点，所有全局和共享内存访问由 this_grid().sync() 之前的这些线程对所有线程可见 网格。

我怀疑这是否正确，因为在 CUDA 编程指南中，grid.sync(); 下面的几行有以下语句：

为了保证线程块在 GPU 上的共同驻留，需要仔细考虑启动的块数。

这是否意味着如果我使用这么多线程以至于没有线程块的共同驻留，我最终可能会陷入线程可能死锁的情况？

当我尝试使用coalesced_threads().sync() 时，也会出现同样的问题。以下是正确的吗？

coalesced_threads().sync();

等到 warp 中的所有 活动 线程都有 达到这一点，所有全局和共享内存访问由 coalesced_threads().sync() 之前的这些线程对所有线程可见 活动线程列表。

以下示例是否从 while 循环中退出？

auto ct = coalesced_threads();
assert(ct.size() == 2);
b = 0; // shared between all threads
if (ct.thread_rank() == 0)
while (b == 0) {
    // what if only rank 0 thread is always taken due to thread divergence?
    ct.sync(); // does it guarantee that rank 0 will wait for rank 1?
}
if (ct.thread_rank() == 1)
while (b == 0) {
    // what if a thread with rank 1 never executed?
    b = 1; 
    ct.sync(); // does it guarantee that rank 0 will wait for rank 1?
}

为了清楚上面的例子，没有ct.sync()是不安全的，可以死锁（无限循环）：

auto ct = coalesced_threads();
assert(ct.size() == 2);
b = 0; // shared between all threads
if (ct.thread_rank() == 0)
while (b == 0) {
    // what if only rank 0 thread is always taken due to thread divergence?
}
if (ct.thread_rank() == 1)
while (b == 0) {
    // what if a thread with rank 1 never executed?
    b = 1; 
}

Answer 1

所以，我的问题是正确的：

this_grid().sync();

等到网格中的所有线程都达到这一点，并且这些线程在 this_grid().sync() 之前进行的所有全局和共享内存访问对网格中的所有线程都是可见的。

是的，这是正确的，假设您有适当的合作启动。适当的合作发布意味着许多事情：

1. 协同启动属性在您运行的 GPU 上为真
2. 您已使用正确形成的合作发射进行发射
3. 您已满足合作启动的网格大小要求
4. 合作上线后cudaGetLastError()返回cudaSuccess

是否意味着如果我使用了这么多线程，以至于没有线程块的共同驻留

如果您违反了合作启动的要求，您就是在探索未定义的行为。试图明确回答这些问题是没有意义的，只是说行为是未定义的。

关于您关于合并线程的陈述，它们是正确的，但必须仔细理解措辞。特定指令的活动线程与合并线程相同。

在您的示例中，您正在创建一个非法案例：

auto ct = coalesced_threads();
assert(ct.size() == 2); //there are exactly 2 threads in group ct
b = 0; // shared between all threads
if (ct.thread_rank() == 0) // this means that only thread whose rank is zero can participate in the next instruction - by definition you have excluded 1 thread
while (b == 0) {  
    // what if only rank 0 thread is always taken due to thread divergence?
    // it is illegal to request a synchronization of a group of threads when your conditional code prevents one or more threads in the group from participating
    ct.sync(); // does it guarantee that rank 0 will wait for rank 1?
}

两个不同的.sync() 语句，在代码的不同位置，不能满足单个同步屏障的要求。它们各自代表一个单独的障碍，必须正确满足其要求。

由于非法编码，这个例子也有未定义的行为；相同的 cmets 适用。