GPU 硬件与资源共享

硬件架构 · CUDA 执行模型 · 数据搬运 · 资源共享 · 性能诊断 · 调度对比

Module Switcher

GPU 内容模块

内容模块

GPU硬件直觉

基础★★★⏱ 30 min

核心概念

CPU 定位：低延迟 · 强控制流 · 通用计算 — 分支预测、乱序执行、大缓存，适合 OS 调度、业务逻辑、IO 编排、串行控制
GPU 定位：高吞吐 · 数据并行 · 数值计算 — 大量轻量执行单元（SM）+ 高带宽显存（HBM），适合矩阵乘、卷积、Attention、batch 并行
CPU-GPU 协作：CPU 做数据准备/控制/调度/kernel launch/IO；GPU 做大规模数值计算/Tensor Core 加速
常见瓶颈：CPU↔GPU 数据搬运、显存带宽、kernel 并行度、多卡通信链路

对比表

维度	CPU	GPU
设计目标	低延迟、强控制流、通用计算	高吞吐、数据并行、数值计算
核心形态	少量复杂核心，每个核心能力强	大量轻量执行单元，组织在 SM 内
控制逻辑	分支预测、乱序执行、复杂流水线、大缓存	控制逻辑更简单，依赖大量线程隐藏延迟
时钟频率	通常更高，强调单线程响应	通常较低，强调整体吞吐
并行方式	线程或进程级并行，规模相对小	SPMD 模型，同一 kernel 映射到大量 thread
内存体系	DDR 容量大，带宽相对中等	HBM 容量较小但带宽极高
延迟隐藏	靠缓存、预测、乱序执行降低单线程等待	靠大量 warp 切换隐藏访存和执行延迟
适合任务	控制流复杂、依赖强、低延迟、IO 密集	矩阵乘、卷积、Attention、向量化和批量计算
不适合任务	海量规则数值计算吞吐不足	串行依赖强、分支复杂、OS/IO 控制任务

为什么 GPU 适合深度学习

深度学习特征	GPU 匹配点	面试表达
大量 GEMM / 卷积 / Attention	Tensor Core 和 CUDA Core 能并行执行大量乘加	模型主要算子可转成高吞吐矩阵计算
数据并行度高	batch、token、矩阵 tile 可以拆给大量线程	单个线程不强，但整体并发规模很大
显存带宽需求高	HBM 带宽远高于普通 DDR	参数、激活、KV cache 的读写需要高带宽
算法结构规则	同一 kernel 可在大量元素上执行类似逻辑	GPU 喜欢规则、可向量化、分支少的工作
可容忍单线程延迟	吞吐比单线程响应更关键	训练和批量推理更关心单位时间完成多少计算

可以把 GPU 的优势概括成三点：

算子形态匹配：深度学习核心算子大多是矩阵乘和张量运算。
并行粒度匹配：样本、token、矩阵块都能拆成大量相似任务。
资源需求匹配：模型参数、激活和 KV cache 对显存带宽要求高。

常见误区

误区	正确理解
GPU core 等于 CPU core	GPU 的 CUDA Core 更像 SM 内的轻量执行单元，不是独立运行复杂程序的通用核心。
GPU 一定比 CPU 快	只有任务有足够并行度、访存模式合理、数据搬运成本可控时，GPU 才有优势。
GPU 可以替代 CPU	GPU 不擅长 OS 调度、复杂控制流、IO 编排和低延迟串行任务，仍需要 CPU 作为 Host。
显存越大性能越强	显存容量决定能不能放下模型和数据，性能还要看带宽、SM、Tensor Core、互联和算子效率。
GPU 利用率高就代表训练快	`GPU-Util` 只是粗指标，还要看 SM Active、Tensor Core 利用率、显存带宽、通信和数据加载。

核心概念

概念	作用	面试抓手
SM（Streaming Multiprocessor）	GPU 的主要计算组织单元，内部包含 warp scheduler、CUDA Core、Tensor Core、寄存器和 shared memory	看 kernel 能不能把 SM 铺满
CUDA Core	常规数值执行单元，处理 FP32/INT 等标量或向量指令	不要直接类比 CPU core
Tensor Core	矩阵乘加专用硬件，支持 FP16/BF16/TF32/FP8/INT8 等不同精度	深度学习 GEMM、卷积、Attention 的核心加速点
HBM	高带宽显存，存放模型权重、激活、梯度、KV cache	Decode、embedding、elementwise 常受 HBM 带宽限制
L2 Cache	多个 SM 共享的缓存层	缓解重复访问，影响访存效率
NVLink / NVSwitch	GPU-GPU 高速互联和单机多卡交换	影响 tensor parallel、pipeline parallel、all-reduce
PCIe	CPU-GPU、GPU-NIC、部分 GPU-GPU 数据路径	H2D/D2H、GPUDirect RDMA、NUMA 亲和都和它相关

主流 GPU 对比

指标	A100 80GB	H100 80GB	H200 141GB
架构	Ampere	Hopper	Hopper
SM 数	108	132	132
FP16/BF16 Tensor Core 峰值	312 TFLOPS	989 TFLOPS	989 TFLOPS
显存	80GB HBM2e	80GB HBM3	141GB HBM3e
显存带宽	2.0 TB/s	3.35 TB/s	4.8 TB/s
NVLink 单向带宽	300 GB/s	450 GB/s	450 GB/s
TDP	400W	700W	700W

关键机制

GPU 性能通常由下面几类资源共同决定：

计算资源：SM 数量、Tensor Core 能力、时钟频率、支持的数据类型。
片上资源：寄存器、shared memory、L1/L2 cache，决定 occupancy 和数据复用能力。
显存资源：HBM 容量决定能放下多大模型和 batch，HBM 带宽决定 memory-bound 算子的上限。
互联资源：PCIe/NVLink/NVSwitch/RDMA 决定多卡和跨机通信效率。
功耗散热：训练集群里功耗墙、降频和散热也会影响实际吞吐。

常见误区

误区	正确理解
TFLOPS 越高训练一定越快	只有算子能用上 Tensor Core 且数据供给跟得上，峰值算力才有意义。
显存越大就是性能越强	容量解决“放不放得下”，带宽和算子效率决定“跑得快不快”。
CUDA Core 数量可以直接比较 GPU 性能	深度学习更要看 SM、Tensor Core、HBM、互联和实际 kernel 效率。
单卡强就代表多卡也强	多卡性能还受 NVLink/NVSwitch、PCIe、RDMA、NCCL 拓扑影响。

图 1：NVLink、PCIe、RDMA 三大互联技术带宽量级对比。NVLink 4.0 单卡聚合带宽 900 GB/s，约为 PCIe Gen5 x16 的 7 倍，约为单口 400G RDMA NIC 的 18 倍。

三大互联技术定位

技术	主要用途	典型通信范围	H100 典型带宽	训练中的角色
NVLink / NVSwitch	GPU-GPU 高速互联	单机内（或 NVLink Switch 扩展域）	900 GB/s 双向/GPU	单机内 AllReduce、张量并行、MoE All-to-All、Pipeline stage 通信
PCIe	通用 I/O 总线	CPU↔GPU、GPU↔NIC、GPU↔NVMe、部分 GPU-GPU P2P	128 GB/s 双向（Gen5 x16）	Host→GPU 拷贝、GPU→NIC、无 NVLink 时的 GPU-GPU fallback
RDMA / InfiniBand / RoCE	跨服务器远程直接内存访问	服务器之间	~50 GB/s（单口 400G NIC）	跨节点梯度同步、参数交换、分布式训练 collective

简单类比：

NVLink：GPU 之间的"高速内环路"——专用、超宽、超低延迟。
PCIe：服务器内部的"通用高速公路"——什么设备都能走，但车道少。
RDMA：服务器之间的"跨城专线"——距离远，但 NIC 可以直接读写对端 GPU 显存。

面试高频追问：为什么 NVLink 带宽比 PCIe 高这么多？

NVLink 是 NVIDIA 专为 GPU-GPU 设计的专有协议，物理层使用更多差分对（H100 有 18 条 NVLink 4.0 链路），不需要兼容通用 PCIe 设备枚举、地址空间映射等 I/O 开销；同时 NVSwitch 芯片提供 in-network reduction（SHARP）硬件加速，AllReduce 类 collective 可以在交换机芯片内完成累加，不需要每张 GPU 都收到完整数据再本地相加。PCIe 是通用 I/O 总线，要服务 CPU、GPU、NIC、SSD 等多种设备，协议栈更重，链路数量也受限于插槽形态。

NVLink 与 NVSwitch：单机 GPU-GPU 的高速通道

NVLink 演进

NVLink 自 2016 年 P100 推出以来已发展四代：

代数	代表 GPU	每条链路带宽	链路数/GPU	聚合带宽/GPU
NVLink 1.0	P100	20 GB/s 双向	4	80 GB/s
NVLink 2.0	V100	25 GB/s 双向	6	150 GB/s
NVLink 3.0	A100	50 GB/s 双向	12	600 GB/s
NVLink 4.0	H100/H200	50 GB/s 双向	18	900 GB/s
NVLink 5.0	B200/GB200	~100 GB/s 双向	18	1.8 TB/s

A100 SXM 单 GPU 有 12 条 NVLink 3.0，聚合 600 GB/s；H100 SXM 增加到 18 条 NVLink 4.0，聚合 900 GB/s，是 PCIe Gen5 x16（~128 GB/s 双向）的 7 倍。

NVSwitch：让所有 GPU 全互联

仅靠 NVLink 点对点连接（NVLink Bridge）只能支持两张 GPU 直连（如 PCIe 版本的 H100 双卡桥接）。要让 8 张 GPU 两两之间都有高带宽通路，需要 NVSwitch 交换芯片。

HGX H100 8-GPU 基板上集成了 4 颗第三代 NVSwitch 芯片，形成全互联无阻塞（non-blocking）交换矩阵：任意 GPU 对之间都能以 ~900 GB/s 的速度同时通信，不会争用带宽。

图 2：HGX H100 8-GPU 基板拓扑概念图。8 颗 GPU 通过 4 颗 NVSwitch 芯片形成全互联，每颗 GPU 的 18 条 NVLink 分布连接到 4 颗 NVSwitch，实现 all-to-all 无阻塞通信。CPU 仅负责控制面（kernel 启动、通信调度），不参与 GPU-GPU 数据搬运。

HGX A100 使用 6 颗第二代 NVSwitch（每颗 GPU 有 2 条 NVLink 连到每颗 NVSwitch），对分带宽（bisection bandwidth）2.4 TB/s；HGX H100 使用 4 颗第三代 NVSwitch（每颗 GPU 有 4~5 条 NVLink 连到 NVSwitch），对分带宽提升到 3.6 TB/s。

第三代 NVSwitch 的一个关键新特性是 in-switch reduction（SHARP 风格）：AllReduce 操作中的加法运算可以在 NVSwitch 芯片内的 ALU 上直接完成，减少数据在 GPU 间的往返，常见 collective（如 all-reduce）有效带宽相比 A100 提升约 3 倍。

PCIe：通用 I/O 总线

PCIe（Peripheral Component Interconnect Express）是服务器内部连接 CPU、GPU、NIC、NVMe SSD 的通用总线。它不是 GPU-GPU 通信的最优路径，但在没有 NVLink 时是 fallback 选择，同时承担 CPU-GPU、GPU-NIC 等关键数据路径。

PCIe 带宽计算

PCIe 带宽 = 代际速率 × lane 数 × 编码开销。常见 GPU 使用 x16 链路：

PCIe 版本	编码前每 lane 速率	x16 单向有效带宽	x16 双向合计	常见对应 GPU
PCIe Gen3	8 GT/s → ~1 GB/s/lane	~16 GB/s	~32 GB/s	V100 PCIe
PCIe Gen4	16 GT/s → ~2 GB/s/lane	~32 GB/s	~64 GB/s	A100 PCIe
PCIe Gen5	32 GT/s → ~4 GB/s/lane	~64 GB/s	~128 GB/s	H100 PCIe
PCIe Gen6	64 GT/s → ~8 GB/s/lane	~128 GB/s	~256 GB/s	B100/未来 GPU

PCIe 是树形拓扑：设备通过 PCIe switch 连接到 CPU root complex。两张 GPU 是否能高效通信，取决于它们在 PCIe 树中的位置：

同 PCIe switch 下的 P2P：路径最短，延迟和带宽相对最好。
跨 PCIe switch / 跨 root complex：路径变长。
跨 CPU Socket：需要经过 CPU 间互联（Intel UPI、AMD Infinity Fabric），延迟显著增加，抖动变大。
Host staging（最差情况）：GPU→CPU pinned memory→GPU，多一次拷贝，占用 CPU 内存带宽。

注意：PCIe P2P（Peer-to-Peer）需要硬件和驱动支持。即使两张 GPU 物理上在同一 PCIe switch 下，也可能因为 ACS（Access Control Services）、IOMMU 配置或驱动限制无法直接 P2P，此时会静默退化为 host staging。排查时可用 nvidia-smi topo -p2p 确认 P2P 是否可达。

RDMA 与 GPUDirect：跨机通信主力

跨服务器时 NVLink 无法直接延伸，必须通过 NIC 和网络。高性能 AI 训练集群使用 InfiniBand 或 RoCE（RDMA over Converged Ethernet） 实现 RDMA（Remote Direct Memory Access）。

RDMA 的核心优势

传统网络通信需要 CPU 参与数据搬运（socket buffer → kernel → NIC），RDMA 允许 NIC 绕过远端 CPU 直接读写远端内存，大幅降低延迟和 CPU 开销。400G InfiniBand（NDR）单口线速约 50 GB/s（400 Gb/s ÷ 8），一台服务器通常配 4~8 张 ConnectX-7 NIC，通过 multi-rail 并行提供数百 GB/s 的跨机总带宽。

GPUDirect RDMA：让 NIC 直接读写 GPU 显存

GPUDirect RDMA 是 RDMA 的关键增强：NIC 可以通过 PCIe 直接 DMA 读写 GPU HBM 显存，不需要先把数据拷贝到 CPU 内存再发送。

图 3：传统跨机通信路径 vs GPUDirect RDMA 路径对比。传统路径在发送端和接收端都需要经过 CPU 内存中转（host staging），共 4 次额外拷贝；GPUDirect RDMA 让 NIC 直接读写 GPU HBM，消除 host staging，CPU 仅负责控制面操作。

传统路径：GPU0 HBM → CPU DRAM → NIC → 网络 → NIC → CPU DRAM → GPU3 HBM（4 次拷贝，CPU 全程参与）。 GPUDirect RDMA 路径：GPU0 HBM → PCIe → NIC → 网络 → NIC → PCIe → GPU3 HBM（零 CPU 内存中转）。

这会带来：

延迟降低：省去 2~4 次内存拷贝。
CPU 开销降低：CPU 仅负责控制面（注册内存、提交 Work Request、轮询 Completion Queue），不参与数据搬运。
带宽提升：避免 CPU 内存带宽成为瓶颈，PCIe 带宽利用率更高。
支持 GPU-Direct Storage（GDS）：同样的思路让 NVMe 直接读写 GPU 显存。

一台 8 卡 H100/A100 服务器内，数据如何流转？

8 卡 SXM/HGX 形态（DGX/HGX 高端训练服务器）

这是大模型训练的主力机型。8 张 H100 SXM GPU 通过 NVSwitch 全互联：

计算数据主要在 HBM 中：模型参数、梯度、激活值、KV Cache 都在 GPU HBM 中（H100 SXM HBM 带宽 3.35 TB/s）。
GPU-GPU 通信走 NVLink/NVSwitch：GPU0 发梯度给 GPU3，数据从 GPU0 HBM → NVLink → NVSwitch → NVLink → GPU3 HBM，CPU 不拷贝数据。
CPU 负责控制面：启动 CUDA kernel、初始化 NCCL communicator、管理进程和内存注册，但不参与数据平面。
PCIe 仍然存在但不是主通道：负责 CPU-GPU 控制命令、Host 到 GPU 的初始数据加载、GPU 到 NIC 的跨机路径、GPU 到 NVMe 的存储 I/O。

8 卡 PCIe 形态（消费级/推理服务器）

PCIe 形态的 8 卡服务器没有 NVSwitch，GPU-GPU 通信依赖 PCIe 拓扑：

同 PCIe switch 下 P2P：路径短，带宽相对可接受，但远低于 NVLink。
跨 CPU Socket：经过 UPI/Infinity Fabric，延迟和抖动大幅增加。
Host staging：如果 P2P 不可用，退化为 GPU→CPU→GPU，性能最差。

这就是为什么 8 卡 PCIe 服务器训练大模型性能通常不如 SXM/NVSwitch 服务器——通信密集型任务（张量并行、MoE 专家并行、大 batch AllReduce）会被互联瓶颈严重限制。

单机 AllReduce 的数据路径

以 8 卡数据并行训练为例，NCCL 根据拓扑自动选择 ring、tree、CollNet、NVLS 等算法。在 SXM/NVSwitch 机器上：

每张 GPU 把梯度切成多个 chunk。
reduce-scatter 阶段：每个 chunk 沿 ring 或 tree 传播，在目标 GPU 上累加——走 NVLink/NVSwitch。
all-gather 阶段：累加后的结果再通过 NVLink/NVSwitch 分发给所有 GPU。
NVSwitch 的 in-switch reduction 可以在交换芯片内完成部分加法，减少 GPU 间数据往返。

跨机数据流转：分层通信

多机多卡训练采用分层通信策略：

单机内 GPU↔GPU

NVLink/NVSwitch（最高优先级，充分利用单机带宽）

本机 GPU↔NIC

PCIe（GPUDirect 直连，同 Socket 最优）

跨机 NIC↔NIC

InfiniBand/RoCE RDMA（多 NIC multi-rail 并行）

远端 NIC↔GPU

PCIe（GPUDirect 直连）

CPU 控制面

仅做内存注册、调度、进程管理，不搬运数据

两台 8 卡服务器做 AllReduce 的典型流程：

每台机器内部先 reduce-scatter：8 张 GPU 通过 NVLink/NVSwitch 做本机内归约，把高带宽用满。
跨机交换：每台服务器的 GPU 数据通过 GPUDirect RDMA 发到对端，多 NIC 并行。
每台机器内部再 all-gather：跨机同步完成后，通过 NVLink/NVSwitch 分发给本机所有 GPU。

这就是 NCCL 做 topology-aware 通信优化的核心思想：先用单机内最快的 NVLink/NVSwitch 做本地聚合，减少必须跨机传输的数据量。

NUMA 与 GPU-NIC 亲和性：容易被忽视的性能杀手

即使都配了 H100 + 400G RDMA，性能也可能差很多——关键在于 GPU 到 NIC 的 PCIe 路径。

理想情况：GPU 和 NIC 同 Socket / 同 PCIe switch

GPU → PCIe switch → NIC（同 Socket 下，路径短、延迟低、无跨 UPI）

糟糕情况：GPU 和 NIC 跨 Socket

GPU → PCIe switch → CPU Socket 0 → UPI/Infinity Fabric → CPU Socket 1 → PCIe root → NIC

跨 Socket 带来的问题：

延迟增加 1~2 μs 甚至更多。
有效带宽降低 30%~50%。
抖动变大（UPI 链路被其他流量竞争）。
CPU socket 间互联带宽被占用，影响其他任务。

调度多机训练任务时，不仅要保证 8 张 GPU 属于同一 NVSwitch 域，还要确保 每张 GPU 到其对应 NIC 的 PCIe 路径是同 Socket 的。nvidia-smi topo -m 输出中的 CPU Affinity 和 NUMA Affinity 列可以直接看到 GPU 与哪些 CPU 核亲和。

通信路径优先级总结

单机内 GPU-GPU（带宽从高到低）

优先级	路径	典型带宽	说明
1	NVLink / NVSwitch	900 GB/s	单机内最优 GPU-GPU 路径，全互联无阻塞
2	PCIe P2P（同 switch）	~50~100 GB/s	无 NVLink 时的较优 fallback
3	PCIe P2P（跨 root / 跨 Socket）	~20~50 GB/s	路径长，延迟和抖动增加
4	GPU → CPU pinned memory → GPU（host staging）	<20 GB/s	最差，占用 CPU 内存带宽，应尽量避免

跨机 GPU-GPU

优先级	路径	说明
1	GPU HBM → PCIe → NIC → RDMA 网络 → NIC → PCIe → GPU HBM	GPUDirect RDMA，理想路径，零 host staging
2	GPU → CPU 内存 → NIC → 网络 → NIC → CPU 内存 → GPU	Host staging，多 4 次拷贝，性能差 30%~50%

大小消息的敏感性差异：以上排序是大消息（tensor/gradient 传输）场景的常见优先级。小消息（如控制信号、barrier）对延迟更敏感，可能 NVLink P2P 比 NVSwitch 多跳延迟更低；实际最优路径还取决于 NCCL 算法选择、通信 overlap、网络拥塞状态等。

调度策略启示

单机 8 卡任务

最优：同一台 HGX/DGX 8 卡机器 + 同一 NVSwitch fabric + GPU-NIC 拓扑均衡。

不建议把一个强通信任务拆成「4 卡机器 A + 4 卡机器 B」——这会把本来可以走 NVLink 的通信变成 RDMA 跨机通信，带宽从 900 GB/s 降到 ~50 GB/s，性能可能差 10 倍以上。

4 卡任务

优先级：

同一 NVSwitch 域内的 4 张 GPU（最优）。
同一 PCIe switch / 同一 Socket 下的 4 张 GPU。
同机但跨 Socket 的 4 张 GPU。
跨机器 2+2 或 1+3（最差）。

多机 8×N 卡任务

调度器需要同时考虑：

每台机器内部是否是完整的 8 卡 NVSwitch 拓扑。
每张 GPU 到 NIC 的距离（NUMA 亲和性）。
NIC 数量和 rail 分布（multi-rail 平衡）。
多台机器是否在同一 leaf switch / 同一网络 pod。
RDMA 网络是否有拥塞（ECN/PFC 配置）。
NCCL topology file 或自动探测是否正确。

常见排查命令

目标	命令 / 工具	看什么
查看 GPU-GPU / GPU-NIC 拓扑	`nvidia-smi topo -m`	GPU 间是否 NVLink 互联（NV#），GPU 到 NIC 是 PIX/PXB/PHB/NODE/SYS 哪种路径，CPU/NUMA Affinity
查看 P2P 可达性	`nvidia-smi topo -p2p r` / `nvidia-smi topo -p2p w`	GPU 对之间是否支持 P2P 读写
查看 NCCL 选择的路径和算法	`NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,GRAPH`	ring/tree/NVLS/IB 路径是否符合预期，是否 fallback 到 TCP
查看 NVLink 状态和速率	`nvidia-smi nvlink -s`	每条 NVLink 的速率、是否 active、有无错误
查看 IB/RDMA 设备	`ibstat`、`ibv_devinfo`	端口速率、链路状态、HCA 能力、Firmware 版本
查看 PCIe 拓扑	`lspci -tv`	GPU、NIC、PCIe switch、root complex 层级关系
查看 NUMA 拓扑	`numactl --hardware`、`lstopo`	CPU、内存、PCIe 设备的 NUMA 亲和性
查看 Fabric Manager 状态	`systemctl status nvidia-fabricmanager`	NVSwitch 是否正常初始化、有无 SXid 错误

排查时不要只看「有没有 8 张 GPU」，要同时确认：

GPU-GPU 是否真的走 NVLink/NVSwitch（NV12/NV18 而非 SYS）。
GPU-NIC 是否同 Socket（PIX/PXB 优于 NODE/SYS）。
RDMA 是否启用 GPUDirect（NCCL 日志中不应出现 Trees are not used 或 TCP fallback）。
是否存在某个慢 rank、慢 NIC 或拥塞路径拖累整体（看 NCCL 日志中各 rank 的 ALG/CHANNEL 信息）。

nvidia-smi topo -m 输出示例解读

典型的 HGX H100 8-GPU 输出：

        GPU0    GPU1    GPU2    GPU3    GPU4    GPU5    GPU6    GPU7
GPU0     X      NV18    NV18    NV18    NV18    NV18    NV18    NV18
GPU1    NV18     X      NV18    NV18    NV18    NV18    NV18    NV18
GPU2    NV18    NV18     X      NV18    NV18    NV18    NV18    NV18
...（全矩阵 NV18，即每对 GPU 间 18 条 NVLink）

Legend 中关键缩写：

NV#：通过 # 条 NVLink 连接（最好）。
PIX：经过单个 PCIe bridge（PCIe 内最好）。
PXB：经过多个 PCIe bridge。
PHB：经过 PCIe Host Bridge（CPU root complex）。
NODE：跨 PCIe Host Bridge，同一 NUMA 节点内。
SYS：跨 NUMA 节点（经过 UPI/Infinity Fabric，最慢）。

面试回答

