别把公式、框架和 GPU 分开背：一次算子会把它们全部串起来
这套题训练一条完整解释链：先手算 Shape 与梯度，接着说明 PyTorch 如何记录图和解释同一块 Storage，最后落到 Warp、内存访问和 Kernel。你不必先会写高性能 CUDA，但要能从正确性一路定位到性能。
先看整套知识的坐标，再逐题理解
这不是背答案清单。先用地图确认概念之间的依赖，再沿“白话直觉 → 核心结论 → 原理与取舍 → 常见误区”阅读每个知识点。
先看关键链路，单个知识点才不会变成碎片
面试官追问的方向，往往就是这张图里的下一根箭头：输入怎样变化、状态存在哪里、哪一步最贵、失败怎样被验证。
换一个关键词，或切回“全部章节”。
先能在纸上算对，才谈得上让框架替你算
面试题常从一个小 Shape 或概率公式开始。回答时先定义维度和求和轴，再解释数值稳定与梯度含义。
Broadcasting 的规则是什么？怎样手算输出 Shape，并发现隐藏的大内存开销？#
先用大白话把它讲明白
把两个 Shape 从最右边开始对齐：对应维相等，或其中一个是 1，才能广播；缺少的左侧维按 1 处理。输出维取两者较大值。
例如 [B,T,D] 加 [D]，后者相当于 [1,1,D]，输出仍是 [B,T,D]；[B,T,1] 乘 [B,1,D] 会得到 [B,T,D]。
广播可理解为用 stride=0 的扩展视图重复读取小张量，不必先完整复制输入，但结果仍可能很大。若把 [B,T] 与 [B,T,1] 相乘，B≠T 时通常直接报维度不兼容；恰好 B=T 时，[B,T] 左补为 [1,B,T]，结果反而合法扩成 [B,T,T]，让 bug 随输入 Shape 出现或消失。
Broadcast 从尾维对齐，每一维必须相等、某方为 1 或某方缺失，输出取较大维。它通常不物化扩展输入，但运算结果按广播后的 Shape 产生，所以我会先写左补 1 后的 Shape，再估结果元素数与字节，防止意外二次扩张。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
广播是元素级索引规则，不是自动做矩阵乘。`expand` 可用 stride 0 表示重复读取同一值；`repeat` 则通常真的复制数据。
in-place 运算还要求目标 Tensor 的 Shape 不能因广播改变，因为底层 Storage 没有空间接收扩大后的结果。
compatible if aᵢ=bᵢ or aᵢ=1 or bᵢ=1；outᵢ=max(aᵢ,bᵢ)工程取舍
广播让代码简洁且可避免输入复制，但可能隐藏巨大输出和重复计算；显式 reshape/unsqueeze 更啰嗦，却能把语义和轴写清。
常见误区
“Broadcast 会先把小 Tensor 完整复制成大 Tensor”通常不对。输入多为 View/索引语义，但结果和某些后续算子仍会按大 Shape 分配。
面试官可能继续问
追问：`expand` 与 `repeat` 的区别？为什么 stride=0 的 View 不适合某些 in-place 写入？[2,1,4] 与 [3,4] 输出什么？
Softmax 与 Cross Entropy 怎样连起来？为什么要减最大值而不是直接 exp？#
先用大白话把它讲明白
Softmax 把一组 logits 变成和为 1 的概率：先取指数，再除以所有指数之和。logit 只看相对差值，加同一个常数不会改变概率。
指数增长很快，若直接算 exp(1000) 会溢出。先减本行最大 logit，最大项变成 exp(0)=1，其他项不大于 1，概率完全不变。
分类 Cross Entropy 取正确类别概率的负对数。PyTorch 对类别索引目标的 CrossEntropyLoss 在数学上等价于 LogSoftmax 后接 NLLLoss，并以稳定形式计算；具体后端和 Shape 是否融合为单个 Kernel 不是 API 保证。
对 logits z，softmaxᵢ=exp(zᵢ−m)/Σexp(zⱼ−m)，m 取最大值避免上溢且不改结果。单个硬标签、无类别权重与标签平滑、尚未做 batch reduction 时，Cross Entropy 等于 logsumexp(z)−z_y，logit 梯度才是 p−one_hot(y)；加权、平滑或 mean reduction 会相应改写或缩放梯度。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
`logsumexp` 用 m + log Σ exp(z−m) 稳定计算归一化项；直接先 softmax 再 log 会在极小概率处丢精度。
CrossEntropyLoss 的 input 应是未归一化 logits，不是已经 softmax 的概率；target 可以是类别索引或满足约束的类别概率。`ignore_index` 只适用于类别索引 target；类别维和 reduction 也要说明。
基础硬标签：CE(z,y)=logΣⱼexp(zⱼ)−z_y；∂CE/∂z = softmax(z)−one_hot(y)工程取舍
Label smoothing 可减少过度置信并改善泛化，但会改变概率校准和极少数确定标签的上限；类别加权能处理不平衡，也会改变梯度尺度与决策阈值。
常见误区
“Cross Entropy 越小就说明输出概率一定校准”错误。它奖励真实标签概率，但分布外、类别不平衡和温度仍会导致过度自信。
面试官可能继续问
追问：为什么给 CrossEntropyLoss 传 softmax 后的值不好？temperature 改变什么？ignore_index 怎样影响 mean reduction 的分母？
链式法则在反向传播里到底算什么？为什么梯度要在分支处相加？#
先用大白话把它讲明白
若 y 由 x 经过多步函数得到，x 轻微变化对 y 的影响，要把沿路径每一步的局部变化率相乘。这就是链式法则。
反向传播从最终标量 loss 的梯度 1 开始，把上游梯度乘当前节点的局部导数，传给输入。它计算的是向量—Jacobian 乘积，不会显式存整张巨大 Jacobian。
要分清两次“相加”：一次 backward 内，同一个 x 的多条下游路径按链式法则求和；多次 backward 之间，叶子 Tensor 的 `.grad` 缓冲区会继续累加，直到训练循环显式清空。后者不是计算图里多出了一条路径。
反向传播是反向模式自动微分：从 dL/dL=1 出发，每个节点用上游梯度乘局部 Jacobian，把 VJP 传给输入；同一变量有多条下游路径时按多元链式法则求和。它避免显式构造 Jacobian，适合参数多、标量 loss 的训练。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
对 y=f(x)、L=g(y)，有 dL/dx=(dL/dy)(dy/dx)。向量情况下乘法方向取决于约定，但框架实现的是高效 VJP。
中间激活常被保存给 backward；checkpointing 通过重算 forward 换显存。图内路径求和是微积分规则，跨 backward 的 `.grad` 累加是叶子梯度缓冲区的更新语义，`zero_grad` 则是训练循环策略。
一次 backward 内：dL/dx = Σ_paths (dL/dy_path)·(dy_path/dx)工程取舍
保存更多激活使 backward 快但占显存；activation checkpointing 降显存却增加重算。Forward-mode 更适合输入少输出多，Reverse-mode 更适合参数多标量 loss。
常见误区
“Backward 是把 Forward 的公式倒着运行”不准确。每个算子有自己的梯度规则，并按计算图拓扑传播 VJP；不可逆函数也可能有定义好的局部梯度。
面试官可能继续问
追问：为什么非标量 Tensor 调 backward 要传 gradient？`retain_graph` 做什么？图内路径求和与梯度累积训练有什么区别？
Tensor 不只是数值，图、Storage、Stride 和 dtype 都影响结果
很多 bug 不是数学公式错，而是框架语义没说清：谁是叶子、谁共享 Storage、哪一步复制、哪一步在低精度执行。
PyTorch Autograd 的动态图、leaf、grad_fn、detach 与 no_grad 分别是什么？#
先用大白话把它讲明白
PyTorch 在执行 Forward 运算时同步记录由 Function 节点组成的图。Tensor 的 `grad_fn` 指向产生它的节点；用户创建且 `requires_grad=True` 的参数通常是 leaf，梯度最终累积在 leaf 的 `.grad`。
`detach()` 返回与原 Tensor 共享数据但从当前图断开的 Tensor；`no_grad` 在代码块内不记录普通运算，常用于推理或参数更新。两者都不等于复制数据。
动态图每次执行 Python 控制流都可构造不同图，灵活但也意味着一次 backward 后图通常被释放；想再次反传要重跑 Forward 或谨慎 `retain_graph`。
Autograd 在 Forward 时按实际执行记录动态 DAG。非叶子结果有 grad_fn，叶子参数接收累计 `.grad`。detach 让结果不再沿当前图反传但通常共享 Storage；no_grad 是执行上下文，不记录反向图。一次 backward 后为反传保存的中间量通常释放。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
原地修改被 backward 保存的 Tensor 会触发 version counter 检查，避免用被篡改的值算错梯度。
`model.eval()` 只切换 Dropout/BatchNorm 等模块行为，不会关闭梯度；`inference_mode` 比 `no_grad` 还会关闭 View tracking 与 version counter bump，开销更低但限制更强，而且同样不会自动调用 `eval()`。
parameter update: θ ← θ − η·θ.grad；then clear or set_to_none before next accumulation window工程取舍
动态图易调试和表达数据依赖控制流；编译捕获可融合和优化，但要处理 graph break、guard 与动态 shape。`retain_graph` 便捷却可能长期占住激活。
常见误区
“调用 model.eval() 就不会建图”错误。它不改变 `requires_grad` 或 Autograd 记录，只影响特定层的训练/推理分支。
面试官可能继续问
追问：为什么非叶子 Tensor 默认不保留 `.grad`？detach 后 in-place 修改可能怎样影响原 Tensor？`inference_mode` 创建的 Tensor 为什么更受限？
View、Reshape、Stride、Contiguous 是什么关系？Transpose 后为什么 view 可能失败？#
先用大白话把它讲明白
Tensor 可以看成 Storage 加一组解释方式：shape 告诉有几维，stride 告诉每个索引在 Storage 里要跳多少元素，storage_offset 告诉从哪里开始。
`transpose` 只交换 shape/stride，通常不搬数据，因此结果可能不连续。`view` 要求新 Shape 能用现有 stride 解释同一块数据；不满足时会报错。
`reshape` 会尽量返回 View，不行时复制成可用布局；`contiguous()` 在已连续时返回自身，否则分配并复制。这个隐式复制可能成为大模型里看不见的带宽开销。
View 按定义共享 Storage，只改 shape/stride/offset；transpose 常产生 non-contiguous layout。view 只有在 stride 满足可折叠条件时零复制成功，reshape 可能 View 也可能复制，contiguous 在需要时把数据重排成连续布局。性能分析要检查是否出现隐式 copy。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
元素地址可由 offset 加各维 index×stride 求出。连续行优先布局的 stride 通常从右往左累乘 size。
高级索引通常返回副本，基础切片多为 View；但写入语义还需看具体操作。共享 Storage 应比较底层 Storage 指针并同时看 storage_offset；只比较 Tensor 首元素的 `data_ptr()`，可能把带偏移的 View 误判为不共享。
storage_index = storage_offset + Σᵢ indexᵢ·strideᵢ工程取舍
保留 View 零复制且省显存，但不规则 stride 可能让后续 Kernel 访存不合并；先 contiguous 多一次复制，却可能让重计算路径更快。需用 profiler 判断摊销。
常见误区
“reshape 永远不复制”错误。官方语义允许它在布局不兼容时创建副本，因此代码不应依赖返回值一定共享 Storage。
面试官可能继续问
追问：为什么不能只用 Tensor 首元素的 data_ptr 判断是否共享 Storage？expand 的 stride 为什么可为 0？什么时候 contiguous 的一次复制值得？
FP16、BF16、FP32 与 AMP 怎样取舍？GradScaler 为什么主要和 FP16 搭配？#
先用大白话把它讲明白
浮点格式像“可表示范围”和“刻度精细度”的取舍。FP32 范围与精度都更高但占 4 字节；FP16 占 2 字节、尾数较细但范围窄；BF16 也占 2 字节，指数范围接近 FP32，尾数更粗。
AMP 的 autocast 按算子选择低精度或保持高精度：大矩阵乘适合低精度，某些归一化、归约和损失更需要稳定性。它不是把全模型粗暴 `.half()`，也不替你决定参数或优化器状态用什么 dtype。
FP16 的小梯度可能下溢为 0，GradScaler 先放大 loss，让 backward 中的梯度落入可表示范围，更新前再缩回并检查 inf/NaN。BF16 指数范围大，通常不靠 scaling 解决下溢，但仍需验证收敛。
FP16 尾数比 BF16 多，但指数范围小；BF16 的指数范围接近 FP32，尾数更短。autocast 只负责按算子选择计算 dtype，参数、梯度与 optimizer state 的 dtype 是另一组策略。FP16 训练常用 GradScaler 抵抗小梯度下溢；发现 inf/NaN 时跳过更新并调小 scale。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
低精度是否加速取决于硬件 Tensor Core、Shape 对齐、Kernel 和带宽；小算子或频繁 cast 可能没有收益。
GradScaler 不承诺 scale 始终大于 1。若 BF16 预训练模型改用 FP16 后产生 overflow，scale 可能持续下降到 1 以下；这时要先检查 dtype 是否与模型数值范围兼容。
scaled_loss = s·loss；unscaled_grad = scaled_grad / s工程取舍
BF16 因指数范围更大而较少 overflow/underflow，但尾数粗可能影响小更新；FP16 在部分硬件吞吐好，却更易超出范围。FP32 稳定但带宽与容量成本高。
常见误区
“AMP 保证结果和 FP32 一样”错误。它是性能—精度策略，仍要用 loss、梯度、收敛和最终任务指标验证。
面试官可能继续问
追问：为什么 optimizer state 常保留 FP32？TF32 属于什么？遇到 GradScaler 持续下降先检查什么？
GPU 快在并行和数据复用，不是单个线程更聪明
先画 Grid/Block/Warp，再画一条代表性的数据访问与复用路径。性能判断要同时看并行度、访存模式和算术强度。
Grid、Block、Thread、Warp、SM 分别是什么？分支发散为什么只在 Warp 内发生？#
先用大白话把它讲明白
启动一个 CUDA Kernel 时，会创建 Grid；Grid 由多个 Block 组成，Block 由线程组成。CUDA 编程模型保证同一 Block 的线程在同一个 SM 上执行，因此才能共享该 SM 上的 Shared Memory 并做块级同步；不要把它扩大成所有硬件时刻都绝不暂停，Dynamic Parallelism 等少见场景可把 Block 状态暂存到内存。
硬件把同一 Block 中连续的 32 个线程组成 Warp。CUDA 的 SIMT（Single Instruction, Multiple Threads，单指令多线程）模型让 Warp 中各 lane 执行同一 Kernel 代码、操作各自数据。
若一个 Warp 内部分线程走 if、另一些走 else，执行某条路径时只有匹配的 lane 活跃，满 Warp 利用率下降。Volta 及以后支持 Independent Thread Scheduling，可在 sub-warp 粒度分歧与重汇合，但并没有让分支发散变成免费。不同 Warp 本就独立调度，不属于同一个 Warp 的发散。
Grid 是一次 Kernel 的全部线程，分为 Block；同一 Block 的线程在同一 SM 上执行并可用 shared memory/块级同步；硬件把线程按 32 个组成 Warp，以 SIMT 执行。Warp 内数据依赖分支会让不同路径只激活部分 lane，常需分路径推进并降低利用率；跨 Warp 走不同分支不叫同一组 divergence。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
Block 大小通常取 32 的倍数以避免最后一个 Warp 大量空 lane，但最优值还受寄存器、shared memory 和占用率限制。
线程通过 blockIdx、blockDim、threadIdx 计算全局索引；越界 guard 必须正确，尾部 Warp 有少量不活跃 lane 很常见。
global_idx = blockIdx.x·blockDim.x + threadIdx.x；warps_per_block=ceil(threads/32)工程取舍
大 Block 可摊调度并提供更多协作线程，却消耗更多每块资源、降低同时驻留 Block；减少分支有利于利用率，但复杂重排也可能增加指令。
常见误区
“GPU 上所有线程严格同时运行”错误。线程分批以 Warp 执行，Block/Warp 由 SM 调度，驻留和进度受资源与 stall 影响。
面试官可能继续问
追问：为什么普通 Block 间不能直接用 `__syncthreads`？Warp divergence 与 memory divergence 有何区别？Block 不是 32 倍数会怎样？
HBM、L2、Shared Memory、Register 怎样分工？Coalescing 和 Tiling 为什么重要？#
先用大白话把它讲明白
Global Memory 是 CUDA 的地址空间，通常由设备 DRAM 承载：数据中心 GPU 常用 HBM，消费级 GPU 也可能用 GDDR，二者不能画等号。L2 为全 GPU 共享，每个 SM 还有 L1/统一数据缓存、Shared Memory 和 Register File。
Coalescing 的直觉是让一个 Warp 的相邻线程访问相邻地址，硬件可用较少内存事务取回整片数据。若地址散乱，会浪费传输；但实际事务数还受访问宽度、对齐与缓存命中影响。
经典 tiled GEMM 常把矩阵块从 Global Memory 搬到 Shared Memory 供 Block 复用，再把操作数放进 Register 计算。它是一条代表性优化路径，不是每个 Kernel 必经的固定楼梯。
不存在所有 Kernel 都固定经过的 HBM→L2→Shared→Register 流水线。普通 Global load 可由 L1/L2 缓存后进入 Register；需要块内复用时才显式用 Shared。合并访问减少内存事务，Tiling 提高目标层级的算术强度。Ampere 的异步 copy、Hopper 及以后的 TMA 还能把 Global→Shared 搬运与计算重叠，并避免普通寄存器中转。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
Shared Memory 生命周期与 Block 一致，线程协作装载后通常需同步再消费；Register 由编译器分配，溢出可能落到 local memory，而 local address space 的数据可驻留在设备内存并由缓存服务。
矩阵乘的高性能来自块内复用：同一 A/B tile 被多个乘加使用；但是否使用 Shared、采用何种缓存策略和异步搬运，要看架构、Kernel 与后端生成代码。
arithmetic_intensity(level) = operations / bytes_moved_at_that_memory_level工程取舍
更大 tile 提高复用，但占更多 Shared/Register，可能降低驻留 Warp；复杂布局能消除 bank conflict，却增加索引、同步和流水化难度。
常见误区
“Local Memory 是片上、所以很快”错误。CUDA 的 local address space 表示线程私有，不等于物理片上；访问可能落到设备内存并由 L1/L2 缓存。
面试官可能继续问
追问：Shared bank conflict 是什么？为什么 transpose 常伤 coalescing？`cp.async`/TMA 为什么能避免普通寄存器中转？
Occupancy 越高越好吗？怎样用 Roofline 判断 Memory-bound 还是 Compute-bound？#
先用大白话把它讲明白
Occupancy 是一个 SM 上活跃 Warp 数占硬件最大值的比例。更多 Warp 能在某个 Warp 等内存时切换别的 Warp，帮助隐藏延迟。
但 100% Occupancy 不保证快：如果每个 Warp 都做低效访问或同样被带宽卡住，再多 Warp 也不会创造额外带宽。某些高性能 Kernel 用更多 Register 降低 Occupancy，却减少访存并更快。
Roofline 把算术强度放横轴、性能放纵轴。低强度区域的上限约是内存带宽×强度；高强度区域才接近峰值计算。它给的是上限与优化方向，不单独证明具体 stall 根因，还要结合 profiler。
Occupancy 是隐藏延迟的手段，不是最终目标。Roofline 上可达性能受 min(峰值算力, 某层级带宽×该层级算术强度) 约束；斜线区的上限偏内存，平台区的上限偏计算。实际点离对应 roof 多远、stall 和吞吐指标是什么，决定下一步该减少字节、提高复用还是优化指令。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
理论 Occupancy 受每 Block 线程、Register、Shared Memory 和硬件上限共同限制；Achieved Occupancy 还受实际调度和工作不均。
层级 Roofline 可分别看 device memory、L2、L1/Shared ceiling；“memory-bound”也要说清是哪一级。
attainable_performance(level) ≤ min(peak_compute, bandwidth(level)·arithmetic_intensity(level))工程取舍
压低寄存器用量可提高 Occupancy，却可能造成 spill 和更多指令；扩大 tile 提高强度，却可能牺牲并行驻留。需要扫参数并比较 profiler baseline。
常见误区
“Occupancy 低就是性能差的根因”错误。它是症状/约束之一；低占用若仍能覆盖延迟且算力满载，强行提高反而更慢。
面试官可能继续问
追问：Ridge point 是什么？为什么 L2 Roofline 与 device-memory Roofline 结论可能不同？如何识别 Register spill？
最后用可复现的最小实验，把正确性和性能问题分开
调优不从“换一个神奇 API”开始。先缩小问题，固定输入与版本，测基线，再一次只改一个变量。
torch.compile 在做什么？Graph Break、Guard 和 Recompile 为什么会让它反而更慢？#
先用大白话把它讲明白
`torch.compile` 尝试捕获 PyTorch 运算图，把多个算子交给后端优化和融合，减少 Python/Kernel launch 开销并生成更贴合输入条件的代码。
默认 `fullgraph=False` 时，遇到无法捕获的 Python 时会先编译已捕获的 FX 区域，回到 Python 执行不支持部分，再从后续可捕获区域继续；Graph Break 多了会丢跨区域融合并增加切换。
每个编译版本带 Guard，例如 dtype、device、Shape 或 Python 值条件。运行时会先找 Guard 匹配的缓存版本；若没有匹配项才可能 Recompile。默认 `dynamic=None` 会在发现 Shape 变化后尝试更动态的版本，而不是简单地为每个 Shape 永久编一个 Kernel。
torch.compile 用 Dynamo 捕获 FX 图，后端如 Inductor 做融合和代码生成。Graph Break 把函数切成编译区与 Python 区；Guard 决定缓存版本能否复用，无版本匹配时才触发 Recompile。PyTorch 2.13 的默认 recompile limit 为 8；默认 `fullgraph=False` 达到上限后回退 Eager，而 `fullgraph=True` 会抛出限制命中错误。收益要同时扣除首次编译、有限次重编译和断图开销。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
`fullgraph=True` 要求整个函数捕获成单图，遇到 Graph Break 会显式报错；`TORCH_LOGS=graph_breaks,guards,recompiles,dynamic` 或 tlparse 可查看原因。
编译缓存按 code object 管理；默认模式达到重编译上限后会回退 Eager，但此前已编译且 Guard 匹配的版本仍可能复用。具体限制可由 `torch.compile(..., recompile_limit=...)` 或全局配置调整，不应靠无限抬高上限掩盖 Shape 抖动。
net_speedup if N_runs·(T_eager−T_compiled) > T_compile + ΣT_recompile + T_region_switch工程取舍
静态/有限动态 Shape 更容易深度优化；动态 Kernel 提高复用，却可能限制优化。局部跳过难编译代码可保住主要热区，但会增加区域边界。
常见误区
“Graph Break 会算错”通常不对。默认模式会在编译区之间执行 Python/Eager 以保持语义，主要代价是丢优化与切换；真正的正确性问题仍需单独测试。
面试官可能继续问
追问：为什么 `.item()` 常触发断图？怎样区分第一次慢和稳态慢？Guard miss 为什么不一定每次都重新编译？
训练遇到 NaN、OOM 或突然变慢，应该按什么顺序定位？#
先用大白话把它讲明白
先让问题可复现：固定代码、数据 batch、seed、硬件和版本，保存第一个异常 step。把“NaN”“OOM”“慢”分开，因为修复路径不同。
NaN 从 loss 和梯度逐层检查 finite、输入范围、学习率、归一化与 AMP scale；训练 OOM 区分参数、梯度、优化器状态、保存的激活、临时 workspace、引用泄漏与 allocator 状态，不把推理阶段的 KV Cache 混进来；慢先分 CPU、数据、通信还是 GPU Kernel。
再做二分与消融：关 AMP、缩 batch/seq、禁编译、单卡、替换模块；使用 anomaly detection、memory snapshot、torch.profiler/Nsight，而不是一次改十个开关。
我先固定最小复现并定位第一个异常 step。NaN 查输入→forward activation→loss→gradient，比较 FP32/AMP；OOM 估静态与峰值内存、查意外广播和引用泄漏；慢用 timeline 分 CPU/data/communication/kernel，再用 profiler 看热点。每次只改一个变量。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
CUDA 默认异步，CPU 墙钟计时若不 synchronize 或不用 CUDA Event 会低估 Kernel；异常也可能在后续同步点才暴露，因此定位阶段要在正确边界同步。
PyTorch memory snapshot 能看到其 CUDA allocator 管理的分配与 OOM 事件，但看不到 NCCL 等直接通过 CUDA API 分配的显存；必要时应把 allocator 统计与设备总使用量对照。
training_memory_peak ≈ parameters + gradients + optimizer_states + saved_activations + temporaries + allocator_overhead工程取舍
更强的 anomaly/check/trace 会明显变慢，只用于小复现；生产训练保留轻量 finite、grad norm、memory 与 step time 监控，异常时提升采样。
常见误区
“CUDA OOM 就清缓存 `empty_cache()`”不是根治。它不会释放仍被 Tensor 引用的内存，也不能改变活跃工作集的理论需求。
面试官可能继续问
追问：怎样正确计时异步 CUDA？为何 loss 正常但 gradient 先 NaN？memory snapshot 为什么可能看不到 NCCL 占用？
编码题：设计一个动态批处理等待队列，怎样保证 deadline、取消、保序与复杂度？#
先用大白话把它讲明白
每个请求带 request_id、到达序号、相对 deadline 和 token/KV 预算。FIFO 容易让长请求挡住快过期请求；纯 EDF（Earliest Deadline First，最早截止期优先）又可能让长 deadline 请求饥饿，而且在到达负载超过容量时，任何队列都不能保证所有 deadline，必须配合准入、拒绝或降级。
可用 dict 做 request_id→state 的期望 O(1) 查找与幂等取消；最小堆 entry 存 `(effective_deadline, arrival_seq, version, request_id)`，相同 deadline 时用序号稳定排序。公平性必须落成明确规则，例如等待越久 effective deadline 越早的 aging，或按租户保留配额；仅仅旁边再放一个 deque 并不会自动公平。
取消或改 deadline 采用 lazy deletion，弹出时用 version/state 跳过旧 entry。每轮在 token/KV budget 内选请求；保序要明确是“单请求 token 顺序”“同优先级稳定顺序”还是“全局响应顺序”，最后一种需要额外 reorder buffer。
我用 request_id→state 哈希表管理取消/完成，用 `(effective_deadline, arrival_seq, version, id)` 最小堆做 EDF+aging，并用租户配额防独占。取消只改 state/version；pop 时跳过旧 entry。每轮在 token/KV budget 内选候选，过期就拒绝并记录原因。系统过载时靠 admission control 守可行 deadline，队列本身不能创造容量。
展开原理与工程取舍FORMULA · TRADE-OFF · FOLLOW-UP
原理拆解
lazy deletion 下，若一次有效 pop 前跳过 k 个 stale entry，本次代价是 O((k+1)log h)，不是固定 O(log n)；但每个 stale entry 只会被弹一次，所以按插入 entry 摊销仍为 O(log h)。
堆内存是 O(n+s)，n 为 live 请求、s 为尚未清理的 stale entry。可在 `heap_size > c·live_size` 时用 live state 重建并 `heapify`，一次 O(n)，把空间恢复到 O(n)。所有相对 deadline 用 monotonic clock，避免墙钟回拨。
enqueue O(log h)；lookup/cancel mark expected O(1)；pop after k stale O((k+1)log h)；memory O(n+s), compact → O(n)工程取舍
EDF 降低紧急请求 miss，却可能饿死长 deadline；aging、租户配额或多级队列改善公平，但可能牺牲最早 deadline。组 batch 若检查 m 个候选，调度开销还要计入 O(m log h) 量级，策略必须与产品 SLO 对齐。
常见误区
“有 deadline heap 再加一个 deque 就同时保证 deadline 与公平”错误。数据结构只提供操作，真正的顺序、aging、配额、过载拒绝和输出契约必须由调度策略定义。
面试官可能继续问
追问：怎样选堆压缩阈值并证明摊销复杂度？多租户怎样防止一个租户占满？如何写 property-based test 验证不重复交付？