4. CUDA内存优化：全局内存、共享内存、寄存器、常量内存的使用与优化策略

核心原理：CUDA内存层次结构与带宽瓶颈

CUDA设备内存呈现典型的金字塔结构：寄存器位于最顶层，访问延迟约1个时钟周期；共享内存和L1缓存次之，延迟约20~30周期；L2缓存约200周期；全局内存位于最底层，延迟高达400~800周期。推理管线中，卷积、矩阵乘、归一化等算子对带宽的消耗远超计算量，因此内存访问模式往往决定内核性能上限。理解各内存的物理特性与作用域，是优化推理内核的第一步。

上图展示了推理内核中各类内存的访问延迟与容量关系。优化策略的核心思路是：将高频复用的数据尽量放置在靠近计算单元的内存层级，同时保证全局内存访问满足合并访问条件（即同一 warp 的 32 个线程访问连续地址）。

方法流程：四步内存优化法

针对推理管线中的典型算子（如 GEMM、LayerNorm、Softmax），推荐以下四步优化流程：

分析访存模式：使用 nvprof --metrics gld_throughput 或 ncu --section MemoryWorkloadAnalysis 统计全局内存吞吐量、缓存命中率、共享内存 bank 冲突次数。
数据分块与共享内存复用：将大矩阵切分为 tile（如 32×32 或 64×64），通过 __shared__ 声明共享内存缓冲区，由线程块协作加载并复用数据块。
寄存器变量显式化：将循环中反复使用的全局内存读取结果存入局部变量（寄存器），减少重复访存；同时注意控制每线程寄存器数量（__launch_bounds__ 可辅助限制）。
常量内存与只读缓存利用：对于权重、缩放因子、查找表等只读数据，使用 __constant__ 声明；若数据量超过 64KB 或访问模式不规则，改用 __ldg() 只读路径。

// 示例：共享内存 tile 加载（矩阵乘核心片段）
__global__ void sgemm_tiled(const float* A, const float* B, float* C, int M, int N, int K) {
    __shared__ float As[32][32];
    __shared__ float Bs[32][32];
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = by * 32 + ty, col = bx * 32 + tx;
    float sum = 0.0f;

    for (int k0 = 0; k0 < K; k0 += 32) {
        // 协作加载 tile 到共享内存（合并访问）
        As[ty][tx] = A[(row) * K + (k0 + tx)];
        Bs[ty][tx] = B[(k0 + ty) * N + col];
        __syncthreads();

        for (int k = 0; k < 32; ++k)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

案例：TensorRT 中 GEMM 内核的共享内存与寄存器协同优化

以 TensorRT 的 Ampere 架构 GEMM 内核为例，其采用 128×128 线程块、每线程计算 8×8 输出子块。关键优化点包括：

共享内存双缓冲：使用两套共享内存缓冲区交替加载下一 tile，隐藏全局内存延迟。
寄存器累加器：每个线程持有 64 个浮点寄存器作为累加器，避免中间结果写回全局内存。
向量化加载：使用 float4 类型实现 16 字节对齐访问，提高内存事务效率。
Bank 冲突消除：共享内存数组填充 padding（如 As[32][32+1]），使同一 warp 访问不同 bank。

优化手段	访存变化	性能提升（相对朴素版本）
仅合并访问	全局内存事务减少 4×	1.8×
+共享内存 tile	全局内存访问减少 90%	3.2×
+双缓冲	隐藏加载延迟	4.5×
+寄存器累加 & 向量化	减少指令数 & 提高带宽利用率	6.1×

误区与复盘：常见内存优化陷阱

误区1：盲目增加共享内存使用量。共享内存过多会降低线程块占用率（occupancy），导致无法隐藏延迟。应通过 cudaOccupancyMaxPotentialBlockSize 验证占用率与性能的平衡点。

误区2：忽略 bank 冲突。共享内存虽快，但若同一 warp 内多个线程访问同一 bank 的不同地址，会被串行化。例如 shared[32][32] 按列访问时，shared[ty][tx] 与 shared[ty+1][tx] 会冲突。解决方案是声明为 shared[32][33]。

误区3：常量内存使用不当。常量内存适合所有线程访问同一地址（广播）的场景；若线程间访问地址不同，常量缓存会退化为串行访问，此时应改用全局内存 + __ldg()。

复盘要点：在一次 LayerNorm 内核优化中，最初使用全局内存直接读取均值与方差，延迟高达 800 周期。改为 __constant__ 存储后，由于所有线程读取相同地址，广播机制将延迟降至约 20 周期，整体内核提速 2.3 倍。但后续将权重也放入常量内存时，因权重按通道索引访问（地址不同），性能反而下降 15%。最终将权重移至全局内存并使用 __ldg()，才获得正确优化效果。

内存优化需要结合具体算子的访存特征与硬件架构，通过 profiling 数据驱动决策。建议在每次修改后对比 gld_throughput、shared_load_throughput 和 achieved_occupancy 三项指标，确保优化方向正确。