AI 基础设施：复习题

早年的节点名大致对应晶体管的某个物理尺寸，今天更像一代工艺的名字：IEEE Spectrum 2020 年的文章指出，“7 nm”晶体管的大部分关键尺寸远大于 7 nm，并介绍了 Wong 等人提出的 LMC 指标（逻辑密度、内存密度和两者之间的连接密度）。厂商的提升数字要先看基准：TSMC 说 N4 是 N5 的增强版，N4P 的“性能高 11%”是相对 N5；N2P 相对 N3E 是同功耗快 18% 或同速度功耗低 36%，两者一般不能同时拿到。NVIDIA 的 4N、4NP 是定制工艺，和 TSMC 公开的 N4 系列差在哪里，没有找到公开说明。

所以节点名只是起点，最后要比同一个 workload 下实测的性能、功耗和成本。真正卡住单块 die 面积的是一次曝光（exposure）的视场：26 mm × 33 mm，即 858 mm²，GH100 的 814 mm² 已经占了 95%。见芯片怎么造出来。

生成一个 token，每个权重只做一次乘加（2 FLOP），BF16 下却要读 2 字节，算术强度约 1 FLOP/字节；H100 SXM5 的屋脊点是 dense BF16 的 989.4 TFLOPS 除以 3,352 GB/s，约 295。7B 模型的 BF16 权重约 $1.4\times10^{10}$ 字节，读一遍至少 4.18 ms，单请求最多约 239 tokens/s，而这些运算本身只要 0.014 ms。

加大 batch 分摊的是权重：只算权重时，强度约等于 batch 大小 $B$；$B$ 小于约 295 时每步时间几乎不变，吞吐随 $B$ 线性上涨。分摊不了的是 KV：每条请求有自己的 KV，attention 这部分的强度只有约 4 FLOP/字节，和 batch 无关。按每条请求 4096 token 上下文算，$B=64$ 时每步下限已是 14.4 ms，吞吐上限约 4,440 tokens/s；$B=128$ 时权重加 KV 约 77 GiB。先碰到的是 KV 的带宽和容量，不是算力。见性能。

Occupancy 是一个 SM 上活跃的 warp 数，除以这个 SM 最多能驻留的 warp 数，由最先用完的那项资源决定。Compute capability 9.0（如 H100）每个 SM 最多 2048 个线程（64 个 warp）、32 个 block、65,536 个 32 位寄存器、228 KB shared memory。每个 block 256 个线程时：每线程 32 个寄存器能驻留 8 个 block，100%；64 个寄存器只剩 4 个，50%；每个 block 用 48 KB shared 也是 4 个，50%；80 个寄存器只剩 3 个，37.5%。寄存器还按每个 warp 256 个向上取整；实际先用 nvcc --resource-usage 看用量，再用 Nsight Compute 或 cudaOccupancyMaxActiveBlocksPerMultiprocessor 核对。

Occupancy 回答的是“有没有足够的 warp 可以切换”，不是“这个 kernel 快不快”。Best Practices Guide 写明，超过某个点再提高 occupancy 不会更快；高性能矩阵乘 kernel 常常每线程用很多寄存器、occupancy 不高，靠指令级并行和异步拷贝（Ampere 起的硬件异步拷贝、Hopper 的 TMA）把等待藏起来。延迟可以藏，带宽上限藏不掉。见 CUDA 和 GPU 架构。

Global memory 以 32 字节为单位传输。一个 warp 执行同一条 load 时，硬件把 32 个线程的地址合并，碰到几个 32 字节段就发几个 transaction。每个线程读一个 float32、起点对齐时：相邻线程读相邻元素，4 个段、128 字节全部用上；起点错开一个元素，5 个段，用上 80%；stride 32 时每个线程独占一段，搬 1024 字节只用上 128 字节，即 12.5%。按列读一个行优先的矩阵就是这种情况。

段数不等于耗时：Best Practices Guide 在 V100 上测“起点错开”的拷贝，带宽约为对齐时的 9/10，而不是按段数推出的 4/5，因为相邻 warp 会用上 cache 里多拉进来的数据。排查时看 Nsight Compute 的 sectors/request：32 位访问的理想值是 4，stride 32 时接近 32。改法是让 threadIdx.x 落在地址的最低维。见 CUDA。

Sparse 数字来自 Ampere 引入的 2:4 结构化稀疏：每 4 个值里最多 2 个非零时，Tensor Core 吞吐翻倍，普通 dense 模型拿不到这个倍数。H100 产品页的 BF16 1,979 TFLOPS 注明 with sparsity，白皮书 V1.04 的 dense 值是 989.4 TFLOPS。但不能一律除以 2：GB300 NVL72 页面的 FP4 是 1,440（sparse）对 1,080（dense）。

