我正在开发一个 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);
}
原子操作导致分支发散的说法正确吗?
最佳答案
在某种程度上,CUDA 中的原子实现可能会因 GPU 架构而异。但特别是对于 GTX 480(费米级 GPU),__shared__
内存原子不是作为单个机器指令实现的,而是实际上由 a sequence of machine (SASS) instructions that form a loop 实现。 .
这个循环本质上是在争夺锁。当特定线程获取锁时,该线程将在识别的共享内存单元上原子地完成请求的内存操作,然后释放锁。
循环获取锁的过程必然涉及到分支发散。在这种情况下,分支发散的可能性在 C/C++ 源代码中并不明显,但如果您检查 SASS 代码,就会很明显。
全局原子通常作为单个(ATOM
或 RED
)SASS 指令实现。然而,如果由扭曲中的多个线程执行,全局原子仍然可能涉及访问的序列化。我通常不会认为这是“分歧”的情况,但我不完全确定分析器将如何报告它。如果你进行一个仅涉及全局原子的实验,我认为它会变得很清楚。
您的案例中报告的差异可能完全是由于共享内存差异(这是预期的)造成的,如上所述。
关于CUDA原子导致分支发散,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/35156157/