NVLink/NVSwitch 是 8 卡 H100/A100 服务器内 GPU-GPU 通信的主路径，H100 SXM 单卡聚合带宽 900 GB/s（18 条 NVLink 4.0 × 50 GB/s），约为 PCIe Gen5 x16（128 GB/s）的 7 倍；PCIe 是 CPU、GPU、NIC、NVMe 间的通用 I/O 总线，带宽低但通用性强；RDMA（InfiniBand/RoCE）是跨服务器通信机制，单口 400G NIC 约 50 GB/s 线速，配合 GPUDirect RDMA 让 NIC 直接读写 GPU HBM、绕过 CPU 内存中转，比传统 host staging 路径延迟低 30%~50%。单机内训练通信应优先走 NVLink/NVSwitch（NCCL 会做 topology-aware 分层 reduce），跨机时使用 GPUDirect RDMA 并保证 GPU-NIC 同 Socket 亲和，避免跨 UPI/Infinity Fabric 导致的带宽下降和抖动。

参考资料

内容模块

CUDA编程模型

系统链路

CPU Host

-> kernel<<<gridDim, blockDim>>>(args)

-> Grid：一次 launch 的全部 block

-> Block：调度到 SM 的基本任务包

-> Thread：程序员看到的逻辑并行单元

-> Warp：硬件调度的 thread 组，通常 32 个 thread

-> SM：执行 block、调度 warp、使用 CUDA Core/Tensor Core/HBM

先建立一张脑图：从一次 kernel launch 到 GPU 硬件执行

CUDA 执行主线是：CPU 端发起一次 kernel launch，CUDA runtime 把它描述成一个 grid；grid 里面有很多 block；block 里面有很多 thread；GPU 硬件把 block 分配到不同的 SM 上执行；SM 内部再把 thread 按 warp 组织和调度。CUDA 官方文档把 CUDA 描述为 NVIDIA 的并行计算平台和编程模型，用来让程序利用 GPU 的计算能力[CUDA Programming Guide]。

你写的是 kernel，启动的是 grid，组织单位是 block 和 thread，硬件执行单位是 SM 和 warp。

层级图：软件概念和硬件概念怎么对应

理解 CUDA 最容易混乱的地方，是把软件层级和硬件层级混在一起。grid / block / thread 是 CUDA 编程模型里的逻辑层级；GPU / SM / CUDA Core / Tensor Core / HBM 是硬件层级；warp 介于两者之间，它是硬件实际调度 thread 的基本方式。

CPU Host 代码
  |
  |  kernel<<<gridDim, blockDim>>>(args)
  v
CUDA Kernel Launch
  |
  v
Grid：一次 kernel launch 对应一个 grid
  |
  +-- Block 0  ---- threads: 0, 1, 2, ...
  +-- Block 1  ---- threads: 0, 1, 2, ...
  +-- Block 2  ---- threads: 0, 1, 2, ...
  |
  v
GPU 硬件调度
  |
  +-- SM 0 执行若干 blocks
  +-- SM 1 执行若干 blocks
  +-- SM 2 执行若干 blocks
       |
       v
     每个 block 内的 threads 被切成多个 warp
     一个 warp 通常包含 32 个 threads

注意这里的“对应”不是一一绑定。一个 grid 可以有很多 block；一个 SM 可以先后执行很多 block；一个 block 在运行期间通常只会被放到一个 SM 上；一个 SM 上也可能同时驻留多个 block。理解这一点后，很多性能问题就能解释清楚：block 太少会导致 SM 吃不饱，block 太大可能导致寄存器或 shared memory 压力过高，warp 内分支发散会降低执行效率。

GPU Core 到底是什么：不要把它等同于 CPU Core

面试里经常会问：“GPU 有很多 core，是什么意思？”这里最重要的是：不要把 GPU core 直接类比成 CPU core。

CPU core 通常是一个功能完整、控制能力很强的通用核心，擅长：

复杂控制流；
分支预测；
乱序执行；
低延迟响应；
单线程性能。

而 GPU 的所谓 “core”，通常指的是大量相对简单的计算执行单元。它们不是独立运行复杂程序的通用核心，而是被组织在 SM（Streaming Multiprocessor） 里面，用来服务大规模并行计算。

更准确地说：

NVIDIA GPU 的核心调度和资源管理单位是 SM，而不是单个 CUDA Core。

一个 SM 里面会包含多类执行和存储资源，例如：

CUDA Core：执行常规 FP32 / INT 等标量或向量计算；
Tensor Core：执行矩阵乘加类计算，常用于 GEMM、卷积、Attention；
Load / Store Unit：负责访存指令；
Special Function Unit：处理特殊数学函数；
Register File：保存线程寄存器状态；
Shared Memory / L1 Cache：供同一个 SM 内的线程块共享和缓存数据；
Warp Scheduler：以 warp 为单位调度指令。

所以讨论 GPU 性能时，不能只看“有多少 CUDA Core”。更重要的是看：

有多少 SM；
kernel 能不能把 SM 铺满；
每个 SM 里有多少 active warps；
访存是否连续、高效；
Tensor Core 是否被用起来；
显存带宽、L2 cache、通信是否成为瓶颈。

一个好记的类比：GPU 是工厂

可以把 GPU 想成一个工厂：

GPU 组件	类比	作用
GPU	整个工厂	承接大规模并行任务
SM	车间	GPU 的主要调度和执行单元
CUDA Core	普通工位	做常规数值计算
Tensor Core	矩阵乘专用产线	加速 GEMM、卷积、Attention 等
Warp Scheduler	车间调度员	决定哪个 warp 发射指令
Register File	工位旁边的小储物格	保存线程局部数据
Shared Memory / L1	车间内共享缓存	供同一 SM 内线程快速共享数据
L2 Cache	工厂级中转仓	多个 SM 共享的数据缓存
HBM / 显存	大仓库	存放模型参数、激活、输入输出数据

用这个类比来记：

SM 是车间，CUDA Core 是普通工位，Tensor Core 是矩阵乘专用产线。

所以真正关心的不是“工位数量”本身，而是：

任务是否能拆成足够多的并行工作；
数据是否能及时送到工位；
专用产线是否被用起来；
车间之间是否需要频繁等待或通信。

面试官常问

Q：GPU core 和 CPU core 有什么区别？

CPU core 是复杂的通用核心，擅长复杂控制流、分支预测、乱序执行和低延迟单线程任务。GPU core 更轻量，数量更多，主要服务于高吞吐的数据并行计算。NVIDIA GPU 里更核心的组织单位是 SM，CUDA Core、Tensor Core、load/store unit、register file、shared memory 等资源都组织在 SM 内。GPU 不是靠少数强核心跑得快，而是靠大量线程、warp 调度、高带宽显存和专用矩阵计算单元把吞吐做上去。

Q：为什么不能只看 CUDA Core 数量判断 GPU 性能？

因为 CUDA Core 数量只反映了一部分标量计算资源。实际性能还取决于 SM 数量、频率、显存带宽、L2 cache、Tensor Core 能力、kernel 并行度、occupancy、访存模式、warp stall、算子是否能用 Tensor Core 等。尤其在深度学习里，GEMM、卷积和 Attention 是否跑到 Tensor Core 上，往往比单看 CUDA Core 数量更关键。

Q：Tensor Core 和 CUDA Core 有什么区别？

CUDA Core 主要执行常规标量或向量计算，例如 FP32、INT 等指令；Tensor Core 是专门为矩阵乘加设计的硬件单元，适合 FP16、BF16、TF32、INT8 等矩阵计算。深度学习中的 GEMM、卷积、Attention 等如果能使用 Tensor Core，性能会显著提升。但 elementwise、索引、gather/scatter、部分归约类算子不一定能用到 Tensor Core。

Kernel 是什么：你写的一段 GPU 函数

kernel 是运行在 GPU 上的一段函数。普通 C/C++ 函数通常由 CPU 调用并在 CPU 上执行；CUDA kernel 则由 CPU 端发起，但实际在 GPU 上并行执行。你可以把 kernel 理解成“一份要在很多线程上同时执行的程序模板”。每个 thread 执行同一份 kernel 代码，但通过自己的 threadIdx、blockIdx 计算出不同的数据位置，从而处理不同元素。

__global__ void add_kernel(float* a, float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    c[i] = a[i] + b[i];
  }
}

int threads_per_block = 256;
int blocks = (n + threads_per_block - 1) / threads_per_block;
add_kernel<<<blocks, threads_per_block>>>(a, b, c, n);

这段代码里，add_kernel 是 kernel；<<<blocks, threads_per_block>>> 是 kernel launch 配置；blocks 决定 grid 里有多少个 block；threads_per_block 决定每个 block 有多少 thread。CUDA 编程模型要求开发者用 grid、block、thread 这样的层级组织并行工作，官方文档也强调理解 CUDA 编程模型有助于理解 GPU 如何执行代码[CUDA Programming Guide]。

学习者理解

不要把 kernel 理解成“启动一个线程”。一次 kernel launch 往往会启动成千上万个逻辑 thread。每个 thread 做的事情通常很小，比如处理一个数组元素、一个矩阵 tile 的一部分、一个 attention block 的一部分。GPU 的优势来自这些小工作被大规模并行执行。

Grid：一次 kernel launch 的全部工作

grid 是一次 kernel launch 产生的全部 block 的集合。你在代码里写的 kernel<<<gridDim, blockDim>>> 中，gridDim 描述 grid 的形状，可以是一维、二维或三维。为什么需要二维、三维？因为很多数据天然是二维或三维的，比如图像、矩阵、卷积特征图、三维仿真网格。用二维 grid 可以让代码更贴近数据结构。

一维 grid：处理向量
blockIdx.x * blockDim.x + threadIdx.x

二维 grid：处理矩阵
row = blockIdx.y * blockDim.y + threadIdx.y
col = blockIdx.x * blockDim.x + threadIdx.x

三维 grid：处理体数据或更复杂的张量切片
x/y/z 三个方向分别映射到数据维度

grid 是逻辑概念，不是说 GPU 上真的有一个叫 grid 的硬件。GPU driver/runtime 会把这个 grid 中的 block 交给硬件调度，硬件再把 block 分派到 SM 上执行。block 的数量通常应该足够多，这样不同 SM 都能拿到活干，某些 block 等待访存时，SM 还能切换到其他 warp 做计算。

面试官常问

Q：一次 kernel launch、grid、block 三者是什么关系？

可以这样答：一次 kernel launch 会启动一个 grid，grid 是这次启动的全部 block 集合；每个 block 又包含多个 thread。grid 和 block 是软件层面的并行组织方式，GPU 硬件会把 block 调度到 SM 上执行。调优时要让 grid 里有足够多的 block，否则 SM 数量再多也可能吃不满。

Block：资源分配和协作的关键边界

block 是 CUDA 里非常重要的边界。一个 block 里的 thread 可以通过 shared memory 共享数据，也可以用 __syncthreads() 做 block 内同步。不同 block 之间通常不能直接同步，也不能直接共享 shared memory。这个设计让 GPU 可以把不同 block 灵活分配到不同 SM 上，甚至用任意顺序执行，从而提高可调度性。

一个 block 在运行时会占用 SM 上的一部分资源，例如 thread slots、register、shared memory。一个 SM 能同时驻留多少个 block，取决于 block 大小、每个 thread 用多少 register、每个 block 用多少 shared memory，以及 GPU 架构限制。很多性能问题的根源就在这里：block 太小，单个 block 并行度不足；block 太大，资源占用过高，导致 SM 上同时驻留的 block/warp 变少。

怎么理解 block size

常见的 threads_per_block = 128、256、512 不是随便写的。因为 warp 通常是 32 个 thread，所以 block size 一般会取 32 的倍数，避免最后一个 warp 只有部分 thread 有效。比如 256 threads/block 就是 8 个 warp。这个值不是越大越好，要结合 register、shared memory、访存模式和 occupancy 分析。

面试官常问

Q：为什么不同 block 之间不能随便同步？

因为 CUDA 希望 block 之间尽量独立，这样硬件可以自由调度 block 到任意 SM 上。如果 block 之间需要频繁全局同步，调度复杂度和硬件成本都会上升。通常一个 kernel 内只做 block 内同步；如果需要全局同步，最常见的方式是拆成多个 kernel launch，因为不同 kernel 之间天然有顺序边界。当然 CUDA 也有 cooperative groups 等高级能力，但面试基础题先讲清楚 block 独立性即可。

Thread：逻辑上的最小并行工作单元

thread 是 CUDA 编程模型里最小的逻辑执行单元。每个 thread 有自己的 threadIdx，有自己的寄存器和局部变量。写 kernel 时，你通常会让每个 thread 根据自己的编号算出要处理的数据下标。

int global_id = blockIdx.x * blockDim.x + threadIdx.x;
if (global_id < n) {
  // 当前 thread 处理第 global_id 个元素
}

这里的 if (global_id < n) 很常见，因为总线程数通常会向上取整到 block size 的倍数，最后一个 block 可能有一些 thread 超出真实数据范围。这个边界判断能防止越界访问。

学习者理解

thread 不是越多越好。线程太少，GPU 吃不满；线程足够多以后，性能主要取决于访存是否连续、是否有分支发散、是否能复用数据、是否能用上 Tensor Core、是否被通信或同步拖慢。

Warp：硬件真正调度的一组 thread

warp 是理解 GPU 执行效率的核心概念。CUDA 里你写的是 thread，但 NVIDIA GPU 硬件通常按 warp 调度执行，一个 warp 通常包含 32 个 thread。也就是说，SM 不是完全独立地一个 thread 一个 thread 执行，而是把一组 thread 作为调度单位。

这会带来两个非常重要的性能现象。第一是 分支发散：如果同一个 warp 里的不同 thread 走了不同的 if/else 分支，硬件通常需要分批执行不同路径，等价于一部分 thread 先闲着，另一部分执行，然后再反过来，所以效率会下降。第二是 访存合并：如果同一个 warp 里的 thread 访问连续内存，硬件可以更高效地合并内存请求；如果访问很散，带宽利用率会变差。

理想情况：同一个 warp 的 threads 访问连续地址
thread 0 -> a[0]
thread 1 -> a[1]
thread 2 -> a[2]
...
thread31 -> a[31]

较差情况：同一个 warp 的 threads 访问分散地址
thread 0 -> a[0]
thread 1 -> a[1024]
thread 2 -> a[17]
...
thread31 -> a[99999]

面试官常问

Q：什么是 warp divergence？为什么会影响性能？

可以这样答：warp divergence 指同一个 warp 内的 thread 因为条件分支走了不同路径。由于 warp 是硬件调度单位，不同路径往往需要被串行执行，导致部分 thread 暂时 inactive，所以有效并行度下降。优化时要尽量让同一个 warp 内 thread 执行相似路径，或者把分支改写成更规则的数据布局和计算方式。

SM：block 被放到哪里执行

SM 是 Streaming Multiprocessor，是 GPU 上最关键的计算组织单元。一个 GPU 有多个 SM；一次 kernel launch 产生很多 block；GPU 会把这些 block 分配到各个 SM 上。每个 SM 内部有 warp scheduler、register file、shared memory、L1/cache 相关结构、CUDA Core、Tensor Core 等资源。

一个 block 被调度到某个 SM 后，它里面的 thread 会被切成多个 warp，由 SM 内部的 warp scheduler 调度执行。当某个 warp 等待 HBM 访存时，SM 可以切换去执行另一个 ready 的 warp，用大量并发来隐藏访存延迟。这也是为什么 GPU 需要很多 thread/warp：不是每个时刻所有 thread 都在算，而是通过足够多的可调度 warp 把硬件流水线填满。

学习者理解

SM 像车间，block 像一个被派到车间的任务包，warp 像车间班组，thread 像班组里的工人。车间要高效运转，需要任务包足够多、班组排班合理、材料供应及时、不要有太多工人因为分支或等待内存而闲着。

面试官常问

Q：block 和 SM 是什么关系？

block 是 CUDA 编程模型里的逻辑任务块，SM 是 GPU 硬件里的执行单元。运行时，block 会被调度到 SM 上执行；一个 block 在运行期间通常不会跨多个 SM；一个 SM 可以同时驻留多个 block，具体数量受 thread 数、register、shared memory 等资源限制。

CUDA 内存层级：为什么“数据怎么搬”比“怎么算”还重要

很多 GPU 性能问题不是算力不够，而是数据搬不动。CUDA 程序里常见的内存层级包括 register、shared memory、L1/L2 cache、global memory/HBM。越靠近计算单元，速度越快、容量越小、使用约束越多；越远离计算单元，容量越大、延迟越高。

按距离理解

register 是每个 thread 私有的最快存储，适合放局部变量，但使用太多会降低 occupancy。shared memory 是 block 内 thread 共享的片上存储，适合做 tile 和数据复用。L2 cache 是多个 SM 共享的缓存，可以缓解重复访问。global memory / HBM 容量最大，模型参数、activation、KV cache 大多在这里，但访问延迟和带宽压力也最大。

面试官常问

Q：为什么 shared memory 能加速矩阵乘？

因为矩阵乘有大量数据复用。如果每个 thread 都直接从 HBM 反复读取元素，会浪费显存带宽。更好的做法是把矩阵分块，把一个 tile 先搬到 shared memory，block 内多个 thread 复用这块数据，再进行多次乘加。这样可以减少 HBM 访问，把更多时间花在计算上。

把概念串起来：一个向量加法 kernel 怎么跑

假设我们要计算 c[i] = a[i] + b[i]，数组长度是 1,000,000。CPU 端决定每个 block 256 个 thread，于是需要大约 3907 个 block。一次 kernel launch 会创建一个包含 3907 个 block 的 grid。GPU 运行时把这些 block 分派给多个 SM。每个 SM 同时驻留若干 block；每个 block 的 256 个 thread 会被切成 8 个 warp；warp scheduler 选择 ready 的 warp 执行 load、add、store。

向量加法的执行链路：

1. CPU 调用 add_kernel<<<3907, 256>>>(a, b, c, n)
2. CUDA runtime 创建一个 grid，里面有 3907 个 block
3. GPU 把 block 分配给不同 SM
4. 每个 block 的 256 个 thread 被组织成 8 个 warp
5. 每个 thread 计算自己的 global_id
6. 每个 thread 从 HBM 读取 a[i] 和 b[i]
7. CUDA Core 执行加法
8. 结果写回 c[i]
9. 所有 block 完成后，这次 kernel 结束

这个例子也说明：不是所有 GPU kernel 都能很好利用 GPU。向量加法每个元素只做一次加法，却要读两个数、写一个数，所以很容易是 memory-bound；矩阵乘每读入一块数据可以做很多乘加，更容易提高算术强度并利用 Tensor Core。

面试标准题：CUDA kernel 的执行模型是什么？

标准回答：CUDA kernel 由 CPU Host 端发起，一次 kernel launch 会创建一个 grid。grid 由多个 block 组成，block 由多个 thread 组成。block 是被调度到 SM 上执行的基本单位；thread 是程序员看到的逻辑执行单元；硬件实际通常以 warp 为单位执行，一个 warp 通常包含 32 个 thread。SM 内部的 warp scheduler 会选择 ready warp 发射指令，通过让多个 warp 同时驻留来隐藏访存延迟。

层级	属于什么	作用	面试重点
Host	CPU 端	分配显存、准备参数、发起 kernel launch	`kernel<<<gridDim, blockDim>>>(args)`
Grid	CUDA 逻辑层级	一次 kernel launch 的全部 block 集合	grid 不是硬件，而是任务组织方式
Block	CUDA 逻辑层级 + 调度边界	一组 thread，调度到某个 SM 上执行	block 内可用 shared memory 和 `__syncthreads()`
Thread	CUDA 逻辑执行单元	运行同一份 kernel 代码，处理不同数据	用 `blockIdx`、`blockDim`、`threadIdx` 计算全局下标
Warp	硬件执行单位	通常 32 个 thread 一组执行同一条指令	关注 warp divergence 和 memory coalescing
SM	GPU 硬件执行单元	承载 block，调度 warp，提供 register/shared memory/CUDA Core/Tensor Core	关注 occupancy、warp stall、register pressure、shared memory

Host 启动 kernel，kernel 形成 grid，grid 拆成 block，block 调度到 SM，block 内 thread 被组织成 warp，SM 以 warp 为单位执行并通过多 warp 并发隐藏内存延迟。

Q: CUDA kernel 的执行模型是什么？

回答主线

从 kernel launch 开始讲：CPU Host 发起 kernel，CUDA runtime 创建 grid；grid 中有多个 block；block 中有多个 thread；block 被调度到 SM 上执行；SM 把 thread 组织成 warp，并由 warp scheduler 调度 ready warp 执行。

为什么 block 是关键边界

block 内 thread 可以共享 shared memory，也可以用 __syncthreads() 同步。不同 block 默认不能直接同步，因为 block 可能被调度到不同 SM，执行顺序也不确定。如果需要全局同步，通常拆成多个 kernel launch。

为什么 warp 是性能关键

硬件实际按 warp 调度。一个 warp 通常 32 个 thread，同一 warp 内如果分支不同，会出现 warp divergence；如果访存地址连续对齐，可以 memory coalescing，显著提升有效带宽。

如何联系性能优化

理解执行模型后，优化方向就很清楚：让 grid/block 足够多以铺满 SM；控制 register/shared memory 避免 occupancy 过低；减少 warp divergence；让 global memory 访问 coalesced；用 shared memory 做数据复用；能用 Tensor Core 的算子尽量满足 dtype 和 shape 条件。

面试口径：CUDA 执行模型的关键不是“线程很多”，而是 grid/block/thread 的逻辑组织如何映射到 SM/warp 的硬件执行。

面试官怎么问：从概念题到排障题

如果面试官问 CUDA 层级，通常不是为了考你背定义，而是想确认你能不能从执行模型解释性能现象。建议回答时使用这个顺序：先讲 kernel launch，再讲 grid/block/thread，再讲 warp/SM，再讲内存层级，最后联系性能问题。

Q: 请你解释一次 CUDA kernel launch 之后 GPU 上发生了什么。

一次 kernel launch 会从 CPU host 端发起，形如 kernel<<<gridDim, blockDim>>>(args)。这次 launch 对应一个 grid，grid 由多个 block 组成，每个 block 又由多个 thread 组成。GPU 会把 block 调度到 SM 上执行；block 内 thread 被组织成 warp，warp 是硬件调度的重要单位。每个 thread 运行同一份 kernel 代码，但通过 blockIdx 和 threadIdx 计算不同数据下标，从而并行处理不同数据。

答题结构：CPU launch → grid → block → thread → block 调度到 SM → thread 组成 warp → 并行处理数据。

Q: grid、block、thread、warp、SM 哪些是软件概念，哪些是硬件概念？

grid、block、thread 是 CUDA 编程模型中的逻辑概念，是程序员组织并行任务的方式。SM 是 GPU 硬件上的计算单元。warp 是硬件实际调度 thread 的重要执行单位，连接了逻辑 thread 和硬件执行。回答时可以补充：block 会被调度到 SM 上，thread 会被组成 warp 执行，一个 SM 上可以驻留多个 block/warp。

Q: 为什么 block size 通常设置成 32 的倍数？

因为 NVIDIA GPU 通常以 warp 为单位调度 thread，一个 warp 通常是 32 个 thread。如果 block size 不是 32 的倍数，最后一个 warp 可能只有部分 thread 有效，造成执行资源浪费。当然 block size 不是只看 32 的倍数，还要看 register 使用、shared memory 使用、occupancy、访存模式和具体 kernel 的计算特征。

Q: 什么是 occupancy？是不是越高越好？

occupancy 可以粗略理解为 SM 上活跃 warp 数相对理论最大 warp 数的比例。它反映 SM 有多少可调度工作可以用来隐藏访存或执行延迟。但 occupancy 不是越高越好：如果一个 kernel 已经受 HBM 带宽限制，提高 occupancy 未必有用；如果为了提高 occupancy 降低单线程寄存器使用，反而导致重复访存或指令变多，也可能变慢。所以 occupancy 是诊断指标，不是最终目标。

Q: 为什么同一个 warp 内分支不同会变慢？

因为 warp 是硬件调度单位。同一个 warp 里的 thread 如果走不同分支，硬件通常需要分别执行不同路径：走 A 分支时，走 B 分支的 thread 暂时 inactive；再执行 B 分支时，A 分支的 thread 暂时 inactive。这样虽然逻辑上还是并行程序，但有效并行度下降了，这就是 warp divergence。

Q: 如果一个 CUDA kernel 很慢，你会从哪些层级分析？

我会按执行链路分层分析。第一看 launch 配置：grid/block 是否足够让 SM 吃满；第二看 warp 层面：是否有严重分支发散、访存是否连续；第三看 SM 资源：register 和 shared memory 是否限制 occupancy；第四看内存层级：是否频繁访问 HBM、是否可以用 shared memory 做复用；第五看计算路径：是否用上 Tensor Core、shape 是否适合高性能 kernel；最后看端到端是否被 CPU 数据准备、kernel launch overhead 或多卡通信拖慢。

