CUDA优化：Bank Conflict
使用到的测试代码：github – cuda_perf
1. Bank Conflicts (Shared Memory)
针对Shared Memory的访问，CUDA使用bank机制，将shared memory的访问（读/写）映射到不同的bank，以实现并行访问。bank以4字节为单位，共32个bank。这样，一个时钟周期内，可以并行访问32个不同的bank，即访问128字节的数据。映射公式bank index = (address /4) % 32。
图示transaction：
1
2
3
4
5
6
7
Thread（在CUDA Core中）
↓
访问Shared Memory
↓
[Bank系统处理] ← Transaction在这里发生
↓
返回数据到Thread
举例：warp中定义的shared memory如何映射到banks：
1
__shared__ float s[64];
如上变量，其映射如下：
在一次transaction的时候，如果，当warp中的不同线程访问到同一个bank中的不同地址时，就会产生Bank Conflicts，导致访问串行化：需要分成多次transaction。有N个线程访问同一个bank，称为N-way Bank Conflicts。
Bank Conflicts影响仅限一个warp内的一次transaction，且内存为shared memory。
如下情况，会产生Bank Conflicts：
- 一次
transaction中，warp中的多个线程，访问同一个bank中的不同地址； - 一次
transaction中，warp中的多个线程，访问shared memory的下一个128字节，被映射到同一个bank。此时也是属于上一种情况：同一个bank中的不同地址。
如下情况，Bank Conflicts不会产生：
warp中的线程，访问地址唯一对应到bank簇的每个bank，不论是顺序，还是错位；warp中的多个线程，访问同一个bank中的相同地址–使用boardcast分发相同地址数据到多个线程；warp中的线程，单个线程一次访问多个bank，但其他线程不访问这些bank。此时，生成多次transaction。
2. Bank Conflicts 示例
如下示例，产生32路Bank Conflicts：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void all_conflicts() {
__shared__ float s[32][32];
[[maybe_unused]] int warp_id = threadIdx.y;
int lane_id = threadIdx.x; // thread 在 warp 中的 id
float *ptr = &s[lane_id][0];
int addr = (int)(uintptr_t)ptr & 0xFFFF;
[[maybe_unused]] float r1; // 声明输出变量
for (int j = 0; j < num_iters; j++) { // num_iters 定义为 100'000
asm volatile("ld.volatile.shared.f32 %0, [%1];" : "=f"(r1) : "r"(addr));
}
}
// launched withall_conflicts<<<1, dim3(32, 8)>>>();
// Gride size: 1（即只有一个Block）
// Block size: dim3(32, 8)（即有8个warp，每个warp 32个线程）
使用
Nsight Compute分析bank conflicts（Metrics full），分析步骤见使用Nsight Compute分析Bank Conflict。
2.1. conflict free 代码参考
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void conflict_free_kernel() {
__shared__ float s[8][32];
int warp_id = threadIdx.y;
int lane_id = threadIdx.x;
float *ptr = &s[warp_id][lane_id];
int addr = (int)(uintptr_t)ptr & 0xFFFF;
[[maybe_unused]] float r1; // 声明输出变量
for (int j = 0; j < num_iters; j++) {
asm volatile("ld.volatile.shared.f32 %0, [%1];" : "=f"(r1) : "r"(addr));
}
}
// conflict_free_kernel<<<1, dim3(32, 8)>>>();
// Gride size: 1（即只有一个Block）
// Block size: dim3(32, 8)（即有8个warp，每个warp 32个线程）
3. 矢量读写指令
使用矢量指令ld.shared.v4可以读取4个连续的32位数据。如下代码，一个线程读取s[4*i], s[4*i+1], s[4*i+2], s[4*i+3]（分为四个transaction）。会产生四路Bank Conflicts：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
__global__ void vectorized_loads() {
__shared__ float sh[8][128];
int warp_id = threadIdx.y;
int lane_id = threadIdx.x;
float4* ptr = reinterpret_cast<float4*>(&sh[warp_id][lane_id * 4]);
int addr = (int)ptr & 0xFFFF;
float4 r;
for (int j = 0; j < num_iters; j++) {
asm volatile ("ld.volatile.shared.v4.f32 {%0,%1,%2,%3}, [%4];"
: "=f"(r.x), "=f"(r.y), "=f"(r.z), "=f"(r.w)
: "r"(addr));
}
}
上图只给出了一半的线程访问情况。编号(
Lane)为0的线程与编号为8的线程，访问的shared memory中的数据映射到了同一个bank 0；同时，编号为16，以及24，同样映射到了bank 0。 不过，由于其每个线程一次访问4个32位数据，其平均访问时间折算下来，与32位加载相当。
要想避免Bank Conflicts，可以错开（interleave）冲突的线程访问的顺序，比如：
- 线程0：s[0] -> s[1] -> s[2] -> s[3]
- 线程8：s[33] -> s[34] -> s[35] -> s[32]
3. 避免 Bank Conflicts 的方法
3.1. Padding
warp内多个线程访问同一bank会引发冲突，导致串行化访问。 通过在二维共享内存数组的列数上 +1 padding，可打破冲突模式。示意图：
示例代码：
1
2
3
4
5
6
7
__shared__ float sData[BLOCKSIZE][BLOCKSIZE + 1]; // +1 避免bank冲突
int x = threadIdx.x;
int y = threadIdx.y;
sData[x][y] = matrix[y * col + x];
__syncthreads();
matrixTest[y * col + x] = sData[x][y];