GPU bank conflict

背景

最近在学习cuda的一些基础知识。学习CNN时，有bank冲突的测试。在cpu中，有一个很相似的问题，cpu cacheline false sharing。这两个问题都回到了体系结构上。

比较

概念上，bank是指将CPU的高速缓存（Cache）或物理内存（RAM）划分为多个独立的逻辑或物理存储区块。多个并发请求就可以同时访问不同的bank，提高访存带宽。bank conflict就是多个请求打到了同一个bank上，退化为串行访问。

bank在概念上更接近硬件端，cpu中对等的概念是cache中的bank conflict，但是cpu中有大量的bank访问端口，削弱了访问冲突导致的性能损失，平时基本也不会遇到这种问题。cpu中提到更多的其实是cacheline false sharing。cacheline层级上是“打包的bank”概念。

查阅一些资料后，我比较认可的说法是cpu和gpu涉及哲学的不同。cpu没有设计很强的“访存并发能力”，核心中的ld/st端口即使满载，一个cycle内能访问的数据量也不大。AX512这样的指令，一个cycle内也只需要16*32个float的数据量，这就天然导致了bank冲突很少发生。而gpu一个warp中就有32个thread，这样一次并发的吞吐就已经超过了CPU，且warp还只是gpu的基本执行单元。

示例

教学示例中共享内存优化的cnn写法如下, 仅截取写入share mem的部分：

// Load input data to shared memory
    for (int c = 0; c < inputChannels; c++) {
        // Each thread loads multiple elements to cover the tile with padding
        for (int dy = 0; dy < tileSizeWithPadding; dy += tileSize) {
            for (int dx = 0; dx < tileSizeWithPadding; dx += tileSize) {
                int in_y = in_y_base + ty + dy;
                int in_x = in_x_base + tx + dx;
                
                // Check bounds and apply padding
                float value = 0.0f;
                if (in_y >= 0 && in_y < inputSize && in_x >= 0 && in_x < inputSize) {
                    value = input[
                        b * inputChannels * inputSize * inputSize +
                        c * inputSize * inputSize +
                        in_y * inputSize + in_x
                    ];
                }
                
                // Store in shared memory if within tile bounds
                if (ty + dy < tileSizeWithPadding && tx + dx < tileSizeWithPadding) {
                    sharedInput[
                        c * tileSizeWithPadding * tileSizeWithPadding +
                        (ty + dy) * tileSizeWithPadding + (tx + dx)
                    ] = value;
                }
            }
        }
    }

使用ncu采集bank conflict相关事件：

bash

ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe
# 输出如下
convolutionSharedKernel(float *, float *, float *, int, int, int, int, int, int, int, int) (4, 4, 16)x(8, 8, 8), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
-------------------------------------------------------- ----------- ------------
Metric Name                                              Metric Unit Metric Value
-------------------------------------------------------- ----------- ------------
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                    67200
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                     6144
-------------------------------------------------------- ----------- ------------

可以发现有大量的bank conflict产生。必然是读写share mem时发生了conflict。

根据基本知识：share mem会被划分为32个bank，每个bank通常为4byte（float）,一个warp中的32个thread如果顺序访问bank，是不会产生conflict的。那有了这个理论的“锤子”，可以回头来考究一下上面的代码。

读取share mem的index为c * tileSizeWithPadding * tileSizeWithPadding + (ty + dy) * tileSizeWithPadding + (tx + dx), 对于初学者来说，这个索引由于高维和一维的转换而产生了复杂的index，不利于理解，可以做简化，比如设定通道数c先为0，input数据的维度dy和dx也为0，就得到了简化后的索引公式ty * tileSizeWithPadding + tx.

假设线程块设置是 blockDim.x = 16, blockDim.y = 16。这意味着一个线程块有 256 个线程。GPU 会把它们切成 Warp，切分的顺序是先看 tx，再看 ty。比如：

线程 0 到 15：ty = 0，tx = 0
线程 16 到 31：ty = 1，tx = 0

本示例中的block维度为dim3 blockDim(8, 8, 1);，那么warp的分配策略为：

ty=0, tx=0~7
ty=1, tx=0~7
ty=2, tx=0~7
ty=3, tx=0~7

本示例中的卷积核的维度为5，所以tileSizeWithPadding=8+5-1=12。那么根据简化后公式：

