CONEGICT_FREE_OFFSET宏用于GPU Gems 3的并行前缀算法

Question

首先,这里是算法的链接http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html

为了避免存储体冲突,每个NUM_BANKS(即,对于可计算性2.x的设备为32)元素将填充添加到共享存储器阵列.这是通过(如图39-5)完成的:

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

我不明白ai/NUM_BANKS是如何与宏等效的:

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))  

不等于

n >> LOG_NUM_BANKS

任何帮助表示赞赏.谢谢

Answer 1

我编写了该代码并共同撰写了该文章,并且我要求您仅使用本文来了解扫描算法,并且不要使用其中的代码.它是在CUDA是新的时候写的,我是CUDA的新手.如果您在CUDA中使用现代扫描实现,则不需要任何银行冲突避免.

如果您想以简单的方式进行扫描,请使用thrust::inclusive_scan或thrust::exclusive_scan.

如果您真的想要实现扫描,请参阅最近的文章,例如[1].或者对于具有更快代码的实际操作,但这需要更多的研究,这一个[2].或者阅读Sean Baxter的教程(尽管后者不包括对扫描算法的开创性工作的引用).

[1] Shubhabrata Sengupta,Mark Harris,Michael Garland和John D. Owens."用于多核GPU的高效并行扫描算法".在Jakub Kurzak,David A. Bader和Jack Dongarra,编辑,科学计算与多核和加速器,Chapman和Hall/CRC计算科学,第19章,第413-442页.Taylor&Francis,2011年1月.http://www.idav.ucdavis.edu/publications/print_pub？ pub_id = 1041

[2] Merrill,D.和Grimshaw,A.Parallel Scanures for Stream Architectures.技术报告CS2009-14,弗吉尼亚大学计算机科学系.2009年12月