4、全局内存Bank冲突：虽然全局内存没有Bank，但理解其与共享内存的协同至关重要

说实话，我第一次听到「全局内存Bank冲突」这个说法时，也愣了一下。

全局内存哪来的Bank？它又不是共享内存。但后来我明白了——大家真正想聊的，是全局内存访问模式如何引发共享内存的Bank冲突。这两者之间的协同关系，才是性能优化的关键。

4.1 先搞清楚：全局内存到底有没有Bank？

没有。全局内存的硬件结构里，不存在Bank这个概念。

Bank是共享内存和寄存器文件的专属。全局内存的访问单元是内存控制器和缓存行。一个全局内存请求，会被拆成128字节的缓存行来服务。

那为什么大家总把「全局内存」和「Bank冲突」放在一起说？

因为你的数据从全局内存加载到共享内存后，存储方式决定了后续的Bank冲突情况。我见过太多人，全局内存加载写得飞快，结果数据一进共享内存，Bank冲突直接让性能腰斩。

4.2 一个让我印象深刻的案例

几年前我优化过一个矩阵转置的kernel。全局内存的合并访问做得很好，每个线程连续读取128字节，效率拉满。

但问题出在写入共享内存之后。

每个线程把读到的数据写到共享内存的对应位置。然后，另一个方向的线程来读取。你猜怎么着？读取的时候，16个线程同时访问同一个Bank。

嗯，16路Bank冲突。性能直接掉了60%。

我当时盯着profiler看了半天，心想：全局内存明明没问题啊，怎么整体这么慢？后来才意识到，是共享内存这头堵死了。

4.3 全局内存到共享内存的映射关系

咱们来拆解一下这个过程。假设你有32个线程，每个线程从全局内存读一个float。

// 全局内存加载（合并访问，没问题）
float data = global_mem[threadIdx.x];

然后你把它存到共享内存：

__shared__ float shared_data[32];
shared_data[threadIdx.x] = data;

这一步看起来人畜无害。但问题出在后续的读取。

假设另一个方向的线程要按列读取：

// 按列读取，假设矩阵是32x32
float val = shared_data[threadIdx.y * 32 + threadIdx.x];

如果threadIdx.y是0，那所有线程读的都是第0行。第0行的32个元素，分布在32个Bank上吗？

不一定。

共享内存的Bank宽度是4字节。一个float正好占一个Bank。32个float，如果连续存放，正好占满32个Bank——每个Bank一个元素，完美无冲突。

但如果你存的是struct，或者数据有padding，或者访问步长是32的倍数……那Bank冲突就来了。

4.4 常见的坑：步长与Bank冲突

我整理了一个表格，方便你对照：

访问步长（以float为单位）	Bank冲突情况	实际带宽利用率
1（连续访问）	无冲突	100%
2	2路冲突	50%
4	4路冲突	25%
8	8路冲突	12.5%
16	16路冲突	6.25%
32	32路冲突（最坏情况）	3.125%

你看，步长32的时候，所有线程都在抢同一个Bank。这比全局内存的非合并访问还惨。

4.5 如何避免？我常用的三种方法

方法一：padding

在共享内存数组的每一行末尾加一个额外的元素。比如原本是shared_data[32][32]，改成shared_data[32][33]。这样每行多出一个float，Bank的映射就被打乱了。

__shared__ float shared_data[32][33];  // 加padding
// 写入时还是按行写
shared_data[threadIdx.y][threadIdx.x] = data;
// 读取时按列读
float val = shared_data[threadIdx.x][threadIdx.y];  // 注意索引交换

方法二：改变数据排布

在写入共享内存时，就按后续读取的方式排布。比如后续要按列读，那写入时就按列写。说白了，就是提前转置。

方法三：使用向量化加载

用float4一次加载4个float。这样每个线程处理4个元素，Bank冲突的粒度就变了。

float4 data = reinterpret_cast<float4*>(global_mem)[threadIdx.x];
shared_data[threadIdx.x * 4 + 0] = data.x;
shared_data[threadIdx.x * 4 + 1] = data.y;
shared_data[threadIdx.x * 4 + 2] = data.z;
shared_data[threadIdx.x * 4 + 3] = data.w;

4.6 全局内存与共享内存的协同策略

我个人习惯把这个问题拆成三步：

1. 全局内存加载阶段：保证合并访问。每个线程连续读取，warp内的线程地址连续。
2. 共享内存存储阶段：按需排布。想清楚后续怎么读，现在就怎么存。
3. 共享内存读取阶段：检查Bank冲突。用profiler跑一遍，看有没有超过2路的冲突。

我曾经在一个项目里，把这三步画成了流程图，贴在工位上。每次写新kernel都对照着检查一遍。效果很好，基本没再出过Bank冲突的问题。

4.7 总结一下

全局内存没有Bank，但你的数据最终要进共享内存。共享内存的Bank冲突，根源往往在全局内存的访问模式上。

说白了，这是一个数据排布的问题。你在全局内存怎么放，决定了共享内存怎么存，进而决定了有没有Bank冲突。

我建议你每次写kernel时，都画一张数据流图。从全局内存到共享内存，再到计算单元。每一步都问自己：这一步有没有冲突？有没有更好的排布方式？

嗯，养成这个习惯后，你会发现很多性能问题，其实在写代码之前就能预判到。

---

公众号:蓝海资料掘金营，微信deep3321