1. 背景知识
-
常说的 bankConflict 指的在共享内存的 store load 过程中产生。
-
共享内存位于 L1TexCache 上，使用场景通常为从 GlobalMemeory store to SharedM，sharedM load to register。
-
共享内存的硬件实现：为实现高带宽，共享内存被划分大小相等的存储体，这些模块可以被同时访问。
-
共享内存的访存原理：如果 n 个数据请求落在n 个不同的存储体(逻辑上叫bank)，可以被同时处理，相当于获得单个存储体带宽 n 倍的整体带宽。
2. BankConflict
-
首先一个warp 的所有线程对共享内存访问的一个内存事务可以是非连续的地址；
-
如果一次内存事务中不同线程访问数据在一个存储体（包含一个bank不同字节的情况）就会引发串行访问，此时发生 m 路 bankConflict，需要 m 个内存事务请求。
-
如下图，一个warp 非连续访存不同 bank(存储体)，只需要一个访存事务：
- 如下图，如果一个warp 访存同一个bank(存储体)，则会产生 bankConflict，在第18号存储体发生串行访问，需要 4 个内存事务：
3. 解决方案
- padding
- bank_B 的元素在 padding 后的数组中列坐标计算：
- A: 共享内存的原数据每行占用 bank 数量；
- P: 行末 Padding bank 数量；
- B: 数据在原数组中的 bank 坐标；
- Yi: 元素在 padding 后数组列坐标；
- i：数据的 bank 标记循环次数（i≥0）。
\[Y_{i} = (i*32 + B) \bmod (A+P)
\]
- Padding 参数 P 的选择：无论对共享内存连续访问还是非连续访问，通常设置为最小值 1；
- swizzle：混合布局。将共享内存原来的存放地址进行转换，得到新的物理存放地址。避免bankconflict。
- 算法原理：将地址行列坐标做异或运算的结果覆盖原坐标中的列坐标数据。
- 将矩阵每8个连续元素分块, 比如 16 * 64大小的矩阵，如下图：
- addr 地址末三位为tile 内的位置，一行8个元素，只需要定位到 tile 起始位置，所以一定为 000；中间三位为块的列坐标，前三位为块的行坐标。
- 对于gaddr 在共享内存的原地址addr: B 表示列需要的二进制位数，M 表示一个块内的元素索引需要的二进制位数，S 表示 addr 地址按块划分的行列坐标需要位移的二进制位数=M。
template<uint32_t B,uint32_t M,uint32_t S>
__device__ __forceinline__ uint32_t swizzle(uint32_t srcAddr)
{//行列坐标值取后三位进行异或运算//掩码用来获取行坐标uint32_t mask = (1 << B - 1) << M;uint32_t addr = ((srcAddr >> S) & mask) ^ srcAddr;return addr;
}
