GPU 执行模型 3：运算与内存子系统

首先我们对商品 GPU 种类、GPU 的整体结构、GPU 核心结构、SM 和 partition 结构进行简单介绍。然后我们我们详细介绍运算子系统和内存子系统。

背景知识

商品 GPU 种类

PCIe 卡

一张独立的卡，有自己的 PCB、散热器、供电模组，竖着插在主板的 PCIe 插槽里。整张卡是一个完整的独立单元，可以随时拔插。
HGX

HGX 板上装好了 8 个 SXM GPU 模块和 NVSwitch，但不包含 CPU、机箱和其他配件。

Dell、Supermicro、联想这些 OEM 厂商买 HGX 板子，然后自己搭配 CPU、机箱、电源、散热做成完整的服务器卖给客户。
- SXM 模块
  
  GPU 核心和 HBM 显存封装在一个扁平的小模块上（大概巴掌大小），模块底部是密密麻麻的触点。
  
  这个模块从上方按压到 HGX 板上的专用插座里，有点像 CPU 装进 LGA 插座的感觉。
  
  HGX 板上有 8 个这样的插座，每个插座下面的 PCB 走线直接连到 NVSwitch 芯片，实现 GPU 之间的 NVLink 全互联。
DGX

NVIDIA 的「整机解决方案」，把 GPU、CPU、NVLink/NVSwitch、InfiniBand 网卡、NVMe 存储、散热系统、操作系统和软件栈全部整合好，出厂即用。

GPU 整体结构

包括 A40 在内的 PCIe GPU 由以下部分组成：

GPU 核心：A40 的核心为 GA102，628 平方毫米 die
显存颗粒：显存颗粒占了很大面积。
- A40 有 48GB GDDR6。这些显存是一颗颗独立的 DRAM 芯片，围绕在 GPU chip 周围焊在 PCB 上。
- A100 内存类型是 HBM。HBM 是垂直堆叠的芯片，通过硅中介层（silicon interposer）跟 GPU die 紧密贴在一起，虽然面积比 GDDR6 小但高度更高。
供电模组（VRM）：供电模组也占了不少板面积。

GPU 动辄几百瓦，需要大量的电压调节器把 12V 的输入电压转换成 GPU 核心需要的低电压大电流。这些供电电路包括 MOSFET、电感、电容，在 PCIe 卡上通常占据了背面和侧面的大片区域。
散热

散热方案虽然不在 PCB 上，但决定了卡的整体尺寸。
- PCIe 卡上方那个巨大的散热器加风扇占了卡的大部分体积。
- SXM 模块虽然没有自带散热器，但模块本身也需要为冷板留出接触面积。
走线的空间

PCB 走线本身就需要空间：
- GPU die 到每颗显存之间需要高速信号线（比如 384-bit 总线意味着几百根高速差分线），
- 还有 PCIe 接口的走线、
- 供电的铜层、
- NVLink 接口的走线，
这些都需要面积来布局和布线。
其他芯片

其他芯片也占一些空间，比如视频输出接口的驱动芯片、BIOS ROM、温度传感器、风扇控制器等。

所以 GPU die 本身只是整张卡上最核心但面积占比并不大的一部分。大量的板面积留给了喂饱这颗 die 所需要的供电和显存。

GPU 核心结构

GPU 前端
- Host Interface Unit
- FIFO 缓冲区
- Command Processor
- （Syncpoint 寄存器）
GPU 后端
- GigaThread Engine：处理 kernel launch 的引擎，把 kernel 计算任务分发到 SM。
  - Streaming Multiprocessors：具体执行计算任务的单元。
- Copy Engine：负责内存 DMA 搬运的引擎。

SM 和 Partition 结构

1 GA102 is composed of:

PCI Express 4.0 Host Interface
GigaThread Engine
7 Graphics Processing Clusters (GPCs)
- 1 Raster Engine per GPC
- 2 Raster Operators (ROPS) partitions per GPC
  - 8 ROP units per ROPS partition
- 6 Texture Processing Clusters (TPCs) per GPC
  - 2 Streaming Multiprocessors (SMs) per TPC, 84 SMs in total
  - 1 PolyMorph Engine per TPC
