本文介绍了Cuda原子导致分支分歧的处理方法,对大家解决问题具有一定的参考价值,需要的朋友们下面随着小编来一起学习吧!
问题描述
我正在开发一个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);
}
说原子操作导致分支发散正确吗?
推荐答案
在某种程度上,CUDA中的原子实现可能会因GPU架构而异。但特别是对于GTX 480(费米级图形处理器),__shared__
内存原子不是作为单个机器指令实现的,而是实际上通过a sequence of machine (SASS) instructions that form a loop实现的。
此循环实质上是在争夺锁。当锁被特定线程获取时,该线程将在标识的共享内存单元上以原子方式完成请求的内存操作,然后释放锁。
循环获取锁的过程必然涉及分支发散。在这种情况下,分支分歧的可能性在C/C++源代码中并不明显,但如果您检查SASS代码,则会很明显。
全局原子通常实现为单个(RED)sass指令。然而,如果由WARP中的多个线程执行,全局原子仍然可能涉及访问的串行化。我通常不会认为这是"分歧"的情况,但我不完全确定分析器将如何报告它。如果你做了一个只涉及全球原子公司的实验,我想这一点就会变得很清楚。您的案例中报告的差异可能完全是由于如上所述的共享内存差异(这是预期的)。
这篇关于Cuda原子导致分支分歧的文章就介绍到这了,希望我们推荐的答案对大家有所帮助,也希望大家多多支持!