4、访存模式优化：合并访问、Bank Conflict与数据预取策略

访存优化，说白了就是让GPU的显存带宽别闲着。

我见过太多项目，计算单元利用率不到30%，但显存带宽已经跑满了。你想想看，GPU算力再强，数据喂不进去也是白搭。这一章，咱们就聊聊怎么把数据搬得又快又稳。

4.1 合并访问：让数据像火车一样整齐

GPU的显存控制器有个特点——它喜欢连续的数据。就像火车，一整列拉走效率最高，零散的小包裹反而麻烦。

什么是合并访问？

当一个warp（32个线程）同时访问显存时，硬件会把这些请求合并成尽量少的事务（transaction）。如果32个线程访问的地址是连续的，那一个事务就能搞定。如果地址是乱序的，那就得多发几个事务，带宽利用率直线下降。

举个例子，假设我们有一个float数组，每个线程读一个元素：

// ❌ 糟糕的访问模式：按列访问
__global__ void bad_access(float *data, int width) {
    int col = threadIdx.x;
    int row = blockIdx.x;
    float val = data[row * width + col];  // 同一warp内，col不连续
}

// ✅ 好的访问模式：按行访问
__global__ void good_access(float *data, int width) {
    int row = threadIdx.x;
    int col = blockIdx.x;
    float val = data[row * width + col];  // 同一warp内，row连续
}

我在项目中遇到过一个问题：一个图像处理算法，按列处理像素，结果带宽利用率只有15%。改成按行处理后，直接飙到85%。嗯，有时候就是换个思路的事。

4.2 Bank Conflict：共享内存的隐形杀手

共享内存虽然快，但它也有自己的脾气。共享内存被分成32个bank（存储体），每个bank每周期只能响应一个请求。

什么是Bank Conflict？

当同一个warp内的多个线程访问同一个bank的不同地址时，就会发生bank conflict。这些访问会被串行化，导致性能下降。

举个例子：

__shared__ float s_data[32][32];

// ❌ 有bank conflict：同一列的数据在同一个bank
int idx = threadIdx.x;
float val = s_data[idx][0];  // 所有线程访问第0列

// ✅ 无bank conflict：同一行的数据在不同bank
float val = s_data[0][idx];  // 每个线程访问不同列

为什么会这样？因为共享内存的地址到bank的映射是：bank_id = (byte_address / 4) % 32。所以连续地址会落在不同bank上。

常见的解决策略：

• Padding：在数组每行末尾加几个空元素，改变地址对齐
• 改变访问模式：让线程访问不同的bank
• 使用向量化类型：比如float4，一次读4个元素

访问模式	Bank Conflict次数	实际带宽利用率
无冲突	0	100%
2路冲突	1	50%
4路冲突	3	25%
32路冲突	31	3.125%

4.3 数据预取：让数据提前到位

预取，就是提前把数据从慢速存储搬到快速存储。GPU里最常见的预取场景，就是把数据从全局内存搬到共享内存或寄存器。

为什么需要预取？

因为全局内存的延迟是几百个周期，而计算只需要几个周期。如果不预取，计算单元就得干等着。

常见的预取策略：

1. 软件预取：手动用__pld()指令或ldg()提前加载
2. 双缓冲：用两个共享内存缓冲区，一个计算，一个加载
3. 循环展开+预取：在循环中提前加载下一轮的数据

来看一个双缓冲的例子：

__global__ void double_buffer_kernel(float *input, float *output, int N) {
    __shared__ float s_buf[2][BLOCK_SIZE];
    
    int tid = threadIdx.x;
    int global_id = blockIdx.x * blockDim.x + tid;
    
    // 预取第一块数据
    s_buf[0][tid] = input[global_id];
    __syncthreads();
    
    for (int i = 0; i < N / BLOCK_SIZE - 1; i++) {
        int current = i % 2;
        int next = (i + 1) % 2;
        
        // 预取下一块数据
        if (global_id + (i + 1) * BLOCK_SIZE < N) {
            s_buf[next][tid] = input[global_id + (i + 1) * BLOCK_SIZE];
        }
        
        // 计算当前数据
        output[global_id + i * BLOCK_SIZE] = s_buf[current][tid] * 2.0f;
        
        __syncthreads();
    }
    
    // 处理最后一块
    output[global_id + (N / BLOCK_SIZE - 1) * BLOCK_SIZE] = 
        s_buf[(N / BLOCK_SIZE - 1) % 2][tid] * 2.0f;
}

4.4 综合实战：一个矩阵乘法的优化案例

咱们把上面三个策略串起来，看看一个矩阵乘法是怎么优化的。

原始版本：

// 每个线程算一个元素，直接读全局内存
__global__ void matmul_naive(float *A, float *B, float *C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    for (int k = 0; k < K; k++) {
        sum += A[row * K + k] * B[k * N + col];
    }
    C[row * N + col] = sum;
}

这个版本的问题很明显：每次循环都读全局内存，而且A的访问是连续的，但B的访问是跨步的。

优化后版本：

// 使用共享内存 + 合并访问 + 预取
__global__ void matmul_optimized(float *A, float *B, float *C, 
                                  int M, int N, int K) {
    __shared__ float s_A[TILE_SIZE][TILE_SIZE];
    __shared__ float s_B[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int tile = 0; tile < K / TILE_SIZE; tile++) {
        // 合并访问：连续读A和B
        s_A[threadIdx.y][threadIdx.x] = A[row * K + tile * TILE_SIZE + threadIdx.x];
        s_B[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        // 无bank conflict：每个线程读不同列
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += s_A[threadIdx.y][k] * s_B[k][threadIdx.x];
        }
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

这个版本做了三件事：

• 合并访问：每个tile内，A和B的读取都是连续的
• 消除bank conflict：共享内存的访问模式是行读列，没有冲突
• 数据预取：用共享内存做缓冲，一次加载多次使用

4.5 避坑指南

最后，分享几个我踩过的坑：

• 别盲目用共享内存：共享内存只有几十KB，用多了反而降低occupancy
• 注意对齐：128字节对齐的访问效率最高，不对齐会触发多个事务
• 小心原子操作：原子操作会串行化，能避免就避免
• 别忽略寄存器溢出：寄存器不够用时会 spill 到本地内存，性能暴跌

我曾经在一个项目中，为了消除bank conflict，把共享内存数组从[32][32]改成[32][33]（padding）。结果因为多占用了共享内存，occupancy从50%降到了33%，性能反而下降了。嗯，优化就是这样，牵一发而动全身。

记住一句话：没有银弹，只有权衡。每个优化策略都有代价，关键是要找到当前瓶颈的平衡点。