第四章：内存访问优化（二）——Shared Memory的使用与Bank Conflict

好，咱们接着聊内存优化。上一章讲了全局内存的合并访问，说白了就是让数据排好队进寄存器。但光靠这个还不够——全局内存再快，延迟也是几百个周期起步。这时候就得请出Shared Memory这位老将了。

Shared Memory是什么？它是GPU芯片上的一块SRAM，每个SM独享。延迟只有几十个周期，比全局内存快了一个数量级。我刚开始做CUDA优化时，总觉得Shared Memory容量太小（也就几十KB），懒得用。后来被一个矩阵乘法的性能瓶颈逼得没办法，才认真研究它。嗯，真香。

4.1 Shared Memory的基本用法

Shared Memory用__shared__关键字声明。它的生命周期是整个block，block内的所有线程都能读写。注意，不同block之间的Shared Memory是隔离的，互不干扰。

__global__ void shared_mem_example(float *input, float *output, int N) {
    __shared__ float shared_data[256];  // 每个block独享256个float
    
    int tid = threadIdx.x;
    int global_idx = blockIdx.x * blockDim.x + tid;
    
    // 从全局内存加载到Shared Memory
    if (global_idx < N) {
        shared_data[tid] = input[global_idx];
    }
    __syncthreads();  // 同步！确保所有线程都加载完毕
    
    // 在Shared Memory上做计算...
    // 比如做一次规约操作
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        if (tid % (2 * stride) == 0) {
            shared_data[tid] += shared_data[tid + stride];
        }
        __syncthreads();
    }
    
    if (tid == 0) {
        output[blockIdx.x] = shared_data[0];
    }
}

这里有个关键点：__syncthreads()。为什么需要它？因为Shared Memory没有自动同步机制。线程A写数据，线程B马上读，如果不加同步，B读到的可能是旧数据。我在项目中遇到过这种bug，排查了整整一天才发现是同步没加对位置。

4.2 Bank Conflict：Shared Memory的隐形杀手

Shared Memory虽然快，但有个坑——Bank Conflict。Shared Memory被分成32个Bank（对应warp的32个线程），每个Bank的带宽是32位（4字节）。如果同一个warp里的多个线程访问同一个Bank的不同地址，就会发生冲突，访问会串行化。

举个例子：

__shared__ float shared[32][32];
int tid = threadIdx.x;
int row = tid / 32;
int col = tid % 32;

// 情况1：按列访问（冲突！）
float val = shared[row][col];  // 同一warp的线程访问同一列，Bank冲突

// 情况2：按行访问（无冲突）
float val = shared[col][row];  // 同一warp的线程访问不同列，无冲突

为什么会这样？因为Shared Memory的地址映射规则是：地址除以4字节后，对32取模，得到Bank编号。按列访问时，同一列的数据在内存中是连续存储的，相邻线程访问相邻地址，但地址对32取模后可能落在同一个Bank上。

核心原则：同一个warp内的线程，尽量访问不同的Bank。如果必须访问同一个Bank，尽量让所有线程访问同一个地址（广播机制，无冲突）。

4.3 Padding技术：解决Bank Conflict的利器

Padding说白了就是给数组加个填充列，改变地址对齐方式。比如一个32x32的矩阵，按列访问时冲突严重。我们改成32x33：

// 原始：32x32，按列访问冲突
__shared__ float shared_no_pad[32][32];

// 加Padding：32x33，按列访问无冲突
__shared__ float shared_pad[32][33];  // 多了一列

为什么加一列就能解决？因为地址偏移了。原来第0列第0行的地址是0，第0列第1行的地址是32（按行存储）。加了Padding后，第0列第0行的地址是0，第0列第1行的地址是33。33对32取模是1，和0不同，所以Bank冲突消失了。

我曾经在一个图像处理项目里，用Padding把Shared Memory的带宽利用率从40%提升到了95%。效果立竿见影。

小技巧：Padding的大小不一定是1列。你可以根据访问模式灵活调整。比如按2的幂次访问时，加1列通常就够了。如果访问步长是2，可能需要加2列。

4.4 实战：矩阵乘法优化（Tiling）

好了，理论讲完了，咱们来点实战。矩阵乘法是CUDA优化的经典案例，也是检验Shared Memory使用水平的试金石。

