内存访问模式优化：全局内存合并访问、共享内存bank冲突、常量内存与纹理内存

内存访问，说白了就是GPU性能的命门。

我见过太多项目，kernel写得花里胡哨，结果一跑profile，发现大部分时间都卡在内存访问上。你想想看，计算单元再快，数据喂不进去也是白搭。今天咱们就聊聊内存访问模式优化的几个核心点：全局内存合并访问、共享内存bank冲突，还有常量内存和纹理内存的妙用。

全局内存合并访问

全局内存是GPU里最大的存储空间，但也是访问最慢的。不过，如果你能让线程束（warp）里的32个线程访问连续的内存地址，硬件就会把这些访问合并成一次或几次大的内存事务。这就是所谓的“合并访问”。

我个人习惯在写kernel之前，先画个内存访问的草图。举个例子：

// 不合并的写法（糟糕）
__global__ void bad_access(float *A, float *B, int N) {
    int tid = threadIdx.x;
    for (int i = 0; i < N; i += blockDim.x) {
        B[tid] = A[tid * N + i];  // 跨步访问，不连续
    }
}

// 合并的写法（好）
__global__ void good_access(float *A, float *B, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        B[tid] = A[tid];  // 连续访问，完美合并
    }
}

我在项目中遇到过最典型的坑是：二维数组按行访问和按列访问的性能差异。按行访问时，相邻线程访问相邻地址，合并访问效率极高。按列访问时，每个线程跳着读，性能直接掉一个数量级。

共享内存bank冲突

共享内存是片上的SRAM，速度比全局内存快两个数量级。但它有个“脾气”——bank冲突。

共享内存被分成32个bank（对应一个warp的32个线程）。每个bank每个时钟周期只能响应一个访问请求。如果同一个warp里多个线程访问同一个bank的不同地址，就会发生bank冲突，访问会被串行化。

嗯，这里要注意：如果多个线程访问同一个bank的同一个地址，反而没事，会触发广播机制。

我举个例子说明bank冲突：

// bank冲突的写法
__shared__ float s_data[32][32];
int tid = threadIdx.x;
s_data[tid][tid] = ...;  // 所有线程访问同一列，bank冲突严重

// 无bank冲突的写法
s_data[tid][0] = ...;    // 连续访问，无冲突

我曾经在一个图像处理项目里，因为共享内存bank冲突，性能死活上不去。后来把数据布局从行主序改成列主序，再配合padding（在每行末尾加几个无用元素），冲突就解了。

解决bank冲突的常用套路：

• padding：在共享内存数组的每一行末尾加几个元素，错开bank映射
• 改变访问模式：让连续线程访问连续地址
• 使用向量化类型：比如float4，一次读4个float，减少访问次数

常量内存与纹理内存

这两种内存比较“专”，用对场景效果拔群，用错场景就是浪费。

常量内存

常量内存只有64KB，但它的厉害之处在于：当一个warp里所有线程访问同一个常量地址时，硬件会广播，只发一次请求。我习惯把那些“所有线程都用到的参数”放常量内存，比如卷积核的权重、矩阵乘法的维度等。

// 声明常量内存
__constant__ float c_filter[256];

// 主机端赋值
cudaMemcpyToSymbol(c_filter, host_filter, sizeof(float) * 256);

// kernel中使用
__global__ void conv_kernel(float *input, float *output) {
    float val = c_filter[threadIdx.x];  // 所有线程读同一个地址，广播
}

纹理内存

纹理内存最初是为图形学设计的，但它的一些特性在通用计算里也很有用：

• 空间局部性优化：纹理缓存对2D空间访问做了优化，适合图像、矩阵等二维数据
• 硬件插值：可以自动做线性插值，适合需要浮点坐标访问的场景
• 边界处理：支持多种寻址模式（clamp、wrap等），省去手动边界判断

我记得有一次做图像金字塔，需要频繁访问不同尺度的图像数据。用纹理内存后，缓存命中率大幅提升，性能提升了30%以上。

知识体系总览

下面这张图是我自己总结的内存访问优化决策流程，你可以照着这个思路来排查性能问题：

说白了，内存访问优化就是“让数据离计算更近，让访问模式更规整”。全局内存合并访问是基本功，共享内存bank冲突是进阶技巧，常量内存和纹理内存是锦上添花的利器。

嗯，今天就聊到这儿。这些技巧我在实际项目中反复验证过，确实能带来实实在在的性能提升。你可以在自己的代码里试试看，跑个profile对比一下，效果立竿见影。