企业官网建设流程全解析

学习目标

学完本节你将能够：

• 理解 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）为例，给出数量级参考值，帮助建立性能直觉：

架构差异备注：Ada、Hopper、RDNA 等不同架构的缓存容量和带宽数值会有差异，但延迟的数量级关系不变，核心性能规律通用。消费级 GPU（如 RTX 3060）跑代码时数值可能对不上，但寄存器 < 共享内存 < L1/L2 < 全局内存的性能排序始终成立。

关键结论：

1. 寄存器和共享内存的延迟比全局内存快两个数量级以上。
2. 一次全局内存访问（约 600 周期）的代价，以 1.5 GHz 主频计算相当于 400 ns。同一时间 SM 可以执行数百条算术指令，这就是为什么内存优化如此关键。
3. L2 缓存是全局内存访问的最后一道缓存防线，命中率对实际延迟影响巨大。

3. L1 与 L2 缓存的角色

3.1 L1 缓存

• 位于 SM 内部，与共享内存共享物理存储（可配置划分）。
• 缓存全局内存和局部内存的访问。
• 对空间局部性敏感：连续地址访问命中率高。
• Volta 及之后的架构上，L1 缓存对全局内存的合并访问也有优化作用。

架构差异注意：共享内存与 L1 的物理共享取决于架构：

• Kepler、Maxwell：分开设计，互不影响
• Volta 及之后（含 Ampere、Ada、Hopper）：合并设计，支持可配置划分比例

因此旧架构的调优经验不能直接套用到新 GPU 上。

3.2 L2 缓存

• 位于所有 SM 之间，是 GPU 片上的最后一级缓存。
• 容量远大于 L1（A100 上 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，最终到达寄存器

关键点：

1. 步骤 2（地址合并） 是影响全局内存性能的第一道关卡，下一节会详细展开。
2. 步骤 3‑5 的缓存层级决定了实际延迟。合并访问做得好，缓存命中率就高，事务数就少。
3. 地址合并发生在两个层次：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); regOnly<<<gridSize, 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); sharedMemAccess<<<gridSize, 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); coalescedAccess<<<gridSize, 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); globalMemAccess<<<gridSize, 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 检测合并访问效率