Bank冲突的理解

bank conflict的定义

同一warp内的线程访问同一bank的不同字地址时，会发生 bank conflict，导致这些访问被串行化，从而增加访问延迟。
不同warp不存在bank冲突

不同warp访问同一bank实验

#include <iostream>
#include <vector>
#include "cuda_helper.cuh"

__global__ void func(float* a) {
    __shared__ float shared_data[32][32];
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    shared_data[threadIdx.y][threadIdx.x] = a[idx];
    __syncthreads();
    a[idx] = shared_data[threadIdx.y][31 - threadIdx.x] * 0.5f; // 这里让0-31号线程访问元素31-0，交错一下，避免太简单被编译器搞成不知道什么样子
}

int main() {
    std::vector<float> test_data(32*32);
    std::vector<float> res_data(32*32);
    for (int i = 0; i < 32*32; i++) {
        test_data[i] = static_cast<float>(i);
    }
    float * d_test_data{};
    CHECK_CUDA(cudaMalloc(&d_test_data, sizeof(float)*32*32));
    CHECK_CUDA(cudaMemcpy(d_test_data, test_data.data(), sizeof(float)*32*32, cudaMemcpyHostToDevice));
    func<<<1, dim3(32,32)>>>(d_test_data);
    CHECK_CUDA(cudaDeviceSynchronize());
    CHECK_CUDA(cudaMemcpy(res_data.data(), d_test_data, sizeof(float)*32*32, cudaMemcpyDeviceToHost));
    for (int i = 0; i < 32*32; i++) {
        std::cout << res_data[i] << " ";
    }
    return 0;
}

这段代码使用ncu --set full -k func -o m ./main分析时，报告中Details\Memory Workload\Memory Table中显示如下，不存在任何bank冲突

分析

不同分区上并行的warp在访问同一bank不同地址时会怎么样呢？
shared memory有一个Pipeline,分为如下几个阶段：

1. Issue Decode
   由warp scheduler发出一条ld.shared/st.shared指令，经过IF、ID
2. Address Generation
   每个线程的地址由AGU计算
3. Bank Mapping & Conflict Resolution
   将warp发出的请求映射到bank,检测bank conflict，若有，将访问拆分为多个serialized transactions，若无，生成一个32-way transaction
4. SMEM Crossbar Arbitration
   共享内存交叉开关仲裁，有一个队列。
   这里的仲裁是per-bank并行执行的，如果bank busy则等待
5. Bank Access
6. Data Return