先回顾一下矩阵乘法的公式：C = A * B，其中A是MxK，B是KxN，C是MxN。每个C[i][j] = sum(A[i][k] * B[k][j]) for k in 0..K-1。

朴素实现中，每个线程计算C的一个元素，需要从全局内存读取A的一行和B的一列。A的一行有K个元素，B的一列也有K个元素。如果K很大，全局内存访问次数是O(K)，而且每次访问都不连续——A是按行访问（连续），B是按列访问（不连续）。

Tiling的核心思想：把矩阵分块，每个block负责计算一个tile（小块）。先把A和B的tile加载到Shared Memory，然后在Shared Memory上做计算。这样，全局内存访问次数从O(K)降到了O(K / tile_size)。

看代码：

#define TILE_SIZE 16

__global__ void matmul_tiled(float *A, float *B, float *C, 
                              int M, int N, int K) {
    __shared__ float shared_A[TILE_SIZE][TILE_SIZE];
    __shared__ float shared_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 - 1) / TILE_SIZE; tile++) {
        // 加载A的tile
        if (row < M && tile * TILE_SIZE + threadIdx.x < K) {
            shared_A[threadIdx.y][threadIdx.x] = 
                A[row * K + tile * TILE_SIZE + threadIdx.x];
        } else {
            shared_A[threadIdx.y][threadIdx.x] = 0.0f;
        }
        
        // 加载B的tile
        if (col < N && tile * TILE_SIZE + threadIdx.y < K) {
            shared_B[threadIdx.y][threadIdx.x] = 
                B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        } else {
            shared_B[threadIdx.y][threadIdx.x] = 0.0f;
        }
        
        __syncthreads();
        
        // 在Shared Memory上计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += shared_A[threadIdx.y][k] * shared_B[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    // 写回全局内存
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

这段代码有几个要点：

Tile大小选择：16x16是常见选择，因为每个block需要16*16*4*2 = 2KB的Shared Memory（两个tile）。加上其他开销，一个SM可以同时驻留多个block。
边界处理：当矩阵维度不是tile_size的整数倍时，需要做边界检查，越界的元素补0。
同步点：加载完每个tile后需要同步，计算完当前tile后也需要同步（防止下一个tile覆盖数据）。

注意：Shared Memory的容量有限。Fermi架构是48KB，Kepler是48KB可配，Maxwell是64KB，Turing是64KB可配。如果你的tile太大，会导致occupancy下降。我一般建议tile_size不超过32。

4.5 性能对比与调优

咱们做个简单对比。假设M=N=K=1024，tile_size=16：

实现方式	全局内存访问次数	Shared Memory使用	相对性能
朴素实现	1024次/线程	无	1x（基准）
Tiling 16x16	64次/线程	2KB/block	约8-10x
Tiling + Padding	64次/线程	2.1KB/block	约10-12x

看到没？Tiling带来的提升是数量级的。加上Padding后，还能再提升20%左右。

我个人习惯在调优时先跑一遍profiler，看看Shared Memory的利用率。如果利用率低于80%，大概率有Bank Conflict。这时候加Padding试试，往往能立竿见影。

我曾经在一个推荐系统的矩阵乘法里，把tile_size从16改成32，结果性能反而下降了。为什么？因为Shared Memory占用太大，导致occupancy从50%降到了25%。所以调优不是一味地增大tile，要综合考虑occupancy和计算密度。

总结一下本章核心：

Shared Memory是片上SRAM，延迟低但容量小
Bank Conflict会导致访问串行化，用Padding解决
Tiling技术通过分块加载，大幅减少全局内存访问
调优时要平衡tile_size和occupancy

这张图把本章的知识体系串起来了。从全局内存到Shared Memory，再到三个核心问题和对应的解决方案，最后落到矩阵乘法的实战优化。你想想看，是不是每个环节都环环相扣？

好了，本章的内容就到这里。记住，Shared Memory用好了，性能起飞；用不好，反而可能拖后腿。多动手写代码，多跑profiler，慢慢就能找到感觉。

📚 课程目录

第四章：内存访问优化（二）——Shared Memory的使用与Bank Conflict

4.1 Shared Memory的基本用法

4.2 Bank Conflict：Shared Memory的隐形杀手

4.3 Padding技术：解决Bank Conflict的利器

4.4 实战：矩阵乘法优化（Tiling）

4.5 性能对比与调优