12 32-bit memory controllers (384-bit total)
512KB L2 cache paired with each controller, 6144KB L2 cache in total

1 SM is composed of:

4 partitions / processing blocks
- 16384 32-bit register file per partition, 65536 per SM
- 1 L0 i-cache per partition
- 1 warp scheduler per partition
- (12 warp slots, 48 per SM)
- 1 dispatch unit per partition
- 1 3rd-gen Tensor cores per partition, 4 per SM, 336 in total
- 16 FP32 CUDA cores per partition, 64 per SM, 5376 in total
  - capable of half precision arithmetic
- 16 FP32/INT32 CUDA cores, 64 per SM, 5376 in total
  - capable of half precision arithmetic
- 4 Special Function Units (SFUs) per partition
- 8 Load/Store units per partition
2 FP64 units per SM, 168 in total
128KB L1 data (texture) cache/Shared Memory per SM shared by 4 partitions
(16 block slots)
4 Texture units per SM
1 2nd-gen Ray Tracing core per SM

In compute mode, the GA10x SM will support the following configurations for L1 d-cache/Shared Memory:

128 KB L1 + 0 KB Shared Memory
120 KB L1 + 8 KB Shared Memory
112 KB L1 + 16 KB Shared Memory
96 KB L1 + 32 KB Shared Memory
64 KB L1 + 64 KB Shared Memory
28 KB L1 + 100 KB Shared Memory

运算子系统

GA102 的 partition 里，做运算的单元主要有四种：

32 个 CUDA Core
1 个 Tensor Core
4 个 Special Function Unit
8 个 Load/Store Unit

除了 partition 内部的运算单元，每个 SM 还有：

2 个 FP64 unit

写高性能 kernel 的时候理想状态是让这几种单元同时忙起来，不要让任何一种成为瓶颈。

背景知识

FFMA、FADD、FMUL

FFMA 就是 Floating-point Fused Multiply-Add，浮点融合乘加指令。一条指令完成 D = A × B + C，一个时钟周期算两个浮点操作（一次乘法一次加法）。
- FFMA 比先 FMUL 再 FADD 精度更高
  
  FMA 之所以叫「fused」，是因为硬件内部把乘法和加法合并成一步完成，中间结果不做舍入，只在最终输出时舍入一次。跟先算乘法再算加法（两次舍入）相比，fused 版本精度更高。这在科学计算和深度学习中都很重要。
- 相邻的乘法和加法会被自动合并为 FMA
  
  在实际的 GPU kernel 里，编译器会尽可能把相邻的乘法和加法合并成 FMA，因为 FMA 又快又省又精度高。你写 c = a * b + d，nvcc 编译器会自动生成 FFMA 指令，不需要你手动指定。只有那些真的凑不成乘加对的操作才会留下独立的 FADD 或 FMUL。
- FFMA、FADD、FMUL latency 大约都是 4 个时钟周期，没有快慢区别
  
  FFMA 并不是把一个乘法器和一个加法器串联起来。它是一个专门设计的单一电路，内部结构跟独立的乘法器或加法器都不一样。
  
  关键在于普通的 FMUL 做完乘法后要做一次舍入和规格化（把结果调整成标准浮点格式），这一步很耗时。
  
  FADD 在做加法前也要做对阶（把两个浮点数的指数对齐）。如果你串联 FMUL FADD，中间要经历：乘法→舍入→规格化→对阶→加法→舍入→规格化，整个流水线非常长。
  
  FMA 电路的优化是：乘法做完之后跳过中间的舍入和规格化，直接把全精度的中间结果（比标准浮点格式更宽的内部表示）送进加法逻辑，加法完成后才做一次舍入和规格化，省掉了中间那次舍入和规格化的延迟。这也是为什么它叫「fused」：乘法和加法在硬件层面融合成了一个不可分割的操作。
  
  最终结果是，在 NVIDIA GPU 上，FFMA、FADD、FMUL 的 latency 都是大约 4 个时钟周期，throughput 也一样。FMA 并没有比单独的加法或乘法更慢。这是因为 FMA 单元从一开始就是按照「一次完成乘加」来设计流水线深度的，而不是在已有的乘法器上面再接一个加法器。
