4、显存层次结构：全局内存、共享内存、寄存器、L1/L2 Cache、常量内存、纹理内存
显存层次结构，说白了就是GPU里数据流动的「高速公路网」。我刚开始接触CUDA时，总觉得显存就是显存，哪来那么多门道？直到有一次我写了个暴力矩阵乘法，全局内存读写占用了90%的时间，我才意识到——不理解层次结构，性能就是纸上谈兵。
GPU的显存不是铁板一块。它分了好几层，每一层的速度、容量、用途都不一样。你想想看，CPU有L1/L2/L3缓存，GPU其实也有类似的设计，只不过名字和用法更「硬核」一些。
核心原则：越靠近计算单元（SM），速度越快，容量越小。越远离，速度越慢，容量越大。
4.1 全局内存（Global Memory）
全局内存是GPU里最大的存储空间，通常有4GB、8GB甚至80GB（比如A100）。所有线程都能访问它，但代价是——慢。一次全局内存访问大概需要400~800个时钟周期。
我在项目中遇到过一个问题：一个简单的向量加法，带宽利用率只有20%。后来发现是合并访问（coalesced access）没做好。相邻线程访问相邻地址，才能把带宽跑满。
个人经验：写全局内存时，尽量让线程ID和内存地址对齐。比如 threadIdx.x 对应 data[threadIdx.x]，而不是 data[threadIdx.x * 2]。
4.2 共享内存（Shared Memory）
共享内存是SM内部的一块高速存储，所有线程块内的线程都能访问。它的延迟只有几个时钟周期，比全局内存快了两个数量级。
嗯，这里要注意：共享内存是手动管理的。你得自己把数据从全局内存搬进去，用完再搬回去。这听起来麻烦，但正是性能优化的关键。
// 共享内存示例：矩阵分块乘法
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int t = 0; t < N / BLOCK_SIZE; t++) {
// 加载到共享内存
As[threadIdx.y][threadIdx.x] = A[row * N + t * BLOCK_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(t * BLOCK_SIZE + threadIdx.y) * N + col];
__syncthreads();
// 从共享内存计算
for (int k = 0; k < BLOCK_SIZE; k++)
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
__syncthreads();
}
C[row * N + col] = sum;
}避坑指南：我曾经在共享内存上踩过一个坑——bank conflict。当多个线程同时访问同一个bank的不同地址时，访问会被串行化。解决办法是给数组加一个padding，比如 __shared__ float data[32][33] 而不是 [32][32]。
4.3 寄存器（Register）
寄存器是GPU里最快的存储，没有之一。每个线程都有自己的私有寄存器，延迟只有1个时钟周期。但数量有限——每个SM通常只有65536个寄存器，分给所有线程。
你想想看，如果一个线程用了32个寄存器，那每个SM最多只能跑2048个线程（65536/32）。寄存器用多了，并行度就降了。
| 存储类型 | 延迟 | 容量 | 作用域 |
|---|---|---|---|
| 寄存器 | 1 cycle | ~256KB/SM | 单个线程 |
| 共享内存 | ~5 cycles | 48~96KB/SM | 线程块内 |
| L1 Cache | ~5 cycles | 与共享内存共享 | SM内 |
| 常量内存 | ~100 cycles | 64KB | 全局只读 |
| 纹理内存 | ~100 cycles | 取决于显存 | 全局只读 |
| L2 Cache | ~200 cycles | ~40MB | 芯片级 |
| 全局内存 | ~400-800 cycles | GB级别 | 全局 |
4.4 L1/L2 Cache
L1 Cache和共享内存共用同一块物理存储。你可以通过编译选项调整它们的比例：比如 --cuda-gpu-arch=sm_80 配合 cudaFuncSetAttribute 设置。
L2 Cache是所有SM共享的，容量大概几十MB。它的作用是缓存全局内存的访问，减少对HBM（高带宽显存）的直接访问。
我的习惯：如果数据访问模式是「一次读多次用」，我会优先用共享内存手动缓存。如果数据访问模式是「随机且稀疏」，那就靠L2 Cache自动缓存，省心。
4.5 常量内存（Constant Memory）
常量内存只有64KB，但它的厉害之处在于广播能力。当一个warp里的所有线程访问同一个常量地址时，硬件只需要一次读取，然后广播给所有线程。
我记得有一次做物理模拟，需要把一组物理常量（比如重力、摩擦系数）传给所有线程。用常量内存后，性能提升了15%。
// 常量内存声明
__constant__ float gravity[3];
// 主机端设置
float h_gravity[3] = {0.0f, -9.8f, 0.0f};
cudaMemcpyToSymbol(gravity, h_gravity, sizeof(float) * 3);
// 内核中使用
__global__ void physics_kernel(float *pos, float *vel, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
vel[idx * 3 + 1] += gravity[1] * dt; // 所有线程读同一个地址
}
}4.6 纹理内存（Texture Memory）
纹理内存最初是为图形渲染设计的，但它对某些计算场景特别有用。它支持硬件插值和边界处理，而且对二维空间局部性有优化。
我个人觉得，纹理内存最适合做图像处理或者需要「邻居访问」的算法。比如图像滤波、卷积操作，用纹理内存比全局内存快不少。
关键区别：纹理内存有硬件缓存，而且对二维访问模式做了优化。如果你的数据访问是「行优先」的，用全局内存+合并访问就够了。如果是「随机二维坐标」，纹理内存更合适。
4.7 如何选择？我的经验法则
- 寄存器：放线程私有的临时变量，比如循环计数器、中间结果。
- 共享内存：放线程块内需要共享的数据，比如矩阵分块中的子矩阵。
- 常量内存：放所有线程都读同一个值的只读数据，比如物理常量、查找表。
- 纹理内存：放二维/三维空间访问模式的数据，比如图像、地形数据。
- 全局内存：放所有其他数据，但要注意合并访问。
说白了，显存层次结构就是让你用最快的存储处理最频繁的数据。我见过太多人一股脑把所有数据扔进全局内存，然后抱怨GPU不够快。其实不是GPU慢，是你没用好它的层次结构。
最后提醒：别为了优化而优化。先跑一遍profiler，看看瓶颈到底在哪。我曾经花了两天优化共享内存，结果发现瓶颈在CPU到GPU的数据传输上——白忙活了。