07 > 存储层次

GPU 并不是只有一种内存，而是有一整架 “梯子”。靠近核心的几级很小、很快，离得远的那一级很大、很慢。这一课在 L40S 上沿着这架梯子往下走，看每一级谁能访问、有多大、有多快，并算一算哪里能放下多少数据。

一架梯子，三个问题

第 00 课介绍过寄存器、共享内存和 L2 缓存，第 06 课讲了显存带宽为什么会限制 GPU。这一课把它们按顺序排好，这个顺序就叫存储层次。

对每一级，都问三个问题。

• 它在哪里？在 SM 内部，在 GPU 芯片的其他位置，还是在旁边单独的显存芯片上。
• 谁能访问它？一个线程，同一个线程块里的所有线程，还是整个 GPU 上的所有线程。
• 它有多大、多快？小的层级快，大的层级慢，没有哪一级又大又快。

“快” 有两种含义。带宽是一级存储每秒能送出多少字节。延迟是一次加载要等多久数据才到，用时钟周期来计。

寄存器

寄存器就在每个 SM 内部，紧挨着核心。每个线程都有自己的寄存器，用来放局部变量，别的线程读不到。用寄存器不需要额外等待，因为核心在执行指令时直接读取它。

L40S 一个 SM 的寄存器文件有 65,536 个 32 位寄存器，也就是 256 KB。一个线程最多能用 255 个。H100 每个 SM 也是 256 KB。

[!NOTE] 把整块芯片加起来，寄存器其实不少。L40S 有 142 个 SM × 256 KB = 36,352 KB 寄存器，约 35.5 MB，比它的 142 × 100 KB = 14,200 KB 共享内存还多。

共享内存和 L1

每个 SM 还有一块高速的片上存储，分给两种用途。

• 共享内存，由核函数自己管理。同一个线程块里的所有线程都能读写它，所以它们用它来共享数据。其他线程块的线程看不到它。
• L1 缓存，由硬件管理。它把最近从全局内存加载的数据留在 SM 附近，同样的数据再读一次就很快。

L40S 是 CC 8.9 的 GPU，在它上面，这块存储每个 SM 有 128 KB。核函数可以选择怎么分，这叫 carveout。共享内存取 0、8、16、32、64 或 100 KB，剩下的给 L1。CUDA 会给每个线程块预留 1 KB，所以一个线程块最多能用 99 KB。H100 上这块存储每个 SM 有 256 KB，其中最多 228 KB 可作共享内存。

已发表的微基准测试在 RTX 4090 上做了测量，它和 L40S 用的是同一种 AD102 芯片。结果是共享内存加载大约 30 个周期，L1 命中大约 43 个周期。

L2 缓存

L2 缓存在 GPU 芯片上，但在 SM 之外。所有 SM 共用它，所以它服务整个 GPU。对全局内存的每一次读写都要经过它。

L40S 有 96 MB 的 L2，H100 有 50 MB。同一组微基准测试在 RTX 4090 上测得 L2 命中大约 273 个周期，约为共享内存加载的 9 倍。

全局内存

全局内存是 GPU 的主内存，也就是 L40S 上 48 GB 的 VRAM。它在单独的显存芯片上，所有线程块的所有线程都能访问它，CPU 也会把数据拷进拷出。全局内存里的数据在多次核函数启动之间一直保留。

L40S 用的是 GDDR6，带宽 864 GB/s。H100 SXM 有 80 GB 的 HBM3，带宽 3.35 TB/s。全局内存是最慢的一级，在 RTX 4090 上每次加载大约 541 个周期，是 L2 命中的 2 倍，约为共享内存加载的 18 倍。

常量内存和纹理内存

全局内存还有两种特殊的访问方式，各自带一个小缓存。

常量内存是 64 KB 的只读数据，每个 SM 缓存其中 8 KB。当一个线程束的所有线程读同一个地址时，读一次就广播给全部 32 个线程。如果它们读不同的地址，这些读取就只能一个接一个地完成。

纹理内存是一条经过 L1 缓存的只读通路，为图形而设计，因为图形里相邻的线程往往读相邻的像素。在 Ada 和 Hopper 上，纹理缓存和 L1 是同一个单元，所以大多数 CUDA 代码直接读全局内存，让 L1 去缓存就行。

本地内存和寄存器溢出

本地内存和寄存器一样只属于一个线程，但它其实在全局内存里。它会经过 L1 和 L2 缓存，可一旦未命中，代价就和任何一次全局内存加载一样。每个线程最多能用 512 KB。

寄存器不够用时，编译器就把数据放到本地内存，这叫寄存器溢出。如果线程用运行时才知道的值去索引一个数组，也会这样，因为寄存器不能这样索引。

[!WARNING] “本地” 的意思是私有，不是离得近。本地内存在芯片外，和全局内存一样慢。编译时加上 -Xptxas -v，编译器会报告每个核函数用了多少寄存器，以及溢出写入和溢出读取各多少字节。

一览表

这些周期数是在 RTX 4090 上测得的。H800 是和 H100 一样的 Hopper GPU，在它上面测得的结果几乎一样，分别是 29、41、263 和 479 个周期。

算一算哪里能放下多少

一个 float 占 4 个字节。在 CUDA 的表格里，1 KB 是 1,024 字节，1 MB 是 1,024 KB。

L40S 的一个 SM 最多有 100 KB 共享内存。

• 100 × 1,024 = 102,400 字节。
• 102,400 / 4 = 25,600 个 float。
• 一个 32 × 32 的 float 分块占 32 × 32 × 4 = 4,096 字节 = 4 KB，所以一个 SM 能放 100 / 4 = 25 个这样的分块。
• 一个线程块最多用 99 KB，也就是 99 × 1,024 / 4 = 25,344 个 float。

