学习目标学完本节你将能够理解 GPU 内存子系统的物理架构SM 内部存储 vs 显存了解 L1/L2 缓存的层级关系和在内存访问路径中的角色掌握各级存储的带宽、延迟数量级差异建立性能直觉理解从 Warp 发起内存请求到数据返回的完整路径认识局部内存的本质及其性能陷阱1. GPU 内存子系统总览GPU 的存储体系分为片上On‑Chip和片外Off‑Chip两大部分┌─────────────────────────────────────────────────────┐ │ SM流式多处理器 │ │ ┌──────────────┐ ┌──────────────────────────────┐ │ │ │ 寄存器文件 │ │ 共享内存 / L1 缓存 │ │ │ │ (256KB/SM) │ │ (可配置通常 128KB) │ │ │ └──────────────┘ └──────────────────────────────┘ │ │ ┌──────────────┐ ┌──────────────────────────────┐ │ │ │ 常量缓存 │ │ 纹理缓存 │ │ │ └──────────────┘ └──────────────────────────────┘ │ └─────────────────────────────────────────────────────┘ │ │ L2 缓存片外但靠近显存控制器 │ ┌─────────────────────────────────────────────────────┐ │ 显存Global Memory │ │ (HBM2/HBM3容量大延迟最高) │ └─────────────────────────────────────────────────────┘关键区分片上存储寄存器、共享内存、常量缓存、纹理缓存。位于 SM 内部速度快但容量小生命周期与线程/Block 绑定。片外存储全局内存显存、L2 缓存。位于 GPU 芯片外部或边缘容量大但延迟高。补充说明局部内存Local Memory局部内存不是独立的物理存储而是线程栈内存物理上落在全局显存中。延迟与全局内存一致最高级是寄存器溢出后的隐藏性能杀手。每个线程都有独立的局部内存空间但访问速度远慢于寄存器。补充说明常量缓存Constant Cache用于只读常量广播适合存放权重、超参数、滤波系数等。当同一 Warp 内所有线程读取相同地址时广播机制使访问几乎无延迟。如果 Warp 内线程读取不同地址性能急剧下降串行化。2. 各级存储的带宽与延迟定量对比以下数据以 Ampere 架构A100为例给出数量级参考值帮助建立性能直觉存储层级延迟时钟周期带宽近似容量每 SM / 整卡寄存器~0极高SM 内部全速256 KB/SM共享内存~1‑5极高与 L1 同级164 KB/SM可配置常量缓存~1‑5广播时极高64 KBL1 缓存~20‑30高128 KB/SM与共享内存共享L2 缓存~200高~7 TB/s40 MB/卡全局内存显存~400‑600高~1.5‑2 TB/s HBM240‑80 GB/卡局部内存~400‑600同全局内存受显存容量限制架构差异备注Ada、Hopper、RDNA 等不同架构的缓存容量和带宽数值会有差异但延迟的数量级关系不变核心性能规律通用。消费级 GPU如 RTX 3060跑代码时数值可能对不上但寄存器 共享内存 L1/L2 全局内存的性能排序始终成立。关键结论寄存器和共享内存的延迟比全局内存快两个数量级以上。一次全局内存访问约 600 周期的代价以 1.5 GHz 主频计算相当于 400 ns。同一时间 SM 可以执行数百条算术指令这就是为什么内存优化如此关键。L2 缓存是全局内存访问的最后一道缓存防线命中率对实际延迟影响巨大。3. L1 与 L2 缓存的角色3.1 L1 缓存位于 SM 内部与共享内存共享物理存储可配置划分。缓存全局内存和局部内存的访问。对空间局部性敏感连续地址访问命中率高。Volta 及之后的架构上L1 缓存对全局内存的合并访问也有优化作用。架构差异注意共享内存与 L1 的物理共享取决于架构Kepler、Maxwell分开设计互不影响Volta 及之后含 Ampere、Ada、Hopper合并设计支持可配置划分比例因此旧架构的调优经验不能直接套用到新 GPU 上。3.2 L2 缓存位于所有 SM 之间是 GPU 片上的最后一级缓存。容量远大于 L1A100 上 40 MB带宽接近 HBM 的 5 倍。缓存全局内存访问服务所有 SM。对多 SM 共享数据和重复访问有显著加速效果。3.3 缓存与合并访问的关系合并访问不仅决定了一次内存事务读取多少数据还直接影响 L1/L2 的缓存利用率。非合并访问会导致大量无效数据被读取到缓存浪费带宽。缓存行利用率低有效数据占比小。严重时会引发Cache Thrashing缓存颠簸频繁踢出有效数据。4. 内存访问路径从 Warp 请求到内存事务当一个 Warp 内的线程执行一次全局内存读取时完整的硬件流程如下1. Warp 发起读取请求 │ ├── 32 个线程各自给出访问地址 │ 2. 地址合并Coalescing │ 硬件分析 32 个地址的分布 │ ├── 连续且对齐 → 合并为最少的内存事务 │ └── 分散/跨步 → 拆分为多个内存事务 │ │ 注实际硬件还会在 L2 层对来自不同 Warp 的相邻请求 │ 进行二次合并进一步提高显存带宽利用率 3. L1 缓存查询 │ ├── 命中 → 直接返回数据快 │ └── 未命中 → 继续向下查询 4. L2 缓存查询 │ ├── 命中 → 返回数据 │ └── 未命中 → 访问显存 5. 显存访问 │ HBM2/HBM3 读取数据块 │ 一个内存事务通常读取 32/64/128 字节 6. 数据返回 │ 依次回填 L2、L1最终到达寄存器关键点步骤 2地址合并 是影响全局内存性能的第一道关卡下一节会详细展开。步骤 3‑5 的缓存层级决定了实际延迟。合并访问做得好缓存命中率就高事务数就少。地址合并发生在两个层次Warp 内合并决定事务数量和L2 跨 Warp 合并提高带宽利用率。5. 代码演示测量各级存储的延迟差异以下程序通过简单的时间测量直观对比四种访问模式纯寄存器计算基线共享内存访问合并全局内存访问非合并全局内存访问跨步取模典型坏访存#include cstdio #include cuda_runtime.h // 1. 纯寄存器计算基线 __global__ void regOnly(float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; float val 0.0f; for (int i 0; i 1000; i) { val idx * 0.001f; // 纯寄存器计算无内存访问 } if (idx N) out[idx] val; } // 2. 共享内存访问 __global__ void sharedMemAccess(float *out, int N) { __shared__ float tile[256]; int idx blockIdx.x * blockDim.x threadIdx.x; int tid threadIdx.x; tile[tid] tid * 1.0f; __syncthreads(); float val 0.0f; for (int i 0; i 1000; i) { val tile[tid]; // 共享内存访问片上快 } if (idx N) out[idx] val; } // 3. 合并全局内存访问对照 __global__ void coalescedAccess(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; float val 0.0f; for (int i 0; i 1000; i) { val in[idx]; // 连续访问同一 Warp 内线程读相邻地址 } if (idx N) out[idx] val; } // 4. 非合并全局内存访问跨步取模 __global__ void globalMemAccess(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; float val 0.0f; for (int i 0; i 1000; i) { // 跨步 7 并取模完全无空间局部性无法合并 // 且频繁跳跃导致 Cache Thrashing缓存颠簸 val in[(idx * 7) % N]; } if (idx N) out[idx] val; } int main() { const int N 1 16; const int blockSize 256; const int gridSize (N blockSize - 1) / blockSize; float *d_in, *d_out; cudaMalloc(d_in, N * sizeof(float)); cudaMalloc(d_out, N * sizeof(float)); cudaEvent_t start, end; cudaEventCreate(start); cudaEventCreate(end); // 寄存器版本 cudaEventRecord(start); regOnlygridSize, blockSize(d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); float ms_reg; cudaEventElapsedTime(ms_reg, start, end); printf(Register only: %f ms\n, ms_reg); // 共享内存版本 cudaEventRecord(start); sharedMemAccessgridSize, blockSize(d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); float ms_shared; cudaEventElapsedTime(ms_shared, start, end); printf(Shared memory: %f ms\n, ms_shared); // 合并全局内存版本 cudaEventRecord(start); coalescedAccessgridSize, blockSize(d_in, d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); float ms_coalesced; cudaEventElapsedTime(ms_coalesced, start, end); printf(Global (coalesced): %f ms\n, ms_coalesced); // 非合并全局内存版本 cudaEventRecord(start); globalMemAccessgridSize, blockSize(d_in, d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); float ms_global; cudaEventElapsedTime(ms_global, start, end); printf(Global (non‑coalesced): %f ms\n, ms_global); cudaFree(d_in); cudaFree(d_out); return 0; }预期结果趋势实际数值因 GPU 而异Register only: ~1.2 ms Shared memory: ~1.5 ms Global (coalesced): ~2.0 ms Global (non‑coalesced): ~8‑15 ms结果解读寄存器版本最快因为完全没有内存访问是计算性能的上限。共享内存与寄存器差距很小因为都在 SM 内部延迟极低。合并全局内存比共享内存慢但差距可控因为合并访问提高了缓存和带宽利用率。非合并全局内存比合并版本慢 4‑7 倍比寄存器慢 10 倍以上。根本原因就是前文所述每次访问都要等待约 600 个时钟周期的显存延迟且缓存被大量无效数据占满无法掩盖延迟。6. 课后练习练习1定性排序将以下存储按访问延迟从低到高排序并说明你的理由L2 缓存寄存器共享内存全局内存显存L1 缓存局部内存练习2合并访问初步观察修改本节的globalMemAccessKernel将跨步系数从 7 改为 1即in[idx]再分别测试跨步 2、7、100 的性能。记录数据并分析为什么跨步越大越慢跨步 1 和跨步 2 的差异有多大练习3共享内存与 L1 的关系查阅你的 GPU 架构文档确认共享内存和 L1 缓存是否共享物理存储Volta 及之后架构支持可配置划分。如果支持用cudaFuncSetAttribute调整划分比例配合--ptxas‑options-v观察共享内存大小变化。注意对全局内存访问性能的影响通常很小不需要追求性能差异。练习4计算缓存命中率影响假设一个 Kernel 反复访问一个 32 MB 的全局内存数组而你的 GPU L2 缓存为 40 MB。分析这个 Kernel 的缓存行为哪些访问会命中 L2哪些会穿透到显存如果数组大小增至 80 MB超过 L2 容量性能会如何变化如何调整访问顺序提高 L2 命中率练习5绘制内存层次图独立绘制一张 GPU 内存层次图标注每级存储的物理位置、典型容量和延迟数量级。用这张图向他人解释 GPU 内存子系统的工作原理。7. 下一步下一节将深入全局内存合并访问Coalesced Access的硬件机制你将学习内存事务的粒度32/64/128 字节Warp 内地址分布如何决定事务数量如何编写符合合并访问规则的 Kernel非合并访问的典型模式及其性能损失使用 Nsight Compute 检测合并访问效率