首先,这里是算法的链接:
GPU Gems 3,39章:Parallel Prefix Sum (Scan) with CUDA。
为了避免存储体冲突,每个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
任何帮助都是非常感谢的。谢谢
发布于 2012-03-14 09:25:34
这段代码是我写的,也是这篇文章的合著者之一,我请求您使用这篇文章仅用于学习扫描算法,并且不要使用其中的代码。它是在CUDA刚出现的时候写的,我也是CUDA的新手。如果您在CUDA中使用scan的现代实现,则不需要任何银行冲突避免。
如果你想以一种简单的方式进行扫描,使用thrust::inclusive_scan
或thrust::exclusive_scan
。
如果您真的想实现扫描,请参考更多最近的文章,比如这篇[1]。或者对于具有更快代码的真实作品,但这将需要更多的研究,这一个[2]。或者阅读Sean Baxter's tutorial (尽管后者不包括对scan算法的开创性工作的引用)。
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月。
2 Merrill,D.和Grimshaw,A。流架构的并行扫描。技术报告CS2009-14,弗吉尼亚大学计算机科学系。2009年12月。
https://stackoverflow.com/questions/9689185
复制相似问题