CUDA原子导致分支发散

Question

我正在开发一个 CUDA 内核来计算图像的直方图 (NVIDIA GTX 480)。我注意到使用 cuda profiler 发现了 82.2% 的分支分歧。探查器指示以下函数作为差异源，位于名为 device_functions.h 的文件中(特别是包含 return 语句的行)。

static __forceinline__
unsigned int __uAtomicAdd(unsigned int *p, unsigned int val)
{
  return __nvvm_atom_add_gen_i((volatile int *)p, (int)val);
}

原子操作导致分支发散的说法正确吗？

Answer 1

在某种程度上，CUDA 中的原子实现可能会因 GPU 架构而异。但特别是对于 GTX 480(费米级 GPU)，__shared__ 内存原子不是作为单个机器指令实现的，而是实际上由 a sequence of machine (SASS) instructions that form a loop 实现。 .

这个循环本质上是在争夺锁。当特定线程获取锁时，该线程将在识别的共享内存单元上原子地完成请求的内存操作，然后释放锁。

循环获取锁的过程必然涉及到分支发散。在这种情况下，分支发散的可能性在 C/C++ 源代码中并不明显，但如果您检查 SASS 代码，就会很明显。

全局原子通常作为单个(ATOM 或 RED)SASS 指令实现。然而，如果由扭曲中的多个线程执行，全局原子仍然可能涉及访问的序列化。我通常不会认为这是“分歧”的情况，但我不完全确定分析器将如何报告它。如果你进行一个仅涉及全局原子的实验，我认为它会变得很清楚。

您的案例中报告的差异可能完全是由于共享内存差异(这是预期的)造成的，如上所述。