CUDA原子导致分支发散

Question

我正在开发一个 CUDA 内核来计算图像的直方图 (NVIDIA GTX 480)。我注意到使用 cuda 分析器发现了 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（Fermi-class GPU），__shared__ 内存原子不是作为单个机器指令实现的，而实际上是由 a sequence of machine (SASS) instructions that form a loop.

实现的

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

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

全局原子通常作为单个（ATOM 或 RED）SASS 指令实现。但是，如果由 warp 中的多个线程执行，则全局原子仍可能涉及访问的序列化。我通常不会将此视为 "divergence" 的情况，但我不完全确定探查器将如何报告它。如果你运行一个只涉及全局原子的实验，我认为它会变得清晰。

您报告的差异可能完全是由于共享内存差异（这是预期的），如上所述。

CUDA原子导致分支发散

CUDA atomics causes branch divergence

profiler

cuda

atomic

nvidia