ty=0, tx=0~7。映射bank范围0~7
ty=1, tx=0~7。映射bank范围12~19
ty=2, tx=0~7。映射bank范围24~31
ty=3, tx=0~7。计算范围36~43，%32后实际范围4~11，和ty=0，tx=4~7就发生了conflict。概率4/32=0.125.

稍微深入研究一下

似乎发现，某些算法似乎天生就会产生bank conflict。比如上面示例中，block的维度是（8，8），卷积核的size（5，5），就会有12.5%的冲突概率。如果卷积核改为（3，3）大小，计算冲突概率为18.75%.

在实际的测试中，卷积核（3，3）实际的冲突计数反而少了：

bash

convolutionSharedKernel(float *, float *, float *, int, int, int, int, int, int, int, int) (4, 4, 16)x(8, 8, 8), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
-------------------------------------------------------- ----------- ------------
Metric Name                                              Metric Unit Metric Value
-------------------------------------------------------- ----------- ------------
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                    32257
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                     4096
-------------------------------------------------------- ----------- ------------

这是因为小的卷积核会有更少的访存次数，所以纯看冲突计数不能佐证我们的发现。这里单纯的ncu的指标似乎不能计算出“冲突比例”这样的指标。

这里如果要计算“冲突比例”，首先需要了解ncu提供的指标的意义：

l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld， ncu中的解释为of shared memory data bank conflicts generated by LDS, LD, 3D。我理解为LDS, LD, 3D（这个不是指令）指令执行时发生bank冲突事件的次数。
sm__sass_inst_executed_op_shared_ld, ncu中解释为of warp instructions executed: LDS, LD，LDS, LD指令（读取shared mem）执行的次数。需要注意的是指令是以warp为单位执行的，一个warp中的指令会被32个thread同时执行。

那就可以计算出“冲突比例”这样一个指标。

设定kernel size为3时：

bash

convolutionSharedKernel(float *, float *, float *, int, int, int, int, int, int, int, int) (4, 4, 16)x(8, 8, 8), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
-------------------------------------------------------- ----------- ------------
Metric Name                                              Metric Unit Metric Value
-------------------------------------------------------- ----------- ------------
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                    32256
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                     4096
sm__sass_inst_executed_op_shared_ld.sum                         inst        32256
sm__sass_inst_executed_op_shared_st.sum                         inst        12288
-------------------------------------------------------- ----------- ------------

设定kernel size为5时：

bash

convolutionSharedKernel(float *, float *, float *, int, int, int, int, int, int, int, int) (4, 4, 16)x(8, 8, 8), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
-------------------------------------------------------- ----------- ------------
Metric Name                                              Metric Unit Metric Value
-------------------------------------------------------- ----------- ------------
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                    67200
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                     6144
sm__sass_inst_executed_op_shared_ld.sum                         inst        89600
sm__sass_inst_executed_op_shared_st.sum                         inst        12288
-------------------------------------------------------- ----------- ------------

采集计数器后，根据对指标含义的理解，kernel size=3时，每读一次share mem，就发生一次bank conflict。kernel size=5时，每读一次share mem，就发生0.75次bank conflict。按照之前的思路计算conflict ratio，kernel size为5时的ratio确实比kernel size为3时的小。

inst是以wrap为单位执行的，作为分母是偏小的，但是即使粗暴推理乘上一个32，也和理想的10%左右的值相去甚远。而且我怀疑conflict的计数也并不是简单一个线程的冲突一次记录一次，gpu应该不会设计很精细的性能计数器。这里ncu又没有提供其它更合适的性能事件，以ratio的角度解释就有些奇怪了。这里更合理的解释应该是每发生一次share mem的访问，会带来1个cycle的bank conflict的性能损耗。

how to fix

最终的目的仍然是减少bank conflict的发生。查阅一些资料后有以下方法：

padding 填充

回到上面推导的的冲突的原因

ty=0, tx=0~7。范围0~7
ty=1, tx=0~7。范围12~19
ty=2, tx=0~7。范围24~31
ty=3, tx=0~7。计算范围36~43，%32后实际范围4~11，和ty=0，tx=4~7就发生了conflict。概率4/32=0.125.

