首先，这是算法的链接:

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

感谢任何帮助。谢谢

最佳答案

我写了那段代码并合写了这篇文章，我要求你只将这篇文章用于学习扫描算法，不要使用其中的代码。是刚接触CUDA的时候写的，我刚接触CUDA。如果您在 CUDA 中使用现代扫描实现，则不需要任何银行冲突避免。

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

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

[1] Shubhabrata Sengupta、Mark Harris、Michael Garland 和 John D. Owens。 “多核 GPU 的高效并行扫描算法”。 Jakub Kurzak、David A. Bader 和 Jack Dongarra，编辑，Scientific Computing with Multicore and Accelerators，Chapman & Hall/CRC Computational Science，第 19 章，第 413-442 页。 Taylor & Francis，2011 年 1 月。

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

关于cuda - GPU Gems 3 的并行前缀算法中使用的 CONFLICT_FREE_OFFSET 宏，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/9689185/