流水线 latency 和 throughput 的区别

假设 FFMA 执行单元是 4 级流水线，那么一条 FFMA 指令从输入到结果出来需要 4 个周期（latency = 4）。

但是因为流水线的每一级在每个周期都可以接受新的指令，所以当流水线填满之后，每个周期都有一条指令完成（throughput = 1条/周期）。

所以说「每周期做一次浮点操作」描述的是 throughput，不是 latency。算峰值 FLOPS 的时候用的也是 throughput。

但是 latency 在写 kernel 的时候也很重要：如果一个 warp 的下一条指令依赖上一条 FFMA 的结果（数据依赖），那它必须等 4 个周期才能继续。这时候 warp scheduler 就会切换到另一个 ready 的 warp 去发射指令，用别的 warp 的工作来填充这 4 个周期的空档。

这就是为什么 GPU 需要大量的 warp 来掩藏 latency：如果一个 partition 上只有一个 warp，每条有依赖的指令都要空等 4 个周期，throughput 直接掉到 1/4。有足够多的 warp 轮流上，流水线才能保持满载，throughput 才能达到理论峰值。

CUDA Core

同时支持 FP32 或 INT32 的 CUDA Core

每个 partition 里有 16 个 FP32 core 和 16 个同时支持 FP32 或 INT32 的 core。
每 SM 每 cycle 完成的操作数

CUDA C Best Practices Guide 给出了数据类型以及相应的每 SM 每 cycle 完成的操作数：
- 16-bit floating-point add, multiply, multiply-add (2-way SIMD): 64
  - 16 x 4，一个 partition 只有 16 个 CUDA Core 在干活
- 32-bit floating-point add, multiply, multiply-add: 128
  - 32 x 4，一个 partition 32 个 CUDA Core 一起干活
- 32-bit integer multiply, multiply-add, extended-precision multiply-add: 64
  - 16 x 4，一个 partition 只有 16 个 CUDA Core 在干活
峰值性能

ga102 白皮书给出了以下类型的 peak TFLOPS/TOPS：
- FP16/BF16：37.4
  - 16 x 4 x 84 x 1740 M = 9.35 T，
  - FMA 一条指令两次操作，乘 2
  - 一条指令对两份数据操作得到两份结果（SIMD），再乘 2
- FP32：37.4
  - 32 x 4 x 84 x 1740 M = 18.7 T，
  - FMA 一条指令两次操作，再乘 2
- INT32：18.7
  - 16 x 4 x 84 x 1740 M = 18.7 T，
  - FMA 一条指令两次操作，再乘 2

Tensor Core：矩阵运算

Tensor Core 只支持矩阵运算

它不是传统意义上的标量 ALU，而是一个专门做小矩阵乘加运算的硬件单元。
支持的数据类型

CUDA Programming Guide 给出了支持的数据类型：
```
FP64  TF32  BF16  FP16  FP8  FP6  FP4  INT8  INT4
NO    YES   YES   YES   NO   NO   NO   YES   YES
```
Ampere 是第三代 Tensor Core，相比 Volta 和 Turing 新增了 BF16 和 TF32 数据格式的支持。TF32 比较特别，它用 FP32 一样的指数位宽但截断了尾数，让你可以不改代码（还是用 float 类型）就能享受接近 Tensor Core 的加速，只是精度略有损失。
峰值性能

ga102 白皮书给出了以下类型的 peak TFLOPS/TOPS：
- FP16 Tensor TFLOPS with FP16 Accumulate：149.7
  - 128 x 4 x 84 x 1740 M = 74.8 T，
  - FMA 一条指令两次操作，乘 2
- FP16/BF16 Tensor TFLOPS with FP32 Accumulate：149.7
- TF32 Tensor TFLOPS: 74.8
  - 64 x 4 x 84 x 1740 M = 37.4 T，
  - FMA 一条指令两次操作，乘 2
- INT8 Tensor TOPS: 299.3
  - 256 x 4 x 84 x 1740 M = 149.7 T，
  - FMA 一条指令两次操作，乘 2
- INT4 Tensor TOPS: 598.7
  - 512 x 4 x 84 x 1740 M = 299.3 T，
  - FMA 一条指令两次操作，乘 2

