4. CUDA内存管理：全局内存、共享内存、常量内存、寄存器

CUDA程序性能的瓶颈往往不在计算单元，而在内存访问。GPU拥有庞大的并行线程规模，但若内存访问模式低效，数万个线程会同时阻塞在访存路径上，导致SM（流式多处理器）空转。理解CUDA内存层次结构，是编写高性能内核的必修课。本章聚焦四种最常用的内存：全局内存、共享内存、常量内存和寄存器，剖析其特性、适用场景与典型误区。

4.1 核心原理：CUDA内存层次结构

CUDA内存模型从物理上分为板载显存（全局内存、常量内存、纹理内存）和芯片内存储（寄存器、共享内存、L1/L2缓存）。从编程视角看，它们具有不同的作用域、生命周期和访问速度。下图展示了线程、线程块与各内存的访问关系：

从上图可见：寄存器是每个线程私有的“工作台”，速度最快但容量极小；共享内存是线程块内所有线程共享的“高速暂存区”，用于线程间协作和数据复用；全局内存是GPU显存，容量大但延迟高，所有线程均可访问；常量内存是全局内存的“只读缓存变体”，适合广播式访问。理解这些差异，是优化内存访问的第一步。

4.2 方法流程：内存分配与访问模式

在CUDA中，内存管理遵循“分配→拷贝→内核访问→拷贝回→释放”的标准流程。以下代码展示了全局内存的典型使用方式：

// 1. 分配主机内存并初始化
float *h_data = (float*)malloc(N * sizeof(float));
for (int i = 0; i < N; i++) h_data[i] = i * 1.0f;

// 2. 分配设备全局内存
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 3. 主机→设备拷贝
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 4. 内核访问（合并访问示例）
__global__ void scale_kernel(float *data, float factor, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) data[idx] *= factor;  // 相邻线程访问相邻地址
}

// 5. 设备→主机拷贝
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// 6. 释放资源
cudaFree(d_data);
free(h_data);

共享内存的使用则需在内核内部声明，并通过__syncthreads()同步线程。典型模式是“加载→同步→计算→写回”：

__global__ void tile_reduce(const float *in, float *out, int n) {
    __shared__ float tile[256];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;

    // 加载到共享内存
    tile[tid] = (idx < n) ? in[idx] : 0.0f;
    __syncthreads();

    // 块内归约（树形加法）
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) tile[tid] += tile[tid + stride];
        __syncthreads();
    }

    // 写回结果
    if (tid == 0) out[blockIdx.x] = tile[0];
}

常量内存的声明方式特殊，需在文件作用域使用__constant__修饰符，并通过cudaMemcpyToSymbol进行拷贝：

__constant__ float c_coeff[4];

// 主机端
float h_coeff[4] = {1.0f, 2.0f, 3.0f, 4.0f};
cudaMemcpyToSymbol(c_coeff, h_coeff, 4 * sizeof(float));

// 内核中直接访问
__global__ void apply_coeff(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) data[idx] *= c_coeff[idx % 4];
}

4.3 案例/场景：矩阵乘法中的内存优化

以经典的矩阵乘法C = A × B为例，展示不同内存策略的性能差异。朴素实现中，每个线程直接读取全局内存中的A和B，会产生大量重复访问。以256×256矩阵为例，每个输出元素需要读取A的一行（256个元素）和B的一列（256个元素），全局内存访问次数高达256³ = 16,777,216次。

采用共享内存分块（Tiling）策略后，每个线程块负责计算一个16×16的输出子块。先将A和B的对应子块加载到共享内存，再进行计算。这样，每个元素从全局内存只读取一次，共享内存访问次数大幅增加，而全局内存访问量降低到原来的1/16。实测在NVIDIA A100上，分块版本比朴素版本快约8~12倍。

策略	全局内存访问次数	共享内存使用	相对性能
朴素实现	16.7M次	0	1.0×
16×16分块	1.05M次	2×16×16×4B = 2KB	8~12×
32×32分块+向量化	0.26M次	2×32×32×4B = 8KB	15~20×

常量内存在此场景中适合存储固定系数矩阵（如卷积核权重），当所有线程访问同一地址时，常量缓存能提供接近寄存器的速度。但需注意：若线程访问不同地址，常量内存会串行化，性能反而下降。

4.4 误区/复盘：常见内存管理陷阱

误区1：过度使用共享内存。共享内存容量有限（通常几十KB），若每个线程块使用过多，会降低SM中同时驻留的线程块数量，导致占用率下降。应通过cudaOccupancyMaxPotentialBlockSize计算最优块大小。

误区2：忽略合并访问。全局内存访问必须满足“相邻线程访问相邻地址”才能获得最大带宽。若线程访问跨步地址（如按列访问二维数组），带宽会降低数倍。应使用行主序存储并让线程索引与内存地址线性对应。

误区3：寄存器溢出。当内核中局部变量过多，超出寄存器可用数量时，编译器会将变量“溢出”到局部内存（实际在全局内存中），导致性能骤降。可通过--maxrregcount编译选项限制寄存器使用，或重构代码减少变量。

误区4：常量内存分支发散。常量内存适合“所有线程读同一地址”的广播场景。若线程根据索引访问不同常量地址，硬件会串行处理每个不同地址的访问，造成严重性能损失。此时应改用全局内存或共享内存。

复盘建议：在优化内存时，应遵循以下排查顺序——先用nvprof或Nsight Compute分析内存吞吐量，确认瓶颈是带宽受限还是延迟受限；再检查全局内存访问是否合并；然后评估共享内存使用是否合理；最后检查寄存器溢出情况。实际项目中，矩阵转置、归约求和、卷积计算等场景的内存优化策略各不相同，需要结合具体算法特征选择合适的内存类型。

掌握这四种内存的特性与配合方式，是CUDA性能优化的核心能力。下一层次的学习中，你将接触到更高级的异步拷贝、内存池和统一虚拟内存等特性，但所有高级功能都建立在对本章基础内存模型深刻理解之上。