常见误区

误区	正确理解
kernel launch 等于启动一个线程	一次 launch 会创建一个 grid，包含大量 block/thread。
grid/block/thread 是硬件实体	它们是 CUDA 逻辑层级，SM/CUDA Core/Tensor Core 才是硬件资源。
thread 是硬件实际独立调度单位	NVIDIA GPU 通常以 warp 为单位调度，warp 内 thread 分支不同会降低效率。
block 可以跨多个 SM 执行	一个 block 在运行期间通常驻留在一个 SM 上。
occupancy 越高越好	occupancy 是诊断指标，不是最终性能目标。

系统链路

硬件视角：基本单位是 SP（Streaming Processor，也叫 CUDA Core），每个 SP 有自己的 register 和 local memory（片下内存，应对寄存器不足），只能被自己访问；多个 SP 和一块 shared memory 构成一个 SM，shared memory 被 SM 内线程共享；多个 SM 和一块全局内存构成 GPU，global memory 被所有线程访问。

软件视角对应关系：

软件	硬件	可见的内存
thread	SP	私有 register + local memory
block	SM	block 内共享 shared memory，可用原子操作和 barrier 同步
grid（device）	GPU	所有 thread 共享 global memory

关键机制

block size 选择：

范围通常是 1 到 1024。
取 32（warp 大小）的倍数，避免最后一个 warp 只有部分线程有效。
考虑 occupancy：block size 要让 SM 上能挂足够多 warp。
考虑 register 数量：block 太大可能占用过多寄存器，降低同时驻留的 block 数。
常见起点是 128、256、512，但最终要 profile。

grid size 选择：

Element-wise 程序通常让 grid size = 数据量 / block size 向上取整，并保证 block 数足够多铺满所有 SM。

int block = 256;
int grid = (n + block - 1) / block;   // 向上取整
kernel<<<grid, block>>>(...);

## 资源模型

指标	定义	含义
SM 利用率 occupancy	有效活跃线程数 / SM 最大线程数	单个 SM 内 warp 是否够多，能否隐藏访存延迟
GPU 利用率 utilization	有效 SM 数 / 总 SM 数	kernel 是否把所有 SM 都铺满

两者结合才能判断 GPU 是否真用满：block 数太少会导致 GPU utilization 低，即 SM 吃不满；block 内 warp 太少或寄存器占用过高会导致 occupancy 低，即单个 SM 隐藏延迟能力差。 ## 内存墙内存墙指处理器速度与内存访问速度的不匹配：处理器算力增长远快于内存访问速度，导致内存成为瓶颈。在 CUDA 里，内存墙通常指 global memory 的高延迟、相对低带宽会拖慢 GPU。缓解手段是利用 shared memory 把频繁访问的数据缓存到片上，减少对 global memory 的访问次数，例如矩阵乘 tiling 和 FlashAttention 的分块。这也是很多 GPU kernel 优化的核心思路：不是算得更快，而是少搬数据。 ## PyTorch 自定义 CUDA 算子三个步骤： 1. 编写 CUDA 算子和对应的 launch 调用函数。 2. 编写 torch cpp 函数建立 PyTorch 与 CUDA 的联系，用 pybind11 封装。 3. 用 PyTorch 的 cpp 扩展库编译调用。

方式	说明	入口
JIT 编译	Python 运行时再编译 cpp/cuda 文件	`from torch.utils.cpp_extension import load`
SETUP 编译	setup.py 提前编译	`from torch.utils.cpp_extension import BuildExtension, CUDAExtension`
CMAKE 编译	编译生成 .so，运行时加载	`torch.ops.load_library("build/libxxx.so")` → `torch.ops.xxx.torch_launch_xxx()`

## Tensor Core 数据路径 Tensor Core 是专门做矩阵乘加（GEMM、卷积、Attention）的硬件单元。在 Volta / Turing / Ampere 架构上，计算所需输入输出数据都放在与 CUDA Core 共享的寄存器上；在 Hopper 架构上，为了获得更好的带宽，计算所需输入数据可以直接存放在共享内存上。 ## 常见误区

误区	正确理解
Occupancy 越高越好	Occupancy 只是隐藏延迟的条件，过高可能牺牲寄存器和数据复用。
block size 越大越好	block 太大会增加资源占用，降低同时驻留 block/warp 数。
shared memory 一定更快	shared memory 快但容量有限，还要注意 bank conflict 和同步开销。
global memory 慢就完全不用	大数据必须在 global/HBM 中，优化目标是提高复用和访问连续性。

Q: 为什么 block size 常取 128 / 256 / 512？

一是必须是 32 的倍数（warp 大小），避免尾部 warp 浪费；二是要能整除 SM 最大活跃线程数以拿到高 occupancy（主流架构公约数 512）；三是 block 太大占用寄存器过多会减少同时驻留 block 数，所以常折中到 128 或 256。最终值要结合寄存器/shared memory 用量和实际 profiling 决定。

Q: 什么是内存墙，CUDA 里怎么缓解？

内存墙是处理器算力增长远快于内存访问速度造成的瓶颈。CUDA 里 global memory 延迟高、带宽相对有限，容易成为瓶颈。缓解办法是把频繁复用的数据搬到 shared memory（片上）做缓存复用，减少 global memory 访问次数，例如矩阵乘 tiling 和 FlashAttention 的分块计算。

诊断入口

GPU timeline 有空洞或 step time 抖动

-> 看 DataLoader / CPU preprocessing / I/O

-> 看 Nsight Systems 的 Memcpy HtoD / DtoH / stream sync

-> 查 PyTorch 里的 .item() / .cpu() / .numpy() / aten::to

-> 看 pinned memory、non_blocking、prefetch、stream 重叠是否生效

-> 多 GPU 场景检查是否绕 CPU，优先 NCCL / NVLink / GPUDirect

指标解释

现象	可能原因	典型证据
GPU-Util 低且 timeline 有空洞	DataLoader、CPU preprocessing 或 H2D 太慢	kernel 之间有大空洞、HtoD 前 CPU 等待
CPU 线程频繁阻塞	D2H 或显式同步	`cudaStreamSynchronize`、`.item()`、`Memcpy DtoH`
拷贝很多但带宽低	大量小 tensor 小拷贝	timeline 上许多短小 memcpy
多 GPU 扩展差	数据绕回 CPU 或拓扑差	PCIe 流量异常、NCCL 路径不理想

Host-to-Device / Device-to-Host Copy 是什么

CPU 与 GPU 数据搬运的核心是识别触发点、使用 pinned memory、合并小拷贝，并通过 profiler 诊断同步与带宽损失。

Host-to-device copy（H2D） 是把数据从 CPU 内存拷贝到 GPU 显存；device-to-host copy（D2H） 是把数据从 GPU 显存拷贝回 CPU 内存。

术语	含义	典型资源	方向	常见场景
Host	主机端	CPU、系统内存	-	DataLoader、Python 进程、业务服务
Device	设备端	GPU、GPU 显存	-	CUDA kernel、模型前向/反向、推理引擎
H2D	CPU 内存到 GPU 显存	Host memory -> Device memory	CPU -> GPU	把 batch 输入拷到 GPU
D2H	GPU 显存到 CPU 内存	Device memory -> Host memory	GPU -> CPU	`.cpu()`、`.numpy()`、`.item()`、日志和后处理

H2D 是 CPU 内存到 GPU 显存，D2H 是 GPU 显存到 CPU 内存。

为什么叫 Host 和 Device

CUDA / GPU 编程里，一般把 CPU 侧叫 Host，把 GPU 侧叫 Device。Host 负责控制流程、数据读取、任务提交和 kernel launch；Device 负责执行大规模并行计算。

Host memory   = CPU 内存
Device memory = GPU 显存

H2D: Host memory   -> Device memory
D2H: Device memory -> Host memory

这个命名也解释了为什么 CUDA API、Profiler 和 Nsight 里经常看到 Memcpy HtoD、Memcpy DtoH。

训练流程里的 H2D / D2H

训练时 H2D 通常不可避免，因为样本先在 CPU 侧被读取和预处理，GPU 计算前必须把 batch 放到显存里。D2H 则很多时候可以减少，尤其是频繁把 loss、metric、中间 tensor 拷回 CPU。

1. DataLoader 从磁盘读取图片/样本
2. CPU 做 decode / augmentation / tokenize
3. batch 放在 CPU 内存
4. H2D: batch 从 CPU 拷贝到 GPU
5. GPU 前向计算
6. GPU 反向传播
7. GPU 更新参数
8. 可选 D2H: loss / metric / output 拷回 CPU 打日志或后处理

操作	是否常见	是否可优化	说明
batch 输入 H2D	很常见	可加速/重叠	用 pinned memory、non_blocking、prefetch
loss.item() D2H	很常见	可减少频率	每 step 多次调用会触发同步和拷贝
tensor.cpu().numpy()	常见	应谨慎	会把 GPU tensor 拷回 CPU，常导致同步
GPU -> CPU -> GPU	性能反模式	应避免	能在 GPU 上做就不要绕回 numpy

为什么数据拷贝会影响性能

GPU 计算很快，但 CPU 内存和 GPU 显存之间要经过 PCIe、NVLink 等互联。它们通常比 GPU 内部 HBM 访问慢，也比 GPU 内部计算更容易成为流水线瓶颈。如果每个 step 都频繁做 CPU -> GPU -> CPU -> GPU，GPU 就可能在等数据，而不是在计算。

现象	可能原因	典型证据
GPU utilization 低	DataLoader 或 H2D 太慢	timeline 中 kernel 之间有空洞
step time 抖动	CPU 预处理、磁盘 I/O 或拷贝不稳定	Profiler 里 DataLoader 时间波动
CPU 等 GPU	D2H 触发同步	`cudaStreamSynchronize`、`Memcpy DtoH`
带宽利用差	大量小 tensor 小拷贝	很多短小 `Memcpy HtoD`

性能问题的关键不是“有没有拷贝”，而是拷贝是否频繁、是否小而碎、是否阻塞、是否能和计算重叠。

优化 1：减少不必要的 D2H

D2H 操作经常比你想象中更隐蔽。PyTorch 里的 .item()、.cpu()、.numpy()、打印 GPU tensor，都可能把 GPU 上的数据拷回 CPU，并触发同步。

触发操作	问题	优化方式
`loss.item()`	GPU loss 拷回 CPU，可能同步	每 N step 记录一次，或聚合后再取
`tensor.cpu()`	显式 D2H	只在必要输出/保存时调用
`tensor.numpy()`	需要先回 CPU	尽量用 PyTorch GPU tensor 操作替代 numpy
`print(gpu_tensor)`	为了显示内容可能触发同步	少打印，打印 shape/dtype/device 等元信息

# 不推荐：每步都强制 D2H + 同步
loss_value = loss.item()

# 更好：降低频率，或多个指标聚合后再拷贝
if step % 100 == 0:
    loss_value = loss.detach().item()

优化 2：Pinned Memory + Non-blocking Copy

Pinned memory，也叫 page-locked memory，表示这块 CPU 内存固定在物理内存里，不会被操作系统换出。GPU DMA 从 pinned memory 做 H2D 拷贝更高效，也更容易和计算异步重叠。

loader = DataLoader(
    dataset,
    batch_size=128,
    shuffle=True,
    num_workers=8,
    pin_memory=True,
    persistent_workers=True,
    prefetch_factor=2,
)

for batch, label in loader:
    batch = batch.to("cuda", non_blocking=True)
    label = label.to("cuda", non_blocking=True)
    output = model(batch)

参数	作用	注意点
`pin_memory=True`	DataLoader 把 batch 放到 pinned CPU memory	占用不可换出的物理内存，不能无限开
`non_blocking=True`	允许异步 H2D	通常需要 pinned memory 才更有效
`num_workers`	多进程并行读数据/预处理	过大可能 CPU 争用或内存压力大
`prefetch_factor`	提前准备后续 batch	增大可提升流水线，但会占更多内存
`persistent_workers`	跨 epoch 保持 worker	减少 worker 重启开销

优化 3：让拷贝和计算重叠

理想训练流水线是：GPU 正在计算 batch N，CPU 同时准备 batch N+1，H2D 同时拷贝 batch N+1。这样 GPU 不需要在 batch 边界长时间等待数据。

理想流水线：

时间轴 ---->
CPU/DataLoader:  prepare batch N+1   prepare batch N+2
H2D copy:              copy N+1           copy N+2
GPU compute: compute batch N     compute batch N+1

如果 timeline 里看到 GPU compute 结束后才开始 H2D，并且中间有明显空洞，说明数据准备或拷贝没有很好重叠。

优化 4：批量拷贝，减少小拷贝

大量小 tensor 分别 .to("cuda") 会产生很多小 H2D copy，每次拷贝都有固定开销，带宽利用率也差。更好的方式是先在 CPU 侧合并成 batch，再一次性拷贝到 GPU。

# 不推荐：很多小 H2D
for x in list_of_tensors:
    x = x.to("cuda")

# 更好：先合并，再一次 H2D
batch = torch.stack(list_of_tensors)
batch = batch.to("cuda", non_blocking=True)

优化 5：避免 CPU/GPU 来回转换

下面这种链路非常浪费：

x = x.cuda()
y = x.cpu().numpy()
z = torch.tensor(y).cuda()

它会产生 GPU -> CPU -> GPU，既浪费带宽，又可能触发同步。原则是：能在 GPU 上做的计算就留在 GPU 上；能用 PyTorch GPU tensor 操作，就不要中途转 numpy。

优化 6：预处理流水线和数据格式

很多训练瓶颈不在 H2D 本身，而在 H2D 前面的 CPU decode、augmentation、tokenizer、磁盘 I/O 和小文件读取。优化方向包括更快的数据格式、更好的预取、更少的小文件和更靠近 GPU 的预处理。

瓶颈	优化方式
图片 decode / resize / crop / normalization 慢	NVIDIA DALI、GPU augmentation、离线预处理、cache 结果
小文件太多	WebDataset、TFRecord、mmap、顺序读取格式
Tokenizer 慢	批量 tokenizer、预 tokenize、cache token ids
磁盘/网络存储慢	本地 NVMe cache、数据预热、分布式缓存

多 GPU 场景：避免 Host 中转

多 GPU 通信最好直接走 GPU-GPU 或 GPU-NIC-GPU，而不是绕回 CPU。否则会出现低效路径：

差：GPU0 -> CPU -> GPU1
好：GPU0 -> GPU1
好：GPU -> NIC -> GPU

机制	作用	典型场景
NVLink / NVSwitch	节点内 GPU-GPU 高带宽互联	TP、AllReduce、P2P copy
NCCL	多 GPU collective 通信库	DDP、ZeRO、张量并行
GPUDirect RDMA	GPU 显存和 NIC 直接通信，减少 CPU 参与	跨节点训练、RDMA 网络
P2P copy	GPU 间直接拷贝	同机多卡数据交换

推理服务里的拷贝优化

在线推理也会被拷贝拖慢，尤其是请求小、batch 小、后处理重、输出频繁回 CPU 的场景。优化时要减少输入输出 buffer 抖动，让预处理、推理、后处理尽量流水线化。

请求批处理，合并小输入。
输入预处理尽量批量化或放到 GPU。
输出后处理尽量在 GPU 上完成。
使用 zero-copy 或共享内存减少业务进程和推理进程之间的数据复制。
避免频繁把中间 tensor 拷回 CPU。
使用 TensorRT 管理输入输出 buffer，减少反复分配和拷贝。
使用 CUDA Graph 降低固定 shape 推理的 launch overhead。

如何判断是不是数据拷贝瓶颈

如果 GPU utilization 低、CPU utilization 高、DataLoader 慢、每个 step 中间有明显空洞，就要怀疑数据加载或 H2D/D2H 拷贝瓶颈。

工具	看什么
Nsight Systems	timeline 中的 `Memcpy HtoD`、`Memcpy DtoH`、kernel 空洞、stream 同步
PyTorch Profiler	DataLoader 时间、`aten::to`、`cudaMemcpyAsync`、CPU/GPU timeline
nvidia-smi / DCGM	GPU 利用率、显存、PCIe/NVLink 相关指标
torch.cuda.Event	测量 GPU 侧时间，区分 CPU wall time 和 GPU elapsed time
TensorBoard Profiler	训练 step 拆解和输入 pipeline 分析

重点关注：
Memcpy HtoD
Memcpy DtoH
cudaStreamSynchronize
cudaDeviceSynchronize
kernel 之间的大空洞

Q: 什么是 host-to-device / device-to-host copy？怎么优化？

标准回答

Host 指 CPU 侧，Device 指 GPU 侧。Host-to-device copy 是把数据从 CPU 内存拷贝到 GPU 显存，例如训练时把 batch 输入拷到 GPU；device-to-host copy 是把 GPU 上的结果拷回 CPU，例如 .cpu()、.numpy()、.item()、日志和后处理。

为什么影响性能

这类拷贝需要走 PCIe、NVLink 等互联，相比 GPU 内部计算和 HBM 访问慢很多。而且很多 D2H 操作会触发同步，破坏 CUDA 异步执行，让 CPU 等 GPU 或 GPU 等数据。

优化思路

第一，减少不必要的拷贝，尤其避免频繁 .cpu()、.numpy()、.item()；第二，用 pinned memory 和 non_blocking=True 加速 H2D；第三，用 DataLoader 多 worker、prefetch、persistent workers 让数据准备和 GPU 计算重叠；第四，尽量把预处理和后处理放在 GPU 上；第五，合并小 tensor 做批量拷贝；第六，多 GPU 场景使用 NCCL、NVLink、GPUDirect，避免 GPU 通信绕回 CPU。

少拷贝、批量拷贝、异步拷贝、拷贝计算重叠，避免 GPU -> CPU -> GPU 来回搬。

核心概念

Stream 和 Event 用来把 H2D/D2H 拷贝、kernel 与 CPU 准备组织成异步流水线。

GPU 有自己的显存（HBM），CPU 也有自己的内存（DDR）。数据在两者之间搬运，就是 H2D（Host to Device）和 D2H（Device to Host）拷贝。

因为 H2D/D2H 走的是 PCIe 或 NVLink，带宽远低于 GPU 内部 HBM 带宽。一次不注意的同步拷贝可能让 GPU 空等数百微秒，这在推理低延迟场景下尤其致命。

路径	典型带宽	延迟	常见场景
CPU → GPU（PCIe 4.0 x16）	~25 GB/s（双向 ~50 GB/s）	~10 μs 起步	输入数据、模型加载
GPU → CPU（PCIe 4.0 x16）	~25 GB/s	~10 μs 起步	读取推理结果、日志
GPU → GPU（NVLink 3.0）	~300 GB/s	~5 μs	张量并行、NCCL 通信
GPU → GPU（PCIe）	~25 GB/s	~10 μs	无 NVLink 的卡间通信

系统链路

同步拷贝（cudaMemcpy）：CPU 发起后阻塞等待完成，GPU 和 CPU 都在等。适合初始化、模型加载等不在乎延迟的场景。

异步拷贝（cudaMemcpyAsync）：CPU 发起后立即返回，拷贝在 GPU 端异步执行。必须配合 CUDA stream 使用，才能实现计算和拷贝重叠。

// 同步：CPU 阻塞等待
cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice);

// 异步：CPU 立即返回，拷贝在 stream 上执行
cudaMemcpyAsync(d_ptr, h_ptr, size, cudaMemcpyHostToDevice, stream);

## 关键机制

策略	原理	效果
Pinned Memory（页锁定内存）	用 `cudaMallocHost` 分配 CPU 端不可换页内存，DMA 直接传输	异步拷贝带宽可提升 2-3×
计算与拷贝重叠	用多个 CUDA stream，一个算当前 batch，另一个拷下一个 batch	隐藏拷贝延迟
减少不必要的 D2H	尽量让数据留在 GPU，减少中间结果回传 CPU	减少拷贝次数和延迟
批量传输	多次小拷贝合并为一次大拷贝	减少启动开销
Unified Memory	`cudaMallocManaged` 让 CPU/GPU 共享地址空间，按需迁移	编程简单，但性能不如手动管理

Pinned Memory 为什么更快： 普通 malloc 分配的内存是可换页的，GPU DMA 无法直接访问，需要先拷贝到临时 pinned buffer 再传输，多了一次隐式拷贝。cudaMallocHost 分配的内存被锁定在物理页上，GPU DMA 可以直接传输，省掉中间环节。但 pinned memory 不可换页，会占用实际物理内存，不能无限分配。 CUDA stream 是 GPU 上的任务队列。同一个 stream 内的操作按提交顺序串行执行；不同 stream 之间的操作可以并行执行（如果硬件资源允许）。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// stream1：拷贝输入 + 计算
cudaMemcpyAsync(d_in1, h_in1, size, cudaMemcpyHostToDevice, stream1);
kernel_a<<<grid, block, 0, stream1>>>(d_in1, d_out1);

// stream2：同时拷贝另一组输入 + 计算
cudaMemcpyAsync(d_in2, h_in2, size, cudaMemcpyHostToDevice, stream2);
kernel_b<<<grid, block, 0, stream2>>>(d_in2, d_out2);

// 等待两个 stream 都完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

## 资源模型推理或训练中，经典优化是让数据拷贝和计算重叠。三缓冲（triple buffering）是最常见的模式：

Buffer A: GPU 正在计算（kernel 执行中）
Buffer B: H2D 正在拷贝下一批数据（DMA 传输中）
Buffer C: CPU 正在准备再下一批数据

下一个时间步：
  A 计算完成 → D2H 取结果
  B 拷贝完成 → 开始计算
  C 准备完成 → 开始 H2D 拷贝

这样 GPU 不用等 CPU 准备数据，CPU 也不用等 GPU 算完，三者流水线化。实现要点： - 每个缓冲区用独立的 CUDA stream。 - 输入数据用 pinned memory 分配。 - 用 cudaMemcpyAsync 而不是 cudaMemcpy。 - 用 cudaStreamSynchronize 或 event 等待完成。 - 缓冲区数量通常 2 或 3，多了收益递减。 ## 常见误区

误区	正确理解
用了多个 stream 就一定并行	并行取决于硬件资源、依赖关系、默认 stream 语义和 kernel 是否占满 SM。
`cudaMemcpyAsync` 一定异步	H2D 异步拷贝通常需要 pinned memory；pageable memory 可能触发隐式同步或 staging copy。
默认 stream 可以直接用于重叠	默认 stream 可能和其他 stream 隐式同步，做流水线应使用自定义 stream 和 event。
stream 越多越好	stream 过多会增加调度复杂度，资源冲突时收益有限。

Q: 默认 stream（stream 0）和自定义 stream 有什么区别？

默认 stream（NULL / 0）是同步流——它会和所有其他 stream 隐式同步。如果你把操作提交到默认 stream，它会等所有其他 stream 完成；其他 stream 也会等它完成。自定义 stream 之间没有这种隐式同步，所以可以实现真正的并行。要做计算拷贝重叠，必须用自定义 stream。

Q: 两个 stream 一定能并行吗？

不一定。并行取决于硬件资源是否足够。如果两个 stream 的 kernel 都需要全部 SM，它们会被串行执行；如果一个在拷贝（使用 DMA engine / copy engine），一个在计算（使用 SM），它们可以并行，因为用的是不同硬件单元。这也是为什么计算+拷贝重叠比计算+计算重叠更容易实现——DMA engine 和 SM 是独立的硬件资源。

内容模块

资源共享与调度

一条完整主线

业务需求

独占、强 SLA、开发测试、细粒度显存或算力

隔离目标

故障、显存、SM、带宽、时延分别要隔离到什么程度

选择机制

整卡 / MIG / MPS / Time-Slicing / HAMi

暴露资源

Device Plugin、GPU Operator 或 HAMi 向 Kubelet 注册资源

调度放置

kube-scheduler 或 HAMi Scheduler 选择节点和物理 GPU

运行时执行

MIG 硬件实例、MPS server、CUDA 时间片或 HAMi-Core 执行限制

监控兜底

DCGM、进程显存、P99、Xid、OOM 与共享密度

先分清五个层次

层次	机制	真正被切分的对象	隔离强度
设备直通	整卡 / PCIe Passthrough	整张物理 GPU	最强，但粒度最粗
硬件分区	MIG	GPU Instance / Compute Instance，对应 SM、显存切片和缓存等硬件资源	强
驱动级并发	MPS	多个 CUDA client 通过 MPS server 共享执行资源	中等，仍共享故障域和部分硬件
时间复用	Time-Slicing	多个进程交错获得 GPU 执行时间	弱，不提供显存和性能硬隔离
K8S 软件虚拟化	HAMi	用调度账本分配显存/算力，再在容器内执行软限制	取决于设备后端，NVIDIA 常见路径是软件隔离

核心机制对比

