时间戳

首页 > 解决方案 > 与 cuda 合作组的僵局

cuda - 与 cuda 合作组的僵局

问题描述

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

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

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

void __syncthreads();

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

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

this_grid().sync();

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

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

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

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

当我尝试使用coalesced_threads().sync(). 以下是正确的吗？

coalesced_threads().sync();

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

以下示例是否从 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; 
}

标签： cuda

解决方案

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

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()的语句，在代码的不同位置，不能满足单个同步屏障的要求。它们各自代表一个单独的障碍，必须适当满足其要求。

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

推荐阅读