Tensor core 的硬件实现

脉动阵列的定义

数据被一个全局时钟「一拍一拍地泵」过整个处理单元网络，就像血液被心脏泵过全身一样。

一个严格的脉动阵列，要同时满足这几条：
- 规则的、只做局部连接的 PE 网格。一维线阵、二维方阵、六边形阵都行,但每个处理单元（PE）只跟它的直接相邻 PE 通信，没有全局互连。
- 每个 PE 内部很简单，通常就是一个乘累加（MAC）单元加几个本地寄存器。
- 数据在相邻 PE 之间按节拍同步流动，整个阵列在单一全局时钟下锁步（lockstep）运行，每「拍」所有 PE 同时计算一次、同时把数据往邻居推一格。
- 只有边界访存。数据从阵列边缘流入，在内部逐级传播，结果从边缘流出。内部 PE 永远不碰主存/大寄存器堆。
- 每个数据一旦从存储器取出，会在它流经的每个 PE 上被复用一次。这是 Kung 当年的核心卖点：用空间上的复用，绕开 von Neumann 访存瓶颈。
常见的两种脉动阵列
- 输出驻留：部分和留在 PE 原地累加，两个输入矩阵都流入
  
  这是最自然的脉动阵列。PE 网格的大小就是 mxn。从开始到右下角的元素最后一个算完，latency 为 k + m-1 + n-1。如果算上取出结果，从横轴或者纵轴取，还要再多花 m 或者 n 个 cycle。
- 权重驻留：XW=Y，权重 W 留在 PE，激活 X 流过，部分和流出。
  
  经典例子是 Google TPU v1 的 MXU。PE 网格大小是 kxn。每个 PE 的输出 = 上面 PE 输出 + 当前 PE 收到的 x 乘上当前 PE 上的权重 w。
  
  从开始到获得最后一个元素，latency 也是 m+n+k-2。如果算上装填 W，那么还需要 k 或者 n 个 cycle。
```
    * * *   *   *   *
  * * *     *   *   *
* * *       *   *   *

cycle 1    Y11
cycle 2    Y21 Y12
cycle 3    Y31 Y22 Y13
cycle 4        Y32 Y23
cycle 5            Y33
```
Tensor core 实际上不是脉动阵列

根据英伟达专利 US10338919B2，tensor core 实际上不是脉动阵列，而采用了点积引擎的设计。

根据 FP16/BF16 峰值数据，我们可以得到一个 tensor core 一个周期做 128 次 FMA。我们需要明确，tensor core 的实际点积引擎构造是很难获得的，而且我们也不需要获得，128 这个数据已经够用。但是我们可以进行一下合理的猜想，128 = 4x4x8。即一个 tensor core 有 16 个单元，每个单元有 8 个物理乘法器和一棵加法树 (Adder Tree)。

n 个周期内，8 个乘法器同时把长度为 8 的向量 A 和长度为 8 的向量 B 进行相乘，产生 8 个乘积。加法树接着把这 8 个乘积连同标量 C 瞬间加总成一个标量结果 D。共计 8 次乘法 8 次加法，即 8 次 FMA。8 x 4 x 4 = 128 次 FMA。

对应的矩阵乘法表达式应该是 4x8 x 8x4 + 4x4 = 4x4。mnk 为 4x4x8。

注意，n 个周期开始的时候，tensor core 的矩阵数据寄存器的值是当前要计算的矩阵 4x8 和 8x4。n 个周期要结束的时候，4x4 网格已经把这个结果算好了。tensor core 不用像常规脉动阵列那样，让矩阵的每个元素流过网格。

另外注意，流水线级数 n 英伟达并没有公开。

如何调用 tensor core

WMMA API
- load_matrix_sync
- store_matrix_sync
- fill_fragment
- mma_sync
PTX 指令
- WMMA 指令（黑盒）
  
  WMMA 指令直接对应 WMMA API。
  
  在这一套指令里，数据被装在一个叫做「不透明寄存器组（Opaque Register File）」里。你作为程序员，不需要知道（官方也刻意不让你知道）Warp 里的 0 号线程到底拿了矩阵的哪几个元素。
  
  你只需要提供一个全局内存（或共享内存）的指针，调用 wmma.load，硬件和编译器就会自动帮你把数据切碎、排布好，然后送进 Tensor Core。
