📚 课程目录第3章 CUDA内存管理与性能优化

3. CUDA内存管理与性能优化

在GPU上写kernel，性能瓶颈往往不在算术指令，而在数据搬运。CUDA把存储资源分成若干层次，每一层的容量、带宽、延迟和访问规则都不同。理解这些差异，是写出高性能kernel的前提。本节以矩阵转置为贯穿案例，说明全局内存的合并访问、共享内存的bank冲突、常量内存的广播机制以及寄存器的使用约束，并给出可复现的优化路径。

一、四类存储资源的定位

寄存器是线程私有的最快存储，容量极小（每个SM通常65536个32位寄存器），由编译器分配，生命周期与线程一致。共享内存（shared memory）是线程块内共享的片上存储，容量约48KB～164KB（依架构而定），延迟远低于全局内存，但需要显式管理，且存在bank冲突问题。常量内存（constant memory）容量64KB，位于片外但带常量缓存，同一warp内所有线程访问同一地址时可广播，速度接近寄存器；若地址分散则退化为串行。全局内存（global memory）容量最大、延迟最高（数百个周期），通过L1/L2缓存缓解，其性能关键在于访问是否合并（coalesced）。

一个常被忽略的事实是：全局内存的访问以32字节或128字节的transaction为单位。若一个warp的32个线程访问连续且对齐的地址，硬件可将其合并为少数几个transaction；若地址跨步或分散，则transaction数量成倍增加，有效带宽骤降。这正是矩阵转置案例的核心矛盾。

二、矩阵转置案例：从朴素实现到合并访问

设矩阵为N×N的float数组，按行主序存储，即元素(i,j)位于索引 i*N+j。转置要求输出out(j,i)=in(i,j)。朴素实现让每个线程处理一个元素：

__global__ void transposeNaive(float *out, const float *in, int N) {
    int i = blockIdx.y * blockDim.y + threadIdx.y;
    int j = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N && j < N)
        out[j * N + i] = in[i * N + j];
}

读取时，同一warp内线程的j连续，in[i*N+j]是连续地址，读取是合并的。但写入时，out[j*N+i]中j随线程变化，地址跨步为N，同一warp的32个线程写入地址相隔N个float，无法合并，产生32个独立transaction。结果是读快写慢，整体带宽利用率可能只有峰值的1/8到1/4。

解决方案是引入共享内存做中转：每个block先把一块tile从全局内存合并读入共享内存，再从共享内存按转置后的顺序合并写回全局内存。这样全局内存的读写都变成合并访问，代价是共享内存的bank冲突需要处理。

#define TILE 32
__global__ void transposeShared(float *out, const float *in, int N) {
    __shared__ float tile[TILE][TILE + 1];   // +1 填充避免bank冲突
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    if (x < N && y < N)
        tile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();
    x = blockIdx.y * TILE + threadIdx.x;
    y = blockIdx.x * TILE + threadIdx.y;
    if (x < N && y < N)
        out[y * N + x] = tile[threadIdx.x][threadIdx.y];
}

三、bank冲突的成因与消除

共享内存被划分为32个bank，每个bank宽4字节，地址按4字节字映射到bank：bank = (地址/4) % 32。若同一warp内多个线程访问同一bank的不同地址，硬件必须串行处理，称为bank冲突。在转置的写回阶段，线程读取tile[threadIdx.x][threadIdx.y]，当threadIdx.x变化时，地址跨度为TILE+1个float。若没有+1填充，跨度为32，恰好使所有线程落到同一bank，形成32路冲突，共享内存访问退化为32次串行。加上+1后，跨度变为33，33与32互质，线程访问的bank依次错开，冲突降为1路，即无冲突。

这里的关键参数是填充量。填充1列即可打破32的整除关系；填充过多会浪费共享内存容量，降低occupancy。实践中TILE取32、填充1是经典组合，也可用TILE=16、填充1以适配更小的共享内存预算。

四、常量内存与寄存器

常量内存适合存放kernel执行期间不变的小数据，例如转置中的矩阵维度N、缩放系数、滤波器权重。声明为__constant__ float c_weight[256];，用cudaMemcpyToSymbol从主机拷贝。当warp内所有线程读取同一地址时，常量缓存一次读取即可广播给全部线程，延迟接近寄存器；若各线程读取不同地址，则退化为串行，反而比全局内存更慢。因此常量内存的使用前提是“同址访问”。

寄存器方面，编译器为每个线程分配寄存器，数量由kernel复杂度决定。可用nvcc --ptxas-options=-v查看每个kernel的寄存器用量。寄存器越多，单个SM能同时驻留的线程越少，occupancy下降；但寄存器太少又会导致溢出到local memory（实际落在全局内存），性能更差。可用__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)提示编译器控制寄存器上限。转置kernel本身简单，寄存器压力小，通常不是瓶颈，但在融合了计算或多次累加的kernel中需要重点关注。

五、优化路径与验证方法

实际调优建议按以下顺序推进：先用朴素实现跑通并测量基线；再引入共享内存tile，观察带宽提升；然后检查bank冲突，加入填充；最后根据occupancy和寄存器报告微调block尺寸。验证时可用Nsight Compute查看l1tex__data_bank_conflicts_pipe_lsu_mem_shared指标判断bank冲突，用dram__throughput.avg.pct_of_peak_sustained_elapsed判断全局内存带宽利用率。若带宽利用率接近峰值且bank冲突接近零，说明内存路径已优化到位。

转置案例的通用启示：当发现某个方向的访问不合并时，用共享内存做一次“重排”，把不合并的访问转化为片上访问，再以合并方式写回。这一模式在卷积、矩阵乘、直方图等kernel中反复出现。

共享内存填充不是越多越好。填充列会增大共享内存占用，可能降低occupancy。TILE=32、填充1是常见起点，但应根据实际架构的bank数量和共享内存容量做实测调整。

小结：全局内存看合并，共享内存看冲突，常量内存看同址，寄存器看占用。四者并非孤立，而是共同决定kernel的实际带宽与延迟。矩阵转置虽小，却完整覆盖了这四条规则，是理解CUDA内存层次的最佳入口。