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

Question

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

减少#5 通过应用最后的经纱展开来改进先前的减少#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()行之间没有（以及下一行之间）。

我的问题：是否有可能在同一个线程中，一个线程在另一个线程执行A行之前执行B行？（似乎不可能，请任何人确认并详细说明）

标签： cudasynchronizationnvidia

Answer 1

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

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

然而，执行 warp 级别操作的最先进的方式已经改变，这种类型的设计模式在一些最新的架构上可能是不安全的。相反，您应该更喜欢扭曲级别原语来减少而不是隐式扭曲同步。有关更多信息，请参阅此博客文章。