- MMA 指令（手动准备矩阵数据）
  
  MMA 指令没有任何官方的高层 C++ API 对应（只能靠内联汇编手写）。
  
  MMA 需要你自己准备矩阵数据。它不会帮你做 load 的操作。你必须在调用 mma.sync 之前，极其精确地把数据塞进特定的寄存器里。
  
  f16 指令如下：
```
mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype d, a, b, c;

.alayout      = {.row, .col};
.blayout      = {.row, .col};
.dtype .ctype =  .f32, .f32
                 .f32, .f16
                 .f16, .f16

mma.sync.aligned.m16n8k8. row.col.f16.f16.f16.f16 d, a, b, c;
mma.sync.aligned.m16n8k8. row.col.f32.f16.f16.f32 d, a, b, c;
mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 d, a, b, c;
mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 d, a, b, c;
```

PTX 指令的矩阵大小支持

WMMA

WMMA 指令和 WMMA API 支持的矩阵大小是一模一样的。

f16:   16x16x16 8x32x16 32x8x16
bf16:  16x16x16 8x32x16 32x8x16
tf32:  16x16x8
u8/s8: 16x16x16 8x32x16 32x8x16
u4/s4: 8x8x32
b1:    8x8x128

MMA

f16:  8x8x4                              16x8x8 16x8x16
bf16:                                    16x8x8 16x8x16
tf32:                             16x8x4 16x8x8
u8/s8:      8x8x16                              16x8x16 16x8x32
u4/s4:             8x8x32                               16x8x32 16x8x64
b1:                       8x8x128                                       16x8x128 16x8x256

PTX 指令到 SASS 指令

我们只讨论 f16 MMA PTX 指令生成 sm_86 机器码，命令为 nvcc -arch=sm_86 -cubin -keep mma_f16.cu -o mma_f16.cubin。

对于 16x8x8 16x8x16 的 MMA PTX 指令，ptxas 会翻译成对应的 tensor core SASS 指令：

HMMA.16816.F16 R18, R4.reuse, R2.reuse, R8 ;  /* 0x000000020412723c */
                                              /* 0x0c4ff00000000808 */
HMMA.16816.F32 R20, R4, R2, R8 ;              /* 0x000000020414723c */
                                              /* 0x008fe20000001808 */
HMMA.1688.F16 R6, R4.reuse, R2.reuse, R8 ;    /* 0x000000020406723c */
                                              /* 0x0c0fe80000000008 */
HMMA.1688.F32 R12, R4, R2, R8 ;               /* 0x00000002040c723c */
                                              /* 0x000ff60000001008 */

对于 8x8x4 的 MMA PTX 指令，ptxas 会翻译成 CUDA core SASS 指令。

SASS 指令到 tensor core 的映射

只讨论 f16。8x8x4 这个大小忽略，这个形状是专门为 CC 7.0 设计的。这个大小在 CC 8.6 上可以跑，但是只会跑在 CUDA core 上。

f16 剩下的矩阵大小为 8x32x16 16x8x8 16x8x16 16x16x16 32x8x16，挑一个分析一下。其实主要问题是，一次只能计算小矩阵，大矩阵的计算 C = A x B 怎么分解成小矩阵的计算。

SASS 大小是 16x8x16，tensor core 大小假设是 4x4x8

比如大小是 16x8x16。假设 tensor core 的输出矩阵是 4x4，我们只能把 16x8 的结果 C，拆成 4x2 块，每一块 4x4。

在计算一块 4x4 时，我们把 A B 中对应的区域进行矩阵乘法。因为 k 是 16，而 tensor core 的 k 是 8，所以要把 A B 对应的区域拆成两部分，分两个周期算。

一共有 8 块 4x4，所以一共 8 x 2，16 个周期。

分块行为由硬件自动完成。

其实更方便的算法是，16x8x16 / 128（白皮书得来的数据）= 16。这样子就不用假设 tensor core 的构造。
SASS 大小是 16x8x8，tensor core 不假设大小

