存储
NVIDIA GPU 的“存储”需要从两个视角理解：物理存储层级描述数据实际经过哪些硬件，CUDA 内存空间则描述线程如何访问数据、数据对谁可见以及存活多久。两者存在映射关系，但不是一组可以逐项等同的名称。
1. 物理存储层级
从靠近执行单元到设备显存，典型数据路径如下：
SMSP Register File
→ SM Unified L1 Data Cache / Shared Memory
→ GPU L2 Cache
→ Memory Partition / Memory Controller
→ HBM / GDDR越靠近执行单元，容量通常越小、访问延迟越低；越靠近设备显存，容量越大，但需要更多并行请求才能隐藏延迟并利用带宽。
| 物理资源 | 所在范围 | 管理方式 | 主要作用 |
|---|---|---|---|
| Register File | SMSP / SM | 编译器与硬件分配 | 保存线程操作数和执行状态 |
| Unified L1 / Shared Memory | SM | L1 由硬件缓存，Shared Memory 由 Kernel 显式分配 | 保存近期数据或 Block 协作 Tile |
| L2 Cache | 整张 GPU | 硬件缓存，可配置部分持久化策略 | 汇聚全部 SM 的设备内存访问 |
| HBM / GDDR | GPU Device Memory | CUDA Runtime、Allocator 和应用管理 | 保存权重、KV Cache、Activation 与 Workspace |
1.1 Register File
Register File 是计算单元直接使用的片上存储。CUDA C++ 中频繁访问的标量局部值通常会被编译器放入寄存器，但源码变量数量不等于最终寄存器数量；编译器会拆分、合并和复用变量。
寄存器没有缓存层级，也不需要显式 Load 到程序可见的地址空间。代价在于容量固定：每线程寄存器使用越高，一个 SM 能同时驻留的 Warp 越少。若编译器不能把值保留在寄存器中，会生成 Local Memory 访问，即 Register Spill。
1.2 Unified L1 与 Shared Memory
现代 NVIDIA GPU 的 L1 Data Cache 与 Shared Memory 使用同一组 Unified Data Cache 物理资源，但承担两种不同职责：
- L1 Cache：由硬件自动缓存 Global/Local Memory 等访问，Kernel 不直接决定某个 Cache Line 的装入和淘汰。
- Shared Memory：由 Kernel 按 Block 显式分配，程序负责装入数据、同步和复用。
部分架构允许按 Kernel 设置 L1 与 Shared Memory 的 Carveout。统一存储池容量、最大可配置 Shared Memory、单 Block 可用 Shared Memory 是不同口径，不能把产品规格中的 Unified L1/Shared 容量全部视为 Kernel 可分配空间。
Shared Memory 被划分为多个 Bank。一个 Warp 的访问若均匀落到不同 Bank，可以并行完成；若多个不同地址映射到同一 Bank，硬件会拆分请求，形成 Bank Conflict。同一地址被多个线程读取通常可以广播，不属于一般意义上的冲突。
1.3 L2 Cache
L2 Cache 位于 SM 与设备显存之间，由整张 GPU 的 SM 共享。Global Memory、Local Memory Spill 和部分只读数据都可能经过 L2。L2 可以降低重复访问 HBM/GDDR 的流量，但它不是某个 Block 或 Kernel 独占的空间，命中率会受并发 Kernel、数据规模和访问模式影响。
对大模型推理而言，L2 能缓存热点权重、KV Cache 元数据和中间结果，但模型权重通常远大于 L2 容量，不能假设完整模型在多次 Decode 迭代中始终保留在缓存里。
1.4 HBM 与 GDDR
HBM/GDDR 是离散 GPU 的主要设备显存。HBM 通过更宽的接口和多 Stack 组织提供高带宽，常见于数据中心 GPU；GDDR 使用板级显存颗粒，常见于消费级和部分工作站 GPU。
显存访问会分布到多个 Memory Partition 和 Memory Controller。控制器负责把请求映射到 Channel、Bank 和 Row，调度读写并满足 DRAM 时序。它本身不保存模型数据，而是连接 L2 与 HBM/GDDR 的访问路径。
容量和带宽解决不同问题：容量决定权重、KV Cache 和 Workspace 能否驻留在 GPU 上；带宽决定 Memory-bound Kernel 每秒能搬运多少数据。标称带宽是物理上限，实际有效带宽还取决于访问合并、缓存命中、Bank/Partition 分布和并发程度。
2. CUDA 内存空间
CUDA 内存空间按访问范围、生命周期和读写语义分类。一个空间可以映射到多个物理层级，同一物理存储也可以承载不同 CUDA 空间的数据。
| CUDA 空间 | 典型物理位置 | 可见范围 | 生命周期 | 主要特点 |
|---|---|---|---|---|
| Register | Register File | 单个 Thread | Thread | 最低访问开销，由编译器分配 |
| Local Memory | HBM/GDDR，经 L1/L2 缓存 | 单个 Thread | Thread | 逻辑上私有，物理上通常不在片上 |
| Shared Memory | SM Unified Data Cache | Thread Block；Cluster 扩展下可跨 Block 访问 | Block | 程序显式管理、低延迟、需要同步 |
| Global Memory | HBM/GDDR，经 L1/L2 缓存 | Device 上全部 Thread，Host 可通过 API 管理 | Allocation | 容量最大，是模型张量的主要存储空间 |
| Constant Memory | Device Memory + Constant Cache | Device 上全部 Thread，只读 | Allocation | 同一 Warp 广播相同地址时效率高 |
| Texture Memory | Device Memory + Texture Cache | Device 上全部 Thread，读取接口为主 | Allocation | 提供专用缓存、寻址与格式处理能力 |
2.1 Register 与 Local Memory
普通自动变量可能被编译器放入 Register，也可能因数组动态索引、变量过多或其他限制进入 Local Memory。Local 描述的是线程私有可见性，不表示数据位于线程附近的片上 SRAM。
Local Memory 访问具有设备显存级别的基础延迟，并通过 L1/L2 缓存。大型线程私有数组和 Register Spill 会增加 Local Load/Store，应结合编译器输出与 Profile 判断，而不是只根据 CUDA C++ 源码推测。
2.2 Shared Memory
Shared Memory 使用 __shared__ 声明，按 Block 分配：
__global__ void tiled_kernel(const float* input, float* output) {
extern __shared__ float tile[];
int tid = threadIdx.x;
tile[tid] = input[blockIdx.x * blockDim.x + tid];
__syncthreads();
output[blockIdx.x * blockDim.x + tid] = tile[tid];
}它适合让 Block 内线程重复使用同一 Tile、重排数据或完成归约。Shared Memory 并不会自动缓存 Global Memory；Kernel 必须显式加载数据，并在生产者与消费者之间正确同步。
2.3 Global Memory
Global Memory 是 CUDA Kernel 存放大规模张量的主要地址空间。对于离散 GPU，cudaMalloc 分配通常驻留在设备显存；Managed Memory、Mapped Host Memory 和多 GPU 虚拟内存会引入更复杂的数据位置，但 Kernel 仍通过统一的指针和地址空间访问。
Global Memory 的关键优化是 Coalescing。一个 Warp 中相邻线程访问相邻且对齐的数据时，硬件能用较少的内存事务服务请求；跨大步长或离散访问会搬运更多未使用字节，降低有效带宽。
2.4 Constant 与 Texture Memory
Constant Memory 面向只读且访问模式一致的数据。当同一 Warp 的线程读取同一地址时，Constant Cache 可以广播结果；若线程访问大量不同地址，请求可能被序列化。
Texture Memory 通过 Texture Object/Reference 提供专用缓存和寻址语义，包括特定数据格式转换、边界处理和空间局部性优化。对于普通线性张量，现代 CUDA Kernel 更常直接使用 Global Memory Load；是否使用 Texture 路径应由访问模式和 Profile 决定。
3. QA
几组容易混淆的关系如下：
- Register 与 Register File：前者是编译器为线程提供的存储语义，后者是承载它的物理硬件。
- Shared Memory 与 L1 Cache：两者共享物理池，但 Shared Memory 由程序显式管理，L1 由硬件自动管理。
- Global Memory 与 HBM/GDDR：Global Memory 是 Device Code 看到的地址空间，HBM/GDDR 是离散 GPU 上最常见的物理承载。
- Local Memory 与 HBM/GDDR：Local Memory 虽然线程私有，却通常落在设备显存中；“Local”描述可见范围，不描述物理距离。
4. 参考
- NVIDIA CUDA Programming Guide：GPU Memory
- NVIDIA CUDA C++ Programming Guide：Memory Hierarchy
- NVIDIA CUDA Best Practices Guide：Device Memory Spaces
- NVIDIA CUDA Best Practices Guide：Coalesced Access
- NVIDIA CUDA Best Practices Guide：Shared Memory
- NVIDIA CUDA Programming Guide：Configuring L1/Shared Memory Balance
- NVIDIA CUDA Programming Guide：Unified and System Memory
- NVIDIA Hopper Architecture In-Depth