2

我最近问了一个关于 CUDA 中一个块的线程之间的同步问题的问题，在这里：提前退出线程会破坏块中 CUDA 线程之间的同步吗？我的问题的其中一条评论提供了一个类似线程的链接，该线程引用了 PTX 指南中关于 CUDA 屏障 (__syncthreads()) 指令的以下内容：

屏障是在每个 warp 的基础上执行的，就好像一个 warp 中的所有线程都处于活动状态一样。因此，如果 warp 中的任何线程执行 bar 指令，就好像 warp 中的所有线程都执行了 bar 指令。warp 中的所有线程都将停止，直到屏障完成，并且屏障的到达计数按 warp 大小（而不是 warp 中的活动线程数）递增。在有条件执行的代码中，只有当已知所有线程都以相同的方式评估条件时才应使用 bar 指令（warp 不会发散）。由于屏障是在每个 warp 的基础上执行的，因此可选的线程数必须是 warp 大小的倍数。

我仍然对这句话中解释的机制有点困惑。它说如果我们在条件代码中使用屏障，并且如果某些线程通过在条件代码处采用不同的路径而无法到达屏障指令，那么它可能会导致未定义的行为甚至死锁。问题是，我不明白这种机制如何导致死锁。（即使线程数不是 warp 大小的倍数，也是危险的。）该文档说，即使单个线程执行 bar 指令，它也会被视为 warp 中的所有线程都执行了 warp 指令并且到达计数器由经纱中的线程数更新。可能 CUDA 架构通过检查这个到达计数器来确定是否所有线程都已被同步；通过将其与块中的实际线程数进行比较。如果它是基于每个线程更新的，那么这可能会导致死锁，因为计数器永远不会达到最大值。编号。线程，因为其中一些采用不包含 bar 指令的条件路径。但是在这里，这个数字是用经纱中的线程数更新的。所以，我并不完全理解这里的底层机制。

我的另一个问题是关于整体条件语句。我知道warp中的所有线程在给定时间执行相同的指令，在if子句的情况下，采用if和else分支的线程通过保持空闲并在条件结束时再次同步来相互等待。所以对于这样的条件码有一个隐式的同步机制。现在，这将如何在如下代码中工作：

int foundCount=0;
for(int i=0;i<array1_length;i++)
{
    for(j=0;j<array0_length;j++)
    {
        if(i == array0[j])
        {
            array1[i] = array1[i] + 1;
            foundCount++;
            break;
        }
    }

     if(foundCount == foundLimit)
        break;
}

这是我当前任务中的一段代码；对于array1 的每个成员，我需要检查当前array1 索引是否包含在array0 中。如果是这样，我会增加 array1 中当前索引的元素，并且由于它已经包含在 array0 中，所以我使用 break 语句退出内部循环。如果 array1 包含的索引总数达到限制，我们不需要继续外循环，我们也可以退出它。这对于 CPU 代码来说很简单，但我想知道 CUDA 的扭曲机制如何处理这种嵌套的条件情况。想象一下，有 32 个线程正在处理这段代码，其中一些可以处理内部循环，而其中一些已经退出它，其中一些甚至可以退出外部循环。在这种情况下，架构如何组织线程的工作，它是否包含线程当前“等待点”的列表？在如此复杂的情况下，它如何确保同一 warp 中的线程确实处理同一行代码？

4

1 回答 1

3

条件分支是通过让 warp 中的所有线程执行所有分支来实现的。那些不跟随分支的线程执行相当于一个空操作。这通常被称为屏蔽执行，这也是部分扭曲可以适应的方式：部分扭曲包含永久屏蔽的线程。还有直接条件执行指令可用于实现诸如三元运算符之类的东西而无需分支。

这些机制不适用于标准barPTX 指令。正如您所注意到的，这是使用简单的计数器递减方案实现的，如果块中的所有线程都不将计数器递减为零，则会导致死锁。

于 2012-07-27T12:41:41.503 回答