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

说到CUDA编程，内存模型这块儿，我敢说它是新手最容易踩坑的地方，也是性能优化的核心。你想想看，GPU里有那么多核，数据怎么搬、怎么存、怎么读，直接决定了你的程序是飞起来还是趴窝。今天咱们就把这六种内存掰开揉碎了聊一聊。

3.1 全局内存（Global Memory）

全局内存，说白了就是GPU里最大的那个仓库。所有线程都能访问，容量最大，但速度最慢。我刚开始做CUDA时，习惯把数据一股脑全塞进全局内存，结果发现性能惨不忍睹。后来才明白，全局内存的延迟大概在400-800个时钟周期，这可不是闹着玩的。

这里有个重要的概念叫「合并访问」。什么意思呢？就是相邻的线程最好访问相邻的地址。我曾经在项目里遇到过一个坑：线程访问模式是跳跃的，结果全局内存带宽利用率只有可怜的10%。后来改成连续访问，性能直接翻了8倍。

// 好的访问模式：合并访问
__global__ void good_access(float* data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] = data[idx] * 2.0f;  // 相邻线程访问相邻地址
    }
}

// 坏的访问模式：非合并访问
__global__ void bad_access(float* data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[threadIdx.x * N + blockIdx.x] = 1.0f;  // 跨步访问
    }
}

3.2 共享内存（Shared Memory）

共享内存是我个人最喜欢的。它就在SM内部，延迟只有几个时钟周期，比全局内存快了两个数量级。但容量很小，每个SM通常只有几十KB到上百KB。

共享内存有个坑叫「bank conflict」。嗯，这里要注意，共享内存被分成了32个bank，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。我曾经在优化一个卷积算子时，就因为bank conflict浪费了整整两天时间。

// 共享内存使用示例：矩阵转置
__global__ void transpose(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;
    
    // 加载到共享内存
    if (x < width && y < width) {
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    }
    
    __syncthreads();
    
    // 从共享内存转置写出
    int new_x = blockIdx.y * TILE_SIZE + threadIdx.x;
    int new_y = blockIdx.x * TILE_SIZE + threadIdx.y;
    
    if (new_x < width && new_y < width) {
        out[new_y * width + new_x] = tile[threadIdx.x][threadIdx.y];
    }
}

3.3 寄存器（Register）

寄存器是速度最快的内存，没有之一。每个线程私有的，编译器会自动分配。但数量有限，每个SM的寄存器总数是固定的（比如Kepler架构是65536个），分给每个线程的数量取决于线程数。

怎么控制寄存器使用量？可以用 __launch_bounds__ 或者编译选项 -maxrregcount。我个人建议，在写kernel时尽量复用变量，减少不必要的临时变量。

3.4 本地内存（Local Memory）

本地内存听起来像是「本地」的，其实它物理上就在全局内存里。只是每个线程私有的，编译器在寄存器不够用时，会把一些变量「溢出」到本地内存。你想想看，这速度能快吗？

哪些情况会用到本地内存？数组索引访问、大型结构体、或者编译器觉得寄存器不够用的时候。我曾经调试过一个kernel，发现性能比预期慢很多，一查profile，发现本地内存使用量高达几百字节——这就是典型的寄存器溢出。

3.5 常量内存（Constant Memory）

常量内存，顾名思义，存的是只读数据。它有个特点：所有线程读同一个地址时，速度极快（一个周期就能广播给所有线程）。但如果每个线程读不同的地址，那就惨了，会串行化。

我习惯把一些查找表、滤波器系数、或者不会变的参数放到常量内存里。比如在做图像处理时，把卷积核放到常量内存，效果就很好。

// 常量内存声明
__constant__ float filter[256];

// 主机端初始化
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 设备端使用
__global__ void apply_filter(float* image, int width, int height) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    // 所有线程读同一个filter值，速度很快
    float val = image[y * width + x] * filter[0];
}

3.6 纹理内存（Texture Memory）

纹理内存，这玩意儿最初是为图形学设计的。它有一些特殊能力：硬件插值、边界处理、缓存优化。对于有空间局部性的访问模式，纹理内存比全局内存快不少。

我记得在做一个医学图像处理项目时，需要对3D体数据进行三线性插值。用纹理内存的硬件插值功能，代码量减少了一半，性能还提升了30%。

3.7 各内存对比总结

内存类型	位置	访问范围	延迟	容量	缓存
寄存器	SM内部	线程私有	1 cycle	~256个/线程	无
共享内存	SM内部	线程块内	~5 cycles	~48KB/SM	无
全局内存	显存	所有线程	~400 cycles	GB级别	L1/L2
本地内存	显存	线程私有	~400 cycles	有限	L1/L2
常量内存	显存	所有线程	~1 cycle(广播)	64KB	常量缓存
纹理内存	显存	所有线程	~100 cycles	显存大小	纹理缓存

最后说一句，内存模型的选择没有银弹。你得根据具体场景来权衡：数据量大的用全局内存，需要频繁共享的用共享内存，只读且访问模式统一的用常量内存，有空间局部性的用纹理内存。我这些年做过的项目，没有一个是一上来就选对内存类型的，都是在profile和调优中慢慢摸索出来的。嗯，这就是实战的乐趣所在。