机制	K8S 如何表达	能保证什么	不能保证什么	典型场景
MIG	每个 MIG profile 作为独立扩展资源，例如 `nvidia.com/mig-1g.10gb`	固定规格的硬件资源与故障隔离，性能更稳定	不能动态借用相邻实例的空闲资源；profile 组合会产生碎片	强 SLA 推理、多租户生产、稳定小训练
MPS	官方 Device Plugin 按 `replicas` 暴露共享访问，MPS daemon 管理 client	多进程 kernel 并发；可限制 client 的执行资源和显存	不等于 MIG 级隔离；共享故障域，干扰仍需监控	可信 workload、小 kernel、多路推理
Time-Slicing	Device Plugin 为每张卡创建多个逻辑引用	让多个 Pod 获得同一张 GPU 的共享访问权	`replicas=N` 不代表固定 `1/N` 算力，也不隔离显存	Notebook、开发测试、低优离线任务
HAMi	`nvidia.com/gpu` + `nvidia.com/gpumem` + `nvidia.com/gpucores`	细粒度显存/算力请求、设备感知调度、异构加速器统一管理	软件限制不自动获得 MIG 的硬件故障隔离；能力依赖设备后端	私有云、多团队研发、细粒度 vGPU、国产异构设备
CUDA VMM	不是 K8S 资源类型，由应用或运行时调用 CUDA Driver API	虚拟地址与物理显存页分离，支持按需映射和弹性显存池	不负责多租户算力隔离，也不是 GPU 调度器	KV Cache、模型权重缓存、弹性显存

为什么默认 Kubernetes 不够

默认 NVIDIA Device Plugin 把 GPU 注册为整数扩展资源，例如 nvidia.com/gpu: 8。kube-scheduler 主要知道“还有几份资源”，不知道：

两个逻辑资源是否来自同一张物理 GPU。
当前还剩多少真实显存、SM 和 HBM 带宽。
两个 workload 放在一起会产生多大干扰。
一个共享进程崩溃是否会影响同卡其他进程。

因此，GPU 虚拟化至少需要两部分：

资源表达和放置：Device Plugin、GPU Operator、HAMi Scheduler 等告诉 Kubernetes 能分配什么。
执行和隔离：MIG、MPS、CUDA driver、HAMi-Core 等真正执行硬件或软件限制。

只改 scheduler 的资源数量，不会自动得到显存、算力和故障隔离。

选型决策树

是否要求强故障隔离和稳定 SLA

是 -> 整卡或 MIG

是否需要细粒度显存/算力任意配比

是 -> HAMi，或厂商 vGPU 方案

是否是可信 workload 且希望 kernel 并发

是 -> MPS

是否只想让低利用率任务共享访问

是 -> Time-Slicing

是否主要解决 KV Cache / 权重显存弹性

是 -> CUDA VMM，不要把它当 GPU 算力虚拟化

Q: MIG、MPS、Time-Slicing 和 HAMi 最本质的区别是什么？

核心含义

MIG 在硬件层切设备，MPS 在驱动层并发执行，Time-Slicing 在时间上共享访问，HAMi 在 Kubernetes 调度和容器运行时层做细粒度软件配额。

判断标准

不要只背“隔离强弱”，还要追问隔离对象：MIG 能提供固定硬件实例；MPS 可以限制 client 的执行资源和显存但仍共享故障域；Time-Slicing 只提供共享访问；HAMi 的显存/算力限制依赖软件拦截或设备后端能力。

Q: GPU 虚拟化是不是就是把一张卡显示成很多张卡？

不是。“显示成多份”只是资源表达。真正要看底层有没有硬件实例、执行资源限制、显存限制和故障隔离。Time-Slicing 可以把一张卡上报成多个逻辑 slot，但这些 slot 仍然共享整张卡；MIG 的每一份则对应真实硬件实例。

常见误区

误区	正确理解
`replicas: 4` 就是每个 Pod 固定获得 25% GPU	它只表示 4 个共享访问名额；Time-Slicing 不保证固定算力和显存份额。
MPS 和 MIG 都是“切卡”	MIG 创建硬件实例；MPS 让多个 client 通过 server 并发使用同一 GPU。
HAMi 等于开源 MIG	HAMi 是 Kubernetes 虚拟化中间层，NVIDIA 常见实现是软件调度与容器内限制，不等于硬件 MIG。
CUDA VMM 能隔离多个租户的 SM	VMM 管虚拟地址和物理显存映射，不负责 SM 公平或故障隔离。
kube-scheduler 会自动避免共享干扰	默认 scheduler 不理解真实 SM/HBM 干扰，需要扩展调度、画像或在线监控。

资料来源

先把四个“Runtime”概念分开

名词	典型实现	在 GPU 接入中的职责
CRI 容器运行时	`containerd`、CRI-O	接收 kubelet 创建 Pod/Container 的请求，管理镜像、sandbox 和容器生命周期
低层 OCI Runtime	`runc`、`crun`	按照 OCI spec 真正创建 Linux namespace、cgroup 和进程
NVIDIA Container Toolkit / Runtime	`nvidia-ctk`、`nvidia-container-runtime`、CDI hook/spec	把 GPU device node、驱动库和必要环境注入容器；它不负责 Kubernetes 调度
Kubernetes RuntimeClass	`runtimeClassName: nvidia`	让 Pod 选择 containerd/CRI-O 中某个 runtime handler；它只是 Kubernetes 选择入口，不是另一套容器引擎

GPU Operator v25.10.0 及以后默认启用 CDI（Container Device Interface），用于标准 GPU workload 的设备注入。通过 Device Plugin 正常申请 nvidia.com/gpu 的 Pod 通常不需要显式写 runtimeClassName；旧版或手工 legacy runtime 路径如果没有把 NVIDIA runtime 设为默认，则需要配置并选择 RuntimeClass。不要把不同版本的教程拼成一份配置。

一张图看懂“系统识别 GPU”

主机 PCIe 枚举 GPU

lspci 能看到 NVIDIA 设备

NVIDIA Kernel Driver 绑定设备

nvidia-smi、/dev/nvidia* 正常

Container Toolkit 配置 containerd/CDI

GPU 能被正确注入容器

Device Plugin 向 kubelet 注册

/var/lib/kubelet/device-plugins/ 下建立 gRPC socket

kubelet 上报健康设备

Node Capacity/Allocatable 出现 nvidia.com/gpu

kube-scheduler 选择节点

Pod limits 请求整数 GPU

Device Plugin Allocate

返回 device、mount、env 或 CDI device

容器运行 CUDA canary

vectorAdd/Test PASSED

DCGM/告警接管

持续观测 Xid、ECC、温度、功耗与利用率

这里最重要的边界是：

Driver 让宿主机内核和 CUDA 用户态能够访问 GPU。
Container Toolkit/CDI 让容器能获得设备节点和匹配的宿主机驱动库。
Device Plugin 让 Kubernetes 看见、分配和健康管理 GPU。
GPU Operator 自动部署和维护前面这些组件，本身不是 CUDA 数据面。
NFD/GFD 负责硬件与 GPU 属性标签，帮助 Operator 和 scheduler 识别节点类型。
DCGM Exporter 负责指标，不参与设备分配。

NVIDIA GPU Operator 将 GPU 节点软件栈从手工安装转为自动化管理的对比图

NVIDIA 官方 GPU Operator 概念图：左侧需要分别维护 Driver、Runtime、Device Plugin 和监控，右侧由 Operator 统一部署和收敛。该图用于理解责任边界；生产版本与组件兼容性仍以当前 GPU Operator Platform Support 为准。图片来源：NVIDIA Technical Blog。

接入前检查：先隔离，再部署

1. 节点基础面

新节点首先要完成普通 Kubernetes worker 的基线：

固件、BIOS、BMC、GPU 型号和 PCIe 拓扑符合机器验收单。
OS、kernel、containerd、kubelet 版本符合当前集群支持矩阵。
hostname、DNS、NTP、镜像仓库、软件源和控制面网络可达。
kubelet 的 cgroup driver 与 containerd 配置一致。
如果需要 RDMA/GPUDirect，再单独核对 NIC、OFED、IOMMU、NUMA 和拓扑；不要把“GPU 能跑”当成“多机训练网络已就绪”。

NODE=gpu-node-01

kubectl get node "$NODE" -o wide
kubectl describe node "$NODE"

# 先阻止普通业务进入，canary Pod 会单独容忍这个 taint
kubectl taint node "$NODE" gpu-bootstrap=true:NoSchedule --overwrite

如果节点还没有加入集群，应先走平台既有的 kubeadm、Cluster API 或云厂商节点池流程。不要把包含 token、CA hash 的一次性 kubeadm join 命令写死到公共 Runbook。

2. 宿主机硬件与驱动基线

lspci -nn | grep -i nvidia
uname -r
lsmod | grep '^nvidia'
nvidia-smi -L
nvidia-smi --query-gpu=index,uuid,name,pci.bus_id,driver_version,memory.total \
  --format=csv
nvidia-smi topo -m

验收标准：

lspci 数量和机型 BOM 一致。
nvidia-smi -L 列出的 UUID 数量一致且没有掉卡。
Driver 版本属于平台锁定的 GPU Operator/Container Toolkit/CUDA 兼容矩阵。
拓扑与预期 PCIe/NVLink/NVSwitch 连接一致。
Secure Boot、kernel module 签名、nouveau 冲突、Xid/ECC 错误已经处理。

若 lspci 都看不到 GPU，应先检查硬件、虚拟机 passthrough 或 BIOS；此时安装 Device Plugin 没有意义。

路径 A：已有 GPU Operator，推荐生产使用

1. 先判断 Operator 是否已经是集群能力

helm list -A | grep gpu-operator
kubectl get clusterpolicy
kubectl get pods -n gpu-operator

GPU Operator 是集群级控制器，不需要为每个新节点再 helm install 一次。默认情况下，NFD 发现 NVIDIA PCI vendor ID 后会产生：

feature.node.kubernetes.io/pci-10de.present=true

Operator 根据这个标签把 driver、toolkit、device-plugin、GFD、DCGM 和 validator 等 operand 部署到 GPU worker。应等待真实 NFD 探测结果，不要为了“让 Pod 跑起来”手工伪造硬件标签。

kubectl get node "$NODE" \
  -o jsonpath='{.metadata.labels.feature\.node\.kubernetes\.io/pci-10de\.present}{"\n"}'

kubectl get pods -n gpu-operator -o wide \
  --field-selector spec.nodeName="$NODE"

kubectl get events -n gpu-operator --sort-by='.lastTimestamp'

如果 GPU 节点有业务 taint，需要确认 Operator 管理的 DaemonSet 具有对应 toleration，否则控制器识别到了节点，节点侧 operand 仍然落不下来。

2. 明确 Driver 与 Toolkit 由谁管理

节点镜像现状	Operator 配置原则	注意事项
Driver、Toolkit 都未预装	Operator 默认管理 Driver、Toolkit 和 Device Plugin	GPU worker 的 OS/kernel 需要符合 Driver Container 支持矩阵
已预装 Driver	集群安装时使用 `driver.enabled=false`	驱动版本由镜像/节点运维体系负责升级
已预装 Driver + Toolkit	使用 `driver.enabled=false`、`toolkit.enabled=false`	必须提前验证 containerd/CDI 或 NVIDIA runtime 配置

这些 Helm value 通常影响整个 ClusterPolicy，不是“临时只给一台节点改一下”的开关。混合 OS、混合 Driver 版本或不同责任边界应通过节点池、Driver CRD/nodeSelector 等受控方案管理，不能让同一节点同时被主机包管理器和 Driver Container 争抢内核模块。

如果集群还没有 Operator，安装时必须先从官方 Platform Support/Component Matrix 锁定版本，再写入 GitOps values：

GPU_OPERATOR_VERSION=<经过平台验证并锁定的版本>

helm repo add nvidia https://helm.ngc.nvidia.com/nvidia
helm repo update
helm upgrade --install gpu-operator nvidia/gpu-operator \
  --namespace gpu-operator \
  --create-namespace \
  --version "$GPU_OPERATOR_VERSION" \
  --wait

生产环境还需要把 registry mirror、imagePullSecret、proxy、toleration、priorityClass、驱动类型和版本写入 values 文件，而不是长期依赖命令行 --set。

3. 验证 Operator 状态

kubectl get clusterpolicy cluster-policy \
  -o jsonpath='{.status.state}{"\n"}'

kubectl get pods -n gpu-operator -o wide \
  --field-selector spec.nodeName="$NODE"

kubectl logs -n gpu-operator \
  -l app=nvidia-operator-validator \
  --all-containers --tail=100

目标节点至少要看到对应的 Driver（若由 Operator 管理）、Container Toolkit、Device Plugin、GFD、DCGM Exporter 和 Validator 状态正常。Validator 中 CUDA/Device Plugin 校验失败时不能直接移除接入 taint。

路径 B：手工管理 Driver、Runtime 和 Device Plugin

这条路径适合不可变 GPU 节点镜像、离线环境或平台已经有成熟的 OS 配置管理体系。它的优点是版本可控，缺点是升级、回滚和节点漂移都要自己负责。

1. 安装并锁定 Host Driver

优先使用发行版/企业镜像的软件包管理方式，不建议在自动化节点上长期混用 .run installer。安装后至少验证：

