第二章：CUDA内存架构——全局内存、共享内存、寄存器、常量内存、纹理内存的层次结构与特性
大家好，欢迎来到第二章。这一章我们聊聊CUDA的内存架构。说实话，我刚开始接触CUDA时，最头疼的就是这堆内存类型——全局、共享、寄存器、常量、纹理……名字都记不全。但后来我发现，搞懂它们，你的程序性能就能翻几倍。今天我就带你捋一遍。
2.1 为什么需要这么多内存？
你想想看，GPU要处理海量数据，如果所有数据都放在一个地方，那访问速度肯定慢得离谱。所以NVIDIA设计了一套分层的内存体系，就像CPU的L1/L2缓存一样，但更复杂。
说白了，就是让数据尽量靠近计算单元。越靠近计算单元，访问越快，但容量越小。这是个经典的trade-off。
核心原则：频繁访问的数据放近处，偶尔访问的数据放远处。别把热数据扔在全局内存里，那是自找麻烦。
2.2 全局内存（Global Memory）
全局内存是GPU里最大的存储区域，通常有几十GB。所有线程都能读写它，但速度最慢。延迟大概在400-800个时钟周期。
我在项目中遇到过一个问题：一个简单的向量加法，数据全放在全局内存里，结果带宽利用率只有30%。后来发现是访问模式不对——没有合并访问。
关键点：
- 所有线程都能访问，但速度慢
- 需要合并访问才能达到高带宽
- 生命周期由程序员控制（cudaMalloc/cudaFree）
我的经验：全局内存的访问模式决定了带宽利用率。连续访问比随机访问快10倍以上。写代码时尽量让相邻线程访问相邻地址。
2.3 共享内存（Shared Memory）
共享内存是块上的存储，同一个块内的线程可以共享。延迟只有几个时钟周期，比全局内存快两个数量级。
但容量很小——每个SM通常只有几十KB到上百KB。你得精打细算。
特性：
- 块内线程共享，块间隔离
- 延迟低（约20-30个周期）
- 需要手动管理数据搬入搬出
- 有bank conflict问题，需要注意
避坑指南：我曾经在矩阵乘法里用了共享内存，但没处理好bank conflict，结果性能反而下降了。记住：共享内存不是银弹，用不好就是拖累。
2.4 寄存器（Register）
寄存器是离计算单元最近的内存，速度最快——零延迟。每个线程都有自己的寄存器，但数量有限（通常每个线程最多255个）。
如果你在核函数里声明太多局部变量，编译器会把多余的变量“溢出”到本地内存（local memory），那速度就慢了。
要点：
- 每个线程私有，速度最快
- 数量有限，过多会导致寄存器溢出
- 编译器自动管理，但你可以通过
__launch_bounds__提示
我建议：写核函数时，尽量少用局部变量。如果必须用，考虑用共享内存替代。我见过有人一个核函数用了300个寄存器，结果occupancy掉到20%。
2.5 常量内存（Constant Memory）
常量内存是只读的，所有线程都能访问。它有自己的缓存（constant cache），如果所有线程访问同一个地址，速度非常快。
但有个限制：常量内存只有64KB。而且如果线程访问不同地址，性能会急剧下降。
适用场景：
- 存储不会改变的参数（如滤波器系数、查找表）
- 所有线程访问相同地址
我的经验：在图像处理中，我经常把卷积核放在常量内存里。所有线程读同一个核，缓存命中率100%，速度飞快。
2.6 纹理内存（Texture Memory）
纹理内存是专门为图形处理设计的，但也能用于通用计算。它有硬件支持的插值、边界处理等功能。
纹理内存有缓存，对空间局部性好的访问模式很友好。而且它支持一维、二维、三维数据。
特点：
- 只读，有缓存
- 支持硬件插值和边界处理
- 对二维访问模式特别友好
- 延迟比全局内存低，但比共享内存高
避坑指南：纹理内存不是万能的。如果你的访问模式是随机的，纹理缓存反而会成为瓶颈。我曾经在稀疏矩阵计算里试过纹理内存，结果比全局内存还慢。
2.7 内存层次结构总览
下面这张图展示了CUDA内存的层次结构。你可以看到，从寄存器到全局内存，延迟和容量都在增加。
2.8 如何选择合适的内存？
这个问题没有标准答案，但我可以给你一些经验法则：
| 场景 | 推荐内存类型 | 原因 |
|---|---|---|
| 数据被频繁读写 | 共享内存 | 低延迟，块内共享 |
| 数据只读且所有线程访问相同地址 | 常量内存 | 缓存命中率高 |
| 二维/三维空间局部性访问 | 纹理内存 | 硬件缓存和插值支持 |
| 数据量大且访问模式随机 | 全局内存 | 容量大，但注意合并访问 |
| 临时变量，生命周期短 | 寄存器 | 零延迟，但注意数量限制 |
核心建议：先分析你的数据访问模式，再选择内存类型。别一上来就用共享内存，先问问自己：数据真的需要共享吗？访问模式适合吗？
2.9 实战：一个简单的例子
假设我们要做矩阵乘法。数据在全局内存里，但我们可以把子矩阵搬到共享内存里加速。
__global__ void matMulShared(float* A, float* B, float* C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int row = by * TILE_SIZE + ty;
int col = bx * TILE_SIZE + tx;
float sum = 0.0f;
for (int t = 0; t < N / TILE_SIZE; ++t) {
// 搬数据到共享内存
As[ty][tx] = A[row * N + t * TILE_SIZE + tx];
Bs[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
__syncthreads();
// 计算
for (int k = 0; k < TILE_SIZE; ++k) {
sum += As[ty][k] * Bs[k][tx];
}
__syncthreads();
}
C[row * N + col] = sum;
}这个例子展示了如何用共享内存减少全局内存访问。每个线程只从全局内存读两次（A和B的子块），然后从共享内存读多次。性能提升很明显。
我的经验：共享内存的大小要选好。TILE_SIZE=16或32通常不错。太大可能降低occupancy，太小又浪费带宽。我一般会做profiling来确定最优值。
2.10 总结
CUDA的内存架构其实不复杂，关键是要理解每个层次的特性：速度、容量、作用域。写代码时，多想想数据应该放在哪里。
记住：没有最好的内存，只有最合适的内存。全局内存慢但容量大，共享内存快但容量小，寄存器最快但数量有限。根据你的需求选，别盲目。
好了，这一章就到这里。下一章我们会深入聊聊共享内存的bank conflict问题，那是个坑，但也是优化的关键。我们下次见。