冲突_free_offset在GPU Gems 3的并行前缀算法中使用的宏

Question

首先，这里是算法的链接：

gpu gems 3，第39章：并行前缀总和（扫描）与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

.

任何帮助都得到了赞赏。谢谢

Answer 1

我写了该代码并共同编写了这篇文章，我请求您只使用该文章来学习扫描算法，并且不使用其中的代码。当CUDA是新的，我是新的，我是CUDA的新手。如果您在CUDA中使用现代实施扫描，则不需要任何银行冲突避免。

如果要做简单的方式，请使用 thrust::inclusive_scan 或 thrust::exclusive_scan 。

如果您真的想要实现扫描，请参阅更新的文章，例如此 [1] 。或者对于具有更快的代码的真实opus，但需要更多的研究，这是一个 [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.并行扫描用于流架构。弗吉尼亚大学计算机科学系CS2009-14技术报告。 2009年12月。