16x8x8 / 128 = 8。

注意，一个 SASS 指令从进入 tensor core 流水线到算出最后所有结果，除了上面算出来的周期数之外，还要加上 tensor core 本身的流水线级数。

Special Function Unit

专门做超越函数运算的，比如 sin、cos、exp、log、rsqrt 这些。这些函数如果用 ALU 去算需要很多个时钟周期做迭代，SFU 用硬件查找表加插值的方式几个周期就能出结果。

你写 CUDA 时调用 __expf() 或 __sinf() 这些 intrinsic 走的就是 SFU。

对 LLM 来说，GELU 和 Softmax 里的 exp 运算就会用到 SFU。

Load/Store Unit

ldst unit 严格说不算运算单元，它只负责地址计算和内存访问的执行，

kernel 每次读写 global memory、shared memory 都要经过这些单元。它们的数量直接影响 memory throughput。

参考资料

@INPROCEEDINGS{raihan,
  author={Raihan, Md Aamir and Goli, Negar and Aamodt, Tor M.},
  booktitle={2019 IEEE International Symposium on Performance Analysis of Systems and Software (ISPASS)},
  title={Modeling Deep Learning Accelerator Enabled GPUs},
  year={2019},
  volume={},
  number={},
  pages={79-92},
  keywords={Graphics processing units;Registers;Instruction sets;Computer architecture;Hidden Markov models;Deep learning;Tensor Core;Tesla Titan V;Turing RTX 2080;CUTLASS library;GPGPU-Sim},
  doi={10.1109/ISPASS.2019.00016}}

