Ste*_*onà 2 profiler cuda atomic nvidia
我正在开发一个 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);
}
Run Code Online (Sandbox Code Playgroud)
原子操作导致分支发散的说法正确吗?
在某种程度上,CUDA 中的原子实现可能会因 GPU 架构而异。但特别是对于 GTX 480(费米级 GPU)来说,__shared__
内存原子并不是作为单个机器指令来实现的,而是实际上是通过形成循环的一系列机器(SASS)指令来实现的。
这个循环本质上是在争夺锁。当特定线程获取锁时,该线程将在识别的共享内存单元上原子地完成请求的内存操作,然后释放锁。
循环获取锁的过程必然涉及到分支发散。在这种情况下,分支发散的可能性在 C/C++ 源代码中并不明显,但如果检查 SASS 代码,就会很明显。
全局原子通常作为单个 (ATOM
或RED
) SASS 指令来实现。然而,如果由扭曲中的多个线程执行,全局原子仍然可能涉及访问的序列化。我通常不会认为这是“分歧”的情况,但我不完全确定分析器将如何报告它。如果你进行一个仅涉及全局原子的实验,我认为它会变得很清楚。
您的情况中报告的差异可能完全是由于如上所述的共享内存差异(这是预期的)造成的。
归档时间: |
|
查看次数: |
371 次 |
最近记录: |