nvidia-smi
nvidia-smi -L
ls -l /dev/nvidia* /dev/nvidia-caps/* 2>/dev/null

nvidia-smi 顶部显示的 “CUDA Version” 是 该 Driver 可支持的最高 CUDA API 版本信息，不等于节点已经安装了同版本 CUDA Toolkit。

2. 安装 NVIDIA Container Toolkit 并配置 containerd

containerd --version
nvidia-ctk --version

sudo nvidia-ctk runtime configure --runtime=containerd
sudo systemctl restart containerd
sudo systemctl is-active containerd

nvidia-ctk 会为 containerd 写入 NVIDIA runtime/drop-in 配置。修改后必须重启 containerd，并确认 kubelet 没有因为 CRI socket 或配置格式错误进入 NotReady。

legacy 模式下有两种选择：

使用 --set-as-default 把 nvidia handler 设为默认低层 runtime。
保留普通 runc 为默认，为 GPU Pod 配置 RuntimeClass：

apiVersion: node.k8s.io/v1
kind: RuntimeClass
metadata:
  name: nvidia
handler: nvidia

当前 CDI 路径应按所锁定的 Container Toolkit、containerd 和 Device Plugin 版本统一配置。标准 Device Plugin workload 的 CDI 注入通常对业务 YAML 透明，不应为了照抄旧教程同时开启多套 hook。

3. 部署 NVIDIA Device Plugin

Device Plugin 一般以 DaemonSet 运行在 GPU 节点。生产上使用 Helm/GitOps 锁定 chart 和镜像版本，并配置 GPU nodeSelector、toleration、MIG strategy 与共享策略；官方 static DaemonSet 更适合快速实验。

helm repo add nvdp https://nvidia.github.io/k8s-device-plugin
helm repo update

DEVICE_PLUGIN_VERSION=<经过验证并锁定的版本>
helm upgrade --install nvidia-device-plugin nvdp/nvidia-device-plugin \
  --namespace nvidia-device-plugin \
  --create-namespace \
  --version "$DEVICE_PLUGIN_VERSION"

Device Plugin 启动后会在 /var/lib/kubelet/device-plugins/ 下通过 gRPC 向 kubelet 注册资源名和健康设备。kubelet 再把设备数量写入 Node Status；scheduler 只根据 Capacity/Allocatable 做放置，不会直接执行 nvidia-smi。

kubectl get pods -n nvidia-device-plugin -o wide \
  --field-selector spec.nodeName="$NODE"

kubectl logs -n nvidia-device-plugin \
  -l app.kubernetes.io/name=nvidia-device-plugin \
  --tail=200

kubectl get node "$NODE" \
  -o jsonpath='{.status.capacity.nvidia\.com/gpu}{" / "}{.status.allocatable.nvidia\.com/gpu}{"\n"}'

全链路验收：不要只看 `nvidia-smi`

第 1 关：Node 正常加入集群

kubectl get node "$NODE"
kubectl get node "$NODE" \
  -o jsonpath='{.status.conditions[?(@.type=="Ready")].status}{"\n"}'

目标：Ready=True，CNI、containerd、kubelet 正常。

第 2 关：Kubernetes 已发布 GPU 资源

kubectl get node "$NODE" \
  -o jsonpath='{.status.capacity.nvidia\.com/gpu}{" / "}{.status.allocatable.nvidia\.com/gpu}{"\n"}'

kubectl describe node "$NODE" | grep -A8 -E 'Capacity:|Allocatable:'

整卡模式下，空闲节点的 Capacity/Allocatable 数量应与健康物理 GPU 数量一致。若使用 MIG、Time-Slicing、MPS 或 HAMi，资源名和数量要按对应策略解释，不能继续用整卡数量验收。

第 3 关：运行绑定到新节点的 CUDA canary

apiVersion: v1
kind: Pod
metadata:
  name: gpu-node-canary
spec:
  restartPolicy: Never
  nodeSelector:
    kubernetes.io/hostname: gpu-node-01
  tolerations:
    - key: gpu-bootstrap
      operator: Equal
      value: "true"
      effect: NoSchedule
  containers:
    - name: vectoradd
      image: nvcr.io/nvidia/k8s/cuda-sample:vectoradd-cuda12.5.0
      resources:
        limits:
          nvidia.com/gpu: 1

应把 nodeSelector、镜像 tag 和 runtimeClass（如果走 legacy 非默认 NVIDIA runtime）替换为集群实际值。

kubectl apply -f gpu-node-canary.yaml
kubectl get pod gpu-node-canary -o wide
kubectl describe pod gpu-node-canary
kubectl logs gpu-node-canary

验收标准：

Pod 明确调度到目标新节点。
Events 没有 FailedScheduling、FailedCreatePodSandBox 或 Device Plugin Allocate 错误。
日志输出 Test PASSED。
Pod 运行期间，kubectl describe node 的 Allocated resources 能看到 GPU request；status.allocatable 表示节点可分配容量上限，不会因为 Pod 占用而动态减一。

第 4 关：监控与稳定性

kubectl get pods -n gpu-operator -o wide \
  --field-selector spec.nodeName="$NODE" | grep -E 'dcgm|validator'

nvidia-smi -q -d ECC,POWER,TEMPERATURE
journalctl -k --since '-30 min' | grep -E 'NVRM|Xid'

还要确认 Prometheus 已抓到该节点的 DCGM 指标，至少覆盖 GPU 利用率、显存、温度、功耗、ECC/Xid 与 PCIe/NVLink 相关指标。一次 vectorAdd 成功只能证明基本功能，不能证明长时间训练或通信稳定。

第 5 关：解除隔离并清理 canary

kubectl delete pod gpu-node-canary
kubectl taint node "$NODE" gpu-bootstrap:NoSchedule-
kubectl get node "$NODE" --show-labels

只有前四关都通过才能移除 gpu-bootstrap taint。如果平台还使用 workload=gpu:NoSchedule 等永久 taint，应保留并让正式 GPU workload 显式 toleration。

接入共享方案前必须先过整卡基线

整卡 nvidia.com/gpu canary 通过

-> 选择该节点池的唯一 GPU 管理方案

-> MIG：启用 MIG mode，配置 GI/CI 和 mig.strategy

-> Time-Slicing：加载 Device Plugin replicas 配置

-> MPS：加载 MPS sharing 配置和隔离参数

-> HAMi：切换到 HAMi scheduler/device-plugin/core 链路

-> 再按新资源名和隔离语义跑一轮 canary

MIG：验收 nvidia.com/mig-* profile、MIG Manager 状态和实例拓扑。
Time-Slicing：验收 shared resource 数量，但不能把 replicas 当成固定算力份额。
MPS：验收并发、显存/active thread 限制以及 MPS server 故障边界。
HAMi：验收 Webhook、HAMi Scheduler、Device Plugin、HAMi-Core 和显存/算力配额。

同一张物理 GPU 不应同时被 NVIDIA 官方 Device Plugin 与 HAMi/其他 vGPU Device Plugin 重复注册。生产上应通过独立节点池、nodeSelector 和 DaemonSet 调度范围明确所有权。

分层故障定位

现象	最可能故障层	首先检查
`lspci` 看不到卡	硬件/BIOS/虚拟化 passthrough	BMC、PCIe slot、VM 配置、固件
`lspci` 有，`nvidia-smi` 失败	Host Driver	kernel module、Secure Boot、nouveau、版本、Xid
宿主机正常，GPU 容器创建失败	Toolkit/CDI/runtime	containerd config、CDI spec、runtime handler、kubelet/containerd 日志
容器路径正常，但 Node 没有 `nvidia.com/gpu`	Device Plugin/kubelet	DaemonSet、plugin 日志、gRPC socket、设备健康状态
Allocatable 正常，Pod 一直 Pending	scheduler/策略	limits、taint/toleration、affinity、quota、已有占用
Pod 已运行，但 CUDA 报错	镜像/Driver 兼容或设备注入	Driver-CUDA 兼容、容器内库、可见 UUID、Allocate 结果
运行中 GPU 从 Allocatable 消失	设备健康/Xid	Device Plugin ListAndWatch、kernel log、DCGM/Xid 告警

生产变更与回滚原则

新节点保持 bootstrap taint，失败不会影响业务。
Driver、Toolkit、Device Plugin、GPU Operator 版本写入 GitOps/镜像清单，不使用 latest。
改 Driver 或 runtime 前先 cordon/drain，预留可能重启节点的窗口。
保存 containerd、Toolkit、Operator values 和节点标签基线，确保可回退。
先在单节点 canary 池验证，再扩到同型号节点池。
除基本 CUDA canary 外，根据业务补充 NCCL、RDMA、MIG/共享和长稳测试。

Q: 面试官问“一个新 GPU 节点来了，怎么让 Kubernetes 用起来”，应该怎么回答？

我会先给节点加 bootstrap taint，完成普通 worker 的 OS、containerd、kubelet、CNI 和网络基线；然后从下往上打通 GPU 链路：确认 PCIe 枚举，安装并验证 NVIDIA Driver；配置 NVIDIA Container Toolkit，使 containerd 能通过 CDI 或 NVIDIA runtime 注入设备；部署 Device Plugin，让它通过 gRPC 向 kubelet 注册健康 GPU，最终在 Node Status 中出现 nvidia.com/gpu。生产上我优先用 GPU Operator统一管理 Driver、Toolkit、Device Plugin、GFD、DCGM 和 Validator。最后必须在目标节点跑一个真实申请 nvidia.com/gpu 的 CUDA canary，检查调度、Allocate、容器内 kernel 和 DCGM 指标，全部通过后才移除 taint。若还要上 MIG、MPS、Time-Slicing 或 HAMi，会在整卡基线通过后再切换，并重新按对应资源语义验收。

Q: 为什么宿主机 nvidia-smi 正常，Kubernetes 仍可能看不到 GPU？

nvidia-smi 只证明 Host Driver 能访问设备。Kubernetes 还依赖 Container Toolkit/CDI 打通容器设备注入，依赖 Device Plugin 向 kubelet 注册并持续上报健康设备。Runtime 配错时 Pod 可能创建失败；Device Plugin 未运行或设备被标记 unhealthy 时，Node Status 不会出现可用的 nvidia.com/gpu。

资料来源

MIG 到底切了什么

MIG 把一张支持该能力的 GPU 划分为若干 GPU Instance（GI），每个 GI 再包含可运行 CUDA workload 的 Compute Instance（CI）。profile 名称中的两部分分别表达计算切片和显存规格，例如 1g.10gb；具体 profile、数量和名称取决于 GPU 型号。

NVIDIA 官方 MIG 图：A100 40GB 被抽象为 8 个 5GB 显存切片和 7 个计算切片。来源：NVIDIA MIG User Guide · Partitioning。

一个 1g.5gb GPU Instance 将 1 个计算切片与 1 个 5GB 显存切片组合成固定规格。新型号的 profile 名称和显存容量会变化，面试时应讲清“固定计算份额 + 固定显存份额”，不要死背 A100 数字。来源：NVIDIA MIG User Guide · GPU Instances。

物理 GPU

-> 开启 MIG mode

-> 按 profile 创建 GPU Instance

-> 为 GPU Instance 创建 Compute Instance

-> NVIDIA driver 枚举 MIG device UUID

-> Device Plugin 上报 MIG 扩展资源

-> Pod 申请一个 MIG profile

MIG 的价值是性能和故障隔离更稳定，但资源形状固定。一个 1g workload 空闲时，旁边的 3g workload 不能自动借走它的硬件切片。

裸机怎么用

下面以单张 GPU 为例。命令必须在没有业务进程占用该 GPU 时执行；profile 仅作为示例，要先查询本机实际支持项。

# 1. 确认 GPU 与 MIG 状态
nvidia-smi -L
nvidia-smi -i 0 --query-gpu=name,mig.mode.current,mig.mode.pending --format=csv

# 2. 启用 MIG mode
sudo nvidia-smi -i 0 -mig 1

# 3. 查询可用 GPU Instance profile 与 placement
nvidia-smi mig -lgip
nvidia-smi mig -lgipp

# 4. 示例：创建两个 3g.20gb GI，并同时创建对应 CI
sudo nvidia-smi mig -cgi 3g.20gb,3g.20gb -C

# 5. 验证实例与 UUID
nvidia-smi mig -lgi
nvidia-smi mig -lci
nvidia-smi -L

仅执行 -mig 1 不够：没有创建 GI/CI 时，CUDA workload 还没有可使用的 MIG device。创建后的 MIG device 可以通过 MIG-<UUID> 选择，例如容器运行时可把该 UUID 放进 NVIDIA_VISIBLE_DEVICES。

Kubernetes + GPU Operator 怎么用

1. 选择 MIG strategy

GPU Operator / NVIDIA Device Plugin 常见两种策略：

策略	资源表达	适用场景
`single`	节点上的 MIG 设备采用一致几何，使用方式更统一	同一节点全部做同规格推理池
`mixed`	不同 profile 以独立资源名上报，例如 `nvidia.com/mig-1g.10gb`	同节点需要多种规格

安装时启用 MIG Manager 的示意命令：

helm repo add nvidia https://helm.ngc.nvidia.com/nvidia
helm repo update

helm install --wait gpu-operator nvidia/gpu-operator \
  -n gpu-operator --create-namespace \
  --set mig.strategy=mixed

实际部署时应锁定经过验证的 GPU Operator 版本，不要在生产集群直接跟随 latest。

2. 给节点应用 MIG geometry

MIG Manager 监听节点的 nvidia.com/mig.config label。变更前先清空业务 workload：

NODE=gpu-node-1

kubectl cordon "$NODE"
kubectl drain "$NODE" --ignore-daemonsets --delete-emptydir-data

# profile 名称以该节点的 MIG Manager ConfigMap 为准
kubectl label node "$NODE" \
  nvidia.com/mig.config=all-1g.10gb --overwrite

kubectl get node "$NODE" \
  -L nvidia.com/mig.config,nvidia.com/mig.config.state

只有状态进入 success，并且 Device Plugin 重新上报资源后，节点才重新开放：

kubectl describe node "$NODE" | grep -E 'nvidia.com/(gpu|mig-)'
kubectl uncordon "$NODE"

3. Pod 申请 MIG 实例

mixed strategy 下，Pod 申请具体 profile；资源名必须以 kubectl describe node 的实际结果为准。

apiVersion: v1
kind: Pod
metadata:
  name: mig-demo
spec:
  restartPolicy: Never
  containers:
    - name: cuda
      image: nvcr.io/nvidia/cuda:12.8.0-base-ubuntu22.04
      command: ["bash", "-lc", "nvidia-smi && sleep 3600"]
      resources:
        limits:
          nvidia.com/mig-1g.10gb: 1

验证：

kubectl apply -f mig-demo.yaml
kubectl get pod mig-demo -o wide
kubectl exec mig-demo -- nvidia-smi -L
kubectl exec mig-demo -- nvidia-smi

容器应该只看到被分配的 MIG device 及其显存规格，而不是整张物理卡。

重配与回收

裸机删除顺序通常是先删 CI，再删 GI，最后按需关闭 MIG mode：

sudo nvidia-smi mig -dci
sudo nvidia-smi mig -dgi
sudo nvidia-smi -i 0 -mig 0

GPU Operator 环境优先让 MIG Manager 管理，不要一边手工执行 nvidia-smi mig，一边让 controller 按 label 重建几何。需要恢复整卡时，先 drain 节点，再应用 all-disabled 等实际配置中存在的 profile。

生产风险

风险	原因	处理方式
MIG 重配导致任务中断	MIG Manager 会停止 GPU clients，部分环境还需要重启	cordon/drain，维护窗口重配
profile 碎片	剩余 slice 组合无法满足新 profile	按业务规格建资源池，不频繁在线改 geometry
申请资源名错误	不同 GPU/strategy 暴露的资源名不同	以 Node Allocatable 为准，不硬编码猜测
只看 GPU-Util 误判	MIG 需要实例级监控和归因	使用支持 MIG 的 DCGM 指标和业务吞吐
把 MIG 当弹性共享	MIG 实例之间不会自动借资源	需要弹性时考虑 MPS/HAMi，或在维护窗口重配

Q: 为什么强 SLA 推理优先考虑 MIG？

因为 MIG 的资源边界来自硬件实例，性能和故障隔离比 MPS/Time-Slicing 更可预测。代价是规格固定、资源不能自动借用，并且重配 geometry 会影响节点上的 GPU workload。

资料来源

MPS 执行链路

MPS control daemon

接收管理命令并创建 MPS server

MPS server

持有 GPU 调度资源，接收多个 client 的 CUDA work

CUDA clients

通过 pipe 连接 server，提交 kernel 和内存申请

GPU

允许来自不同 client 的 kernel 并发执行

监控

nvidia-smi / DCGM / MPS ps 观察 client、显存和吞吐

MPS 的目标是降低多 CUDA context 独立运行的切换成本，并增加小 kernel 的并发机会。它不是 MIG：不同 client 仍然共享物理 GPU、部分缓存、带宽和故障域。

裸机怎么用

1. 启动 MPS daemon

下面是单用户、单 GPU 的最小实验流程。生产环境需要独立服务账号、目录权限和进程生命周期管理。

export CUDA_VISIBLE_DEVICES=0
export CUDA_MPS_PIPE_DIRECTORY=/tmp/nvidia-mps
export CUDA_MPS_LOG_DIRECTORY=/tmp/nvidia-log

mkdir -p "$CUDA_MPS_PIPE_DIRECTORY" "$CUDA_MPS_LOG_DIRECTORY"
nvidia-cuda-mps-control -d

# 查看 server 和 client
echo ps | nvidia-cuda-mps-control

随后在同一组 MPS pipe 配置下启动两个 CUDA 进程，它们会连接到 MPS server。应用必须在创建 CUDA context 之前设置资源限制。

2. 限制执行资源

# 该 client 新建 CUDA context 时最多使用约 30% 的可用执行资源
export CUDA_MPS_ACTIVE_THREAD_PERCENTAGE=30
python inference_worker.py

CUDA_MPS_ACTIVE_THREAD_PERCENTAGE 限制的是 client 可用的执行线程比例，最终值会按硬件支持粒度取整。它不是“把每个时刻的 GPU-Util 精确锁死在 30%”，也不直接保证业务吞吐等于整卡的 30%。

MPS 还支持通过控制命令设置未来 client 的默认执行比例：

echo 'set_default_active_thread_percentage 30' \
  | nvidia-cuda-mps-control

3. 限制显存

# 限制未来 MPS client 在 GPU 0 上可分配的 device memory
echo 'set_default_device_pinned_mem_limit 0 8G' \
  | nvidia-cuda-mps-control

# client 侧还可以进一步收紧，不能突破 server 的更低限制
export CUDA_MPS_PINNED_DEVICE_MEM_LIMIT='0=6G'
python inference_worker.py

超出限制时，CUDA 内存分配会返回 OOM。显存限制要预留 CUDA context、库 workspace 和通信 buffer，不要只按模型权重大小计算。

4. 停止 MPS

echo ps | nvidia-cuda-mps-control
echo quit | nvidia-cuda-mps-control

退出前先停止 client workload。异常退出后还要检查残留 server、pipe 目录和 GPU 上下文。

Kubernetes 怎么用

NVIDIA Device Plugin 的 MPS sharing 用 replicas 定义等份共享资源。官方文档仍将该能力标为实验性；MPS 与 Time-Slicing 互斥，并且当前不能用于已启用 MIG 的设备。

1. 创建共享配置

apiVersion: v1
kind: ConfigMap
metadata:
  name: nvidia-device-plugin-config
  namespace: gpu-operator
data:
  mps-4: |-
    version: v1
    flags:
      migStrategy: none
    sharing:
      mps:
        renameByDefault: true
        resources:
          - name: nvidia.com/gpu
            replicas: 4

这表示每张 GPU 暴露 4 个 nvidia.com/gpu.shared，MPS control daemon 为每个共享资源限制为大约总显存和执行资源的四分之一。它与裸机手工设置任意百分比不同：这里是按 replicas 等份管理。

2. 让 Device Plugin 加载配置

GPU Operator 环境可以让 ClusterPolicy 引用 ConfigMap，并通过节点 label 选择配置：

kubectl apply -f nvidia-device-plugin-config.yaml

kubectl patch clusterpolicy/cluster-policy \
  --type merge \
  -p '{"spec":{"devicePlugin":{"config":{"name":"nvidia-device-plugin-config","default":"mps-4"}}}}'

kubectl label node gpu-node-1 \
  nvidia.com/device-plugin.config=mps-4 --overwrite

不同 GPU Operator / Device Plugin 版本的 chart 字段可能变化，上线前要以所锁定版本的配置 schema 为准。

3. Pod 申请与验证

apiVersion: v1
kind: Pod
metadata:
  name: mps-demo
spec:
  restartPolicy: Never
  containers:
    - name: cuda
      image: nvcr.io/nvidia/cuda:12.8.0-base-ubuntu22.04
      command: ["bash", "-lc", "nvidia-smi && sleep 3600"]
      resources:
        limits:
          nvidia.com/gpu.shared: 1

kubectl describe node gpu-node-1 \
  | grep -E 'nvidia.com/(gpu.shared|gpu.sharing-strategy|mps.capable)'

kubectl apply -f mps-demo.yaml
kubectl exec mps-demo -- nvidia-smi

验证不能只看 Pod Running，还要同时运行多份实际 workload，比较：

单独运行与共置运行的吞吐、P50/P99 延迟。
每个 client 的显存峰值和 OOM 行为。
SM Active、DRAM throughput、L2、PCIe 等 DCGM 指标。
任一 client 异常时，其他 client 是否受到影响。

什么时候适合 MPS

适合	不适合
同一团队的多个小模型推理	互不信任租户、需要强故障隔离
单任务 kernel 很小，整卡利用率低	单任务已经接近打满 SM 或 HBM
计算型与 I/O 型 workload 互补	严格 P99 SLO 且无法容忍共享抖动
能持续采集干扰指标并自动降级	只有“多放几个 Pod”而没有运行时保护

Q: MPS 设置 30%，是不是这个进程永远只能获得 30% GPU 性能？

不是。active thread percentage 约束 client 能使用的执行资源上限，不直接等价于业务性能比例。实际吞吐还受 kernel 并行度、显存带宽、cache、CPU/IO、其他 client 负载和硬件取整影响，必须用真实 workload 基准测试。

常见误区

误区	正确理解
启用 MPS 后多个进程就完全隔离	MPS 改善并发和可施加资源限制，但仍共享物理 GPU 与故障域。
Active thread 30% 等于 GPU-Util 30%	它限制执行资源，不是业务层的恒定利用率。
只限制 SM 就不会互相影响	显存容量、HBM 带宽、L2、PCIe、CPU/IO 仍可能成为共享瓶颈。
MPS 可以直接和 MIG sharing 同时打开	官方 Device Plugin 当前不支持在 MIG-enabled device 上启用 MPS sharing。

资料来源

系统链路

Device Plugin 读取 timeSlicing 配置

-> 为每张物理 GPU 创建 N 个逻辑引用

-> Kubelet 把逻辑数量写入 Node Capacity / Allocatable

-> kube-scheduler 按普通扩展资源分配 Pod

-> 多个 Pod 获得同一 GPU UUID 的访问权

-> CUDA driver 在进程之间交错执行

kube-scheduler 不知道这些 slot 来自同一张物理卡，也不会为每个 slot 保留固定 SM 或显存。Time-Slicing 解决的是“能不能共享访问”，不是“怎么保证共享性能”。

Kubernetes 怎么配置

1. 创建 ConfigMap

apiVersion: v1
kind: ConfigMap
metadata:
  name: nvidia-device-plugin-config
  namespace: gpu-operator
data:
  time-slicing-4: |-
    version: v1
    flags:
      migStrategy: none
    sharing:
      timeSlicing:
        renameByDefault: true
        failRequestsGreaterThanOne: true
        resources:
          - name: nvidia.com/gpu
            replicas: 4

关键参数：

参数	含义	为什么重要
`replicas: 4`	每张物理 GPU 创建 4 个共享访问名额	控制最大共享密度，不是算力比例
`renameByDefault: true`	资源名改为 `nvidia.com/gpu.shared`	防止用户把共享 slot 当独占 GPU
`failRequestsGreaterThanOne: true`	拒绝单容器一次申请多个共享 slot	强调申请的是访问权，不是成比例算力

2. 让 Device Plugin 加载配置

GPU Operator 环境示例：

kubectl apply -f nvidia-device-plugin-config.yaml

kubectl patch clusterpolicy/cluster-policy \
  --type merge \
  -p '{"spec":{"devicePlugin":{"config":{"name":"nvidia-device-plugin-config","default":"time-slicing-4"}}}}'

kubectl label node gpu-node-1 \
  nvidia.com/device-plugin.config=time-slicing-4 --overwrite

官方 Device Plugin 的共享方式按节点配置：同一节点上的 GPU 使用相同 sharing method，不能让 GPU0 用 Time-Slicing、GPU1 用 MPS。需要不同策略时，应拆成不同节点池或至少使用不同节点配置。

3. 确认资源已经上报

假设节点有 8 张 GPU、每张设置 4 个 replicas，Node 应显示 32 个共享资源：

kubectl describe node gpu-node-1 \
  | grep -E 'nvidia.com/(gpu.shared|gpu.replicas|gpu.sharing-strategy)'

kubectl get node gpu-node-1 \
  -o jsonpath='{.status.capacity.nvidia\.com/gpu\.shared}{"\n"}'

如果仍然只看到 8 张独占 GPU，检查 ConfigMap namespace、ClusterPolicy 引用、节点 label 和 Device Plugin Pod 日志。

Pod 怎么申请

apiVersion: v1
kind: Pod
metadata:
  name: timeslice-demo
spec:
  restartPolicy: Never
  containers:
    - name: cuda
      image: nvcr.io/nvidia/cuda:12.8.0-base-ubuntu22.04
      command: ["bash", "-lc", "nvidia-smi -L && sleep 3600"]
      resources:
        limits:
          nvidia.com/gpu.shared: 1

kubectl apply -f timeslice-demo.yaml
kubectl get pod timeslice-demo -o wide
kubectl exec timeslice-demo -- nvidia-smi -L
kubectl exec timeslice-demo -- nvidia-smi

Pod 通常会看到整张物理 GPU 的 UUID 和显存容量。多个 Pod 可能看到同一 UUID，这正是它们共享同一张卡的证据。不能因为 nvidia-smi 显示整卡显存，就让每个进程都按整卡容量分配。

如何验证真实边界

不要只创建 4 个 sleep Pod。至少进行下面四组实验：

单任务基线：记录单独运行的吞吐、P50/P99 和显存峰值。
并发压力：同时运行 2/4 个计算密集 workload，观察性能是否近似均分以及上下文切换开销。
显存冲突：让两个 workload 的显存总需求超过物理显存，确认其中一个可能 OOM，证明 replicas 不隔离显存。
异常影响：记录 Xid、Pod 重启和其他同卡 workload 的影响，验证共享故障域风险。

kubectl get pods -o wide
kubectl logs -n gpu-operator -l app=nvidia-device-plugin-daemonset --tail=200
nvidia-smi pmon -s um -c 10

适用与不适用

适合	不适合
Notebook、交互式开发、低负载实验	延迟敏感在线推理
低优离线推理和评测	互不信任租户
可以接受性能抖动的教学/研发集群	显存峰值大、容易 OOM 的训练
希望低成本提高 Pod 密度	要求固定 GPU 百分比计费或 SLA

Q: replicas=4 时，一个 Pod 实际拿到多少 GPU？

它拿到的是“访问一张共享 GPU”的权利，而不是静态 1/4 卡。如果只有一个 CUDA 进程，它可能使用大部分 GPU；多个满负载进程同时运行时，驱动会在它们之间交错执行，但不提供固定算力、显存或 P99 保证。

常见误区

误区	正确理解
每个 Pod 会只看到 1/4 显存	通常仍看到整卡，多个 Pod 的显存申请可能互相挤压。
申请 2 个 shared GPU 就有两倍算力	多申请共享 slot 不代表获得成比例算力，建议开启 `failRequestsGreaterThanOne`。
Time-Slicing 能让训练和推理安全混部	它缺少强隔离，在线 SLO 容易被训练 kernel 和显存争用影响。
GPU-Util 变高就说明收益更好	还要检查吞吐、JCT、P99、OOM 和上下文切换代价。

资料来源

NVIDIA Kubernetes Device Plugin：CUDA Time-Slicing

HAMi 解决什么问题

官方 NVIDIA Device Plugin 的 Time-Slicing/MPS 主要按 replicas 做等份共享。平台如果希望用户声明“我要一张物理 GPU 上的 3GiB 显存和 30% 算力”，默认 kube-scheduler 无法理解这种二维设备容量。

HAMi 为 Pod 增加细粒度资源：

resources:
  limits:
    nvidia.com/gpu: 1
    nvidia.com/gpumem: 3000
    nvidia.com/gpucores: 30

nvidia.com/gpu：需要分配的 GPU 设备数量。
nvidia.com/gpumem：每个设备的显存需求，常用单位为 MiB。
nvidia.com/gpucores：设备计算资源百分比，1 表示 1%。

HAMi 还支持多种 GPU/NPU/MLU/DCU 等后端，但每种设备的共享和隔离能力不同，不能只看“支持列表”就假设所有硬件都有相同能力。

核心架构

Pod 提交细粒度 GPU request

-> Mutating Webhook 识别 HAMi resource，并设置 hami-scheduler

-> Scheduler Extender 维护全局设备视图，执行 Filter / Score / Bind

-> 选定物理 GPU UUID，把分配结果写入 Pod annotation

-> HAMi Device Plugin 在 Allocate 阶段挂载设备和运行时配置

-> HAMi-Core 注入容器，执行显存检查与算力节流

-> Monitor / Metrics 持续观测分配和使用情况

组件	职责	关键状态
Mutating Webhook	识别 HAMi 资源并选择调度入口	`spec.schedulerName`
Scheduler Extender	维护剩余显存/算力，选择节点和物理设备	Node/Pod annotation、设备全局视图
HAMi Device Plugin	注册设备、读取调度结果、向容器注入设备与环境	GPU UUID、Allocate 结果
HAMi-Core	在容器内拦截 CUDA/NVML 调用，执行显存上限和算力节流	显存上限、SM limit、运行时计数器

这条链路解释了为什么 HAMi 不能只靠 Device Plugin：标准 scheduler 在调度阶段看不到每张 GPU 剩余多少显存和算力，而 Device Plugin 的 Allocate 又发生在节点已经选定之后。HAMi 需要 Scheduler Extender 先做设备级放置，再通过 annotation 把结果交给节点侧执行。

Helm 部署

1. 前置检查

GPU 节点应先满足：

NVIDIA driver 和 NVIDIA Container Toolkit 可用。
容器运行时已正确配置 NVIDIA runtime。
Kubernetes、Helm、CUDA/driver 版本符合当前 HAMi release 的兼容矩阵。
同一批节点没有另一个 Device Plugin 同时注册和管理相同 GPU 资源。

HAMi、Volcano vGPU Device Plugin 和 NVIDIA 官方 Device Plugin 不应同时管理同一节点上的同一类 GPU。生产上应通过节点池和 label 划分职责。

2. 安装

# 标记由 HAMi 管理的 GPU 节点
kubectl label node gpu-node-1 gpu=on --overwrite

helm repo add hami-charts https://project-hami.github.io/HAMi/
helm repo update

helm install hami hami-charts/hami \
  -n kube-system

部分 HAMi 版本要求 scheduler 镜像 tag 与 Kubernetes 版本对应。安装前应阅读所锁定 chart 版本的 values 和兼容说明，不要直接复制其他集群的版本参数。

3. 验证控制面与节点组件

kubectl get pods -n kube-system \
  | grep -E 'hami-(scheduler|device-plugin)'

kubectl logs -n kube-system deploy/hami-scheduler --tail=200
kubectl get daemonset -n kube-system | grep hami-device-plugin

至少确认：

hami-scheduler 正常运行并可访问 API Server。
每个目标 GPU 节点都有 hami-device-plugin Pod。
Node 上已经注册 HAMi 管理的 GPU 数量与设备信息。
原 NVIDIA Device Plugin 没有同时争抢相同资源名。

Pod 怎么申请细粒度 GPU

apiVersion: v1
kind: Pod
metadata:
  name: hami-vgpu-demo
  annotations:
    hami.io/node-scheduler-policy: "binpack"
spec:
  restartPolicy: Never
  containers:
    - name: cuda
      image: nvcr.io/nvidia/cuda:12.8.0-base-ubuntu22.04
      command: ["bash", "-lc", "nvidia-smi && sleep 3600"]
      resources:
        limits:
          nvidia.com/gpu: 1
          nvidia.com/gpumem: 3000
          nvidia.com/gpucores: 30

kubectl apply -f hami-vgpu-demo.yaml
kubectl get pod hami-vgpu-demo -o wide
kubectl get pod hami-vgpu-demo \
  -o jsonpath='{.spec.schedulerName}{"\n"}'
kubectl describe pod hami-vgpu-demo
kubectl exec hami-vgpu-demo -- nvidia-smi

预期现象：

Webhook 将任务路由到 HAMi scheduler。
Pod annotation 中记录选中的 GPU UUID、显存和算力分配结果。
容器内 nvidia-smi 看到的显存信息受到 HAMi-Core 的 NVML 拦截影响。
超出显存配额的 CUDA 分配返回 OOM，而不是无限挤占同卡其他 Pod。
算力限制依靠运行时节流，nvidia-smi 的瞬时 GPU-Util 仍可能波动。

注意：Node Capacity/Allocatable 可能主要显示 nvidia.com/gpu，而 gpumem/gpucores 的剩余容量由 HAMi scheduler 的设备视图和 annotation 管理，不一定像普通 Extended Resource 一样直接出现在 Capacity 中。

调度策略

策略	含义	适合场景
`binpack`	优先把任务放到已经使用的节点/GPU，保留完整设备	提高整卡可用性、便于缩容、减少碎片
`spread`	把任务分散到不同节点/GPU	降低热点和共享干扰，提高故障分散

选择不能只看利用率：小模型推理可以 binpack，延迟敏感或带宽密集 workload 更适合 spread。多 GPU 训练还要考虑 NVLink/PCIe 拓扑，不能只按显存余量拼卡。

HAMi 与其他方案怎么选

维度	NVIDIA Time-Slicing	NVIDIA MPS Sharing	HAMi	MIG
分配粒度	共享访问名额	按 replicas 等份	显存 MiB + 算力百分比 + 设备数	固定硬件 profile
调度感知	主要看逻辑数量	主要看逻辑数量	感知单卡剩余显存/算力和设备策略	profile 作为独立资源
隔离方式	时间复用，弱	MPS server 限制	软件/设备后端限制	硬件隔离
异构设备	NVIDIA	NVIDIA	面向多厂商扩展	支持 MIG 的 NVIDIA GPU
系统复杂度	低	中	较高，增加 Webhook/Scheduler/Core	中，需要 profile 生命周期管理

HAMi 还提供 dynamic MIG 等扩展能力，但它们依赖支持的 GPU、驱动和 HAMi 版本。面试时应把“HAMi 软件 vGPU”和“HAMi 调用硬件 MIG 动态创建实例”区分开。

排障路径

Pod 未被调度

看 schedulerName、Webhook、HAMi scheduler 日志

资源不足

看 Pod annotation、设备注册信息、显存/算力剩余视图

Pod 卡在 ContainerCreating

看 Device Plugin Allocate、runtime、libvgpu 注入

显存限制没生效

看 HAMi-Core、LD_PRELOAD、CUDA/NVML 兼容性

性能抖动

看 gpucores、binpack 密度、HBM/SM 指标和同卡 workload

设备重复或数量异常

排查是否与 NVIDIA/Volcano Device Plugin 冲突

Q: HAMi 是怎么做到“申请 3GiB GPU 显存”的？

调度阶段，HAMi Scheduler 先从全局设备视图中选择剩余显存足够的物理 GPU，并把 UUID 和 3GiB 配额写入 Pod annotation；Allocate 阶段，Device Plugin 注入设备、环境变量和 HAMi-Core；容器内 HAMi-Core 拦截 CUDA/NVML 的显存查询和分配调用，超过配额时返回 OOM。它是调度账本和容器内执行限制的组合。

Q: HAMi 和 MIG 最大的区别是什么？

MIG 的边界来自硬件实例，规格固定、故障和性能隔离更强；HAMi 的优势是显存/算力比例更灵活、能做设备感知调度和异构统一管理，但 NVIDIA 常见 vGPU 路径主要依赖软件拦截与节流，不能直接宣称获得 MIG 同等级硬件隔离。

资料来源

一条生产决策链

给 workload 分类

在线推理、离线推理、训练、Notebook、系统任务

定义不可妥协约束

P99、故障隔离、显存上限、吞吐、可抢占性

建立单任务画像

显存峰值、SM Active、HBM、PCIe、CPU/IO、运行时间

选择虚拟化机制

整卡 / MIG / MPS / Time-Slicing / HAMi

选择调度策略

独占、binpack、spread、干扰感知、拓扑感知

设置运行时保护

显存/算力上限、共享密度、在线降级与抢占

灰度与复测

单任务基线 -> 两任务共置 -> 压力/异常 -> 扩大范围

场景选型表

场景	首选	原因	不建议
核心在线推理，严格 P99	整卡或 MIG	性能和故障边界最可预测	直接 Time-Slicing；无监控的 MPS
同团队多个小模型推理	MPS	kernel 可并发，能限制显存和执行资源	把 MPS 当硬隔离卖给不可信租户
Notebook / 开发实验	Time-Slicing	配置简单，提高共享访问密度	承诺固定 1/N 性能
多团队细粒度显存/算力	HAMi	支持 MiB/百分比声明和设备感知放置	忽略 HAMi-Core 兼容性与插件冲突
大训练或通信密集训练	整卡、拓扑感知分配	需要稳定 SM/HBM/NVLink，减少同卡干扰	为了表面利用率强行切碎
KV Cache / 权重弹性	CUDA VMM + 运行时准入	解决虚拟地址和物理显存页弹性	把 VMM 当算力虚拟化

生产架构怎么分层

层次	组件	需要维护的状态
设备发现	GPU Feature Discovery、DCGM、HAMi Device Plugin	型号、显存、MIG 状态、健康、拓扑
资源表达	NVIDIA Device Plugin、MIG profile、HAMi resources、DRA	资源名、容量、共享关系、设备属性
准入与调度	kube-scheduler、Scheduler Plugin、HAMi Scheduler	quota、剩余显存/算力、优先级、干扰画像
运行时执行	MIG、MPS server、CUDA driver、HAMi-Core、CUDA VMM	硬件实例、client 限制、内存映射
监控与回退	DCGM Exporter、业务指标、告警、抢占/迁移	P99、吞吐、SM/HBM、OOM、Xid、降级状态

面试时可以说：Device Plugin 解决“怎么把设备交给 Pod”，不等于解决“共享后性能是否安全”；虚拟化方案必须和调度、监控、回退形成闭环。

上线前验证矩阵

验证项	要证明什么	失败时怎么处理
单任务基线	虚拟化本身的固定开销	检查 context、driver、MPS/HAMi 注入
两任务共置	SM/HBM/显存是否互补	调低共享密度或改 spread
显存超限	限制是否真的生效	检查 MPS/HAMi runtime，必要时改 MIG
延迟压力	P95/P99 是否可接受	在线 workload 改整卡/MIG
异常退出	是否共享故障域	加自动摘卡、重建 daemon、隔离节点池
节点重启	MIG geometry/共享配置能否恢复	用 Operator/controller 声明式重建
监控归因	能否定位到 Pod/MIG/client	补 DCGM、Pod annotation 和业务标签

论文项目怎么映射

论文部分只回答“为什么需要在基础机制上再加系统策略”。详细背景、算法和实验仍放在<a href="../../../content/ai-infra/papers/index.html">论文工作</a>页面。

项目	使用的基础机制	机制本身缺什么	论文补了什么
DeepShare	NVIDIA MPS + DCGM	MPS 能并发，但不知道哪两个作业适合共置，也不知道何时该保护欠保障租户	QAD、Random Forest 干扰预测、动态准入阈值、连续窗口在线降级
ElastiCo	MPS 执行资源限制 + CUDA VMM	固定 SM/显存比例无法追随训练和推理阶段变化	资源形态变换、影子定价、训推阶段感知的动态调整
Maestro	CUDA VMM	VMM 只提供映射能力，不知道每个 Agent stage 应预留多少 KV Cache	输出长度预测、准入控制、按需物理页映射和降级策略

这里不要混淆：

DeepShare 的 GPU 共置依赖 MPS，但贡献不是“发明 MPS”，而是让共置服从租户保障和干扰闭环。
ElastiCo 组合了算力和显存弹性，但 MPS/VMM 只是执行原语，策略来自资源形态与价格信号。
Maestro 使用 CUDA VMM 管 KV Cache，不是把一张 GPU 虚拟成多个 Kubernetes GPU。
HAMi 是可落地的开源平台方案，和论文里的自研策略可以互补，但当前项目没有宣称论文原型直接基于 HAMi。

面试设计题答法

Q: 让 10 个团队共享 100 张 GPU，你会怎么设计虚拟化方案？

先分池

核心在线推理和大训练保留整卡/MIG 池；Notebook、离线推理进入共享池；不要让所有 workload 共用一种策略。

再选机制

强 SLA 用 MIG；可信小 workload 用 MPS；低优研发用 Time-Slicing；需要任意显存/算力配比和异构统一管理时评估 HAMi。

最后闭环

调度器维护 quota、拓扑和共享密度，DCGM 与业务指标检测干扰，超过阈值自动迁移、降级到独占或抢占低优任务。

Q: 为什么不全部使用 HAMi 或全部使用 MIG？

因为它们优化的目标不同。MIG 用固定硬件边界换稳定性，但会产生 profile 碎片；HAMi 用软件调度和运行时控制换细粒度与灵活性，但系统复杂度和隔离边界不同。生产平台通常按 workload 分池组合使用，而不是追求单一方案统一全部场景。

常见误区

误区	正确理解
利用率最高的方案就是最好方案	还要同时看有效吞吐、P99、JCT、OOM、故障域和运维成本。
一个集群只能选一种 GPU 虚拟化	可以按节点池组合，但同一节点/设备的管理插件和 sharing method 要避免冲突。
论文机制可以替代底层 GPU 共享	论文策略建立在 MIG/MPS/VMM 等执行原语上，解决的是准入、调度和保障。
有显存配额就不会干扰	SM、HBM、cache、PCIe、CPU/IO 仍可能竞争。

系统链路

Linux CFS

OS 决定哪个进程/线程拿到 CPU 时间片

CUDA Stream/Event

程序员组织 H2D、kernel、D2H 等 GPU 任务的顺序和依赖

CUDA Block 调度

GPU 把 grid 里的 block 分配到 SM

Warp Scheduler

SM 内部选择 ready warp 发射指令

为什么这个问题容易混淆？

因为“线程”这个词在 CPU 和 GPU 里含义不同。

名词	CPU/Linux 语境	CUDA/GPU 语境
Thread	OS 可调度实体，可能对应内核线程或用户线程	CUDA 逻辑执行单元，运行同一份 kernel 代码
Scheduler	Linux 内核调度器，决定哪个 task 运行在 CPU core	GPU 硬件/驱动/runtime 多层机制，把 block/warp 放到硬件上执行
Context switch	保存/恢复 CPU 寄存器、地址空间、内核栈等	GPU block/warp 不是按 CFS 时间片做频繁 OS 式切换
Fairness	多进程共享 CPU 时间，要尽量公平	GPU kernel 更关心吞吐、occupancy 和 latency hiding
Preemption	CFS 可以抢占当前 CPU task	block 一旦驻留 SM 通常运行到完成，GPU 任务抢占粒度更粗、代价更高

所以面试回答时要先分层：

CPU OS 调度层：进程/线程如何共享 CPU。
CUDA 任务提交层：stream/event 如何组织 GPU work。
GPU kernel 执行层：block/warp 如何映射到 SM。
集群资源调度层：K8s/Volcano/YARN 等如何分配 GPU 设备给任务。

不要把这四层混成一个“GPU 调度器”。

Linux CFS：公平地分 CPU 时间

CFS 全称是 Completely Fair Scheduler，目标不是让某个任务跑到最快，而是在多个 runnable task 之间尽量公平地分配 CPU 时间，同时兼顾交互响应。

CFS 调度对象

CFS 调度的是 Linux 内核里的 task_struct。对用户来说，它可以表现为：

一个进程；
一个线程；
一个容器里的某个线程；
一个 cgroup 下的一组任务。

在 Linux 里，线程和进程最终都可以作为调度实体参与 CPU 调度。CFS 不关心这个任务是不是 AI 训练、数据加载、推理服务、日志线程，它只看到“这个 runnable task 需要 CPU”。

vruntime：谁“欠 CPU 时间”最多

CFS 的核心指标是 vruntime，可以理解成“加权后的虚拟运行时间”。

任务真实运行越久，vruntime 越大。
nice 值越低、优先级越高，权重越大，同样运行一段真实时间，vruntime 增长越慢。
调度器倾向于选择 vruntime 最小的任务运行，因为它看起来“拿到的公平份额最少”。

直觉上：

谁的 vruntime 小，说明谁相对更“饿”，应该优先给 CPU。
谁的 vruntime 大，说明谁已经跑得比较多，可以先等一等。

这和传统时间片轮转不一样。Round Robin 更像“每个人轮流拿固定时间片”；CFS 更像“持续维护每个人已经拿到的公平份额”。

nice、weight 和公平份额

Linux 的 nice 值会影响任务权重。

nice 值	权重直觉	结果
更小，例如 `-10`	权重更大	同样时间内获得更多 CPU 份额
默认 `0`	默认权重	普通任务
更大，例如 `10`	权重更小	更愿意让出 CPU

CFS 不是简单地“高优先级永远先跑”。它通过权重影响 vruntime 增长速度，让高权重任务在长期上获得更多 CPU 时间，但仍然允许其他任务运行。

runqueue 和红黑树

每个 CPU 通常有自己的 runqueue。CFS 会把 runnable task 按 vruntime 组织起来，经典实现使用红黑树。

任务变为 runnable

进入当前 CPU 的 CFS runqueue

按 vruntime 排序

vruntime 小的任务排在更靠左的位置

选择最左任务

调度器选择 vruntime 最小的 task

运行一段时间

task 的 vruntime 增加

重新入队或继续运行

根据抢占、阻塞、唤醒和时间粒度决定

这个结构的意义是：调度器能快速找到“最应该运行”的任务。

抢占和上下文切换

CFS 是抢占式调度。当前任务正在 CPU 上运行时，如果出现一个更应该运行的任务，调度器可以触发抢占。

常见触发点包括：

周期性 tick 或调度时钟更新；
当前任务主动阻塞，例如等待 IO、锁、网络；
新任务唤醒，例如交互请求到来；
当前任务运行超过合理粒度；
更高优先级或更小 vruntime 的任务需要运行。

CPU 上下文切换通常涉及：

保存当前任务寄存器状态；
切换内核栈；
切换或更新地址空间相关状态；
更新调度统计；
恢复下一个任务的执行上下文。

上下文切换不是免费的。线程数过多、锁竞争严重、频繁唤醒阻塞，都可能导致 CPU 时间花在调度和切换上，而不是有效计算上。

CUDA Stream/Event：任务级异步调度

对 CUDA 程序员来说，Stream/Event 是控制 GPU work 提交顺序和依赖的工具。

Stream 是 GPU 任务队列：

同一个 stream 内的操作按提交顺序执行；
不同 stream 的操作可以并发执行，但前提是硬件资源允许；
常见操作包括 H2D 拷贝、kernel launch、D2H 拷贝、event record/wait。

Event 是依赖和计时工具：

可以记录某个 stream 上的进度点；
可以让另一个 stream 等待这个 event；
可以用于测量 GPU 端耗时；
可以避免 CPU 端粗暴 cudaDeviceSynchronize()。

CPU 提交任务

把 H2D、kernel、D2H 放进一个或多个 stream

Stream 保证顺序

同一 stream 内先提交的先执行

Event 表达依赖

一个 stream 可以等待另一个 stream 的完成点

GPU runtime/driver 派发

把 ready 的 work 交给 GPU 执行

硬件执行 kernel

进入 block/SM/warp 层级

这里要强调：stream 不等于 SM 调度器。Stream 决定的是 kernel、memcpy 等任务之间的顺序和并发机会；block 和 warp 怎么在 SM 里执行，是更底层的硬件执行机制。

CUDA Block / Warp 调度：吞吐优先

一次 kernel launch 会产生一个 grid，grid 里包含很多 block。GPU 的工作是把这些 block 分配到 SM 上执行。

block 是调度到 SM 的基本单位

block 也常被称为 CTA（Cooperative Thread Array）。一个 block 里的 thread 可以：

共享 shared memory；
使用 __syncthreads() 做 block 内同步；
通过 thread/block index 处理不同数据。

GPU 硬件会把 block 调度到某个 SM。一个 block 一旦驻留到 SM 上，通常不会像 CPU task 那样被 CFS 时间片频繁抢占并迁移到另一个 SM，而是运行到完成。

这带来两个重要结论：

不同 block 之间默认没有执行顺序保证。
block 数量要足够多，否则 SM 可能吃不满。

block residency：为什么不是想放多少就放多少

一个 SM 能同时驻留多少个 block，受多种资源限制：

限制项	为什么影响 block 驻留
每个 block 的 thread 数	SM 可同时容纳的 thread/warp 数有限
每个 thread 的 register 数	register file 容量有限，register pressure 高会降低 occupancy
每个 block 的 shared memory	shared memory 被 block 独占，使用多会减少可驻留 block
架构上限	每个 SM 最大 block 数、warp 数、thread 数有硬件限制

所以 CUDA 调优不是“block 越大越好”。block 太小，单个 block 并行度不足；block 太大，可能占用太多 register/shared memory，导致 SM 上可同时驻留的 warp 变少。

warp scheduler：隐藏内存延迟

block 被放到 SM 后，block 内 thread 会被组织成 warp。NVIDIA GPU 上一个 warp 通常是 32 个 thread。SM 内部的 warp scheduler 会在多个 ready warp 之间选择并发射指令。

它的关键目标不是公平，而是吞吐：

某个 warp 等 HBM 访存时，切换到另一个 ready warp；
某个 warp 遇到长延迟指令时，让其他 warp 填补流水线；
通过足够多 active warp 隐藏内存和执行延迟。

Block 驻留 SM

block 占用 register、shared memory、thread slots

Thread 组成 warp

通常 32 个 thread 一组

Warp 等待访存

当前 warp 可能 stalled

选择 ready warp

warp scheduler 发射另一个可执行 warp

提高吞吐

用并发 warp 隐藏延迟，而不是追求单线程低延迟

这就是 GPU 和 CPU 的核心差异之一：CPU 用复杂控制逻辑优化单线程延迟；GPU 用大量 warp 并发隐藏延迟，追求整体吞吐。

CFS vs CUDA 调度：核心对比

维度	Linux CFS	CUDA block/warp 调度
所在层级	操作系统内核	GPU runtime/driver + GPU 硬件
调度对象	进程/线程，也就是 task	grid 中的 block/CTA；SM 内部的 warp
目标	公平性、响应性、CPU 时间共享	吞吐、occupancy、隐藏访存延迟
核心指标	`vruntime`、nice/weight、调度延迟	active warps、occupancy、warp stall、SM utilization
抢占方式	可以抢占当前 CPU task	block 通常运行到完成，GPU 任务抢占粒度更粗
切换代价	CPU 上下文切换较频繁，但相对可控	GPU kernel/上下文抢占代价更高，不适合频繁时间片化
公平性	明确追求公平份额	不以线程公平为主，强调硬件利用率
程序员控制	nice、cgroup、affinity、priority、policy	grid/block 配置、stream/event、kernel 设计
典型问题	线程过多、上下文切换、锁竞争、CPU 抢占	block 太少、occupancy 低、warp divergence、访存不合并

面试中最容易犯的错是说“GPU 也像 CPU 一样靠 OS 调度每个 thread”。这不对。CUDA thread 是 GPU kernel 内的逻辑执行单元，不是 Linux CFS 直接调度的 OS thread。

Stream/Event vs Thread Block/Warp：不同层次

对 CUDA 程序员来说，最需要区分的是这两层：

层次	你控制什么	典型 API / 概念	解决什么问题
任务级调度	kernel、memcpy、依赖、并发机会	`cudaStream_t`、`cudaEvent_t`、`cudaMemcpyAsync`	计算/拷贝重叠，多 kernel 并发，减少 CPU 等待
kernel 内执行	block 数、thread 数、shared memory、访存模式	`gridDim`、`blockDim`、`threadIdx`、warp、SM	铺满 SM，减少 divergence，提高 occupancy 和访存效率

例子：

stream1: H2D batch0 → kernel batch0 → D2H result0
stream2: H2D batch1 → kernel batch1 → D2H result1

这是 stream 层的并发组织。每个 kernel batchX 内部又会被拆成 grid/block/thread，并由 GPU 把 block 分配到 SM、把 thread 组织成 warp。

换句话说：

Stream/Event 解决“多个 GPU 任务之间怎么排队、等待、重叠”。
Block/Warp 调度解决“一个 kernel 内部怎么并行执行”。

和 AI Infra 的关系

这个对比不是纯 CUDA 八股。它能解释很多系统现象。

为什么 GPU 训练任务不适合像 CPU 一样频繁抢占？

CPU 线程抢占和上下文切换相对常见，CFS 正是为共享 CPU 时间设计的。但 GPU 训练任务通常有：

大量显存状态：模型参数、梯度、optimizer state、activation；
NCCL communicator 和多 rank 同步；
kernel 执行和通信 overlap；
CUDA context、缓存分配器、通信 buffer；
checkpoint 恢复成本。

如果频繁像 CPU 时间片一样抢占 GPU 训练任务，可能会导致上下文切换、显存迁移、通信重建和 checkpoint 回滚成本远大于收益。所以集群调度里 GPU 抢占往往要 checkpoint-aware、gang-aware，而不是简单时间片轮转。

为什么 Time Slicing 和 MPS 不等于 CUDA block 调度？

Kubernetes / NVIDIA device plugin 里的 GPU time-slicing、MPS 属于多进程/多 Pod 共享 GPU 的资源管理机制。

Time Slicing：多个进程或容器按时间片共享 GPU。
MPS：多个 CUDA 进程通过 MPS server 更高效共享 GPU 执行资源。
CUDA block/warp 调度：一个 kernel 内部 block 和 warp 如何映射到 SM。

它们层次不同。不能说“开了 time-slicing 后，K8s 会调度 CUDA thread block”。K8s 只调度 Pod；NVIDIA device plugin/driver 负责 GPU 资源共享；kernel 内部 block/warp 仍由 GPU 硬件调度。

为什么 CPU 线程很多会拖慢 GPU 任务？

GPU kernel 虽然在 GPU 上执行，但 GPU 程序仍依赖 CPU：

CPU 负责 dataloader、预处理、tokenization；
CPU 发起 kernel launch；
CPU 提交 H2D/D2H；
CPU 管理 CUDA runtime、NCCL、RPC；
CPU 处理服务端请求、排队和调度。

如果 CPU 侧线程过多、上下文切换严重、NUMA 绑定不合理或 dataloader 卡住，GPU 可能出现空洞：SM 没活干、GPU Util 周期性下降。此时问题不是 CUDA block 调度，而是 CPU 供给链路和 OS 调度压力。

高频面试问答

Q: Linux CFS 和 CUDA thread block 调度有什么本质区别？

Linux CFS 是 CPU 上的操作系统调度器，调度对象是进程或线程，目标是公平性、响应性和 CPU 时间共享。它用 vruntime、nice 权重、抢占和上下文切换决定哪个 task 在 CPU core 上运行。CUDA thread block 调度是 GPU kernel 内部的执行机制，调度对象是 grid 中的 block/CTA，GPU 把 block 分配到 SM 上；block 内 thread 再组成 warp，由 SM 的 warp scheduler 选择 ready warp 发射指令。CPU 调度强调公平和低延迟，GPU 调度强调吞吐和隐藏内存延迟。

面试口径：CFS 调 OS task，CUDA block 调 kernel 内工作单元；前者公平，后者吞吐。

Q: CUDA Stream/Event 和 Thread Block/Warp 调度是什么关系？

它们处在不同层次。Stream/Event 是任务级异步调度和依赖管理工具，用来组织 H2D、kernel、D2H 等 GPU work 的顺序、等待和重叠。同一个 stream 内顺序执行，不同 stream 可以并发；event 可以表达跨 stream 依赖。Thread Block/Warp 调度是 kernel 内部的硬件并行执行机制：一个 kernel launch 产生 grid，grid 中的 block 被调度到 SM，block 内 thread 被组织成 warp，SM 的 warp scheduler 在 ready warp 之间切换。

Stream 管 kernel 之间，block/warp 管 kernel 里面。

Q: GPU block 一旦被调度到 SM 后，会像 CPU 线程一样被频繁抢占吗？

通常不会。CPU 线程是 OS 调度实体，CFS 可以通过时间片和抢占频繁切换 runnable task。CUDA block 是 kernel 内部的工作单元，一旦驻留到某个 SM 上，通常运行到完成；SM 内部通过切换 ready warp 来隐藏延迟，而不是像 OS 一样把 block 按时间片迁移来迁移去。现代 GPU 支持某些粒度的抢占能力，但相对 CPU task 抢占更粗、更贵，AI 训练平台也不会把它当作常规公平时间片机制使用。

Q: 为什么 GPU 调度更强调吞吐而不是公平？

GPU 的设计目标是把大量相似计算并行执行，尽量提高 SM、Tensor Core、HBM 带宽等资源的利用率。kernel 内部的 thread/warp 往往属于同一个计算任务，不需要像多用户 CPU 进程那样做强公平分配。SM 的 warp scheduler 更关心哪个 warp ready、能不能填满流水线、能不能隐藏访存延迟。公平性通常出现在更上层，例如多 Pod 共享 GPU、集群队列 quota、租户配额，而不是单个 kernel 内每个 CUDA thread 的公平时间片。

Q: 如果 GPU 利用率周期性掉低，这和 CFS 有关系吗？

可能有间接关系。GPU 利用率掉低可能是 GPU kernel 本身太小、block 不足、访存低效，也可能是 CPU 侧没有及时喂数据。CPU 侧 dataloader、tokenizer、RPC worker、kernel launch 线程都受 Linux 调度影响。如果 CPU 线程过多、上下文切换严重、NUMA 远端访问、锁竞争或 IO 阻塞，GPU 就会等输入或等 launch，表现为周期性空洞。因此排查时要同时看 GPU timeline 和 CPU 侧 perf/top/线程状态，而不是只看 CUDA kernel。

Q: CFS 的 vruntime 可以类比 GPU 的什么指标？

严格来说没有一一对应。vruntime 是 CFS 为了公平分配 CPU 时间定义的虚拟运行时间，用来决定哪个 OS task 更应该运行。GPU kernel 内部没有为每个 CUDA thread 维护类似 vruntime 的公平指标。GPU 更常看的指标是 active warps、occupancy、warp stall、SM Active、memory coalescing、HBM bandwidth 等，它们衡量的是硬件吞吐和延迟隐藏效果，而不是公平份额。

内容模块

性能指标与诊断

核心概念

指标族	代表指标	回答什么问题	常见误区
算力	TFLOPS、TOPS、Tensor Core Util	计算单元的理论上限和实际吞吐	理论峰值不等于模型实际速度
显存带宽	HBM GB/s、Memory Throughput	数据能否及时喂给计算单元	显存容量和显存带宽不是一回事
显存容量	VRAM used、peak active memory	模型、激活、KV cache 能不能放下	显存占用高不代表 GPU 利用率高
利用率	GPU-Util、SM Active、Occupancy	时间上是否有活、空间上是否铺满	GPU-Util 高不代表忙得有效
互联	PCIe、NVLink、IB/RDMA、NCCL 时间	多卡和跨机通信是否拖慢	单卡指标无法解释多卡扩展效率
能耗	Power、tokens/J、TFLOPS/W	成本、供电、散热和能效	只看吞吐不看 TCO

GPU 性能指标全景

GPU 性能优化不要只看一个指标。面试中最好按“算力、显存、利用率、互联、能耗”五条线回答，然后用 Roofline 判断瓶颈。

算力

TFLOPS / TOPS

衡量计算峰值和实际计算吞吐。关注理论峰值、实际 TFLOPS、Tensor Core 是否被用上。

显存带宽

GB/s · TB/s

衡量 HBM 读写速度。Softmax、LayerNorm、KV Cache 读取通常更受带宽限制。

显存容量

决定模型权重、优化器状态、激活值和 KV Cache 能不能放下。

利用率

SM Active · Occupancy

用于判断 GPU 是否在忙，但不能单独代表模型效率或吞吐。

互联

NVLink · PCIe · IB

决定多卡训练、张量并行、流水并行和跨节点通信效率。

能耗

Watt · tokens/J

数据中心场景必须关注功耗、散热、供电和单位吞吐成本。

指标之间的关系图

可以用下面这条链路理解性能指标：模型工作负载先决定计算量和访存量，再被 GPU 的算力、带宽、容量和互联约束，最终表现为吞吐、延迟和成本。

计算性能：TFLOPS 不是实际速度

TFLOPS 是每秒浮点运算次数。厂商宣传的 A100 FP16 312 TFLOPS、H100 FP16 989 TFLOPS 都是理想条件下 Tensor Core 的理论峰值，真实模型通常达不到。

精度	A100 理论峰值	H100 理论峰值	典型场景
FP64	9.7 TFLOPS	34 TFLOPS	科学计算，AI 训练较少使用
FP32	19.5 TFLOPS	67 TFLOPS	通用 CUDA 计算
TF32	156 TFLOPS	494 TFLOPS	Tensor Core 加速的 FP32 训练
FP16/BF16	312 TFLOPS	989 TFLOPS	混合精度训练、主流 LLM 训练
FP8	不支持	1979 TFLOPS	H100 Transformer Engine，推理和部分训练
INT8	624 TOPS	3958 TOPS	量化推理，吞吐优先

怎么估算理论峰值

简化理解：理论峰值 ≈ SM 数量 × 每个周期可完成的矩阵运算量 × 频率。Tensor Core 是矩阵乘加专用单元，所以 FP16/BF16/FP8 峰值远高于普通 FP32 CUDA Core。

为什么实际值低

真实模型会受到 HBM 带宽、kernel launch、算子碎片、通信等待、数据依赖、batch/shape 不规则和 Tensor Core 对齐条件影响。实际 TFLOPS 通常需要用 Nsight Compute、框架 profiler 或 MFU 估算。

Roofline 模型：判断 kernel 是缺算力还是缺数据

Roofline 把 kernel 的 FLOPs、Bytes、实际吞吐和硬件峰值放在一张图里，判断瓶颈是算力、显存带宽还是其他因素。

要看什么	指标来源	说明
Arithmetic Intensity	FLOPs / Bytes	低于 ridge point 通常偏 memory-bound
Achieved FLOP/s	Nsight Compute / profiler	距离 compute roof 有多远
Memory Throughput	Nsight / DCGM	是否接近 HBM roof
Roofline 位置	Nsight Compute Roofline chart	判断优化方向：减少访存还是提高计算吞吐

Roofline 和 GPU Utilization 的关系

nvidia-smi GPU-Util 只能告诉你 GPU 上有没有 kernel 在跑，Roofline 更进一步问：这个 kernel 的性能离硬件理论上限有多远？它是被内存带宽限制还是被计算峰值限制？所以 GPU Utilization 高不代表接近 Roofline。可能出现 GPU-Util = 100% 但 kernel 在 Roofline 图上离屋顶很远，原因可能是访存不合并、cache miss、occupancy 低、warp stall、Tensor Core 没用上、指令依赖、分支发散、kernel 太小等。

怎么获取 Roofline？

Nsight Compute 支持 Roofline 分析，常见做法：

ncu --set full ./your_program
# 或使用带 Roofline 的 section
ncu --section SpeedOfLight_RooflineChart ./your_program

然后在 Nsight Compute UI 里看 Roofline 图，它通常会展示 arithmetic intensity、achieved FLOP/s、memory roof、compute roof、kernel 点位，以及 FP32 / FP64 / Tensor Core 等不同 roof。也可以手动估算：Arithmetic Intensity = FLOPs / Bytes，Achieved Performance = FLOPs / Kernel Time，Memory Bandwidth Used = Bytes / Kernel Time，然后对比硬件峰值。

Q: 什么是 Roofline 模型？面试怎么回答？

先估 FLOPs 和 Bytes，算 Arithmetic Intensity，再和硬件 ridge point 比较，判断热点 kernel 是 memory-bound 还是 compute-bound。

Roofline 把 GPU 性能指标收敛成“缺数据还是缺算力”。

Q: 为什么 Transformer 的 Attention 经常是瓶颈？

Self-Attention 的 QK^T 计算量随序列长度平方增长，同时需要频繁读写 Q、K、V 和 attention score。FlashAttention 的核心价值不是只减少 FLOPs，而是通过分块、在线 softmax 和重计算减少 HBM 读写。

GPU 利用率：高 Util 不等于高效率

nvidia-smi 的 GPU Utilization 表示采样窗口内 GPU 是否有 kernel 活跃，不能代表 Tensor Core 利用率、真实 TFLOPS 或端到端吞吐。

指标	主要来源	真正含义	常见误区
GPU Utilization	`nvidia-smi`	采样周期内至少有 kernel 在执行的时间比例	100% 可能只是小 kernel 很密集，不代表算力打满
SM Active	DCGM / Nsight	SM 有活跃 warp 的时间比例	SM 忙不等于 Tensor Core 高效工作
Tensor Core Util	Nsight Compute	Tensor Core 管线使用程度	FP32 或 shape 不对齐时可能很低
Occupancy	Nsight Compute	每个 SM 上可驻留 warp 与理论上限的比例	高 occupancy 不一定高性能，低 occupancy 也可能是合理的寄存器/共享内存权衡
Memory Throughput	Nsight / DCGM	HBM 带宽使用情况	高带宽可能说明 memory-bound，不一定是好事

Q: nvidia-smi 显示 GPU Util 100%，但训练很慢，为什么？

可能是小 kernel 频繁启动、Tensor Core 没用上、CPU 数据加载慢、H2D/D2H 拷贝等待、多卡 NCCL 通信等待，或者模型本身 memory-bound。排查时要结合 Nsight Systems 看 timeline，再用 Nsight Compute 看单个 kernel 的 Tensor Core、SM、HBM 指标。

多卡互联性能指标

多 GPU 训练/推理时，互联决定扩展效率。节点内优先看 NVLink/NVSwitch，节点间看 InfiniBand/RoCE 和 NCCL 拓扑。

互联方式	典型带宽	拓扑	主要影响
NVLink 3.0	A100 约 300 GB/s	GPU 点对点或经 NVSwitch	单机多卡 all-reduce、张量并行
NVLink 4.0	H100 约 450 GB/s	GPU 点对点或经 NVSwitch	Hopper 单机扩展效率
NVSwitch	高聚合带宽	单节点内近似全互联	降低拓扑不均衡影响
PCIe 4.0 x16	约 32 GB/s 单向	树形，经 CPU 或 PCIe Switch	CPU-GPU 拷贝、无 NVLink 卡间通信
PCIe 5.0 x16	约 64 GB/s 单向	树形	新一代 CPU-GPU 互联
InfiniBand NDR	400 Gbps = 50 GB/s	跨节点网络	多节点 all-reduce、参数同步

扩展效率

线性扩展效率 = 多卡实际吞吐 /（单卡吞吐 × 卡数）。节点间扩展效率通常低于单节点，因为跨节点带宽和延迟远不如 NVLink/NVSwitch。

面试高频：如何判断瓶颈

SM Active 高 + HBM 高GPU 忙且大量访存，进一步看 Tensor Core 和 Roofline。

SM Active 低 + HBM 高典型 memory-bound，优化访存、融合算子、减少 KV/激活读写。

SM Active 低 + HBM 低GPU 没被喂饱，查 CPU、DataLoader、I/O、网络、调度等待。

GPU Util 高 + TFLOPS 低可能是小 kernel、Tensor Core 未使用、shape 不对齐或通信 kernel 占比高。

单卡快 + 多卡慢查 NCCL、拓扑、IB/RDMA、梯度 bucket、通信计算 overlap。

显存够但仍慢容量不是瓶颈，可能是带宽、NUMA、PCIe、page cache 或数据 pipeline。

GPU 性能面试题

Q: 如何计算一个模型训练需要多少显存？

回答思路：先把显存拆成“模型状态 + 激活值 + 临时开销”，再说明不同优化策略会减少哪一部分。

1. 模型状态

包括参数、梯度和优化器状态。以 Adam + FP16 混合精度为例，常用粗估是每个参数约 16 bytes：FP16 参数 2 bytes、FP16 梯度 2 bytes、FP32 master weight 4 bytes、Adam 一阶动量 4 bytes、二阶动量 4 bytes。

2. 激活值

激活值和 batch size、sequence length、hidden size、层数强相关，训练时需要保留中间结果用于反向传播。长序列 LLM 里，激活值经常比直觉中更大。

3. 临时开销

还要预留 CUDA workspace、attention score / KV cache、中间 tensor、通信 buffer、内存碎片等。实际跑训练时不能把显存算到 100%，通常要留安全余量。

4. 优化手段

ZeRO/FSDP 主要切分参数、梯度和优化器状态；activation checkpointing 通过重计算降低激活值；offload 把部分状态放到 CPU 或 NVMe；混合精度和量化降低单个元素字节数。

面试口径：显存 ≈ 模型状态 + 激活值 + 临时 workspace/通信 buffer。Adam + FP16 可先按参数量 × 16 bytes 粗估模型状态，再叠加激活值和安全余量。

Q: A100 和 H100 的主要区别是什么？

回答思路：不要只背参数表，建议按“架构、算力、显存/带宽、互联、软件特性、部署代价”来讲。

1. 架构代际

A100 是 Ampere 架构，H100 是 Hopper 架构。H100 面向大模型训练和推理做了更多专用优化。

2. 计算能力

H100 的 FP16/BF16 Tensor Core 峰值明显高于 A100，并新增 FP8 Transformer Engine。对 LLM 训练和推理来说，FP8 能在合适场景下提升吞吐和降低显存占用。

3. 显存与带宽

A100 80GB 使用 HBM2e，H100 80GB 使用 HBM3，H100 显存带宽更高。带宽提升对 memory-bound 算子、attention、KV cache 读取很重要。

4. 互联与扩展

H100 支持更高代际的 NVLink/NVSwitch，单机多卡和多机训练的通信效率更好，适合大规模并行训练。

5. 工程代价

H100 性能更强，但功耗、散热、供电和采购成本也更高。数据中心部署时要看机房能力、网络拓扑和整体 TCO，而不是只看单卡峰值。

面试口径：H100 相比 A100 的核心升级是 Hopper 架构、FP8/Transformer Engine、更高 Tensor Core 算力、更高 HBM 带宽和更强互联；代价是功耗和部署要求更高。

Q: Tensor Core 为什么比 CUDA Core 快？

回答思路：本质区别是 CUDA Core 做通用标量/向量运算，Tensor Core 是矩阵乘加专用单元，专门加速深度学习里的 GEMM。

1. 硬件定位不同

CUDA Core 更通用，适合执行标量或向量 FMA；Tensor Core 面向矩阵块乘加，可以在一个周期内完成小矩阵 tile 的大量乘加操作。

2. 工作负载匹配

神经网络里的 Linear、卷积、QKV projection、MLP、attention projection 最终都能转成矩阵乘。Tensor Core 正好针对这些高密度矩阵乘做了硬件加速。

3. 精度与吞吐权衡

Tensor Core 通常使用 TF32、FP16、BF16、FP8、INT8 等精度格式，通过降低数据位宽和专用矩阵管线换取更高吞吐。

4. 使用条件

并不是所有算子都会自动高效使用 Tensor Core。shape 对齐、数据类型、矩阵维度、框架 kernel 选择都会影响 Tensor Core 利用率。

面试口径：Tensor Core 快是因为它不是“更快的 CUDA Core”，而是矩阵乘加专用硬件；深度学习核心算子大量是 GEMM，所以能获得数量级更高的吞吐。

Q: Memory Coalescing 是什么？

回答思路：Memory Coalescing 讨论的是一个 warp 内的内存访问能不能合并成少量连续内存事务。

1. 基本定义

GPU 以 warp 为单位调度线程。如果同一个 warp 里的相邻线程访问连续、对齐的地址，硬件可以把多次访存合并成更少的 memory transaction。

2. 为什么重要

合并访问可以提高有效 HBM 带宽，减少访存事务数量，让 SM 少等数据。很多 memory-bound kernel 的优化重点就是提高访存连续性。

3. 反例

stride 访问、非对齐访问、随机访问、不同线程访问分散地址都会破坏 coalescing，导致同样的数据量需要更多 memory transaction。

4. 优化方向

常见做法包括调整数据布局、让线程映射到连续地址、使用 shared memory 做重排、向量化 load/store，以及减少不规则访问。

面试口径：Coalescing 的目标是让一个 warp 的访存尽量连续且对齐，把零散访存合并成少量事务，从而提升有效带宽。

官方参考

official

NVIDIA Ampere Architecture

A100 架构深度解析，SM、Tensor Core、HBM 详解。

official

NVIDIA Hopper Architecture

H100 架构深度解析，Transformer Engine、FP8、DPX。

official

Nsight Compute Profiling Guide

GPU 性能分析指标详解，Roofline、Occupancy、Memory 分析。

paper

FlashAttention

内存高效 Attention 算法，IO-Aware 优化经典论文。

诊断入口

业务慢或 GPU 看起来异常

-> nvidia-smi / DCGM 粗看 GPU-Util、显存、功耗、温度、进程

-> Nsight Systems 看 timeline：空洞、Memcpy、同步、NCCL、CPU 等待

-> Nsight Compute 看热点 kernel：SM Active、Occupancy、Warp Stall、Throughput

-> 结合业务指标：step time、tokens/sec、QPS、TTFT、TPOT、P99

指标解释

指标	能说明什么	不能说明什么
GPU-Util	采样窗口里 GPU 是否有 kernel 在跑	不能说明 SM 是否铺满、Tensor Core 是否用上
memory.used	显存容量占用	不能说明 GPU 是否在计算
SM Active	SM 空间覆盖度	不能说明每个 SM 内 warp 是否足够
Occupancy	SM 内 active warp 是否够多	不是越高越好，不等于性能
Warp Stall	warp 为什么不能发射	需要结合访存、同步、依赖和代码结构解释
Compute / Memory Throughput	瓶颈更像算力还是显存带宽	不能解释 CPU、I/O、通信等待
Tensor Core Util	AI 矩阵算力是否用起来	只对 GEMM/Conv/Attention 等相关

GPU-Util 高不等于 GPU 真正用得好

判断一个 GPU “利用率高不高”，不要只看 nvidia-smi 里的 GPU-Util。它只能告诉你采样窗口里 GPU 时间上是否有 kernel 在跑，但不能告诉你 SM 是否铺满、warp 是否真的能发射、Tensor Core 是否用起来、瓶颈到底在计算还是显存。nvidia-smi 本身定位是 NVIDIA System Management Interface，用于监控和管理 NVIDIA GPU，并支持查询 GPU、显存、功耗、温度、时钟等系统级信息[nvidia-smi Documentation]。

最推荐的判断链路：GPU-Util 看有没有活；SM Active 看活有没有铺满 SM；Occupancy 看 SM 里 warp 是否足够；Warp Stall 看 warp 是否真能执行；Compute/Memory Throughput 看瓶颈在哪；Tensor Core Util 看 AI 算力是否用上；最后用吞吐和延迟判断 GPU 忙得是否有效。

一条排查主线：从系统级到业务级

你可以把 GPU 利用率诊断想成十层漏斗。越往上越粗，适合快速判断；越往下越细，适合解释为什么慢。面试时不要一上来报一堆指标，而是先说你的排查顺序。

1. GPU 有没有活干？
   看 GPU-Util。

2. 显存只是占着，还是 GPU 真的在计算？
   区分 memory.used 和 utilization.memory。

3. 活有没有铺满 SM？
   看 SM Active。

4. SM 里面挂了多少 warp？
   看 achieved occupancy / active warps。

5. warp 是否真的能发射？
   看 eligible warps、issued warps、stall reasons。

6. 瓶颈在 compute 还是 memory？
   看 compute throughput 和 memory throughput。

7. AI workload 是否用上 Tensor Core？
   看 tensor pipe / tensor instruction。

8. 是否被功耗、温度、时钟限制？
   看 power、clock、temperature、p-state。

9. 端到端 timeline 有没有空洞？
   用 Nsight Systems 看 CPU、CUDA、memcpy、NCCL。

10. GPU 忙得是否有业务价值？
    训练看 tokens/sec、samples/sec、step time、MFU；推理看 QPS、TTFT、TPOT、P99。

这个顺序的好处是：它不会把“设备看起来很忙”和“业务真的跑得快”混为一谈。很多线上问题的本质不是 GPU 不忙，而是 GPU 忙在了不该忙的地方，例如重计算、低效访存、通信等待、碎 kernel、错误 dtype 或无效 batch。

第一层：GPU-Util 只看时间上有没有 kernel

GPU-Util 回答的是一个非常粗的问题：在采样周期内，GPU 上是否有一个或多个 kernel 正在执行。它不是 SM 利用率，不是 CUDA Core 利用率，也不是 Tensor Core 利用率，更不是实际 FLOPS。

怎么获取

# 快速看当前 GPU 状态
nvidia-smi

# 持续刷新，观察利用率、功耗、温度等变化
nvidia-smi dmon

# 查询指定字段
nvidia-smi --query-gpu=timestamp,index,name,utilization.gpu,utilization.memory,memory.used,memory.total,power.draw,temperature.gpu --format=csv -l 1

怎么理解

如果 GPU-Util 长期低于 30%，大概率是 GPU 没喂饱，可能是数据加载慢、CPU 预处理慢、batch 太小、请求量不足、I/O 慢、kernel launch 间隔大或多卡通信等待。如果 GPU-Util 长期高于 80%，只能说明 GPU 时间上比较忙，后面还要继续判断它是不是忙得有效。

典型陷阱

一个很小的 kernel 如果持续运行，也可能让 GPU-Util 接近 100%。例如一张有 108 个 SM 的 A100，如果某个 kernel 只用到少数 SM，但一直不断执行，GPU-Util 可能很高，实际硬件并没有被铺满。

第二层：显存占用不等于 GPU 利用率

显存占用和 GPU 计算利用率是两件事。memory.used 表示显存容量用了多少；utilization.memory 更偏采样周期内 device/global memory 是否在读写。模型参数、KV cache 或缓存 allocator 占着显存，不代表 GPU 正在高效计算。

怎么获取

nvidia-smi --query-gpu=memory.used,memory.total,utilization.memory --format=csv -l 1

怎么判断

显存高但 GPU-Util 低，常见于模型已经加载、KV cache 或 tensor 占着显存，但请求流量不足、CPU 没喂上或程序在等待。显存低但 GPU-Util 高，也很常见，例如某些小模型或小 batch 计算忙但显存占用不大。显存接近满说明 batch size、sequence length 或并发可能受限，也有 OOM 风险，但它不是“利用率高”的证据。

第三层：SM Active，看活有没有铺满 GPU

SM Active 比 GPU-Util 更接近我们想知道的问题：GPU 上的 SM 有没有被广泛使用。可以粗略理解为，在所有 SM 和所有采样时间里，有多少比例的 SM 至少有 active warp。

为什么它比 GPU-Util 更细

GPU-Util 是“这段时间 GPU 有没有 kernel”；SM Active 是“这些 SM 在这些时间里有没有活”。如果某个 kernel 只铺到少数 SM，它仍然可能让 GPU-Util 很高，但 SM Active 会暴露空间利用率不足的问题。

怎么获取

# 用 Nsight Compute 采集完整指标
ncu --set full ./your_program

# 只采集 SM Active 相关指标
ncu --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed ./your_program

怎么判断

SM Active 低，通常说明并行度不足，例如 batch 太小、grid 太小、算子规模太小或 kernel 设计没有铺开。SM Active 高但性能仍低，说明大多数 SM 时间上有活，但 warp 可能在 stall，或者计算/访存效率不好，需要继续看 occupancy、eligible warps 和 throughput。

第四层：Occupancy，看 SM 里有多少 warp 可调度

Occupancy 可以粗略理解为：一个 SM 上 active warps 数量占该 SM 最大支持 active warps 数量的比例。它的作用是判断 SM 里是否有足够多的 warp 用来隐藏访存、依赖和执行延迟。

Theoretical vs Achieved

Theoretical Occupancy 可以根据 launch 参数、block size、register 使用、shared memory 使用和 GPU 架构限制静态估算；Achieved Occupancy 是运行时真实采集到的结果。前者告诉你理论上最多能挂多少 warp，后者告诉你实际 workload 中挂了多少 warp。

怎么获取

ncu --set full ./your_program

# 常看 Nsight Compute 中的这些 section
# Launch Statistics
# Occupancy
# Scheduler Statistics

# 也可以指定 active warps 指标
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./your_program

不要误解

Occupancy 高不等于性能高。它只是说明 SM 上有足够多的 warp 可用于隐藏 latency。如果 kernel 已经 memory-bound，提高 occupancy 可能没有明显收益；如果为了提高 occupancy 而减少寄存器导致更多访存，反而可能变慢。面试时要强调：occupancy 是诊断指标，不是最终目标。

第五层：Warp Stall，看 warp 是否真的能发射

一个 warp 是 active，不代表它当前周期就能执行。更细地看，一个 warp 可能处于 active、eligible、selected 或 stalled 状态。真正影响吞吐的是：有多少 warp 已经 ready，可以被 scheduler 发射。

active warp：
  驻留在 SM 上，还没结束。

eligible warp：
  当前周期已经准备好，可以发射指令。

selected warp：
  当前周期被 scheduler 选中发射。

stalled warp：
  因为内存、依赖、同步、barrier 等原因暂时不能发射。

关键判断

如果 active warps 很多，但 eligible warps 很少，说明 SM 里看起来挂了很多 warp，但大部分都在等。此时 GPU 可能很忙，但忙得不高效。

常见 stall 原因

Long Scoreboard：通常表示在等 global memory / L2 / DRAM。
Short Scoreboard：通常表示在等 shared memory 或短延迟依赖。
Barrier：等待 __syncthreads() 或类似同步。
Wait：指令依赖等待。
Not Selected：有 eligible warp，但本周期没被 scheduler 选中。
Math Pipe Throttle：计算管线压力大。
MIO Throttle：memory input/output 管线压力大。

怎么获取

ncu --set full ./your_program

# 重点看：
# Scheduler Statistics
# Warp State Statistics
# Source Counters

第六层：Compute Throughput，看算力侧是否接近峰值

Compute Throughput 或 SM Throughput 更接近“算力有没有打满”这个问题。它回答的是计算侧吞吐达到硬件峰值的多少，而不是 GPU 时间上有没有活。

怎么获取

ncu --set full ./your_program

# 重点看：
# GPU Speed Of Light Throughput
# Compute Workload Analysis

# 示例指标
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./your_program

怎么判断

Compute Throughput 高，可能是 compute-bound；Compute Throughput 低但 SM Active 高，说明 SM 有活但效率不高，需要看 stall、访存和指令 mix；Compute Throughput 低而 Memory Throughput 高，通常是 memory-bound；两者都低，可能是同步、依赖、分支、launch overhead 或数据不连续。

第七层：Memory Throughput，看是不是卡在显存和访问模式

很多 AI 算子并不是 compute-bound，而是 memory-bound。典型例子包括 LayerNorm、Softmax、Embedding、Gather/Scatter、Elementwise、小 batch inference。它们的共同特点是每读写一批数据，只做相对较少计算，算术强度不高。

怎么获取

ncu --set full ./your_program

# 重点看：
# Memory Workload Analysis
# GPU Speed Of Light Throughput
# L1/TEX、L2、DRAM 相关指标

怎么判断

如果 DRAM throughput 高、compute throughput 低，通常是 memory-bound。如果 L2 hit rate 低，说明访问局部性差。如果 global load/store 效率差，说明访存不合并或访问模式不友好。如果 memory util 高但业务吞吐低，可能是带宽被低效访问消耗掉了。

优化方向

常见优化包括 coalesced memory access、提高 L2/cache hit、减少 global memory 访问、使用 shared memory 做 tile 复用、融合算子减少中间结果读写，以及调整数据布局让同一个 warp 内访问更连续。

第八层：Tensor Core Utilization，看 AI 算力是否用起来

对于 Transformer、GEMM、Conv、Attention 这类深度学习 workload，判断 GPU 是否用好，还必须看 Tensor Core 是否真正参与。模型跑在 GPU 上，不代表它用了 Tensor Core。

怎么获取

ncu --set full ./your_program

# 重点看：
# GPU Speed Of Light Throughput
# Compute Workload Analysis
# Instruction Statistics

# 示例 tensor 相关指标
ncu --metrics sm__inst_executed_pipe_tensor.sum,sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_elapsed ./your_program

Tensor Core 没用起来的常见原因

dtype 不合适，例如使用 FP32 但没有允许 TF32。
矩阵 shape 不友好，无法形成高效 tile。
没有走 cuBLAS、cuDNN、TensorRT、CUTLASS 等高性能实现。
框架配置禁用了 TF32 或混合精度。
算子 fallback 到普通 CUDA Core kernel。
batch 太小，矩阵规模太小。

PyTorch 常见检查

import torch

torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cudnn.allow_tf32 = True

第九层：Power / Clock / Thermal，看是不是被限频

如果 GPU-Util 高、SM Active 高，但性能仍低，还要检查功耗、温度和时钟。GPU 可能因为 power limit、thermal throttle、p-state、MIG/MPS 资源隔离或容器限制，导致看起来忙但实际频率和吞吐上不去。

怎么获取

nvidia-smi -q -d POWER,CLOCK,TEMPERATURE,PERFORMANCE

nvidia-smi --query-gpu=power.draw,power.limit,clocks.sm,clocks.mem,temperature.gpu,pstate --format=csv -l 1

怎么判断

如果 power draw 接近 power limit，可能功耗受限；温度持续很高，可能 thermal throttle；clocks.sm 低，可能被降频；p-state 不理想，可能没有进入高性能状态。训练吞吐突然下降时，除了看代码和通信，也要看这些节点级硬件状态。

第十层：端到端业务指标，看 GPU 忙得是否有价值

GPU 指标高，不代表业务性能好。真正的目标不是让 GPU 看起来忙，而是让训练吞吐、推理吞吐、延迟和成本达到预期。

训练场景看什么

训练中建议看 samples/sec、tokens/sec、step time、MFU、data loading time、communication time。如果 GPU 指标很高但 tokens/sec 低，可能是重计算、通信、低效 kernel 或数据 pipeline 让 GPU 忙在了低价值工作上。

推理场景看什么

推理中建议看 QPS、tokens/sec、TTFT、TPOT、P50/P90/P99 latency、batch size、queue time、KV cache memory。LLM 推理尤其要拆开 prefill 和 decode，因为 prefill 更偏计算密集，decode 常受 KV cache 读写、batch 组织和访存影响。

工具选择：每个工具解决不同层级的问题

不要指望一个工具回答所有问题。系统级状态用 nvidia-smi 和 DCGM；timeline 用 Nsight Systems；单 kernel 细节用 Nsight Compute；框架算子归因用 PyTorch Profiler。

nvidia-smi：系统级粗看

适合看 GPU-Util、显存占用、功耗、温度、时钟和进程。优点是方便快速，缺点是不能解释 kernel 为什么慢，也看不到 SM、warp、Tensor Core 细节。

nvidia-smi
nvidia-smi dmon -s pucvmet
nvidia-smi --query-gpu=timestamp,index,utilization.gpu,utilization.memory,memory.used,memory.total,power.draw,clocks.sm,temperature.gpu --format=csv -l 1

DCGM / Prometheus / Grafana：生产监控

DCGM 提供 GPU 指标采集、健康检查、作业统计、拓扑等能力，适合集群监控和告警[Feature Overview — NVIDIA DCGM Documentation]。

DCGM_FI_DEV_GPU_UTIL
DCGM_FI_DEV_FB_USED
DCGM_FI_DEV_POWER_USAGE
DCGM_FI_DEV_GPU_TEMP
DCGM_FI_DEV_SM_CLOCK
DCGM_FI_DEV_MEM_COPY_UTIL

Nsight Systems：端到端 timeline

Nsight Systems 适合看 CPU 线程、CUDA API、kernel、memcpy、stream、NCCL 等在时间线上如何交错；官方文档也建议通过 NVTX 或 profiler API 聚焦性能关键代码区域，减少无关数据干扰[User Guide — Nsight Systems]。

nsys profile -t cuda,nvtx,osrt,cudnn,cublas -o report ./your_program

Nsight Compute / NCU：单 kernel 深挖

Nsight Compute CLI 可以从命令行 profile 应用，并为选定 kernel 采集指标；它支持通过 section 或 metrics 选择采集内容，也支持按 kernel 名称过滤[Nsight Compute CLI]。

ncu --set full ./your_program
ncu --kernel-name regex:your_kernel_name --set full ./your_program

PyTorch Profiler：框架算子归因

PyTorch Profiler 适合把 Python / PyTorch op 和 CUDA kernel 对齐起来，回答哪个 op 最耗时、CPU 时间多还是 CUDA 时间多、是不是 DataLoader 慢、是不是有大量小 kernel。

import torch
import torch.profiler as profiler

with profiler.profile(
    activities=[
        profiler.ProfilerActivity.CPU,
        profiler.ProfilerActivity.CUDA,
    ],
    record_shapes=True,
    profile_memory=True,
    with_stack=True,
) as prof:
    for step, batch in enumerate(loader):
        output = model(batch)
        loss = output.sum()
        loss.backward()
        if step > 10:
            break

print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))

实战路径：从 nvidia-smi 到 profiler

真实排障时，建议按下面的顺序走。先用低成本工具判断方向，再用 profiler 精确定位。

Step 1：先看 nvidia-smi

nvidia-smi

关注 GPU-Util、Memory-Usage、Power、Temperature 和进程。如果 GPU-Util 长期低，先怀疑数据加载、CPU preprocessing、batch 太小、I/O、请求量不足、kernel launch 太碎或多卡通信等待。如果 GPU-Util 长期高，继续判断是不是高效。

Step 2：用 Nsight Systems 看 timeline

nsys profile -t cuda,nvtx,osrt,cudnn,cublas -o report ./your_program

重点看 GPU timeline 是否有空洞、kernel 是否连续、CPU 是否频繁同步、cudaMemcpy 是否阻塞、NCCL 是否占大量时间。大量小 kernel 往往说明 launch overhead 高或算子太碎。

Step 3：用 Nsight Compute 看核心 kernel

ncu --set full ./your_program

重点看 SM Active、Achieved Occupancy、Compute Throughput、Memory Throughput、Warp Stall Reasons、Tensor Core Utilization。这里才能回答“这个 kernel 为什么慢”。

六种典型结论：如何把指标翻译成人话

面试和排障时，最重要的是把指标组合翻译成清晰结论。

A. GPU 没喂饱

GPU-Util 低
SM Active 低
Nsight Systems 看到 GPU timeline 有空洞

常见原因是 DataLoader 慢、CPU preprocessing 慢、batch 太小、I/O 慢、请求量不足、kernel launch 太碎或 CPU-GPU 同步。

B. GPU 时间上忙，但没铺满硬件

GPU-Util 高
SM Active 低
Achieved Occupancy 低

常见原因是 kernel grid 太小、batch 太小、并行度不足、算子规模小或 block 配置不合理。

C. SM 都有活，但效率低

GPU-Util 高
SM Active 高
Occupancy 高
Compute Throughput 低
Memory Throughput 也低
Warp Stall 高

常见原因是指令依赖、同步、分支发散、访存模式差、barrier 多或 atomic 多。

D. Memory-bound

SM Active 高
Memory Throughput 高
Compute Throughput 低
Long Scoreboard stall 高

优化方向是访存合并、提高 cache hit、减少 global memory 访问、使用 shared memory、融合算子、减少中间结果读写。

E. Compute-bound

SM Active 高
Compute Throughput 高
Memory Throughput 不高

优化方向是使用 Tensor Core、优化指令 mix、减少冗余计算、提高 tile 效率、使用更合适 dtype。

F. Tensor Core 没用起来

GPU-Util 高
SM Active 高
Tensor Pipe / Tensor Core Util 低

如果 workload 是 GEMM、Conv 或 Attention，需要检查 FP16/BF16/TF32 是否开启，shape 是否对齐，是否使用高性能库，是否 fallback 到普通 CUDA kernel，batch size 是否太小。

排查路径

GPU-Util：
  时间上有没有 kernel。

SM Active：
  kernel 有没有铺满 SM。

Occupancy：
  SM 里 active warp 是否足够。

Eligible / Issued Warps：
  warp 是否真的能发射。

Warp Stall：
  不能发射是在等什么。

Compute Throughput：
  计算管线是否接近峰值。

Memory Throughput：
  是否被显存带宽或访问模式卡住。

Tensor Core Utilization：
  AI 矩阵算力是否真的用起来。

Power / Clock / Thermal：
  是否被功耗、温度、时钟限制。

业务指标：
  GPU 忙，是否真的换来了 tokens/sec、samples/sec、QPS 和低延迟。

诊断入口

业务指标异常

-> nvidia-smi / DCGM 看 GPU-Util、显存、功耗、PCIe/NVLink

-> Nsight Systems 看 timeline：kernel / memcpy / NCCL / 空洞 / 同步

-> Nsight Compute 看热点 kernel：compute、memory、occupancy、stall

-> 扩展性实验：单卡、多卡、不同 batch、不同 seq_len 对比

指标解释

瓶颈类型	主要证据	常见误判
Compute-bound	Compute Throughput 高、Tensor Core Util 高、Roofline 靠近 compute roof	只看 GPU-Util 高就说 compute-bound
Memory-bandwidth-bound	Memory Throughput 高、Long Scoreboard 高、Roofline 落在 memory roof	把显存容量不足和显存带宽不足混为一谈
Memory-capacity-bound	OOM、batch/seq_len/并发受限、KV cache 放不下	只看显存占用高就说带宽瓶颈
Communication-bound	NCCL/memcpy 占比高、多卡扩展效率差、拓扑敏感	通信 kernel 也会让 GPU-Util 高，容易误判为计算忙

排查路径

计算瓶颈（Compute-bound）

定义：GPU 的计算单元（CUDA Core / Tensor Core）是瓶颈，计算请求已满，但数据供给充足。

特征	表现
Compute Throughput	高，接近硬件峰值
Memory Throughput	相对不高
SM Active	高
Roofline 位置	落在 Compute Roof 附近

典型场景：大矩阵乘法（GEMM）、大 batch 的卷积、大 batch 的 prefill 阶段 Attention。

优化方向：

使用 Tensor Core，确保 dtype、shape 和 layout 对齐。
降低精度，例如 FP16/BF16/FP8/INT8。
优化 tile 大小和指令 mix。
对 decode 类任务考虑 speculative decoding，用小模型生成候选，大模型并行验证。

显存带宽瓶颈（Memory-bandwidth-bound）

定义：GPU 的显存带宽是瓶颈，计算单元在等数据，大量时间花在 HBM 读写上。

特征	表现
Memory Throughput	高，接近 HBM 带宽峰值
Compute Throughput	低
Long Scoreboard Stall	高
Roofline 位置	落在 Memory Roof（斜线区域）

典型场景：LLM decode 阶段、elementwise 算子、LayerNorm、Softmax、Embedding lookup、Gather/Scatter、小 batch 推理。

为什么 LLM decode 常是显存瓶颈：每生成一个 token，需要从 HBM 加载大量模型权重和 KV cache，但新增 token 的计算量相对少，算术强度低，容易落在 Roofline 的 memory-bound 区域。

优化方向：

提高算术强度：增大 batch size、融合算子。
减少 HBM 读写：FlashAttention、shared memory 数据复用、算子融合。
压缩数据量：KV cache 量化、权重量化。
改善访问模式：memory coalescing、提高 L2 cache hit rate。
用 PagedAttention 按需分配 KV cache 物理页，减少显存浪费。

显存容量瓶颈（Memory-capacity-bound）

容量瓶颈是指显存不够放，放不下模型权重、KV cache、activation、optimizer state 或足够的 batch 数据。它和带宽瓶颈不同：容量问题首先表现为 OOM 或并发受限，带宽问题表现为算子在等数据。

典型表现：

OOM。
被迫降低 batch size、sequence length 或并发请求数。
KV cache 容量不足导致 decode 阶段可服务请求数受限。

优化方向：

张量并行：模型分片到多卡。
量化：FP16 到 INT8/INT4，减少权重和 KV cache 占用。
ZeRO / FSDP：切分参数、梯度、优化器状态。
Offload：不活跃数据放到 CPU/NVMe。
PagedAttention + KV cache 压缩：显存超配 + 按需分配。

通信瓶颈（Communication-bound）

定义：GPU 间的数据同步和传输成为瓶颈，GPU 时间花在等待通信完成上。

特征	表现
NCCL 时间占比	高（Nsight Systems timeline 上 NCCL 占比大）
GPU-Util	可能高（通信 kernel 在跑），但 SM Active 可能低
扩展效率	多卡吞吐远低于线性扩展

典型场景：大规模张量并行 all-reduce、流水线并行跨卡激活传递、跨节点训练梯度同步、KV cache 跨卡/跨节点迁移。

优化方向：

通信与计算重叠：NCCL 和 kernel 在不同 stream 上并行。
梯度压缩或量化：减少通信数据量。
拓扑优化：优先利用 NVLink/NVSwitch，再考虑 PCIe/RDMA。
减少通信次数：梯度累积、增大 bucket size。
Ring Attention / 序列并行：减少跨卡通信量。

典型现象

判断方法	计算瓶颈	显存瓶颈	通信瓶颈
nvidia-smi	GPU-Util 高，Power 高	GPU-Util 可能不高，Memory Util 高	GPU-Util 可能高但吞吐低
Nsight Systems	kernel 密集，无空洞	kernel 可能有间隙或很短	NCCL / memcpy 占比大
Nsight Compute	Compute Throughput 高	Memory Throughput 高，Long Scoreboard 高	需结合 Nsight Systems
Roofline	落在水平区域	落在斜线区域	不直接适用
扩展性测试	单卡和多卡都慢	单卡和多卡都慢	单卡快，多卡反而慢

排查顺序：先用 nvidia-smi 看有没有活，再用 Nsight Systems 看谁占时间，然后用 Nsight Compute 看单个 kernel 瓶颈在哪，最后用扩展性测试判断是否通信瓶颈。

Q: GPU-Util 100% 但推理 QPS 很低，怎么排查？

先判断是哪类瓶颈。第一步看 Memory Throughput：如果 HBM 带宽打满、Compute Throughput 低，是显存带宽瓶颈（常见于 LLM decode）。第二步看 Nsight Systems timeline：如果 NCCL 占大量时间，是通信瓶颈。第三步看 Tensor Core Util：如果低，可能没用上 Tensor Core 或 shape 不友好。第四步看 kernel 碎片度：大量小 kernel 导致 launch overhead 高。

Q: 计算瓶颈和显存瓶颈能同时存在吗？

在同一个 kernel 上通常不会——Roofline 模型里一个点要么在斜线区域（memory-bound），要么在水平区域（compute-bound）。但在端到端场景中，不同 kernel 可以是不同瓶颈：prefill 阶段 compute-bound，decode 阶段 memory-bound，某些小算子 launch-bound。所以优化时需要针对不同阶段不同策略，这也是为什么推理引擎会把 prefill 和 decode 分开调度。

内容模块

GPU 面试题库

精通★★★⏱ 45 min

高频追问压缩表

问题	回答主线	易错点
CPU 和 GPU 为什么设计不同？	CPU 优化低延迟控制流和大缓存；GPU 用大量执行单元与高带宽追求吞吐	不要只说“GPU 核多”
SM、Block、Warp、Thread 什么关系？	Grid 包含 Block；Block 被分配到 SM；Thread 按 Warp 执行	Block 不能跨 SM 同时执行
Warp Divergence 为什么慢？	同一 Warp 的分支路径需要分别执行，有效并行度下降	不等于所有分支都会串行
Shared Memory 有什么用？	Block 内可编程片上缓存，用于复用和线程协作	它容量有限，过多会压低 Occupancy
Pinned Memory 为什么更快？	页被锁定，DMA 可稳定访问，也能支持异步拷贝	不是越多越好，会减少 OS 可分页内存
Stream 和 Event 的区别？	Stream 是有序任务队列；Event 用于依赖和计时	多 Stream 不保证自动并行
Occupancy 越高越好吗？	它帮助隐藏延迟，但寄存器/Shared Memory/指令效率同样重要	100% 不是最终目标
GPU-Util 100% 为什么还慢？	继续看 SM、Tensor Core、HBM、Warp Stall 和业务吞吐	GPU-Util 不等于峰值算力利用率
MIG、MPS、Time-Slicing 怎么选？	MIG 强隔离；MPS 提升进程并发；Time-Slicing 低成本复用	Time-Slicing 不保证固定性能份额
NVLink、PCIe、RDMA 分别在哪？	GPU-GPU、Host/设备路径、跨机网络数据面	不要把 NVLink 当跨机网络

关键机制串联

CPU 提交 Kernel

CUDA Runtime/Driver 准备 launch 与参数

Grid/Block 映射 SM

受寄存器、Shared Memory 和 Block 数限制

Warp 发射指令

分支、依赖、访存导致 Stall

数据经过内存层次

多设备互联

PCIe/NVLink/NVSwitch/RDMA

业务指标验收

Step Time、Token/s、TTFT、TPOT、成本

Q: 为什么 GPU 适合深度学习？

深度学习的主要计算是规则的大规模矩阵乘和向量操作，具有高数据并行度。GPU 用大量 SM、Warp 和 Tensor Core 提供吞吐，并用 HBM 提供高带宽；代价是复杂控制流和强串行任务不占优势。工程上还必须通过批处理、合适的矩阵形状和内存复用才能真正利用这些硬件。

Q: CUDA Kernel 从 CPU 到 GPU 是怎么执行的？

Host 代码通过 CUDA Runtime/Driver 发起异步 Kernel Launch，指定 Grid 和 Block。GPU 将 Block 分配到满足资源约束的 SM，SM 再以 Warp 为单位选择就绪指令执行。Kernel 完成不代表 Host 已同步，只有遇到显式同步、相关数据依赖或阻塞 API 时 CPU 才需要等待。

Q: MIG、MPS、Time-Slicing 和 HAMi 的本质区别？

MIG 在支持的 GPU 上做硬件级实例切分，隔离最强但形状固定、重配有成本；MPS 让多个 CUDA 进程共享执行上下文并并发提交，利用率高但故障和性能隔离弱；Time-Slicing 主要轮转上下文并扩大 Kubernetes 可调度份额，不提供显存或固定算力隔离；HAMi 在 Kubernetes 层增加显存/算力感知的调度与容器侧限制，工程灵活但需要额外组件和验证。

通用诊断路径

界定现象

单卡/单机/多机，持续/偶发，从何时开始

确认业务影响

Step Time、Token/s、TTFT、TPOT、失败率

检查设备健康

Xid、ECC、温度、功耗、时钟、掉卡

检查供给链

CPU、DataLoader、磁盘、网络、H2D

分析 Timeline

空洞、Memcpy、同步、NCCL、Kernel 碎片

深挖热点 Kernel

SM、Tensor Core、HBM、Occupancy、Warp Stall

修复并复测

同一输入、同一并发、同一指标窗口

典型场景

现象	优先检查	常见根因
GPU-Util 长期很低	CPU、DataLoader、队列、H2D、请求量	输入供给不足、Batch 小、同步点多
GPU-Util 高但吞吐低	SM/Tensor/HBM、Kernel Timeline	Memory-bound、小 Kernel、未使用 Tensor Core
显存占用高但 GPU 空转	进程状态、队列、模型阶段	权重/KV/缓存常驻，不代表正在计算
CUDA OOM	Active/Reserved、峰值、碎片、Batch	模型真需求超限或缓存分配器碎片
Pod 申请 GPU 一直 Pending	Node Allocatable、taint、affinity、配额	Device Plugin 未注册、资源不足、约束冲突
Pod 内 `nvidia-smi` 失败	Allocate、CDI/runtime、驱动库、device node	Toolkit/CDI 配置或 Driver 注入失败
多卡训练 Hang	rank 状态、NCCL 日志、网络、拓扑	某 Rank 异常、collective 不一致、链路故障
性能突然下降	功耗/时钟/温度、Xid、后台进程	降频、ECC、共享干扰、拓扑变化

常用命令与意义

nvidia-smi -L
nvidia-smi topo -m
nvidia-smi dmon -s pucvmet
nvidia-smi --query-compute-apps=pid,gpu_uuid,used_memory --format=csv
dcgmi diag -r 1

nvidia-smi -L：确认设备和 UUID，不说明 CUDA 工作负载正常。
topo -m：检查 GPU-GPU、GPU-NIC、NUMA 路径，不直接等于实测带宽。
dmon：看一段时间内的利用率、功耗、时钟、温度和错误趋势。
DCGM 诊断：验证健康状态；生产执行更重级别测试前要确认不会影响业务。

Q: GPU 利用率从 90% 掉到 20%，如何排查？

先比较同一业务阶段的 Step Time 或 Token/s，确认不是 Eval、Checkpoint 或流量变化。再判断单卡还是所有 Rank 同时下降：单卡优先查设备健康、进程和局部拓扑；全局下降优先查数据、存储、网络和同步。随后用 Timeline 找空洞或 NCCL/Memcpy 占比，对热点 Kernel 再看 HBM、SM、Tensor Core 与 Warp Stall，最后复现并验证修复。

Q: nvidia-smi 能看到卡，但 Kubernetes Node 没有 nvidia.com/gpu，问题在哪？

Host Driver 已经基本正常，但 Kubernetes 资源注册链未打通。检查 Device Plugin DaemonSet 是否落到该节点、是否能访问设备、是否在 /var/lib/kubelet/device-plugins/ 注册、ListAndWatch 是否上报健康设备，以及 kubelet 是否完成 Node Status 更新。此时 Container Toolkit 仍可能有问题，但它主要影响容器注入，不是 Node 资源完全缺失的第一嫌疑。

Q: 如何区分显存容量瓶颈和显存带宽瓶颈？

容量瓶颈表现为 OOM、可容纳 Batch/KV Cache 受限，核心单位是 GB；带宽瓶颈表现为 Kernel 持续搬数据、算力管线吃不满，核心单位是 GB/s。前者看 Active/Reserved/Peak Memory，后者看 HBM Throughput、算术强度和 Warp Stall。

问题定义

设计一个多租户 GPU 平台，不是简单把 nvidia.com/gpu 加到 Pod limits。系统需要同时解决设备接入、资源抽象、调度放置、共享隔离、作业生命周期、可观测性、故障恢复与成本治理。

目标与约束

目标	需要回答的问题
利用率	整卡、大卡小任务、碎片和空闲资源怎么处理？
性能	如何避免共享干扰，如何保证训练/推理 SLO？
公平性	队列、租户、项目之间如何配额、借用和回收？
可用性	掉卡、Xid、节点故障、作业 Hang 怎么恢复？
可运营性	如何计量 GPU·hour、显存·hour、排队时间和成本？

核心架构

节点资源层

Driver、Container Toolkit/CDI、Device Plugin/DRA、DCGM

资源目录层

型号、显存、MIG Profile、拓扑、健康与租户标签

准入与队列层

默认值、配额校验、优先级、Kueue/Volcano

调度层

Gang、拓扑、Bin Packing、共享干扰、抢占与 Backfill

运行层

Job/Operator、Checkpoint、重试、弹性伸缩

观测与治理层

SLO、GPU 健康、利用率、成本、审计与容量规划

关键权衡

资源抽象

整卡：路径最简单、隔离清楚，适合大训练和稳定推理。
MIG：强隔离，适合规格固定、SLO 明确的小推理；需要管理 Profile 碎片与重配。
MPS/Time-Slicing：提高轻负载利用率，但必须接受共享干扰和故障域扩大。
HAMi 等方案：可表达显存/算力份额，但要验证限制精度、兼容性、升级和故障恢复。

调度策略

先满足硬约束：型号、显存、健康、节点池、拓扑和 Gang；再对可行节点打分：碎片、NVLink/NIC 距离、干扰、能耗和数据位置。训练任务通常重视 Gang、拓扑与 Checkpoint；在线推理更重视 SLO、容量冗余和快速扩容。

配额与抢占

使用“保障额度 + 弹性借用 + 有界回收”。抢占前要计算收益、重启成本和 Checkpoint 新鲜度，避免为了释放一张卡导致多机训练整体重启。

故障与扩展

Device Plugin/Driver 故障：节点隔离，停止新任务进入，保留诊断现场。
Xid/ECC/掉卡：标记设备或节点不健康，驱逐策略要区分训练和在线推理。
训练节点故障：依赖 Job Controller、Checkpoint 和 Rank 重建。
调度器故障：队列状态和资源账本应可重建，绑定操作保持幂等。
指标异常：监控业务 SLO 与硬件指标，避免仅根据 GPU-Util 自动扩缩容。

Q: 如何设计 GPU 资源调度的 Score 函数？

先把不满足型号、显存、健康、拓扑硬约束的节点在 Filter 阶段排除；Score 再组合碎片、拓扑距离、共享干扰、数据局部性和能耗。各项必须归一化并配权重，同时记录每次打分理由，便于解释为什么任务落在某个节点。训练和推理应使用不同 Profile，不能用一套权重解决所有工作负载。

Q: 为什么只使用 Kubernetes 原生整数 GPU 不够？

整数 Extended Resource 能完成基本设备数量调度，但无法天然表达显存容量、型号属性、互联拓扑、共享干扰和动态配置。平台通常需要标签/亲和性、调度插件、队列系统，以及 Device Plugin/DRA 或 GPU 共享方案共同补足这些语义。

Q: 如何衡量 GPU 平台是否做好了？

至少看排队时间与 JCT、GPU·hour 有效利用、任务成功率、SLO 违约率、碎片率、抢占浪费、故障恢复时间和单位有效 Token/训练样本成本。利用率只是一项过程指标，不能替代用户侧吞吐、延迟与成本。