3、CUDA显存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

好，咱们今天来聊聊CUDA的显存模型。说实话，这块内容我当年刚接触时也绕了好一阵子。你想想看，一个GPU里居然有六种不同类型的内存，每种还各有脾气。搞懂了它们，你的程序性能就能起飞；搞不懂，那就等着被带宽瓶颈卡死吧。

我个人习惯把CUDA的内存体系比作一个「数据仓库」。全局内存是那个最大的仓库，啥都能放，但离得远，取货慢。共享内存是工位旁边的小货架，手一伸就能够到。寄存器呢，就是工人手里的工具，最快，但数量有限。嗯，咱们一个一个来看。

核心观点：CUDA程序优化的本质，就是让数据尽量待在离计算单元近的地方。说白了，就是减少对全局内存的访问，多用共享内存和寄存器。

3.1 全局内存（Global Memory）

全局内存是GPU里最大的存储空间，所有线程都能读写。它的容量通常是几GB到几十GB。但代价是什么？延迟高，大概400-800个时钟周期。

我在项目中遇到过一个问题：一个矩阵乘法程序，直接拿全局内存算，结果带宽利用率只有20%左右。后来一查，发现是访问模式不对。全局内存有个硬性要求——合并访问。什么意思？就是相邻的线程要访问相邻的地址。

// 坏的访问模式：线程0访问A[0]，线程1访问A[1000]
// 这会导致多次内存事务，带宽浪费严重
int idx = threadIdx.x;
float val = data[idx * 1024];  // 跨度太大

// 好的访问模式：线程0访问A[0]，线程1访问A[1]
// 一次内存事务就能搞定
int idx = threadIdx.x;
float val = data[idx];  // 连续访问

避坑指南：我曾经在项目里因为没注意全局内存的对齐要求，导致性能直接腰斩。记住：全局内存的访问地址最好是32字节、64字节或128字节对齐的。否则，硬件会帮你「补课」，多读很多没用的数据。

3.2 共享内存（Shared Memory）

共享内存是块上板载的SRAM，位于每个SM内部。它的延迟只有几十个时钟周期，比全局内存快了一个数量级。但容量很小，每个SM通常只有几十KB到一百多KB。

共享内存最大的特点是：同一个线程块内的所有线程都能访问。这就给了我们一个机会——把全局内存里的数据先搬到共享内存里，然后让线程们反复使用。

你想想看，如果每个线程都去全局内存读数据，那带宽很快就撑爆了。但要是先搬一块数据到共享内存，大家轮流用，那效率就高多了。

// 典型的共享内存使用模式：矩阵分块
__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 bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float sum = 0.0f;
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        As[ty][tx] = A[by * TILE_SIZE + ty][t * TILE_SIZE + tx];
        Bs[ty][tx] = B[t * TILE_SIZE + ty][bx * TILE_SIZE + tx];
        __syncthreads();  // 同步，确保所有线程都加载完毕
        
        // 计算子矩阵乘积
        for (int k = 0; k < TILE_SIZE; k++)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();  // 同步，防止覆盖数据
    }
    C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}

个人经验：共享内存用得好不好，直接决定了你的kernel能跑多快。我建议你养成一个习惯：每次写kernel前，先想想哪些数据可以被复用。如果能用共享内存缓存起来，那就别犹豫。

3.3 寄存器（Register）

寄存器是GPU里最快的内存，没有之一。它的访问延迟是0个时钟周期——也就是说，一个时钟周期内就能拿到数据。但寄存器数量极其有限，每个线程最多能用255个（具体取决于GPU架构）。

为什么寄存器这么重要？因为编译器会尽量把局部变量分配到寄存器里。如果寄存器不够用，就会发生「寄存器溢出」，数据被挤到本地内存里，性能瞬间暴跌。

