1. CUDA内存模型全景：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存的层次结构与访问延迟对比
做CUDA开发这些年，我见过太多人一上来就写核函数，结果性能惨不忍睹。其实说白了，GPU编程的核心就一句话：搞清楚数据放在哪，怎么搬最快。今天咱们就把CUDA的内存模型彻底捋一遍。
1.1 为什么内存层次这么重要？
你想想看，GPU里有几千个核心同时干活，如果每个核心都去访问全局内存，那带宽再大也扛不住。我刚开始做项目时，就犯过这个错——把中间结果全存在全局内存里，结果性能还不如CPU快。
CUDA的内存模型，本质上是一个金字塔结构：
核心原则：越靠近计算单元的内存，速度越快，但容量越小。程序员的任务就是把数据安排在合适的位置。
下面这张图是我自己画的，帮你快速建立整体认知：
1.2 各内存类型详解
寄存器（Register）
寄存器是离计算单元最近的地方，访问延迟只有1个时钟周期。每个线程私有的，容量极小——通常每个线程最多255个寄存器。我在做矩阵乘法优化时，就遇到过寄存器溢出问题，性能直接腰斩。
避坑指南：我曾经在核函数里定义了太多局部变量，导致寄存器溢出到本地内存。性能从200GB/s掉到40GB/s。后来我强制限制了每个线程的寄存器数量（使用 __launch_bounds__），才把性能拉回来。
共享内存（Shared Memory）
共享内存是同一个线程块内的线程都能访问的。延迟大约5-30个周期，比全局内存快了一个数量级。我习惯用它来做线程间的数据交换和缓存。
// 共享内存声明示例
__shared__ float s_data[256];
// 使用共享内存做归约
__global__ void reduce_kernel(float *input, float *output) {
__shared__ float s_partial[256];
int tid = threadIdx.x;
s_partial[tid] = input[blockIdx.x * blockDim.x + tid];
__syncthreads(); // 同步！这行很重要
// 树形归约
for (int stride = 128; stride > 0; stride >>= 1) {
if (tid < stride) {
s_partial[tid] += s_partial[tid + stride];
}
__syncthreads();
}
if (tid == 0) output[blockIdx.x] = s_partial[0];
}个人经验：共享内存有bank冲突的问题。我建议你尽量让同一warp内的线程访问不同的bank，或者使用padding技巧来避免冲突。
本地内存（Local Memory）
本地内存听起来像本地存储，其实它物理上就在全局内存里。当寄存器不够用时，编译器会自动把变量溢出到本地内存。延迟和全局内存一样高，所以尽量避免。
怎么判断变量是否被分配到本地内存？我教你一招：编译时加 --ptxas-options=-v 参数，看输出里的 lmem 数值。如果大于0，说明有寄存器溢出了。
全局内存（Global Memory）
全局内存是GPU里最大的存储空间，通常有几十GB。但延迟也是最高的，400-800个周期。访问全局内存时，合并访问是关键——让连续的线程访问连续的内存地址。
核心技巧：全局内存访问要满足128字节对齐，且warp内线程的访问模式要能合并成一次或几次事务。我见过有人把数据结构设计成AoS（数组结构体），结果带宽利用率只有10%。改成SoA（结构体数组）后，直接跑到80%。
常量内存（Constant Memory）
常量内存只有64KB，但它的厉害之处在于：当同一warp内所有线程访问同一个常量地址时，硬件会广播给所有线程，只产生一次内存事务。我常用它来存储滤波器的系数、查找表等只读数据。
// 常量内存声明
__constant__ float filter_coeff[64];
// 主机端拷贝
cudaMemcpyToSymbol(filter_coeff, host_coeff, sizeof(float) * 64);纹理内存（Texture Memory）
纹理内存专门为空间局部性强的访问模式优化。它支持硬件插值、边界处理等功能。我一般在图像处理、科学计算中用到它。不过说实话，现在统一内存模型越来越强，纹理内存的使用场景在减少。
1.3 访问延迟对比表
| 内存类型 | 作用域 | 生命周期 | 延迟（周期） | 典型容量 |
|---|---|---|---|---|
| 寄存器 | 单个线程 | 核函数执行期间 | 1 | 每个线程最多255个 |
| 共享内存 | 线程块内 | 线程块生命周期 | 5-30 | 每个SM 48-164KB |
| L1缓存 | SM内 | 自动管理 | 30-50 | 每个SM 128KB |
| L2缓存 | 芯片全局 | 自动管理 | 200-400 | 40MB（A100） |
| 全局内存 | 所有线程 | 程序运行期间 | 400-800 | 40-80GB |
| 常量内存 | 所有线程 | 程序运行期间 | 广播时≈1 | 64KB |
| 纹理内存 | 所有线程 | 程序运行期间 | 取决于缓存命中 | 取决于设备 |
1.4 我的调优心法
做了这么多年CUDA优化，我总结了一个简单的调优流程：
- 先看全局内存访问模式——是不是合并的？对齐了吗？这一步能解决80%的性能问题。
- 再看共享内存使用——有没有bank冲突？同步点是不是太多了？
- 最后查寄存器使用——有没有溢出？用
__launch_bounds__控制一下。
小技巧：用NVIDIA Nsight Compute工具，它能直接告诉你每个内存操作的延迟和带宽利用率。我每次调优都离不开它。
嗯，内存模型这块就讲这么多。记住一句话：数据在哪，性能就在哪。下一节咱们会深入全局内存的合并访问模式，那才是真正出活的地方。