4. 内存访问优化

大家好，欢迎来到第四讲。

今天聊的话题，是 GPU 性能优化的核心中的核心——内存访问。

我经常跟团队里的新人说一句话：你写的 CUDA 代码快不快，90% 取决于你怎么访问内存。计算本身往往不是瓶颈，数据搬运才是。说白了，GPU 的算力再强，如果数据喂不进去，那也是白搭。

这一讲，我会把内存访问优化的几个关键点串起来讲：合并访问、内存对齐、Bank Conflict 避免、以及用共享内存减少全局内存访问。这些都是我在实际项目中踩过坑、填过坑的地方。

1. 合并访问：让数据像流水一样流进来

先问一个问题：为什么同一个 kernel，换个数据布局，性能能差 10 倍？

答案就在合并访问上。

GPU 的全局内存控制器，一次会从 DRAM 中读取一大块连续的数据（通常是 128 字节）。如果一个 warp（32 个线程）访问的地址是连续的、对齐的，那么硬件只需要一次或少数几次内存事务就能搞定。这就是合并访问。

反过来，如果每个线程访问的地址是跳跃的、分散的，那硬件就得拆成很多次小事务，带宽利用率直接腰斩。

来看一个对比示例：

// ❌ 非合并访问：线程 0 访问 data[0][0]，线程 1 访问 data[1][0]...
// 地址不连续，性能极差
__global__ void bad_access(float* data, int width) {
    int row = threadIdx.x;
    float val = data[row * width + 0];  // 列方向访问
}

// ✅ 合并访问：线程 0 访问 data[0]，线程 1 访问 data[1]...
// 地址连续，一个事务搞定
__global__ void good_access(float* data, int width) {
    int col = threadIdx.x;
    float val = data[col];  // 行方向访问
}

嗯，这里要注意：合并访问不仅适用于全局内存，也适用于共享内存。只不过共享内存的 bank 冲突是另一个故事，我们后面讲。

3. 内存对齐：别让硬件多干活

合并访问的前提是什么？对齐。

如果数据起始地址不是 128 字节对齐的，即使访问模式是连续的，硬件也可能需要两次事务才能读完。为什么？因为内存控制器是按对齐的块来读取的。

我个人习惯在分配内存时，用 cudaMalloc 而不是自己手动 malloc，因为 cudaMalloc 默认会返回对齐的地址。如果你用 malloc 或者自己管理内存池，记得用 posix_memalign 或 aligned_alloc 来保证对齐。

对齐的规则很简单：

• 基本类型（float, int）：4 字节对齐
• 向量类型（float2, float4）：8 或 16 字节对齐
• 结构体：按最大成员对齐，建议手动 padding 到 16 字节

你想想看，如果每个线程读一个 float4（16 字节），32 个线程就是 512 字节，正好是 4 个 128 字节事务。完美对齐，完美合并。

3. Bank Conflict：共享内存里的隐形杀手

共享内存虽然快，但它不是没有代价的。共享内存被划分成 32 个 bank（在大多数架构上），每个 bank 每周期只能响应一个请求。

如果同一个 warp 里的多个线程访问同一个 bank 的不同地址，就会发生 bank conflict，导致请求串行化。

说白了，就是本来可以并行读的，现在得排队了。

举个例子：

// ❌ 有 bank conflict：线程 0 访问 shmem[0]，线程 1 访问 shmem[32]...
// 它们落在同一个 bank 上（因为 32 % 32 == 0）
__shared__ float shmem[256];
int idx = threadIdx.x * 32;  // 步长 32，全部冲突
float val = shmem[idx];

// ✅ 无 bank conflict：线程 0 访问 shmem[0]，线程 1 访问 shmem[1]...
// 每个线程访问不同的 bank
int idx = threadIdx.x;  // 步长 1，无冲突
float val = shmem[idx];

我在项目中遇到过最头疼的情况是矩阵转置。如果直接按列写入共享内存，就会产生严重的 bank conflict。解决方案是加 padding：把共享内存声明为 __shared__ float shmem[BLOCK_SIZE][BLOCK_SIZE + 1]，多出一列来错开 bank 映射。

4. 使用共享内存减少全局内存访问

共享内存最大的价值，就是做数据复用。

很多算法（比如卷积、矩阵乘法、归约）中，同一个数据会被多个线程多次使用。如果每次都从全局内存读，带宽根本扛不住。正确的做法是：先把数据从全局内存搬到共享内存，然后所有线程从共享内存读。

来看一个矩阵乘法的简化示例：

__global__ void matmul_shared(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    
    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE; t++) {
        // 协作加载：每个线程加载一个元素到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col];
        __syncthreads();  // 确保所有线程加载完成
        
        // 从共享内存计算
        for (int k = 0; k < TILE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 确保所有线程用完数据，再加载下一块
    }
    C[row * N + col] = sum;
}

这个模式叫分块（tiling）。每个线程只负责加载一个元素，但所有线程一起把一整块数据搬进共享内存。然后每个线程在计算时，可以反复使用这块数据。

总结一下

这一讲我们聊了四个关键点：

• 合并访问：让连续线程访问连续地址，这是带宽利用率的基石。
• 内存对齐：保证起始地址是 128 字节对齐，避免额外事务。
• Bank Conflict：共享内存的并行访问限制，用 padding 或改变步长来避免。
• 共享内存复用：用分块策略，把全局内存访问次数降到最低。

说实话，这些技巧看起来简单，但真正用好需要大量实践。我每次写新 kernel，都会先跑一遍 profiler，看看内存访问模式是不是最优的。如果发现带宽利用率低于 60%，我就会回头检查上面这四点。

嗯，这一讲就到这里。记住：内存访问优化，是 GPU 性能优化的第一课，也是最后一课。