CUDA 中经纱展开期间的线程同步

Question

我正试图了解 Mark Harris 的 reduction in CUDA 中的减少技术 #5。

Reduction #5 通过应用最后的 warp 展开改进了之前的 reduction #4。

幻灯片 21 提到：“我们不需要 __syncthreads()”，这是我不明白的部分。

下面是主要逻辑的代码：

__device__ void warpReduce(volatile int* sdata, int tid) {
  sdata[tid] += sdata[tid + 32]; // line A
  sdata[tid] += sdata[tid + 16]; // line B
  sdata[tid] += sdata[tid + 8];
  sdata[tid] += sdata[tid + 4];
  sdata[tid] += sdata[tid + 2];
  sdata[tid] += sdata[tid + 1];
}

// later...
for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
  if (tid < s)
    sdata[tid] += sdata[tid + s];
  __syncthreads();
}
if (tid < 32) warpReduce(sdata, tid);

我不明白为什么行 A 和行 B 之间没有__syncthreads()（以及下一行之间）。

我的问题：是否有可能在同一个warp 中，一个线程在另一个线程执行A 行之前执行B 行？ （好像是不可能的，有谁能确认一下并详细说明）

Answer 1

是否有可能在同一个线程中，一个线程在另一个线程执行 A 行之前执行 B 行？

在写这篇文章的时候（大约 10 年前），不可能发生这种情况，因为 Warp 保证会在锁步中执行。请注意，有问题的内存需要声明为volatile，以防止编译器优化在 Fermi 和更新的 GPU 的缩减步骤之间缓存结果。在不需要的原始 Tesla 架构上。

但是，执行 warp 级别操作的最先进方法已经改变，这种类型的设计模式在某些最新架构上可能不安全。相反，您应该更喜欢扭曲级别原语来减少而不是隐式扭曲同步。请参阅this blog post 了解更多信息。