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

各位同学，咱们今天来聊聊CUDA的内存模型。说实话，搞GPU编程，如果你不懂内存，那基本等于开车不看仪表盘——能跑，但迟早要出事。我刚开始学CUDA那会儿，就吃过不少内存的亏，今天把这些经验掰开了揉碎了讲给你听。

4.1 为什么内存模型这么重要？

你想想看，CPU和GPU最大的区别在哪？CPU有强大的缓存和分支预测，而GPU靠的是海量的线程并行。但并行计算有个核心问题：数据在哪？怎么搬？搬得快不快？

CUDA的内存模型，说白了就是一套规则，告诉你数据该放哪、怎么放、谁可以访问。我个人的习惯是，写任何kernel之前，先画一张内存访问图，把数据流理清楚。这一步做好了，后面调优能省一半时间。

4.2 全局内存（Global Memory）

全局内存是GPU里最大的存储空间，所有线程都能访问。但大是有代价的——它最慢。我记得有一次项目里，一个kernel跑得很慢，我查了半天，发现是频繁读写全局内存导致的。后来改成共享内存，速度直接翻了5倍。

全局内存的特点：

• 容量大：通常几个GB到几十个GB
• 延迟高：大约400-800个时钟周期
• 所有线程（包括CPU）都能访问
• 使用 cudaMalloc 分配

这里有个避坑指南：我曾经以为全局内存随便读写就行，结果发现合并访问（coalesced access）是个大坑。如果线程访问的地址不连续，带宽利用率会暴跌。比如这样：

// 好的写法：合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx];  // 连续地址

// 坏的写法：非合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx * 2];  // 地址跳跃，带宽浪费

4.3 共享内存（Shared Memory）

共享内存是我个人最喜欢的存储类型。它位于每个SM（流多处理器）内部，同一个block内的线程可以共享。延迟只有几十个时钟周期，比全局内存快一个数量级。

但共享内存有两个限制：

• 容量小：通常只有几十KB到上百KB
• 生命周期短：只在block存在期间有效

共享内存最常见的用法是做数据缓存。比如矩阵乘法，把子矩阵加载到共享内存里，然后反复使用，避免重复访问全局内存。

__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    for (int t = 0; t < N / BLOCK_SIZE; t++) {
        // 加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + t * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        // 计算
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

4.4 寄存器（Register）

寄存器是GPU里最快的存储，没有之一。每个线程独享，延迟几乎为0。但数量极其有限——每个SM的寄存器总数是固定的，比如1024个或65536个（取决于架构）。

你想想看，如果每个线程用32个寄存器，那一个SM最多能同时运行2048个线程（65536/32）。但如果每个线程用64个寄存器，就只能运行1024个线程。这就是为什么寄存器压力会影响并行度。

我建议你在写kernel时，尽量少用局部变量。比如：

// 寄存器友好
float a = data[0];
float b = data[1];
float c = a + b;

// 寄存器不友好
float arr[100];  // 如果寄存器不够，会溢出到本地内存

4.5 本地内存（Local Memory）

本地内存听起来像“本地”的，但其实它存在全局内存里。当寄存器不够用时，编译器会把一些变量“溢出”到本地内存。说白了，这是寄存器的备胎。

本地内存的特点：

• 每个线程私有
• 物理上在全局内存里，所以延迟高
• 通常用于数组、大结构体、或者寄存器溢出

怎么判断变量被放到了本地内存？用 --ptxas-options=-v 看 lmem 那一项。如果数值很大，说明你的kernel有寄存器压力。

4.6 常量内存（Constant Memory）

常量内存是只读的，所有线程都能访问。它有两个特点：

• 有缓存（constant cache），如果所有线程访问同一个地址，速度极快
• 容量小：只有64KB

常量内存最适合存那些所有线程都用到的只读数据，比如滤波器系数、查找表等。

__constant__ float filter[256];

__global__ void apply_filter(float *input, float *output, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        output[idx] = input[idx] * filter[threadIdx.x % 256];
    }
}

4.7 纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的，但用在通用计算里也有奇效。它也是只读的，有专门的缓存，而且支持硬件插值和边界处理。

纹理内存的优势：

• 对二维/三维空间局部性好的访问模式有优化
• 支持线性插值（比如图像缩放）
• 自动处理边界（比如 clamp 或 wrap）

我记得在做一个图像处理项目时，用纹理内存做双线性插值，比手动实现快了3倍。因为硬件直接帮你算了插值，省了不少指令。

texture<float, 2, cudaReadModeElementType> tex;

__global__ void texture_kernel(float *output, int width, int height) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (x < width && y < height) {
        // 纹理内存自动处理边界
        output[y * width + x] = tex2D(tex, x + 0.5f, y + 0.5f);
    }
}

4.8 总结与对比

好了，咱们把CUDA的六种内存都过了一遍。我整理了一个表格，方便你对比：

内存类型	位置	访问范围	延迟	容量	生命周期
全局内存	显存	所有线程+CPU	高（400-800周期）	大（GB级）	程序运行期间
共享内存	SM内部	同一block	低（几十周期）	小（几十KB）	block运行期间
寄存器	SM内部	单个线程	极低（1周期）	极小（几十个）	线程运行期间
本地内存	显存	单个线程	高（同全局内存）	取决于寄存器溢出	线程运行期间
常量内存	显存（有缓存）	所有线程（只读）	低（缓存命中时）	64KB	程序运行期间
纹理内存	显存（有缓存）	所有线程（只读）	低（空间局部性好时）	大（同全局内存）	程序运行期间

最后说一句：内存模型是CUDA性能优化的基石。你想想看，一个kernel跑得慢，90%的原因都是内存访问没做好。我建议你每次写kernel时，都问自己三个问题：

1. 数据放哪最合适？
2. 访问模式是否连续？
3. 有没有办法减少全局内存访问？

把这三点想清楚，你的CUDA程序性能不会差。好了，这一章就到这里，下一章咱们聊聊线程束（Warp）与线程调度，这可是GPU并行计算的精髓。