关键指标：每个SM的寄存器总数是固定的。比如在A100上，每个SM有65536个寄存器。如果你每个线程用32个寄存器，那每个SM就能同时驻留65536/32=2048个线程。如果每个线程用64个寄存器，那就只能驻留1024个线程。线程少了，隐藏延迟的能力就弱了。

我记得有一次优化一个深度学习算子，发现性能上不去。用nvcc的--ptxas-options=-v选项一查，好家伙，每个线程用了80多个寄存器，溢出严重。后来我手动减少了几个局部变量，把寄存器用量压到48个，性能直接提升了30%。

3.4 本地内存（Local Memory）

本地内存听起来像是「本地」的，其实它就是个幌子。本地内存的数据实际上存储在全局内存里，只不过每个线程有自己的私有空间。它的延迟和全局内存一样高。

什么时候会用本地内存？主要有两种情况：

寄存器溢出：编译器发现寄存器不够用，就把一些变量挤到本地内存里
数组索引不可预测：比如你声明了一个局部数组float arr[10]，但访问它的索引是动态计算的，编译器没法把它分配到寄存器里

避坑指南：我曾经写过一个kernel，里面声明了一个int local_array[32]。我以为它在寄存器里，结果一查PTX代码，发现被放到了本地内存里。性能惨不忍睹。后来我把数组大小改小，或者用__launch_bounds__限制线程数，才解决了问题。

3.5 常量内存（Constant Memory）

常量内存是全局内存的一个特殊区域，容量只有64KB。它有两个特点：一是只读，二是有专门的缓存（常量缓存）。

常量内存最适合存放那些所有线程都读取相同值的数据。比如神经网络的权重、查找表、配置参数等。为什么？因为常量缓存有广播机制——当一个warp里的所有线程都访问同一个地址时，硬件只需要一次内存读取，然后广播给所有线程。

// 常量内存声明
__constant__ float weights[1024];

// 主机端初始化
cudaMemcpyToSymbol(weights, host_weights, sizeof(float) * 1024);

// 设备端使用
__global__ void compute(float *input, float *output) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = input[idx];
    // 所有线程都读取同一个weights[0]
    output[idx] = val * weights[0];  // 常量缓存广播，一次读取搞定
}

个人建议：常量内存虽然只有64KB，但用好了效果惊人。我做过一个图像处理项目，把滤波器的系数放在常量内存里，比放在全局内存里快了将近2倍。原因就是warp内的线程都在读同一个系数，常量缓存的广播机制发挥了作用。

3.6 纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的，但在通用计算里也有用武之地。它也是只读的，有专门的纹理缓存，而且支持一些特殊功能：

硬件插值：对于非整数坐标，硬件会自动做线性插值
边界处理：可以设置边界模式（钳位、包裹等）
数据对齐不敏感：纹理内存对访问模式的要求比全局内存宽松

说实话，纹理内存在深度学习里用得不多，但在图像处理、科学计算（比如有限差分法）里还是很常见的。我记得有一次做流体模拟，用纹理内存来存储速度场，因为需要频繁地做双线性插值，纹理内存的硬件插值功能帮了大忙。

3.7 总结：一张图看懂CUDA内存层次

下面这张SVG图展示了CUDA的内存层次结构。你可以看到，从寄存器到全局内存，容量越来越大，但速度越来越慢。优化的核心就是：把频繁访问的数据放在上层，把不常访问的数据放在下层。

好了，CUDA的六种内存模型就讲完了。每种内存都有自己的脾气和用途。全局内存是主力，共享内存是加速器，寄存器是珍宝，本地内存是陷阱，常量内存和纹理内存是特种兵。搞懂了它们，你写kernel的时候就知道该怎么分配数据了。

想要更多资料，联系可可

📚 课程目录

3、CUDA显存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

3.1 全局内存（Global Memory）

3.2 共享内存（Shared Memory）

3.3 寄存器（Register）

3.4 本地内存（Local Memory）

3.5 常量内存（Constant Memory）

3.6 纹理内存（Texture Memory）

3.7 总结：一张图看懂CUDA内存层次