第三章：CUDA内存模型——全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

说到CUDA编程，内存模型是绕不开的核心话题。我刚开始接触CUDA时，总觉得把数据扔到GPU上就能自动跑得快。结果呢？性能惨不忍睹。后来我才明白，不理解内存模型，就像开车不看仪表盘——能跑，但迟早要出事。

这一章，咱们把CUDA的六种内存掰开揉碎了讲。你想想看，GPU上有成百上千个线程同时干活，数据怎么存、怎么取，直接决定了你的程序是“飞”还是“爬”。

3.1 全局内存——最常用，也最慢

全局内存，说白了就是GPU上的“大仓库”。所有线程都能访问，容量最大（通常几个GB到几十个GB），但速度也是最慢的。

为什么慢？因为全局内存在显存芯片上，离计算单元远。访问一次全局内存，大概需要几百个时钟周期。我做过一个项目，一开始把所有的中间结果都放全局内存，结果性能瓶颈全卡在内存访问上。

这里有个重要的优化技巧——合并访问。什么意思呢？就是让相邻的线程访问相邻的内存地址。举个例子：

// 好的写法：合并访问
__global__ void kernel_good(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = data[idx] * 2.0f;  // 线程0访问data[0]，线程1访问data[1]...
}

// 坏的写法：非合并访问
__global__ void kernel_bad(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx * 32] = data[idx * 32] * 2.0f;  // 线程0访问data[0]，线程1访问data[32]...
}

我曾经在一个图像处理项目里，就因为没注意合并访问，性能差了将近10倍。嗯，这坑我踩过，你们就别再踩了。

3.2 共享内存——线程块的“小团队仓库”

共享内存是CUDA里最“香”的内存之一。它在芯片上，离计算单元非常近，延迟只有几个时钟周期。但容量很小，每个线程块通常只有几十KB。

共享内存的作用是什么？让同一个线程块里的线程可以高效地交换数据。我习惯用它来做数据复用和规约操作。

来看一个矩阵转置的例子，共享内存的经典用法：

__global__ void transpose_shared(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 从共享内存转置写出
    x = blockIdx.y * TILE_SIZE + threadIdx.x;
    y = blockIdx.x * TILE_SIZE + threadIdx.y;
    out[y * width + x] = tile[threadIdx.x][threadIdx.y];
}

这里要注意bank conflict的问题。共享内存被分成32个bank，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。我遇到过最夸张的一次，bank conflict让性能直接腰斩。

3.3 寄存器——最快，但最稀缺

寄存器是每个线程私有的“口袋”，速度最快——零延迟。但数量极其有限，每个线程通常只有几十到几百个。

你写的局部变量，只要编译器觉得合适，就会放到寄存器里。但寄存器溢出（spilling）是个大问题。当寄存器不够用时，数据会被挤到本地内存（后面会讲），速度瞬间掉下来。

怎么查看寄存器使用量？用 --ptxas-options=-v 编译选项，或者用 cudaOccupancyMaxPotentialBlockSize 函数计算。

3.4 本地内存——寄存器的“替补”

本地内存听起来像“本地”的，其实它存在全局内存里，只是每个线程私有的。说白了，就是寄存器不够用时，编译器自动把溢出的变量放到这里。

本地内存的延迟和全局内存一样高，所以尽量避免使用。怎么避免？少声明大数组、减少每个线程的变量数量。

我记得有一次调试一个复杂的算法，发现性能怎么都上不去。用 nvcc --ptxas-options=-v 一看，本地内存用了不少。把一些大数组改成动态分配的全局内存后，性能提升明显。

3.5 常量内存——只读，但带缓存

常量内存是全局内存的一种特殊形式，但有两个特点：只读，并且有专门的缓存（常量缓存）。

常量内存适合存储那些所有线程都要用到的、不会改变的数据。比如物理常数、查找表等。

__constant__ float filter[256];

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

// 主机端设置常量内存
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

常量内存的访问模式很特别：如果同一个warp里的所有线程访问同一个地址，速度极快（一个周期）。但如果每个线程访问不同的地址，就会串行化，性能暴跌。

3.6 纹理内存——为空间局部性而生

纹理内存最初是为图形学设计的，但在通用计算里也很有用。它也是只读的，但缓存策略和常量内存不同——纹理缓存针对2D空间局部性做了优化。

什么意思呢？如果你的程序访问数据时，相邻的线程访问相邻的像素（或网格点），纹理内存的缓存命中率会很高。

我做过一个流体模拟的项目，用纹理内存来存储速度场数据。因为每次计算都要访问相邻格点的值，纹理缓存的优势就体现出来了，比直接用全局内存快了30%左右。

// 声明纹理引用
texture<float, 2, cudaReadModeElementType> tex;

__global__ void kernel_texture(float *out, 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) {
        // 使用tex2D读取纹理内存
        out[y * width + x] = tex2D(tex, x, y) * 2.0f;
    }
}

3.7 内存模型总结与选择策略

好了，六种内存都讲完了。咱们来做个对比：

内存类型	位置	访问范围	速度	容量	生命周期
寄存器	芯片内	单个线程	最快	几十~几百个	线程内
共享内存	芯片内	线程块内	快	几十KB	线程块内
本地内存	显存	单个线程	慢	取决于显存	线程内
全局内存	显存	所有线程+主机	最慢	GB级别	主机控制
常量内存	显存（带缓存）	所有线程（只读）	快（同地址）	64KB	主机控制
纹理内存	显存（带缓存）	所有线程（只读）	快（空间局部性）	取决于显存	主机控制

我个人选择内存的策略很简单：

1. 能放寄存器的，绝不放共享内存——每个线程私有的临时变量，尽量用局部变量
2. 需要线程间通信的，用共享内存——记得加 __syncthreads()
3. 只读且所有线程读同一个值的，用常量内存——比如滤波器系数
4. 有空间局部性的只读数据，用纹理内存——图像、网格数据
5. 其他情况，用全局内存——但一定要保证合并访问

最后说一句：内存优化是CUDA性能调优的重中之重。我见过太多人把精力花在算法复杂度上，结果内存访问模式一塌糊涂，性能还不如CPU。你想想看，GPU的计算能力再强，如果数据喂不进去，一切都是白搭。

下一章，咱们聊聊线程束（warp）的调度和执行模型。这东西理解了，你才能真正明白为什么有些代码跑得快，有些跑得慢。