本文是AI Infra 面试题系列的 Level 5，覆盖 Q55-Q64。重点是Profiling、算术强度、FlashAttention、Triton、occupancy 与图编译。
上一篇：推理引擎与在线服务 · 系列总索引 · 下一篇：集群、可靠性与 RL Infra
6. Level 5：GPU Kernel、编译与性能工程
Q55. 接到“训练吞吐下降 20%”的任务，你的 profiling 流程是什么？
高分回答要点
- 先确认回归真实存在：同硬件、同输入、同 commit/镜像、同 warmup，查看分布而非单点，并定位首次坏版本。
- 自上而下分解 step：data、H2D、forward、backward、optimizer、collective、checkpoint；同时查看 CPU launch、GPU timeline 和各 rank skew。
- 找 critical path 上的最大变化，再进入 kernel：shape、dtype、调用次数、实现、FLOPs、bytes、occupancy、Tensor Core 和 memory throughput。
- 用最小复现与单变量 A/B 验证因果；优化后同时检查 correctness、loss、显存、尾部和其他 shape，防止局部 benchmark 变快而端到端变慢。
- 最后固化 benchmark、trace marker 和回归阈值，避免同类问题重复发生。
评分上限提示：直接打开 profiler 找“最慢 kernel”，但不先确认 workload 和 critical path，最多 2 分。
Q56. 怎样为一个 kernel 估算 arithmetic intensity？
高分回答要点
- 明确计算次数和从目标存储层实际读取/写入的 bytes；不能只按 tensor 理论大小，cache reuse、写回和中间量都要考虑。
- 例如大 GEMM 的每个元素可被 tile 多次复用，intensity 高；elementwise op 通常每元素做少量 FLOPs、至少一读一写，偏带宽受限。
- decode 中权重若每 token/batch 只复用有限次，权重 bytes 主导；batch 增大提高复用和 intensity。
- 将估算点放到对应 GPU 的 HBM/compute roofline，再用 profiler counters 检查实际 DRAM bytes 与 FLOPs。
追问：为什么把多个 elementwise op 融合后 FLOPs 几乎不变，却可能明显加速？
Q57. FlashAttention 的系统原理是什么？
高分回答要点
- 标准实现若 materialize
S x Sscore/probability，会产生大量 HBM 读写；FlashAttention 将 Q/K/V 分块放入更快的片上存储。 - 用 online softmax 维护每行最大值与归一化和，逐块得到与标准 attention 等价的结果，无需保存完整 attention matrix。
- backward 可重算部分统计/score，以更多计算换 HBM IO 和 activation 显存。
- 性能受 head dim、sequence、causal/varlen、dtype、GPU 架构、shared memory/register 和 kernel 实现影响；不是所有 shape 都同样快。
实验锚点：V100 的 SM70 不满足本项目 vLLM FlashAttention-2 backend 要求，运行时回退到兼容 backend；这属于支持矩阵问题，不代表 attention 原理不可运行。
Q58. kernel fusion 的收益和风险是什么？
高分回答要点
- 收益是减少 kernel launch、避免中间 tensor 写回 HBM、提升 producer-consumer locality，例如 bias+activation、RMSNorm、optimizer update。
- 风险是融合后 register/shared-memory 压力上升、occupancy 降低、代码生成膨胀、动态 shape specialization 过多，甚至失去更优库 kernel。
- reduction 与 matmul 融合要考虑跨 block 同步限制；不是把所有算子放进一个 kernel 就最好。
- 用端到端 timeline、HBM bytes、launch 数、occupancy 和多 shape benchmark 判断，不只比较单个融合 kernel。
Q59. 让你用 Triton 写 RMSNorm，你会怎样设计和验证？
高分回答要点
- 每行计算平方和、
rsqrt(mean(x^2)+eps)，再乘输入和权重；先明确 shape、dtype、是否保存统计量和 backward 需求。 - 一行或分块映射到 program，使用向量化 load、mask 非 2 次幂长度，在 FP32 中累积 reduction，再 cast 输出。
- hidden 很大时单 program 的 register 压力可能过高，需要分段 reduction 或两阶段 kernel；权重读取可利用 cache。
- correctness 要覆盖非对齐 shape、极值、不同 dtype，与高精度 reference 比误差；性能扫描 block size/warp 数，并与 PyTorch/已有 fused kernel 在多 shape 比较。
- backward 不能只依赖数值看起来接近；应做 gradient check 和训练片段验证。
Q60. occupancy 越高，kernel 一定越快吗？
高分回答要点
- occupancy 是活跃 warps 相对上限的比例，受 registers/thread、shared memory/block、threads/block 和硬件限制。
- 足够 warps 可隐藏 memory/指令 latency，但超过所需后更高 occupancy 不再提升；降低寄存器以追求 occupancy 可能引发 spill 到 local memory。
- compute-bound kernel 可能靠更高 ILP 和每线程 tile 在较低 occupancy 下更快；Tensor Core pipeline 也不能只用 occupancy 判断。
- 应结合 eligible warps、stall reasons、register spill、memory throughput 和 achieved FLOPs 判断。
Q61. CUDA stream/event 如何实现计算通信 overlap？
高分回答要点
- 不同 stream 的工作只有在数据依赖、硬件资源和 runtime 允许时才可能并发；默认 stream/隐式同步会破坏 overlap。
- producer compute 完成后记录 event，communication stream wait event；consumer 再等待 collective 完成，形成最小依赖而非全设备同步。
- compute 与 NCCL 同时争用 SM、HBM、copy engine 或互连，因此时间区间重叠不等于两者都保持独立峰值。
- tensor 生命周期必须用 stream-aware allocator/
record_stream等机制保护，否则内存可能被过早复用。
追问：为什么插入 cudaDeviceSynchronize() 便于计时，却会改变你要测的 overlap？
Q62. PyTorch 2 图编译链路中，graph break 和 dynamic shape 怎样影响性能？
高分回答要点
- 大体链路包括 Python graph capture、autograd graph/分解、后端代码生成；候选人应说明自己讨论的是哪个层级，而非把
torch.compile当单一黑盒。 - Python side effect、数据依赖控制流、不支持算子和某些 mutation 会 graph break，导致频繁往返 eager、失去跨算子优化。
- dynamic shape 可通过 guards/symbolic shape 支持，但 shape 分布太散会 recompilation、缓存膨胀；服务端常做 bucket/padding 稳定 shape。
- 排查要记录 compile time、graph 数、break reason、recompile/guard failure、生成 kernel、warm latency 和内存。
Q63. NCCL 在多机上如何选择路径？遇到带宽异常看什么？
高分回答要点
- 数据路径可能是 GPU-NVLink/PCIe-NIC，经 GPUDirect RDMA 走 InfiniBand/RoCE；若条件不满足可能经 host staging，性能和 CPU 占用显著变化。
- 检查 GPU/NIC/NUMA 拓扑、驱动/CUDA/NCCL、HCA 与 GID、MTU、RDMA 能力、交换网络、链路速率、rail 绑定和容器设备映射。
- 从小到大分层测试：单 GPU kernel/内存、节点内 pair/all-rank、节点间 point-to-point、collective；用同消息尺寸和并发数复现训练模式。
- 查看 NCCL 拓扑/algorithm/protocol 日志、NIC counters、丢包/ECN/PFC、PCIe 降速和 rank skew；环境变量只应用于验证假设，不是永久“调参玄学”。
当前证据边界：本项目只有单机 PCIe/NCCL 数据。面试时应明确提出上述多机实验计划，而非假装做过 RDMA。
Q64. 怎样建设性能回归测试，而不让结果充满噪声？
高分回答要点
- 分层建立 kernel microbench、单机模型 step、分布式 scaling 和在线 trace replay；每层都有 correctness gate。
- 固定硬件池、时钟/功耗、软件镜像和 workload，做 warmup、多次重复、稳健统计与环境健康检查。
- 阈值同时考虑相对回归、绝对影响和方差；关键指标包括吞吐、p99、显存、compile/startup、错误与数值质量。
- 自动保存 commit、配置、trace 摘要和原始样本；回归发生时可二分版本，并区分代码回归、基础设施噪声和 workload 漂移。
- 不应只守一个平均 tokens/s；例如优化吞吐但让 TTFT p99 或 OOM 率恶化，仍应阻止发布。