4、CUDA内存管理：全局内存、共享内存、常量内存、寄存器、内存访问优化

说到CUDA编程，内存管理绝对是绕不开的核心话题。我刚开始接触GPU计算时，总觉得把数据扔到显存里就能自动跑得快。结果呢？第一次跑金融风险模型，性能还不如CPU。后来我才明白——内存访问模式，才是决定性能的关键。

GPU的内存层次结构，说白了就是一个「速度与容量」的博弈。离计算单元越近，速度越快，但容量越小。咱们得学会在这个层次结构中，把数据放到最合适的位置。

4.1 全局内存：GPU的「大仓库」

全局内存是GPU里容量最大的存储空间，所有线程都能读写。但它的延迟也最高——大概400-800个时钟周期。你想想看，CPU访问L1缓存只要几个周期，这差距有多大。

我在做期权定价模型时，一开始把整个波动率曲面都放在全局内存里。结果每个线程都要随机访问不同位置的波动率数据，性能惨不忍睹。后来我做了两件事：合并访问 + 数据预取，性能直接翻了3倍。

// ❌ 糟糕的访问模式：按列访问
for (int i = 0; i < N; i++) {
    float val = data[threadIdx.x + i * blockDim.x];  // 跨步访问
}

// ✅ 好的访问模式：按行访问
for (int i = 0; i < N; i++) {
    float val = data[threadIdx.x + i];  // 连续访问
}

4.2 共享内存：线程块的「小团队缓存」

共享内存是片上的SRAM，延迟只有几十个周期。它在一个线程块内所有线程共享。说白了，这就是你手动控制的L1缓存。

我记得有一次做矩阵转置，全局内存版本带宽利用率只有15%。加上共享内存做分块转置后，利用率直接飙到80%以上。为什么？因为共享内存允许我们把数据先搬到片上，再按转置后的顺序写回全局内存。

// 共享内存声明
__shared__ float shared_data[256];

// 典型用法：分块矩阵乘法
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE_SIZE; t++) {
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

4.3 常量内存：只读数据的「广播站」

常量内存只有64KB，但它的厉害之处在于——当一个warp里的所有线程访问同一个常量地址时，硬件会把这个值广播给所有线程，只花一个周期！

我在做蒙特卡洛模拟时，把随机数生成器的参数放在常量内存里。因为所有线程都用同一组参数，广播机制让访问几乎零开销。但要注意，如果warp里的线程访问不同的常量地址，那就退化成串行访问了。

// 常量内存声明
__constant__ float params[16];

// 主机端赋值
cudaMemcpyToSymbol(params, host_params, sizeof(float) * 16);

4.4 寄存器：最快的存储，但有限

寄存器是每个线程私有的，访问延迟为0。但每个线程能用的寄存器数量有限——通常是255个（取决于计算能力）。

我曾经犯过一个错误：在一个核函数里声明了太多局部变量，导致寄存器溢出（register spilling）。变量被挤到局部内存里，而局部内存实际上在全局内存上！性能直接掉了60%。

4.5 内存访问优化实战策略

好了，理论说完了。咱们来点实际的。我在金融量化团队里总结了一套「内存优化三步法」：

1. 分析访问模式：用NVIDIA Nsight Compute工具看L1/L2缓存命中率、全局内存负载效率。
2. 选择合适的内存类型：只读数据放常量内存或纹理内存；频繁重用的数据放共享内存。
3. 调整数据布局：结构体数组（SoA） vs 数组结构体（AoS）——GPU更适合SoA。

内存类型	延迟	容量	作用域	典型用途
寄存器	0周期	~256个/线程	线程私有	临时变量、循环计数器
共享内存	~30周期	48KB/块	线程块内	数据重用、规约
常量内存	~1周期(广播)	64KB	全局只读	模型参数、查找表
全局内存	~400周期	GB级别	全局读写	主要数据存储

嗯，最后说一句。内存优化没有银弹。你得根据具体算法、数据规模、GPU型号来调整。我建议你养成一个习惯：每次写完核函数，先用profiler跑一遍，看看内存瓶颈在哪。别像我刚开始那样，闷头写代码，结果优化了半天才发现方向错了。

---