H100 每个 SM 有 228 KB，能放 228 × 1,024 / 4 = 58,368 个 float，是 L40S 的两倍多。

L40S 的 L2 缓存有 96 MB。

• 96 × 1,024 × 1,024 = 100,663,296 字节。
• 100,663,296 / 4 = 25,165,824 个 float，约 2,520 万个。
• 一个 5,000 × 5,000 的 float 矩阵是 25,000,000 个 float，刚好放得下。6,000 × 6,000 的矩阵（36,000,000 个 float）就放不下了。

[!TIP] 一个 SM 上的所有线程共用寄存器。要让一个 L40S SM 跑满 1,536 个线程，每个线程能分到 65,536 / 1,536 ≈ 42.7 个寄存器。硬件按每个线程束 256 个寄存器为一块来分配，所以实际上限是每个线程 40 个，因为 40 × 32 = 1,280 = 5 块，48 个线程束 × 1,280 = 61,440，不超过 65,536。

这对 CUDA 意味着什么

写核函数时，你要为每一份数据选好它放在哪一级。

• 普通的局部变量放在寄存器里。数量要少，否则会溢出到本地内存，跑得和全局内存一样慢。
• 线程块里的线程要反复读取的数据放进共享内存，用 __shared__ 声明。从全局内存只加载一次，之后每次复用只要约 30 个周期，而不是约 541 个。
• 大数组放在全局内存里。读取时让一个线程束的 32 个线程访问相邻的地址，这样 L1 和 L2 就能用几次宽的传输把它们都送到。

一次全局内存加载要几百个周期。GPU 靠切换到其他线程束来掩盖这段等待，这是 第 08 课 的主题。第 09 课 接着讲怎么判断一个核函数是受计算限制还是受内存限制。

术语表

• GPU（Graphics Processing Unit，图形处理器）：这些课讲的处理器，由许多并行运行线程的 SM 组成。
• CPU（Central Processing Unit，中央处理器）：计算机的主处理器，负责把数据拷进、拷出全局内存。
• 存储层次（memory hierarchy）：GPU 各级存储组成的梯子，从又小又快的寄存器到又大又慢的全局内存。
• SM（Streaming Multiprocessor，流式多处理器）：GPU 里的处理单元，有自己的核心、寄存器、共享内存和 L1 缓存。L40S 有 142 个。
• 线程（thread）：一条指令流。每个线程有自己的寄存器和本地内存。
• 线程块（block）：在一个 SM 上运行的一组线程，可以通过共享内存共享数据。
• 线程束（warp）：SM 一起运行的 32 个线程。
• 核函数（kernel）：在 GPU 上运行的函数，由许多线程一起执行。
• 带宽（bandwidth）：一级存储每秒能送出的字节数。
• 延迟（latency）：一次加载要等多久数据才到。
• 时钟周期（clock cycle）：GPU 时钟的一次跳动。本页的延迟都按周期计。
• 寄存器（register）：最快的存储，只属于一个线程。L40S 的一个 SM 有 65,536 个 32 位寄存器。
• 寄存器文件（register file）：一个 SM 的全部寄存器，L40S 和 H100 上都是 256 KB。
• KB（kilobyte，千字节）/ MB（megabyte，兆字节）：在 CUDA 的表格里分别是 1,024 字节和 1,024 KB。
• 片上（on-chip）：直接做在 GPU 芯片里，比如寄存器、共享内存、L1 和 L2。
• 共享内存（shared memory）：每个 SM 里的高速片上存储，同一个线程块的所有线程都能读写。
• L1 缓存（Level 1 cache，一级缓存）：每个 SM 片上存储中由硬件管理的那一部分，把最近用过的数据留在近处。
• carveout（划分比例）：片上存储在共享内存和 L1 之间怎么分，按核函数来选。
• CC（compute capability，计算能力）：GPU 功能的版本号。L40S 是 CC 8.9。
• CUDA（Compute Unified Device Architecture，统一计算设备架构）：NVIDIA 用来编写在其 GPU 上运行的程序的平台。
• 微基准测试（microbenchmark）：只测一件事的小测试程序，比如某一级存储的延迟。
• L2 缓存（Level 2 cache，二级缓存）：所有 SM 共用的片上缓存，L40S 上 96 MB，H100 上 50 MB。
• 全局内存（global memory）：GPU 的主内存，所有线程都能访问，CPU 也往里拷数据，是最慢的一级。
• VRAM（显存）：GPU 卡上存放全局内存的显存芯片。
• GDDR6（Graphics Double Data Rate 6）：L40S 用的显存类型，48 GB，864 GB/s。
• HBM3（High Bandwidth Memory 3，第三代高带宽内存）：和 GPU 芯片封装在一起的堆叠式显存。H100 SXM 有 80 GB，3.35 TB/s。
• GB/s（gigabytes per second，吉字节每秒）/ TB/s（terabytes per second，太字节每秒）：带宽的单位。
• 常量内存（constant memory）：64 KB 只读数据，每个 SM 有 8 KB 缓存，整个线程束读同一个地址时很快。
• 广播（broadcast）：读一次，结果同时送给线程束的全部 32 个线程。
• 纹理内存（texture memory）：一条经过 L1 缓存的只读通路，为图形而设计。
• 本地内存（local memory）：每个线程私有、但放在全局内存里的存储，每个线程最多 512 KB。
• 寄存器溢出（register spill）：寄存器不够用时，把数据从寄存器挪到本地内存。
• float：32 位浮点数，占 4 个字节。
• 分块（tile）：大数组中的一小块方形区域，加载到共享内存里反复使用。