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

说到CUDA的内存模型，我刚开始学的时候也觉得头大。六种内存，各有各的脾气。但说白了，搞懂它们，你的算子性能就能翻几倍。今天咱们一个一个捋清楚。

1. 全局内存（Global Memory）

全局内存是GPU里最大的存储空间，所有线程都能访问。但它有个毛病——慢。延迟大概400-600个时钟周期。你想想看，CPU访问L1缓存才几个周期，这差距有多大。

特点：

容量大：通常几个GB到几十个GB
所有线程可读写
延迟高，需要合并访问才能发挥带宽
生命周期：从cudaMalloc到cudaFree

避坑指南：我曾经在一个图像处理项目里，没注意全局内存的合并访问，结果带宽利用率只有10%。后来改成连续访问，性能直接提升了8倍。记住：相邻线程访问相邻地址，这是黄金法则。

2. 共享内存（Shared Memory）

共享内存是块宝。它在每个SM内部，延迟只有几十个时钟周期。同一个block里的线程可以共享数据。我习惯用它来做数据复用和线程间通信。

关键点：

容量小：通常48KB-164KB per SM
同一个block内线程共享
延迟低，约20-30个周期
需要__syncthreads()同步

// 共享内存使用示例
__global__ void shared_mem_example(float *input, float *output) {
    __shared__ float s_data[256];  // 声明共享内存
    
    int tid = threadIdx.x;
    s_data[tid] = input[tid];     // 加载到共享内存
    __syncthreads();              // 同步，确保所有线程加载完成
    
    // 在共享内存上做计算
    s_data[tid] = s_data[tid] * 2.0f;
    __syncthreads();
    
    output[tid] = s_data[tid];    // 写回全局内存
}

个人经验：我建议把共享内存想象成CPU的L1缓存。但它是软件管理的，你得手动控制数据加载和同步。刚开始容易忘加__syncthreads()，然后debug到怀疑人生。

3. 寄存器（Register）

寄存器是GPU里最快的内存，没有之一。每个线程独享，延迟1个周期。但数量有限——每个线程最多255个寄存器（不同架构有差异）。

为什么重要？

速度最快，编译器优先使用
线程私有，无需同步
数量有限，用多了会spill到本地内存

注意：寄存器溢出（register spilling）是性能杀手。我曾经有个算子，寄存器用了太多，结果编译器把部分变量存到了本地内存，性能直接腰斩。用--ptxas-options=-v可以查看寄存器使用情况。

4. 本地内存（Local Memory）

本地内存听起来像本地存储，其实它存在全局内存里。当寄存器不够用时，编译器会把变量"溢出"到这里。嗯，这里要注意：本地内存虽然名字带"本地"，但延迟和全局内存一样高。

触发条件：

寄存器不够用
数组索引无法在编译时确定
大型结构体或数组

避坑指南：我曾经写过一个复杂的kernel，里面有个大数组，结果编译器把它放到了本地内存。性能惨不忍睹。后来改成共享内存，速度提升了3倍。所以，尽量控制每个线程的变量数量。

5. 常量内存（Constant Memory）

常量内存是只读的，容量64KB。它有个特殊之处——有专门的缓存，当所有线程访问同一个地址时，速度极快。我经常用它来存查找表、系数矩阵这类数据。

使用场景：

存储只读的常量数据
所有线程访问相同地址
查找表、滤波器系数等

// 常量内存声明
__constant__ float filter_coeff[64];

// 主机端初始化
float h_coeff[64] = { ... };
cudaMemcpyToSymbol(filter_coeff, h_coeff, sizeof(float) * 64);

// kernel中使用
__global__ void apply_filter(float *input, float *output) {
    int idx = threadIdx.x;
    // 所有线程访问同一个filter_coeff
    output[idx] = input[idx] * filter_coeff[idx % 64];
}

6. 纹理内存（Texture Memory）

纹理内存最初是为图形学设计的，但用在某些计算场景也很香。它支持硬件插值、边界处理，对二维空间局部性访问特别友好。

优势：

硬件加速的插值操作
自动处理边界情况
对二维访问模式有优化
有独立的纹理缓存

个人经验：我做过一个图像缩放的项目，用纹理内存的硬件插值，比手动实现快了5倍。但要注意，纹理内存有容量限制，而且绑定和解绑有开销。适合数据量大、访问模式不规则的情况。

内存模型对比总结

内存类型	位置	延迟	容量	作用域	生命周期
寄存器	SM内部	1周期	~255/线程	线程私有	kernel执行期间
共享内存	SM内部	~30周期	48-164KB/SM	block内共享	block执行期间
全局内存	显存	~400周期	GB级别	所有线程	程序控制
本地内存	显存	~400周期	有限	线程私有	kernel执行期间
常量内存	显存+缓存	~1周期(缓存命中)	64KB	所有线程	程序控制
纹理内存	显存+缓存	~100周期	显存大小	所有线程	程序控制

内存层次结构图

这张图把六种内存的关系画得很清楚。从上到下，速度越来越慢，容量越来越大。我刚开始做CUDA优化时，就把这张图贴在显示器旁边，时刻提醒自己：数据放对地方，性能就赢了一半。

核心原则：频繁访问的数据放寄存器或共享内存，只读且访问模式规律的数据放常量内存，大块数据放全局内存但注意合并访问。纹理内存适合二维空间局部性访问。本地内存...尽量别用。

好了，内存模型就聊到这儿。每种内存都有它的脾气，摸透了，你的CUDA算子就能跑得飞快。记住我的一句话：数据在哪里，性能就在哪里。

📚 课程目录

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

1. 全局内存（Global Memory）

2. 共享内存（Shared Memory）

3. 寄存器（Register）

4. 本地内存（Local Memory）

5. 常量内存（Constant Memory）

6. 纹理内存（Texture Memory）

内存模型对比总结

内存层次结构图