@misc{jiavolta,
      title={Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking},
      author={Zhe Jia and Marco Maggioni and Benjamin Staiger and Daniele P. Scarpazza},
      year={2018},
      eprint={1804.06826},
      archivePrefix={arXiv},
      primaryClass={cs.DC},
      url={https://arxiv.org/abs/1804.06826},
}

@misc{jiaturing,
      title={Dissecting the NVidia Turing T4 GPU via Microbenchmarking},
      author={Zhe Jia and Marco Maggioni and Jeffrey Smith and Daniele Paolo Scarpazza},
      year={2019},
      eprint={1903.07486},
      archivePrefix={arXiv},
      primaryClass={cs.DC},
      url={https://arxiv.org/abs/1903.07486},
}

@misc{ampere,
      title={Demystifying the Nvidia Ampere Architecture through Microbenchmarking and Instruction-level Analysis},
      author={Hamdy Abdelkhalik and Yehia Arafa and Nandakishore Santhi and Abdel-Hameed Badawy},
      year={2022},
      eprint={2208.11174},
      archivePrefix={arXiv},
      primaryClass={cs.AR},
      url={https://arxiv.org/abs/2208.11174},
}

内存子系统

A40 Memory Hierarchy

结构
- 1 个 partition 8 个 ldst unit
- 1 个 SM 1 个 128KB L1 data (texture) cache / Shared Memory
- 1 个 GPU 1 个 crossbar
- 1 个 GPU 12 个 512KB L2 cache
- 1 个 GPU 多片 DRAM
连接
- 4x8 个 ldst unit 连接 1 个 L1 data (texture) cache / Shared Memory
- 84 个 L1 data (texture) cache / Shared Memory 连接 12 个 L2 cache
- 12 个 L2 cache 通过 12 个 memory controller 访问 DRAM

LDST Unit

当一条 LDG 或 STS 指令带着 32 个线程的地址进入 LD/ST 单元时，它主要负责两件极其核心的工作：

地址合并与冲突检测（Memory Coalescing & Bank Conflict Resolution）：
- 如果是访问 Global Memory，它会检查这 32 个地址，看它们能不能被合并成少数几个 128-Byte 的宽内存请求。
- 如果是访问 Shared Memory，它会检查这 32 个地址是否落在了同一个 Bank 里（计算 Bank Conflict 的冲突惩罚）。
分发请求（Dispatching Requests）：处理完地址后，它将打包好的请求（Request）正式发送给 L1 Cache 控制器或 Shared Memory 硬件。

完成这两步之后，LD/ST 单元的工作就彻底结束了！它会立刻拍拍屁股走人，清空内部的流水线寄存器，准备迎接下一条访存指令。

SM L1 Cache

合并后的请求，第一站就是当前 SM 私有的 L1 Cache。在现代 NVIDIA 架构中，这个部件叫 L1TEX Cache（统一的 L1 数据缓存和纹理缓存）。

命中 (Hit)：如果在 L1 里找到了，数据直接顺着内部总线写回该 Warp 对应的寄存器（GPR）中，延迟极低（通常几十个时钟周期）。
缺失 (Miss)：如果没有找到，这个请求就必须走出 SM 的大门。

(注：Shared Memory 也在这一层，物理上与 L1 共享同一块 SRAM，但 LDG/STG 是针对全局内存的指令，所以只走 L1 路径，完全不经过 Shared Memory)。

Crossbar

芯片内部有一张庞大且极速的交叉开关网络，称为 Crossbar (全互连网络)。

由于 GPU 上有几十上百个 SM，而 L2 Cache 在物理上是被切分成 12 个 Slice (切片) 的。Crossbar 负责根据内存地址的哈希值，把 SM 吐出来的请求，精准路由到对应的那个 L2 Cache 切片上。

Device L2 Cache

请求抵达对应的 L2 切片。

L2 Cache 是全局共享的，这意味着哪怕是不同的 SM 上的线程，只要请求同一个地址，也能在 L2 中实现数据的共享。

全局内存的原子操作（Atomic Operations，如 atomicAdd）通常也是在 L2 这一层由专门的 ALU 直接完成的，这样可以避免数据在不同的 L1 之间来回颠簸（Cache Coherence 开销）。

命中 (Hit)：沿着 Crossbar 原路返回 L1 -> 寄存器。

缺失 (Miss)：向显存发起最终请求。

DRAM

带宽

Shared memory bandwith

A40 Shared memory bandwidth: 每个 SM 的 shared memory 是 32 个 bank，每个 bank 每 cycle 能提供 4 bytes（32 bit）。如果没有 bank conflict（32 个 thread 访问 32 个不同的 bank），那一次访问搬运的数据量是：32 banks × 4 bytes = 128 bytes/cycle/SM

A40 GA102 的 SM 时钟频率为 1.740 GHz，所以：128 bytes × 1.74 × 10^9/s ≈ 223 GB/s per SM

84 个 SM 加起来：223 × 84 ≈ 18 TB/s

对比 DRAM 的 696 GB/s，shared memory 的总聚合带宽是 DRAM 的近 27 倍。而且这 18 TB/s 不存在争抢问题：每个 SM 的 shared memory 是独立的，你用你的我用我的。

这就是为什么把数据搬进 shared memory 复用能突破带宽瓶颈的根本原因：不是因为快（虽然延迟确实低），而是因为总带宽大得多，而且是分布式的不用抢。
DRAM bandwidth:

A40 DRAM bandwidth: Memory interface 12 * 32-bit x Memory clock 14.5 Gbps（白皮书数据）= 696 GB/s

注意这是整块 GPU 所有 84 个 SM 共享的一根管子。平均到每个 SM 只有 696 / 84 ≈ 8.3 GB/s。这就是为什么 DRAM 带宽那么容易成为瓶颈。
DRAM 是半双工双向的

出于芯片引脚数量和主板走线面积的严苛限制，GPU 核心与显存颗粒之间的数据线（Data Pins）是读写共享的。

换句话说，GPU 的 DRAM（如 A40 上的 GDDR6）是「半双工双向」的（Half-Duplex Bidirectional），意思是可以读 DRAM，也可以写 DRAM，但是读写不能同时进行，要么是读操作，要么是写操作。

当显存控制器在「读状态」和「写状态」之间切换时，硬件需要空出几个时钟周期来清空总线上的电信号，这被称为 Read-to-Write Turnaround 延迟。因此，读写极度频繁交替的代码，实际能测到的有效总带宽会比纯读或纯写更低。

所以，一个 kernel 都在读显存时或者都在写显存时，理论最高能达到 696 GB/s。如果是又读又写，读写速度的总和理论上最高也只能达到 696 GB/s。