第四讲：内存优化策略

GPU编程里，内存访问模式常常是性能瓶颈。我自己在做单细胞数据处理时就吃过亏——一个挺简单的核函数跑得比预期慢了好几倍，后来一查，全局内存访问没对齐。你想想看，如果每个线程都东一榔头西一棒子地去读显存，那带宽不就全浪费了吗？

全局内存合并访问

全局内存是显存里最大的区域，但延迟也高。GPU每次访存以32字节为单位，一个warp里的32个线程必须连续地读相邻地址，才能触发合并访问。说白了，就是线程ID和地址偏移要线性挂钩。

我在优化基因表达矩阵计算时遇到过典型反面案例：矩阵按列存储，但线程按行索引，结果每个线程访问的地址跨了很大步长，变成非合并。后来改成按行存储，带宽利用率从不到20%直接飙到85%以上。嗯，这个教训印象很深。

下面看一个合并访问的典型写法：

// 合并访问：线程tid访问data[tid]
__global__ void coalesced_kernel(float *data, float *out, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        out[tid] = data[tid] * 2.0f;
    }
}

// 非合并访问：线程tid访问data[tid * stride]
__global__ void strided_kernel(float *data, float *out, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        out[tid] = data[tid * 8]; // stride = 8，浪费7/8的带宽
    }
}

共享内存的使用

共享内存是片上的SRAM，延迟只有全局内存的十分之一左右，但容量很小，通常一个block也就几十KB。它的典型用途是数据复用和规约操作。单细胞分析里计算基因间相关性的时候，我就习惯把子矩阵先读到共享内存里，再让线程反复访问，这样全局内存的访问次数能减少一个数量级。

我建议把共享内存当作手动缓存来管。不过要注意bank conflict——如果多个线程同时访问同一个bank的不同地址，会被串行化，性能会打折。解决方法是加padding或调整数据布局。

寄存器与本地内存

寄存器是每个线程私有的最快存储，但数量有限。NVIDIA的架构里每个线程最多能分到255个32位寄存器。一旦超出，编译器就会把多余的变量「溢出」到本地内存——其实就是全局内存，不过用局部地址访问。但速度会骤降。

我在优化一个单细胞聚类算法时，发现每个线程里定义了一个很大的局部数组，结果kernel占用率低得可怜。后来把部分临时变量挪到共享内存里，减少寄存器压力，性能直接提升了30%。

nvcc -Xptxas -v kernel.cu   # 编译时会打印寄存器数和shared memory使用量
ncu --metrics launch__registers_per_thread ./myapp

本地内存虽然也叫local memory，实际走的还是全局内存的路径，只不过地址被编译器隐藏了。我曾经踩过这个坑——在核函数里写了一个int arr[256]，觉得是局部数组肯定快，结果性能比预期差很多。一查PTX代码，发现全部被放到了local memory。后来改成共享内存动态分配，问题解决。

总结下来，内存优化的核心就是让数据待在离计算单元最近的地方，并且访问模式要规整。寄存器放最常用的标量，共享内存放线程块内复用的数据，全局内存则确保合并。我自己的习惯是写kernel之前先画一个数据流图，标出每个变量应该放在哪一层，避免事后调优时手忙脚乱。