一、CUDA内存架构全景：层级结构与访问延迟
做GPU加速这些年，我最大的感触就是——内存管理才是性能的命门。很多人一上来就写核函数，结果跑起来发现比CPU还慢，十有八九是内存访问模式出了问题。
今天咱们就把CUDA的内存体系彻底捋一遍。我会结合自己踩过的坑，帮你建立起清晰的内存层级认知。
1.1 为什么需要这么多内存类型？
你想想看，一个GPU上有几千个核心同时干活。如果所有数据都挤在一个地方，那不得堵死？
所以NVIDIA设计了分层内存架构：
- 离计算单元越近，速度越快，但容量越小
- 离计算单元越远，速度越慢，但容量越大
这就像CPU的L1/L2/L3缓存一样，只不过GPU的层次更丰富，而且部分内存是软件可控的。
1.2 六种内存类型速览
先给个全景图，后面咱们逐个细聊。
1.3 逐个拆解：特性与实战
1. 寄存器 (Register)
速度最快，没有之一。寄存器就在SM内部，访问延迟为0个时钟周期。但每个线程能用的寄存器数量有限——一般是255个/线程（取决于计算能力）。
核心要点：
- 每个线程私有，其他线程不可见
- 生命周期：从线程创建到销毁
- 编译器自动分配，但你可以用
__launch_bounds__控制
我的经验： 我曾经在写一个矩阵乘法的核函数时，发现占用率只有33%。一查，每个线程用了128个寄存器。后来我手动拆分循环，把寄存器用量降到64，占用率直接翻倍，性能提升了40%。
记住：寄存器溢出（spilling）会使用本地内存，性能会断崖式下跌。
2. 本地内存 (Local Memory)
说白了，这就是寄存器的"替补"。当寄存器不够用时，编译器会把变量放到本地内存里。但注意——本地内存物理上在全局内存里，延迟高达400-800 cycles。
避坑指南： 我曾经在调试一个核函数时，发现性能异常低。用 --ptxas-options=-v 编译一看，local memory用了8KB。原来是我定义了一个大结构体，编译器放不进寄存器了。
解决方案：把大结构体拆成多个小变量，或者用共享内存手动管理。
3. 共享内存 (Shared Memory)
这是CUDA编程的核心武器。共享内存在芯片上，延迟只有20-30 cycles，比全局内存快20倍以上。同一个block内的线程可以互相访问。
| 特性 | 说明 |
|---|---|
| 容量 | 每个SM 48KB~164KB（取决于架构） |
| 延迟 | ~20 cycles（无bank conflict时） |
| 作用域 | 同一个block内的所有线程 |
| 生命周期 | block执行期间 |
使用模式：
- 数据复用：把全局内存的数据先读到共享内存，然后多个线程反复使用
- 线程间通信：用
__syncthreads()同步后，共享数据 - 避免bank conflict：访问模式要设计成连续或广播
4. 全局内存 (Global Memory)
这是GPU的"主存"，容量最大（GB级），但延迟也最高（400-800 cycles）。所有线程都能访问，但访问模式决定了性能。
我的建议： 全局内存访问一定要合并访问（coalesced access）。也就是说，同一个warp内的线程，访问的地址要连续且对齐。我见过太多人因为随机访问，导致带宽利用率不到10%。
5. 常量内存 (Constant Memory)
常量内存只有64KB，但有个大杀器——它带有一级缓存，而且所有线程访问同一个地址时，只需要一个周期。适合存放查找表、系数等只读数据。
注意： 如果warp内每个线程访问不同的常量地址，那性能会退化成串行。我曾经在物理模拟中用常量内存存粒子参数，结果每个粒子参数不同，性能反而比全局内存还差。
6. 纹理内存 (Texture Memory)
纹理内存最初是为图形学设计的，但用在通用计算上也有奇效。它支持硬件插值和边界处理，而且对二维/三维空间局部性有优化。
我个人习惯在图像处理、卷积操作中使用纹理内存。它的缓存策略对空间局部性友好，比全局内存的随机访问快不少。
1.4 延迟对比一览表
| 内存类型 | 典型延迟 | 容量 | 作用域 | 缓存 |
|---|---|---|---|---|
| 寄存器 | 0 cycle | ~255个/线程 | 线程私有 | 无 |
| 共享内存 | ~20 cycles | 48-164 KB/SM | Block内 | 无（手动管理） |
| 常量内存 | ~200 cycles | 64 KB | 全局 | 有（只读） |
| 纹理内存 | ~200-400 cycles | 取决于分配 | 全局 | 有（空间局部性） |
| 全局内存 | ~400-800 cycles | GB级 | 全局 | L1/L2 |
| 本地内存 | ~400-800 cycles | 同全局内存 | 线程私有 | L1/L2 |
1.5 实战选择策略
说了这么多，到底怎么选？我总结了一个决策流程：
- 先问自己：数据是只读还是读写？如果是只读，考虑常量内存或纹理内存。
- 再看复用性：同一个block内的线程是否反复访问同一批数据？如果是，用共享内存。
- 检查寄存器压力：编译时加
-v参数，看寄存器用量。超过64个就要警惕了。 - 最后优化全局内存：确保访问是合并的，尽量用
float4等向量类型。
一句话总结：
寄存器放最热的变量，共享内存做block内协作，常量内存存只读查找表，全局内存做最后的大容量存储。别搞反了，否则性能会很难看。
1.6 一个简单的代码示例
下面这个例子展示了共享内存和全局内存的配合使用：
__global__ void vector_add_shared(const float* A, const float* B, float* C, int N) {
__shared__ float sA[256];
__shared__ float sB[256];
int tid = threadIdx.x;
int idx = blockIdx.x * blockDim.x + tid;
// 从全局内存加载到共享内存
sA[tid] = A[idx];
sB[tid] = B[idx];
__syncthreads();
// 从共享内存计算，写入全局内存
C[idx] = sA[tid] + sB[tid];
}嗯，这里要注意：__syncthreads() 是必须的，否则线程可能读到未更新的数据。我刚开始写CUDA时，就因为这个同步问题，debug了整整两天。
1.7 小结
CUDA的内存架构，说白了就是用空间换时间，用层次换效率。理解每个内存类型的特点和适用场景，是写出高性能CUDA程序的第一步。
我个人习惯在写核函数之前，先画一张内存访问图，标清楚哪些数据放哪里、怎么流动。这个习惯帮我避免了很多性能陷阱。
下一节我们会深入共享内存的bank conflict问题，那是另一个常见的性能杀手。