使用 NVProf 检测 CUDA kernel 的 bank conflict

NVProf 指令

使用 NVProf 可以对 bank conflict 进行检测:

nvprof --events shared_ld_bank_conflict,shared_st_bank_conflict <app> [args...]

其中:

--events 选项指定的 shared_ld_bank_conflict,shared_st_bank_conflict分别代指从 shared memory 加载(读取)时产生的 bank conflict, 以及向 shared memory 存储(写入)时产生的 bank conflict.
<app> [args...] 即要检测的 CUDA 二进制程序及其参数.

额外说明

值得一提的是, 如果没有从 shared memory 读取的指令, 且没有使用 -G 编译, 则两种 bank conflict 事件都无法检测出来, 即使存在向 shared memory 写入产生的 bank conflict.
(没有读取的 bank conflict 很好理解, 因为都没有从 shared memory 读取数据; 而至于写入的 bank conflict, 应该是编译器做了一定的优化, 即 shared memory 虽被写入但数据没有被读取, 则写入是没有意义的, 这部分代码实际并不执行, 所有写入的 bank conflict 就不会检测到了.)

这个主要作用是, 当我们对自己写的 kernel 的 bank conflict 进行检测的时候, 要确保保留对 shared memory 读取的相关代码或设置 -G 编译选项, 否则可能会影响 bank conflict 的检测.

举例

以下代码是一个很简单的 CUDA kernel 示例, 考虑到 bank conflict 是 warp 层面的问题, 所有 kernel 中我定义了 warp_id, land_id 等变量便于后续 bank conflict 的说明.

#include <iostream>
#include <cstdio>
#include <vector>
#include <cuda.h>using namespace std;constexpr int SIZE_A = 64;
constexpr int SIZE_C = 64;__global__ void kernel(const int* a, int* c) {auto tid = (blockIdx.x * blockDim.x + threadIdx.x);auto lane_id = threadIdx.x & 0x1F;auto warp_id = tid >> 5;auto warp_in_block = threadIdx.x >> 5;__shared__ int shm[SIZE_A];if (tid < SIZE_A) {shm[warp_id * 32 + lane_id] = a[warp_id * 32 + lane_id];}if (tid < SIZE_C) {c[warp_id * 32 + lane_id] = shm[warp_id * 32 + lane_id];}
}int main() {vector<int> a(SIZE_A);for (int i = 0; i < SIZE_A; ++i) {a[i] = i;}int* d_a;cudaMalloc(&d_a, sizeof(int) * SIZE_A);cudaMemcpy(d_a, a.data(), sizeof(int) * SIZE_A, cudaMemcpyHostToDevice);int* d_c;cudaMalloc(&d_c, sizeof(int) * SIZE_C);cudaMemset(d_c, 0, sizeof(int) * SIZE_C);kernel<<<1, 128>>>(d_a, d_c);vector<int> c(SIZE_C);cudaMemcpy(c.data(), d_c, sizeof(int) * SIZE_C, cudaMemcpyDeviceToHost);for (auto x : c) {cout << x << " ";}cout << endl;cudaFree(d_c);cudaFree(d_a);return 0;
}

kernel() 函数完成的功能很简单, 就是想数组 a 中的一部分数据先写至 shared memory shm, 再写入到 c 中. 在没有额外说明时, 不使用 -G 选项编译代码.
很明显的是, 由于 shm 的读写时, 每个 warp 的 32 个线程分片读取不同的 4 字节数据, 因此代码没有 bank conflict.
在这里插入图片描述
使用上述 NVProf 指令检测, 结果也印证了上述推断.

现在将 Kernel 修改如下:

__global__ void kernel(const int* a, int* c) {auto tid = (blockIdx.x * blockDim.x + threadIdx.x);auto lane_id = threadIdx.x & 0x1F;auto warp_id = tid >> 5;auto warp_in_block = threadIdx.x >> 5;__shared__ int shm[SIZE_A];// if (tid < SIZE_A) {//     shm[warp_id * 32 + lane_id] = a[warp_id * 32 + lane_id];// }for (auto i = threadIdx.x; i < SIZE_A; i += blockDim.x) {shm[(i % 2) * SIZE_A / 2 + i / 2] = a[i];}if (tid < SIZE_C) {c[warp_id * 32 + lane_id] = shm[warp_id * 32 + lane_id];;}
}

我们在读取 a 数组到 shared memory 的时候, 进行了一点修改. 可以看到, 对应相邻的两个线程, t 和 t+1 (假设 t % 2 ==0), 则一个写入到 shm[t/2], 一个写入到 shm[SIZE_A/2+(t+1)/2] 即 shm[32+t/2], 由于恰好差了 32 个元素, 因此会访问到相同的 bank, 会触发 bank conflict. 通过 NVProf 检测也得到了证实:
在这里插入图片描述
这里的 2 次, 原因笔者猜测为 SIZE_A 大小为 64, 对应 2 个 warp, 每个 warp 相邻的奇数线程和偶数线程访问同一 bank, 以 warp 为单位, 每个 warp 产生 1 个 bank conflict, 共 2 个.

但如果我们将后面将 shm 写入 c 数组的代码注释掉, 即没有从 shared memory 读取的代码, 则可以看到 NVProf 并不会检测到刚刚的 shared_st_bank_conflict.

    if (tid < SIZE_C) {c[warp_id * 32 + lane_id] = shm[warp_id * 32 + lane_id];;}

在这里插入图片描述

但如果我们在编译的时候使用 -G 选项, 则可以看到刚刚的 shared_st_bank_conflict 有可以被检测到了:
在这里插入图片描述

因此, 可以推断出, 在默认情况下, 编译器对于不读取的 shared memory 的写入操作会进行优化, 实际上并不会执行 shared memory 的写入操作, 而 debug 模式 (带 -G 选项)时, 则不会进行该优化.

如下代码展示了在从 shared memory shm 读取到 c 数组时的 bank conflict.

constexpr int SIZE_A = 64;
constexpr int SIZE_C = 32;__global__ void kernel(const int* a, int* c) {auto tid = (blockIdx.x * blockDim.x + threadIdx.x);auto lane_id = threadIdx.x & 0x1F;auto warp_id = tid >> 5;auto warp_in_block = threadIdx.x >> 5;__shared__ int shm[SIZE_A];if (tid < SIZE_A) {shm[warp_id * 32 + lane_id] = a[warp_id * 32 + lane_id];}if (tid < SIZE_C) {// c[warp_id * 32 + lane_id] = shm[warp_id * 32 + lane_id];c[warp_id * 32 + lane_id] =shm[warp_in_block * 32 + lane_id / 8 + (lane_id % 2) * 32];}
}