1

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

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

任何帮助表示赞赏。谢谢

4

1 回答 1

11

我编写了该代码并共同撰写了这篇文章，我要求您仅将本文用于学习扫描算法，不要使用其中的代码。它是在 CUDA 新的时候写的，我是 CUDA 的新手。如果您在 CUDA 中使用现代扫描实现，则不需要任何银行冲突避免。

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

如果您真的想实施扫描，请参阅最近的文章，例如这篇[1]。或者对于具有更快代码但需要更多研究的真正作品，这个[2]。或者阅读Sean Baxter 的教程（尽管后者不包括对扫描算法开创性工作的引用）。

[1] Shubhabrata Sengupta、马克·哈里斯、迈克尔·加兰和约翰·D·欧文斯。“多核 GPU 的高效并行扫描算法”。载于 Jakub Kurzak、David A. Bader 和 Jack Dongarra，编辑，Scientific Computing with Multicore and Accelerators，Chapman & Hall/CRC 计算科学，第 19 章，第 413-442 页。泰勒和弗朗西斯，2011 年 1 月。

[2] Merrill, D. 和 Grimshaw, A. 并行扫描流架构。技术报告 CS2009-14，弗吉尼亚大学计算机科学系。2009 年 12 月。

于 2012-03-14T01:25:34.483 回答