还要对齐三件事：SKU（同叫 H100，SXM 版 80 GB、3.35 TB/s，NVL 版 94 GB、3.9 TB/s）；计数单位（GB200 superchip 是 1 颗 Grace 加 2 颗 GPU，HGX 页面给 8 卡合计）；初步还是最终规格（2022 年 Hopper 博客的 1,000 TFLOPS、3,000 GB/s 是初步值）。对齐之后，从 V100 的 FP16 到 GB300 的 NVFP4，dense 峰值涨了 120 倍，显存带宽只涨约 9 倍，所以 decode 这类读得多、算得少的工作，代际提速主要跟着带宽走。官方页面之间也会对不上，例如 Rubin 的每卡显存带宽有 22 和 19.2 TB/s 两种写法。见 NVIDIA 各代 GPU。

MFU 由 PaLM 论文提出：观测到的 tokens/s，除以硬件按峰值算力运行时理论上能达到的 tokens/s。理论值只计前向和反向必需的运算（dense decoder 每个 token 约 $6N$，attention 里的矩阵乘再加 $12LHQT$），不计重算。HFU 把重算也算进去，依赖实现，不适合横向比较；论文报告 PaLM 540B 的 MFU 为 46.2%，HFU 为 57.8%。nvidia-smi 的利用率是另一回事：NVML 把它定义为采样周期里至少有一个 kernel 在跑的时间比例。

口径一变，数字就变。同一个教学算例（7B 模型、8 张 H100 SXM5、72,018 tokens/s）：只算 $6N$ 是 38.2%，加上 attention 是 44.1%，全量重算的 HFU 是 50.9%，误把 sparse 峰值当分母只剩 19.1%。报告时要写明是否计 attention、峰值是哪种精度、dense 还是 sparse、tokens/s 有没有算 padding。见性能。

Eager 模式下，y = torch.relu(x @ w + b) 是 3 个 ATen 算子、至少 3 个 kernel：每个算子都要走 Python 调用、dispatcher、分配输出和提交 kernel，逐元素算子还各自把中间结果从显存读写一遍。torch.compile 由 TorchDynamo 抓出 FX graph（带 guard，条件不满足就重编译；遇到不支持的 Python 会 graph break），AOTAutograd 提前抓下反向，TorchInductor 为 GPU 生成 Triton 代码并融合 kernel。

省下的不是乘法。按 4096×4096、BF16 和 2 TB/s 的教学带宽：eager 的 add 和 relu 共读写 128 MiB，约 67 µs；融合成一个 kernel 是 64 MiB，约 34 µs；bias 和 ReLU 都进 GEMM 的 epilogue 就没有额外读写。按 500 TFLOP/s 的教学算力，GEMM 本身约 275 µs，eager 的逐元素部分约相当于它的 24%。把 $M$ 降到 8，字节可以忽略，省的主要是 CPU 提交次数；reduce-overhead 用 CUDA Graphs 把一串提交合成一次。代价是首次编译时间、形状变化引起的重编译和 graph break。见框架。

默认的 kube-scheduler 一次只给一个 Pod 找节点，分布式训练却要所有 rank 到齐才能开始；GPU 又是只能取整数、不能超卖的 extended resource。手算一次：3 台节点各 8 卡；依次来了两个 4 卡任务、一个 16 卡训练 T16（2 个 Pod，各 8 卡）和一个 8 卡任务 T8。默认打分只看 CPU 和内存，两个 4 卡任务可能各占一台；T16 的一个 Pod 拿下第三台，另一个放不下，起来的一半占着 8 张卡空等；剩下两台各空 4 张，也拼不出 T8。24 张卡只有 8 张在干活；再来一个 T16，两个任务可能各占一台、互相等，就是死锁。

成组调度放不下就一个都不绑定，第三台留给 T8，16 张卡在干活；再按 GPU 装箱（MostAllocated 并给 GPU 加权），两个小任务挤进一台，24 张都在干活，轮到 T8 排队。实现方式不止一种：Kubernetes 核心的 gang scheduling 在 v1.37 是 beta、默认关闭；Kueue 用 waitForPodsReady 超时驱逐、重新排队；Volcano 用 PodGroup 的 minMember。另外，Slurm 文档里的 “gang scheduling” 指分时轮流，和这里的全有或全无不是一回事。见 Kubernetes。