可以发现，一些位置的bank是空闲的，比如8~11，20~23。不用经过严格的数学推导，如果stride（就是tileSizeWithPadding）和32有一些数学上的关系，比如倍数，公约数，随着ty的增加，对32取模后，是会固定落在某个范围的bank上的。padding填充的原理就是破坏这种数学关系，操作上就是在tileSizeWithPadding的基础上再加上一些padding，构造出和32互质的数。

比如，多加1个padding，tileSizeWithPadding=8+5-1+1=13。重新计算index：

ty=0, tx=0~7。映射bank范围0~7
ty=1, tx=0~7。映射bank范围13~20
ty=2, tx=0~7。计算范围26~33，映射bank范围26~31, 0~1
ty=3, tx=0~7。计算范围39~46，映射bank范围6~11

可以发现，增加一个padding后，冲突的位置为0~1, 6~7. 还是4个位置发生冲突。这和block的结构有关系，并不能完全消除bank conflict。看起来很难不发生conflict。

拷问AI后，业界的解法是对block做一维展开，比如dim3 blockDim(32, 2, 1); // 或者 (32, 1, 1). 这样天然可以避免bank conflict。

地址重映射/位运算

既然默认的分配策略会导致conflict，那么手动加一层映射层也可以来解决这个问题。

这里再次拷问一下AI，帮忙实现一个转换方法：

cpp

// 辅助内联函数：将 (row, col) 映射为 Swizzled 后的 1D 偏移量
__device__ __forceinline__ int getSwizzledIndex(int channel, int row, int col, int stride_width) {
    // 使用 XOR 操作打破 Bank 的跨行对齐
    int swizzled_col = col ^ row; 
    
    // 计算单层 Channel 的大小 (Height * StrideWidth)
    int channel_offset = channel * stride_width * stride_width;
    
    return channel_offset + row * stride_width + swizzled_col;
}

这里的row=ty+dy，col=tx+dx. 进行异或操作，为了简化，依旧设定dy和dx为0：

ty=0, tx=0~7。按位异或后，索引不变，分配0~7
ty=1, tx=0~7。按位异或后，索引为1, 0, 3, 2, 5, 7, 6
ty=2, tx=0~7。按位异或后，2, 3, 0, 1, 6, 7, 4, 5
ty=3, tx=0~7。按位异或后，这个就不算了，还是在0~7间

乘上stride_width后，和没有进行重映射的conflict行为基本是一样的，也就是重映射只改变了col的映射，stride的影响仍然没有消除。

在拷问AI的时候，似乎当前模型对GPU的理解还不是很透彻，这里的一些推导还是保持谨慎

总结：必须要留个尾巴了

这里其实对于理解gpu的share mem硬件和软件的相关行为已经足够了，再深入就更偏向算法设计层面了。对上面的分析过程和一些必要的沉淀知识总结一下：

bank conflict核心知识：

物理结构：GPU 共享内存（Shared Memory）被物理划分为 32 个独立的存储体（Banks），按 4 字节（32-bit）交错编址。
冲突条件：同一个 Warp 中的 32 个线程，在同一个 clock cycle 内访问了同一个 Bank 中的不同地址（若访问完全相同的地址会触发 Broadcast 机制，不会冲突）。
后果：硬件无法并行满足访存，多线程请求会被串行化（Serialized），导致访问延迟成倍增加。
产生conflict的场景：
- 二维矩阵按列访问/转置（Stride 访问）：当跨步访问步长（Stride）是 2 的幂（如 16, 32）时，极易落入同一 Bank。
- Tile 加载中 Padding 尺寸不当：图像/卷积 Tile 大小加上 Kernel 尺寸后（如 $8+5-1=12$），导致二维索引按一维展开后的行步长与 32 产生非互质因子，引发模重叠。

和CPU的比较：

CPU和GPU都存在bank conflict，只不过CPU中bank有更多的访问端口可以有效避免，但是GPU端口数量少更容易导致串行访问。

解决bank conflict的方法（其实最简单的就是用官方的算子库）：

Padding（填充额外列）
重构block的shape
位异或重映射

GPU bank conflict ​

背景 ​

比较 ​

示例 ​

稍微深入研究一下 ​

how to fix ​

padding 填充 ​

地址重映射/位运算 ​

总结：必须要留个尾巴了 ​

GPU bank conflict

背景

比较

示例

稍微深入研究一下

how to fix

padding 填充

地址重映射/位运算

总结：必须要留个尾巴了