37 道题，考的是
AI Infra
一条完整系统链路

结论

Roofline 能算出的不是一个放之四海而皆准的“最佳 shape”，而是：给定 GEMM 的 K、N、数据类型以及硬件有效算力/带宽后，使算子从 memory-bound 跨到 compute-bound 所需的最小 M。实际最优值还会受 Tensor Core tile、grouped GEMM 调度、L2 命中、量化 scale、kernel launch 和专家负载不均影响，必须用实测校准。

对 C[M,N] = A[M,K] × B[K,N]：

假设 beta=0，A、B 各从 HBM 读一次，C 写一次，则最理想的数据搬运量为：

算术强度：

硬件 Roofline 的转折点（ridge point）：

令 AI=R，解出临界 M：

前提是分母为正；如果分母不为正，说明即使 M→∞，这个 shape 在该数据搬运模型下也达不到计算屋顶。若 A/B/C 都是 BF16，即每元素 2 Byte，则：

当 K、N >> R 时，权重 B 的读取占主导，AI≈M，因此可用非常好记的近似：

H100 + DeepSeek 专家 GEMM 数值例子

假设使用 H100 SXM，不使用结构化稀疏的 dense BF16 峰值约为 989.5 TFLOP/s，HBM 带宽 3.35 TB/s。NVIDIA 页面列出的 1,979 TFLOP/s 带有稀疏加速脚注，所以普通 dense GEMM 应取约一半：

DeepSeek-V3 routed expert 的 hidden size 为 7168，expert intermediate size 为 2048。将 gate/up 两个投影融合后，可近似看成 K=7168,N=4096：

down projection 为 K=2048,N=7168：

所以该理想模型给出的结论是：每个 expert 一次聚到约 300～360 行，才接近 H100 dense BF16 的理论 Roofline 转折点。它不是“线上必须攒到 333 token 才运行”；在线推理通常受延迟约束，会在更小的 M 上运行并处于 memory-bound。

从专家 GEMM 反推 MoE TopK 的 token 数

设：

• E：routed expert 总数；
• k：每个 token 激活的 expert 数；
• T_group：一个 EP group 在本次 MoE 层中汇总的原始 token 数；
• 路由均匀，暂不考虑 capacity padding 和冗余 expert。

总 token-expert assignment 数为 kT_group，每个 expert 的平均 GEMM 行数：

要让平均 expert batch 达到 M*：

DeepSeek-V3 取 E=256、k=8、M*=333：

若 EP=P，每卡进入该 EP group 的本地 token 数相同，则：

例如 EP=32 时，理论上约需每卡 333 个本地 token，才能使平均每个 expert 得到约 333 行。这更像 prefill 大 token batch；decode 的每卡 token batch 往往只有个位数到几十，因此专家 GEMM通常仍是带宽受限。

工程上怎么真正选

1. 用该公式得到候选区间，而不是直接上线。
2. 将理论峰值替换为同精度下实测的有效 Tensor Core 吞吐和 HBM 带宽；通常还要乘利用率系数。
3. 对真实 M 分布跑 grouped GEMM benchmark，至少扫 M={1,2,4,...,512}，同时测 P50/P99，而不只测吞吐。
4. 用真实路由 trace 统计每个 expert 的 token 数分布；用均值会低估长尾。可用 M_p95 或容量系数 c>1 规划 buffer。
5. 将 All-to-All、permute/unpermute 和 padding 一并算入端到端 MoE 时间。单独让 GEMM compute-bound，可能反而因攒 batch 增大排队延迟。

常见追问

• 为什么 AI≈M？ 大模型 decode 的权重矩阵远大于激活，B 的 KN 字节占主导；BF16 下 2MKN/(2KN)≈M。
• FP8 怎么算？ 必须分别写 b_A、b_B、b_C，并加入 scale/反量化流量。不能简单把所有字节都除以 2；输出和累加通常不是 FP8。
• 为什么公式和实测转折点不一致？ 理想模型假设每个元素只过 HBM 一次，且忽略 launch、tile 浪费、寄存器压力、L2 和频率降档。

结论

FlashAttention 是精确 attention，不是稀疏或近似 attention。它没有改变 softmax(QKᵀ)V 的数学结果和 O(N²d) 计算量；核心是 tiling、算子融合和 online softmax，避免把 N×N 的 score/probability 矩阵写回 HBM，从而显著减少 HBM IO 和中间显存。

Flash-Decoding 面向 query length≈1、context 很长、batch 很小 的 decode 场景。普通 FA 只有很少的 Q block，无法占满 GPU；Flash-Decoding 进一步沿 KV 序列维切分并行，最后精确合并各 split 的 softmax 结果。

标准 attention 为什么慢

标准实现通常拆成：

S = QKᵀ        -> 把 N×N 的 S 写 HBM
P = softmax(S) -> 读 S、写 N×N 的 P
O = PV         -> 再读 P

计算量仍是 Θ(N²d)，但 HBM 还要反复读写 Θ(N²) 的中间矩阵。序列很长时，这个 IO 比 Tensor Core 乘法更贵。

FA 的实现：分块 + online softmax

将 Q 切成行块 Q_i，K/V 切成列块 K_j,V_j。对一个 Q tile，只在 SRAM/寄存器中保留当前 score tile，以及每一行的三个状态：最大值 m、归一化分母 l、尚未归一化的输出累积 O。

初始化：

读入一个新的 KV tile 后：

遍历完所有 KV tile 后输出 O/l。当最大值变大时，旧分块的累积量乘 exp(m-m') 即可重新缩放，因此无需保存此前的 scores。这也是不同 KV 分片结果能够精确合并的基础。

实现层面还会：

• 将 QK、mask、scale、softmax、dropout（训练时）和 PV 融在一个 kernel 数据流里；
• 选择适合 shared memory、寄存器和 Tensor Core MMA tile 的块大小；
• backward 不保存整个 attention matrix，而是保存少量归一化统计并重算 score；这是“多做一点 FLOP，少做大量 IO”的典型取舍；
• causal attention 跳过未来 token 对应的 tile；变长 batch 通过 cumulative sequence length/块表处理。

原始论文给出的额外内存是输入输出之外 O(N)，而标准实现需显式保留 O(N²) 中间量。FA 的 HBM 访问复杂度也低于标准 attention。

Flash-Decoding 的三步

1. 把长 KV 序列按 sequence 维切成 s 个 split；切分本身只是 view/索引。
2. 每个 split 独立计算当前 Q 对局部 K_s,V_s 的 partial output，并写出该 split 的 m_s、l_s/O_s 或等价 LSE 状态。
3. 启动一个很小的 reduction kernel，用 log-sum-exp 规则把所有 split 精确归并为最终输出。

若每个 split 返回已归一化 o_s 和 LSE_s，则全局：

什么时候用哪个

• prefill：Q 很长，普通 FA 已有足够多的 Q tiles，主要使用 FlashAttention。
• decode、batch 大且 context 不太长：paged decode kernel 或普通 decode attention 可能更好，split reduction 有额外开销。
• decode、batch 小且 context 很长：Flash-Decoding 通过 KV split 提供更多 thread block，通常更合适。

常见追问

• FA 降低计算复杂度了吗？ dense exact attention 没有，仍为 O(N²d)；主要降低 IO 和中间显存。
• 为什么分块 softmax 仍然精确？ softmax 可以由可结合的 (max, exp-sum, weighted-value-sum) 状态表示并按上述公式合并。
• Flash-Decoding 和 speculative decoding 是一回事吗？ 不是。前者是单步 attention kernel 的并行化；后者是一次验证多个候选 token，减少串行 model forward 次数。

结论

KV Cache 是模型执行 attention 所必需的运行状态：保存当前序列历史 token 在每一层的 K/V（或 MLA latent）。Prefix Cache 是建立在 KV Cache 之上的跨请求复用与保留机制：识别相同 token prefix，让新请求直接引用已经算过的 KV blocks，从而跳过这段 prefix 的 prefill。

所以 Prefix Cache 不是另一种张量，也不能代替 KV Cache。它主要节省重复 prefix 的 prefill 计算；第一次请求、没有共享前缀、或只关心 decode 吞吐时，收益很小。

区别

常规 MHA/GQA 的 KV 容量可估算为：

其中最前面的 2 表示 K 和 V。MLA 则缓存低维 latent 与 RoPE key，容量近似：

哈希块式实现

以 vLLM 的 Automatic Prefix Caching 为例：

1. 将 KV cache 划成固定 token 数的物理 blocks，sequence 的 block table 将逻辑块映射到物理块。

2. 只对完整块建立缓存键。块键通常为：

   hash(parent_hash, block_token_ids, extra_hashes)
   
   extra_hashes 可包含：LoRA_ID、multimodal_hash、cache_salt 等

   parent_hash 使当前块的键包含此前全部 prefix 的语义，而不用把整段 token 反复放进 key。这里要区分“哈希字段”和“缓存命名空间”：vLLM 文档列出的块哈希组件是 parent hash、当前块 token 与上述 extra hashes；模型权重 revision、tokenizer、位置编码配置、KV layout/dtype 等兼容性条件仍必须由独立 engine/cache namespace 或等价隔离保证，不能因为它们没有逐项塞进块 hash 就跨模型共享。

3. 新请求从左向右计算块 hash，在全局 hash table 中寻找最长连续命中；命中的物理 blocks 增加引用计数并直接挂到新请求的 block table。

4. 从第一个未命中的 token 开始做增量 prefill，新写入的完整 block 再进入缓存索引。

5. 多个 sequence 只能共享不可变的缓存块。vLLM 只为完整块建立 Prefix Cache，追加或分叉时通常直接为各自的新尾部申请块，不会回写已命中的完整块；只有某个实现额外共享了仍可写的 partial tail，才需要 copy-on-write。

6. block 的活跃引用计数降到 0 后才可驱逐。常见策略为 LRU；更复杂系统会加入频率、长度、重算成本和传输成本。

多租户场景不能只追求命中率。按 vLLM v0.21.0 文档口径，v0.11 起 sha256 是 Prefix Cache 默认哈希算法；若还要求跨 Python/vLLM 版本、跨语言复现同一个键，可选用规范序列化的 sha256_cbor。该版本文档也列出速度更快的 xxhash / xxhash_cbor，但它们是非加密哈希；只有明确接受碰撞与安全风险时才应选用，不能把它们当成与 SHA-256 等价的安全配置。请求的 cache_salt 会进入首块 hash，把复用限制在同一信任域，降低跨租户命中和基于延迟推断缓存内容的风险。salt/hash 是缓存隔离的一层，不能替代租户鉴权与内存访问控制。

Radix tree 实现

另一条路线是 SGLang 的 RadixAttention：将 token 序列放进 radix tree，树边代表 token span，公共祖先天然对应共享 prefix。它特别适合 Agent、树搜索、few-shot、分支推理等存在嵌套/分叉 prefix 的工作负载。哈希块表更接近页缓存，radix tree 更直接表达层次关系；两者都需要引用计数、块分配和驱逐策略。

常见追问

• 命中 100K prefix 后完全不用读它了吗？ 只是不再重算它的 K/V；后续 decode attention 仍需访问这些历史 KV，除非模型本身使用稀疏/滑窗 attention。
• 为什么只缓存 full block？ 可简化不可变共享、hash 和内存管理；最后一个 partial block 通常继续由当前请求持有。
• Prefix Cache 会改变结果吗？ token、模型权重、位置编码状态、LoRA、多模态输入和 cache dtype 都一致时，复用是等价计算；键缺少这些维度就可能错误命中。

结论

不要把整段压缩 KV 一次性投影成完整多头 K/V 并落到 HBM。核心方案是：按 KV sequence 再分段，load → dequant/proj → attention → online-softmax merge，立刻释放该段的展开激活。进一步可融合 load/projection/attention，让展开后的 K/V tile 只存在于 SRAM/寄存器；显存仍不够时改用 MLA 的 latent/MQA 计算路径或 Ring Attention。

截至 2026-07-15，vLLM Ascend v0.18.0 的官方 Context Parallel 设计对 MLA chunked prefill 正是这样处理：AllGatherKV 对超长 context 会产生明显峰值显存，因此预设每轮最大 KV 量，分段完成 attention，并用 online softmax 顺序合并。这里描述的是该版本后端的实现边界，不代表其他 runtime 自动采用相同路径。

为什么 MLA 反而会冒出大激活

MLA 持久化的不是每头 K/V，而是低秩 latent：

再由 up projection 恢复各头的 no-PE K 和 V。以 DeepSeek-V3 的公开配置为例：

• n_heads=128
• kv_lora_rank=512
• qk_nope_head_dim=128
• qk_rope_head_dim=64
• v_head_dim=128

BF16 持久化 latent+RoPE key 每 token 每层约：

但若一次性展开 no-PE K 和 V：

仅 128K token 的展开临时量就约 8 GiB/层（这是未考虑 TP 分头前的量级示例），而压缩 cache 约 144 MiB/层。MLA 节省的是持久 KV cache；错误的执行计划仍可能制造巨大临时 activation。

分段 projection + online softmax

假设 query chunk 长度为 Tq，历史 KV 长度为 S。把 S 切成 S_tile：

for kv_segment in history:
    load compressed KV segment
    dequant + up-project this segment only
    compute Q × K_segmentᵀ and P × V_segment
    merge (m, l, O) by online softmax
    discard expanded K/V segment

令可用于展开激活的 workspace 为 W，本 rank 负责 h_local 个 head，则粗略上界：

实际还要给 Q、score tile、输出、通信 buffer、量化 scale 和双缓冲留空间，因此不能把 W 全吃满。

可叠加的工程优化

1. Kernel fusion：把 paged KV load、反量化、Wkv_b projection、RoPE 拼接、attention 融合，避免展开 K/V 写 HBM。
2. 双缓冲与流水：segment j 计算时预取 j+1；通信、projection、attention 用不同 stream/event 重叠。
3. 按层 gather、用完即丢：不要 all-gather 全模型所有层的 KV。官方 CP 设计在普通 prefill 中只 gather 当前层，计算后立即释放。
4. MLA latent/MQA 路径（矩阵吸收）：将 K/V up-projection 的权重吸收到 Q 与 output projection，使 attention 直接在 latent cache 上做，避免展开多头 K/V。代价是 latent 维可能大于单头维，计算量和 kernel shape 可能不如 prefill 的 MHA 路径，需按 Tq、S 选择。
5. Ring Attention：每个 rank 只保留一段 KV，边传边算并在线合并，峰值更低；代价是 p 个通信阶段、实现和拓扑调优更复杂。
6. 降低 cache/通信精度：FlashMLA 当前官方支持矩阵明确给出的是 token-level sparse decoding 使用 FP8 KV Cache、以 BF16 做矩阵计算；SM90 dense decoding 路径仍列为 BF16。其每 token FP8 KV 不是简单的 512+64 个低精度元素，而是 656 Bytes：512 个 E4M3 NoPE 值占 512 Bytes，4 个 FP32 分组 scale 占 16 Bytes，64 个未量化 BF16 RoPE 值占 128 Bytes。不能把这一能力泛化成所有 dense/prefill MLA kernel 都支持 FP8，且上线前必须做长上下文精度回归。

常见追问

• 减小 query chunk 是否能解决？ 只能减小 Q/score workspace，无法消除一次性展开整个历史 KV 的 O(S) 峰值；关键还要沿 KV 维分段。
• 分段会改变 softmax 吗？ 使用 (m,l,O) 的精确 online merge 不改变数学结果，只有正常浮点舍入差异。
• 为什么 prefill 有时仍选择展开后的 MHA 路径？ prefill 的 Q 很多，常是 compute-bound；标准 head_dim 的 Tensor Core/FA kernel 可能比 latent 大维度路径更高效。应做运行时分派，而不是一种实现包打天下。

结论

Agent 的典型特征是 system prompt/工具定义长、多轮 append、同一上下文分叉、会话暂停后再恢复。最有效的顺序通常是：先消除重复副本，再提高 GPU KV 利用率，再分层存储，最后才做有损压缩。对应措施是 prefix/session cache、paged blocks + 分支尾块管理、MLA/GQA/低精度 KV、CP 分片，以及 GPU→CPU DRAM→SSD/远端 KV pool 的分层 offload。

1）先把“重复”去掉

• system prompt、tool schema、代码仓库快照等稳定前缀用 Prefix Cache；不要让每个 Agent session 各存一份。
• 多轮会话只对新增 user/tool tokens 做增量 prefill，保留 append-only session KV。
• Agent 做 beam/tree search 或并行 tool plan 时，只读的公共祖先 blocks 可以共享；分支通常新建尾块，只有继续共享半满可写尾块的实现才需要 copy-on-write。PagedAttention 与 radix tree 都适合表达这种共享关系。
• 热 prefix 跨实例复制，冷 prefix 只保留一个远端副本；调度器同时考虑 cache locality 和实例负载。

2）提高 GPU 内存有效容量

• block/page 分配，减少预留和外部碎片；按实际 token 增长分配，而不是按 max length 连续预留。
• KV cache 量化为 FP8/INT8；节省接近 2 倍时，要验证长上下文累积误差、scale 粒度和 kernel 是否真的支持低精度直接计算。
• 使用 GQA/MQA/MLA 架构。它们是模型结构能力，不能在不重新训练/转换验证的情况下对任意 MHA 模型“打开开关”。
• 用 Decode Context Parallel 把 sequence 维 KV 分片到多个设备，消除 TP 域内 KV 副本；代价是每层 attention 多了通信。

3）做分层 KV 存储

建立 GPU HBM、CPU DRAM、本机 NVMe、远端 DRAM/SSD 的多级 cache：

• 活跃会话的最新/最可能访问 blocks 留在 GPU；
• 暂停会话或可复用 prefix 异步 offload 到 CPU/远端；
• 恢复前预取，按层或按 block 与计算重叠；
• 驱逐不只看 LRU，还看复用概率、prefix 长度、重算 FLOP、传输时间和租户优先级。

Mooncake 的 KVCache-centric 架构会利用 GPU 集群中闲置的 CPU、DRAM、SSD，并由全局调度器在 cache reuse、负载和 SLO 之间权衡。这比“GPU OOM 后临时 swap”更适合 Agent 的长生命周期状态。

4）降低必须保留的上下文

以下方法可能改变模型语义，应作为产品/模型策略而非纯 Infra 等价优化：

• 对旧对话做摘要，把可恢复事实写入外部 memory/RAG，再截断原 token；
• 只保留工具结果的结构化摘要，不重复保存大段网页/日志；
• 模型原生支持时使用 sliding-window、local/global 或 sparse attention；
• 基于重要性丢弃 KV 的方法需要任务级精度回归，不能宣称无损。

一个实用容量决策

驱逐分数必须使用一致量纲。不能直接拿“占用字节 × HBM 价格”和“复用概率 × 时间”相除；前者近似容量/货币成本，后者是期望延迟。若优化目标是延迟或 SLO，可对每个 block 在一个调度窗口内估算：

再按每字节保留价值排序：

空间不足时优先驱逐 value_per_byte 小的 blocks。若目标是货币成本，则应把延迟/SLO 违约也换算成同一货币目标，并把预计驻留时间计入 HBM 成本。Agent 场景通常使长 system/tool prefix 的 P_reuse 很高，因此它们虽大，却未必应最先驱逐。

常见追问

• offload 一定增加吞吐吗？ 不一定。若每个 decode step 都同步从 CPU 拉历史 KV，PCIe 会成为灾难；offload 更适合暂停会话/可复用 prefix，或能按层预取并覆盖传输的实现。
• 只开 Prefix Cache 能解决活跃长会话的 decode 读带宽吗？ 不能。它去重和省 prefill，但每步 attention 仍需访问有效历史 KV。
• Agent 最适合 LRU 吗？ 未必。会话可能长时间等待工具后再次命中，普通 LRU 会误删；应加入 session 状态、工作流 DAG 和预计恢复时间。

结论

PD（Prefill/Decode）分离不一定好。它在 prefill 计算密集、decode 带宽密集、两者互相干扰且 TTFT/TPOT SLO 严格时很有价值；代价是两套权重/资源池、额外排队、KV 传输和 cache locality 破坏。Agent 的长 prefix 高命中场景中，真正要算的是“未命中 suffix 有多长”和“prefix KV 在哪里”，而不是原始 prompt 有多长。

如果 100K prefix 已驻留在某个 D 实例，而新一轮只有几十/几百 token，让请求继续在该实例做短增量 prefill，常常比把长 prefix 搬给 P、再把新 KV 搬回 D 更好。

PD 分离为什么有效

• prefill 是大 GEMM/长 Q attention，偏 compute-bound，目标是 TTFT；
• decode 每步 token 少、反复读权重和 KV，偏 memory-bound，目标是 TPOT；
• colocate 时，一个长 prefill 会阻塞 decode iteration，造成 TPOT 抖动；
• 分离后 P、D 可使用不同并行度、batch 和扩缩容策略。DistServe 的核心就是消除 P/D 干扰并按各自 SLO 配资源。

为什么长 prefix cache 可能让 PD 变差

设命中的 prefix 为 L_hit，真正要 prefill 的 suffix 为 L_miss。分离路径近似：

cache-resident colocated 路径近似：

当 L_miss 很小、L_hit 很大时，P 上的计算收益很小，加载/传输长 prefix 的成本却很大。更糟糕的是会话每轮都在 D→P→D 之间来回搬状态。

2026 年的 PPD（Prefill / Prefill-capable Decode）研究专门针对多轮服务验证了这一问题：固定把后续轮次都送到独立 P 节点，会重复制造增量 prefill 与 KV 往返；让一部分 D 节点按负载处理 append-prefill 可降低后续轮 TTFT。这里引用它作为“混合路由值得考虑”的实证，不把论文特定集群上的收益数字外推到所有模型和网络。

工程判断与架构选择

1. cache-aware hybrid routing：
   • cache miss 且 prompt 长：发往 P；
   • D 上命中长 session prefix 且 suffix 短：在 D/mixed worker 增量 prefill；
   • 命中远端共享 cache：比较 reload、recompute 和队列时间再选。
2. 粘性会话 + 负载兜底：优先回到持有 session KV 的 D；当排队收益大于迁移代价才迁移。
3. 全局 KV pool：P、D 不以某张卡为唯一所有者，按 block/层异步装载。Mooncake 就是 KV-centric 调度；但网络和存储仍是实际成本，不是“共享后免费”。
4. 热 prefix 复制：减少所有请求集中到一台 cache-rich worker 的热点；复制数量由命中收益与显存代价决定。
5. layer-wise transfer/overlap：P 算第 l 层时传第 l 层新 KV，避免等完整 prefill 后一次性传输。
6. 保留 mixed 池：流量/命中分布变化大时，纯 P/D 固定切分会资源失衡；一部分 worker 可在 P、D 之间切换或处理短 cached prefill。

什么时候更倾向 PD 分离

• 长 prompt cache miss 多；
• prefill arrival burst 会明显破坏 TPOT；
• 网络带宽高、P/D 拓扑亲和，KV 传输可覆盖；
• TTFT 与 TPOT SLO 都严格，且 P/D 资源需求稳定可预测。

什么时候更倾向 colocate/hybrid

• 多轮 Agent 的 session prefix 命中高、每轮 suffix 很短；
• KV 只能驻留在 D，P 读取它需要跨慢网；
• 低 QPS 下分成两个队列降低 batching；
• P/D 流量比例变化大，固定资源池经常一边空闲一边排队。

常见追问

• PD 分离只需传新 token 的 KV 吗？ 如果 P 已能访问历史 prefix KV，可以只回传新增 KV；如果历史 KV 只在 D，本轮增量 prefill 的 attention 仍需要它，必须加载/传输或改在 D 上算。
• 如何判定迁移？ 估算 剩余服务时间收益 > KV迁移时间 + 目标队列等待 + cache冷启动，并设滞回阈值，避免请求反复抖动。
• 论文说 KV 传输很小，是否可忽略？ 只对论文的模型、拓扑和 placement 成立。DistServe 通过节点亲和让很多传输走 NVLink；跨节点 Agent 长 prefix 需要重新测。

结论

这里先统一术语：CP（Context Parallel）是把 sequence/KV cache 沿上下文维分片到 p 个设备；为了让每个 Q 看见完整历史，可以交换 Q、聚合 KV，或让 KV 沿环流动。下文把 AllGather KV（AG-KV）、AllGather Q（AG-Q）和 Ring Attention 作为三种算法候选比较，不表示任一 runtime、任一版本都同时实现了三条路径。

截至 2026-07-15，vLLM Ascend v0.18.0 文档明确写出：GQA chunked prefill 已实现 AG-Q，MLA chunked prefill 已实现 AG-KV；Ring-Attn 被列为可行/备选算法，但文档没有把它列入这两条已实现路径。回答时应把“算法上可选”与“当前后端已支持”分开。

选择的第一原则是比较本层通信量、峰值显存和可重叠性：

• S 很长、当前 query chunk Tq 很小，尤其 decode/Agent 长 prefix：算法上通常偏 AG-Q 或 Ring，前提是当前后端支持；
• Tq 较大、KV 很小或被 MLA 高度压缩、同节点高速互联：AG-KV 可能更简单、更快；
• AG-KV 的完整 KV 激活放不下：分段 AG-KV + online softmax，或 Ring；
• 不能只看 bytes，还要看 collective latency、额外 reduction、kernel 成熟度以及 CP 是否与 TP 复用同一通信域。

Agent 场景当然还能 AG-KV；只是长 prefix 使它更容易不划算。MLA 压缩 KV、CP 较小、NVLink/HCCL 高带宽、分段后峰值可控，或 query chunk 本身也很大时，AG-KV 仍可能胜出。

三种方案的数据流

设历史 KV 总长度为 S，当前 query 的总 token 数为 Tq，CP 大小为 p，每 token 的 KV/Q/O 元素数分别为 dKV、dQ、dO，元素字节数为 b。为便于比较，下面先采用 sequence-sharded prefill 抽象：每个 rank 初始持有约 S/p 的 KV 与 Tq/p 的 Q。真实实现必须按 collective 前的本地张量字节数重算；例如 vLLM Ascend 的 DCP 与 TP 复用通信域时，GQA 的 Q 是沿 head 维 all-gather，并非这里的 token 维示意。

方案 A：AG-KV

每个 rank 保留本地 Q shard 和 1/p 的 KV，先 all-gather 得到完整 KV，再运行常规 prefill attention：

local Q shard + local KV shard
          ↓ AllGather KV
local Q shard + full KV
          ↓ FlashAttention
local O

若 KV 均匀分片，每 rank 一层大约接收：

优点是计算路径接近成熟的标准 FA，结果不需要跨 rank 合并；缺点是通信随长历史 S 增长，并且每个 rank 都会短暂物化 full KV。

方案 B：AG-Q

每个 rank 保留本地 KV shard，all-gather Q，让每个 rank 用完整 Q 对自己的局部 KV 做 partial attention，再用各 shard 的 max/LSE/output 做精确 softmax merge，最后按需要 reduce-scatter 输出：

local Q + local KV shard
   ↓ AllGather Q
full Q × local KV -> partial (m,l,O)
   ↓ merge / A2A / RS
final O shard

在上述均匀 sequence-sharded 抽象下，本地 Q 为 Tq/p × dQ，所以 Q gather 每 rank 约：

还要加 partial output 和 LSE 的合并通信，粗略为：

当 S dKV >> Tq(dQ+dO) 时，AG-Q 通常更省通信和峰值显存。代价是多一步 partial-softmax merge，且 Q head/TP 布局可能需要额外变换。若 runtime 实际沿 head 维聚合 Q，应以本地 Q head shard 的真实 bytes 替换上式，不能机械套用 Tq/p。

方案 C：Ring Attention

Q 留在本 rank，KV shard 沿 ring 逐跳流动；每收到一段就计算并在线更新 (m,l,O)：

step 0: Q × local KV
step 1: Q × KV from left neighbor
...
step p-1: cover all KV shards

它不物化 full KV，可将通信与计算流水重叠，峰值内存最低；但有 p-1 个阶段，首包/同步延迟、负载平衡和故障处理更复杂。短 query 且每段计算太少时，通信未必能被覆盖。

vLLM Ascend v0.18.0 中 MLA 与 GQA 为什么选不同方案

该版本的官方实现中，GQA chunked prefill 使用 AG-Q，MLA 使用 AG-KV。原因不是“模型名称决定算法”，而是张量布局和通信量不同：

• GQA 的 KV heads 少，而 TP/CP 下 Q heads 的布局使聚合 Q 后对本地 KV 做 partial attention 更自然；
• MLA 持久 KV 是低维 latent，dKV 显著小于展开后的多头 K/V，因此 gather 压缩 KV 可能比想象中便宜，并可复用标准 prefill 路径；
• 但 MLA 的 AG-KV 后若一次性 up-project 整段历史，仍会产生第 4 题的大 activation，所以官方实现对超长 KV 分段处理。

这是一条带版本号的实现案例，不是跨 runtime 的固定规则；Ring 仍可用于推导峰值显存和流水边界，但不能据此声称 vLLM Ascend v0.18.0 已把它作为上述 chunked-prefill 的现行生产路径。

Agent 长 prefix 的选择步骤

1. 用本轮 Attention 实际可见的完整历史 S 和 Tq 代入：S 仍包含命中的 prefix 与新增 suffix，Prefix Cache 命中只省掉重复 prefill，并不会让 decode/CP 少看那段历史。只有模型窗口截断、滑窗或稀疏 Attention 才会缩短这里的有效 S。decode 时 Tq≈batch，通常远小于 S。
2. 估算三种方案每层的 bytes、峰值 workspace 和 collective 次数。
3. 判断链路：同机 NVLink 与跨机 IB/RoCE 的带宽、RTT 差异很大；AG-KV 若跨节点会更敏感。
4. 若 AG-KV 只有峰值显存不满足，先尝试 segmented AG-KV，不必立即完全改 AG-Q。
5. 以端到端 TPOT/TTFT 实测。AG-Q bytes 少，不代表 partial merge kernel 一定更快。

常见追问

• CP 是否一定减少单请求延迟？ 不一定。它减少每卡 KV 容量和局部 attention 工作量，但引入通信；更常见的直接收益是支持更长 context/更大 batch。
• AG-KV 是否意味着永久复制完整 KV？ 不必。可以只 gather 当前层、当前 segment，用完释放；持久 cache 仍保持分片。
• AG-Q 的结果怎么合并？ 每个 KV shard 返回局部 softmax 的 (m,l,O)，按与 FlashAttention 相同的 log-sum-exp 规则精确合并。

结论

长输出的核心问题不是单次算子，而是请求占据 decode slot 和 KV memory 的时间很长，形成 straggler、batch 收缩和实例间负载漂移。应使用iteration-level continuous batching + 按预计剩余 token 工作量调度 + KV-aware 迁移/抢占；同时用 speculative decoding、DCP/KV 分片和低精度 KV 降低单 token 成本。

不能把一条自回归序列的未来 token 直接横向切给多机独立计算，因为 token t+1 依赖 token t。可并行的是同一步的模型/上下文，或一次草拟并验证多个 token。

负载指标不能只看 request 数

一个实例有 10 个即将结束的请求，可能比另一个只有 2 个 100K-context 长推理更空。调度应估算：

其中 Lremain 很难准确预测，可用 max_tokens、任务类型、历史分布、当前 EOS hazard 形成区间估计；对无法预测的请求使用 aging/WFQ 保证公平。

单实例内

1. Iteration-level continuous batching：每个 decode iteration 都移除完成请求并接纳新请求，避免整批等待最长序列。
2. token-budget batching：限制本轮总 token/KV 工作量，而不是固定 request count；将相近 context length 的请求放一起可减少 kernel 尾部和 padding。
3. chunked prefill 调度：给 decode 留固定 token budget，避免长 prefill 一次占满 iteration；严格 TPOT 下优先 decode。
4. 抢占与换出：显存不足时对低优先级/长等待请求 swap 或 recompute。长 context 的重算非常贵，通常优先 offload KV；抢占太频繁说明 admission control 有问题。
5. 公平与 SLO：对超长 generation 设置独立队列、权重或最大并发，防止短请求 head-of-line blocking，但不能让长请求永久饥饿。

多实例之间

• 新请求发往预计 finish time 最小的 D 实例，而不是简单 round-robin。
• 调度成本同时考虑 active tokens、KV bytes、context 分布、cache locality、网络和队列。
• 迁移活跃请求时需搬 KV。只在：

时迁移；否则让旧请求自然 drain 更便宜。

• 对热点 session/prefix 适度复制，避免粘性调度把一台机器压爆。
• 扩容时先把新请求导向新实例（drain-based balancing），比立即迁移所有长请求风险低。

降低长输出的单 token 成本

1. Speculative decoding / MTP：draft 多个 token，由主模型一次验证，减少串行 forward 次数。DeepSeek-V3 原始技术报告在其自身评测中给出第二个 MTP token 接受率约 85%～90%、约 1.8× TPS；这不是跨模型或跨部署的通用倍率，实际收益随采样、任务、batch 与验证实现改变。
2. Flash-Decoding / split-KV decode：小 batch 长 context 时沿 KV 维增加并行度。
3. DCP/sequence-sharded KV：每卡只存/读部分历史，再合并 partial attention；以通信换容量与局部带宽。
4. FP8/INT8 KV、MLA/GQA：减少每步必须读的 KV bytes。
5. 模型支持时使用 sliding-window/sparse attention：把每步读取从全历史改为受限范围；这是模型语义/结构选择，不是任意模型的透明优化。
6. admission control：根据 max_new_tokens 和 KV 配额提前拒绝/降级，避免系统先接收再在中途 OOM。

常见追问

• 为什么迁移长请求很贵？ 其 KV 大小与已生成/输入 context 成正比，可能需要跨网络搬数 GB；迁移收益却只来自剩余 token。
• 能否把一个长 decode 请求拆到两台分别生成一半？ 时间维不能；可用 TP/CP 并行同一个 token step，或 speculative decoding 并行候选验证。
• 如何避免预测输出长度错误？ 使用滚动估计与工作窃取；调度器必须在每个 iteration 更新，而不是只在 admission 时决定一次。

结论

大 EP 的目的主要是分摊专家权重、汇聚更多 token 形成更大的 expert GEMM、扩大可部署模型规模；大 DP 的目的是复制服务实例扩大独立请求吞吐，并为 MoE 路由提供更大的 token pool。EP 不是越大越好：All-to-All 域、跨节点 RTT、负载不均、每卡专家数和通信 buffer 都会成为瓶颈。

对 DeepSeek-V3 的教学例子，若明确假设：

• 256 个 routed experts；
• TopK=8；
• 每个 EP rank 本步有 16 个原始 token；
• 路由完全均匀；
• 希望每个 expert 平均收到 16 个 token；

则 EP=32。这与 DeepSeek-V3 官方 prefill 部署使用 EP32 一致，但后者是端到端 profiling 后的工程配置，不是上述四个数字唯一推导出的普适最优。

推导

设：

• P 为 EP size；
• B_local 为每 rank 进入 MoE 的 token 数；
• E 为 expert 总数；
• k 为 TopK；
• M_target 为希望每个 expert 聚到的平均 token 数。

EP group 总原始 token 为 PB_local，总 assignment 为 kPB_local，因此：

反解：

代入 Mtarget=16、E=256、k=8、Blocal=16：

也可直接看：EP32 共 32×16=512 个 token，产生 512×8=4096 个 assignments，平均分到 256 experts，正好每 expert 16 个。

必须指出的隐含假设

题目只说“单卡 batch size=16”仍然信息不足，除非同时给出或默认 Mtarget=16。若第 1 题 Roofline/实测得出的最佳 expert batch 是 Mtarget=64，同样公式会给出 EP=128；若目标约 333，则需要 EP≈666，已经超过 expert 数并需大量 expert replicas，通常不现实。

此外，线上路由不是完全均匀。若某些 expert 热，平均 16 不代表最热 expert 只处理 16；需要真实 trace、capacity factor 或 redundant expert。

为什么大 EP

• 每卡仅存 E/P 个原始 experts，降低 expert weight 显存；
• EP group 汇总 P 个 rank 的 token，提升每 expert grouped-GEMM 的 M；
• 可以将热门 expert 做冗余副本并动态路由；
• 代价是 dispatch/combine All-to-All 扩大，跨节点小包延迟与拥塞上升，任一热点 rank 会拖慢整个 collective。

为什么大 DP

• attention/dense 部分独立处理不同请求，近线性扩大 aggregate throughput；
• 不增加单请求模型并行通信；
• 配合 DP-attention + EP-MoE，可让 attention 保持小 TP/高效率，同时 MoE 跨 DP ranks 汇总 token；
• 代价是复制非专家权重、占用更多显存，并把 KV/cache locality 分散到更多实例。

DeepSeek-V3 的两份官方部署披露说明了什么

• 2024 技术报告中的原始案例：Prefill 最小部署单元为 4 节点、32 GPU；attention 使用 TP4+SP 与 DP8，MoE 使用 EP32。每卡放 8 个原始 routed experts，另加 1 个冗余 expert。
• 同一技术报告中的 Decode 案例：最小部署单元为 40 节点、320 GPU；attention 为 TP4+SP、DP80，MoE 为 EP320。每卡一个 expert，64 张卡用于 shared/redundant experts。报告还说明 decode 每 expert batch 通常不超过 256，瓶颈偏内存访问，因此使用更少 SM 做 dispatch+MoE+combine，并与另一 microbatch 的 attention 重叠。
• 2025 年 2 月官方在线系统披露：V3/R1 的 Prefill 使用 routed-expert EP32 与 MLA/shared-expert DP32，4 节点中放置 32 个冗余 routed experts，每卡管理 9 个 routed experts 和 1 个 shared expert；Decode 则使用 routed-expert EP144 与 MLA/shared-expert DP144，18 节点中同样放置 32 个冗余 routed experts，每卡管理 2 个 routed experts 和 1 个 shared expert。

两份官方材料已经直接说明：EP32、EP320 或 EP144 都是特定时间、硬件、运行时和负载下的工程配置，不是模型结构唯一推导出的常数。Prefill 和 Decode 的合适 EP 本来就可以不同，而且会随部署演进：前者更重视大 batch 计算效率和双 microbatch overlap，后者更重视单 expert 权重读取、低延迟通信、冗余与 attention overlap。

选 EP 的实际流程

1. 先满足权重容量、E 的可分片性和拓扑边界；优先让高频 A2A 留在 NVLink 域，跨节点再用分层通信。
2. 用路由 trace 算各 P 下的 per-expert M 分布，而不只看均值。
3. 分别 profile grouped GEMM、dispatch、combine，找 T_total(P) 最小点。
4. 检查 P99 expert imbalance、网络拥塞、buffer 峰值和故障域。
5. 热点明显时比较“继续增大 EP”与“复制 hot experts/EPLB”哪个更便宜。

常见追问

• EP 可以大于 expert 数吗？ 可以，但必须有 shared/redundant expert 或其他复制映射；否则多出的 rank 没有独立 expert。DeepSeek-V3 2024 技术报告中的 Decode EP320 就是一个包含冗余/shared 部署的历史案例，不代表后续在线系统仍固定使用 EP320。
• 大 DP 会让单请求更快吗？ 通常不会，主要提高多请求吞吐与排队能力。
• EP32 为什么每卡 8 个 expert？ 256/32=8；另一个冗余 expert 是负载平衡策略，不属于 256 个原始 expert 的均匀切分。

结论

先用 SLO 和显存确定 Decode 的可行 batch，再用真实长度分布 profile 出每个 D replica 的“满足 TPOT 的 goodput”；Prefill 同样用 cache miss 后的 prompt 长度和 TTFT profile 容量。P:D 配比是两个服务中心的容量平衡：

不是按 GPU FLOPS 静态拍一个 1:1 或 1:2。还必须验证 KV 传输带宽、两阶段队列和 burst headroom。

第一步：定 Decode batch size

Decode batch B 同时受三类约束：

1. TPOT SLO：实测 t_step(B,Lctx) 的 P99 必须小于目标 TPOT；连续 batching 下还要算调度间隙。
2. KV 显存：

变长 batch 应按 token/block 总数算，不能简单用 B×平均长度 忽略 P99。 3. 吞吐拐点：

增大 B 到 TPS 增益很小、但 TPOT/排队明显变差时停止。

因此“合适的 B”通常在满足 P99 TPOT 和显存安全的最大值附近。在线低延迟服务还要为 burst 和长 context 留余量；10%～30% 只能作为压测前的演示性起点，不是行业通用常数，最终应由目标流量 trace、故障冗余和 SLO 反推。

第二步：把 D 的 token capacity 转成请求 capacity

若 profile 得到每个 D replica 在目标 SLO 下吞吐 TPS_D，平均输出长度为 E[Lo]：

更准确做法是 trace replay，因为 t_step 随 context 增长，E[Lo] 的均值也掩盖长尾。做第一轮容量草算时可以用 ρ=0.6~0.8 演示安全利用率；它不是固定推荐值，生产取值仍要由到达过程、服务时间分布、扩缩容时延和目标 P99 校准，以免 M/G/k 队列在高利用率下尾延迟陡升：

第三步：profile Prefill 容量

Prefix Cache 开启后要用 Lmiss，不是原始 prompt 长度。若一个 P replica 在目标 TTFT/chunk 策略下能处理 TPS_P 个未命中 prompt tokens/s：

同理：

若已经定了 ND，并让 D 以安全容量 ρD μD 运行，可反推：

但最终仍应取业务到达率 λ 与上式所需的较大者，不能因为 D 暂时配置少就低估 P 的入口 burst。

一个具体例子

假设：

• D 在 B=64 时 P99 step=45 ms，满足 50 ms TPOT；
• TPS_D=64/0.045≈1422 token/s；
• 平均输出 400 token，因此 raw μD≈3.56 req/s；
• 取 ρD=0.7，安全容量约 2.49 req/s；
• P profile 为 50K 未命中 prompt token/s，平均 Lmiss=2000，raw μP=25 req/s；
• 取 ρP=0.7，安全容量 17.5 req/s。

一台 P 理论上可喂约 17.5/2.49≈7 台 D。若部署 8 个 D：

这个 2P:8D 只对上述 profile 和分布成立。若 Agent prefix hit 提高、Lmiss 降低，P 更少；若输出变短，D req/s 提升，反而需要更多 P 才能喂满 D。

别忘了 KV 传输与队列

每秒 KV 传输需求：

要求链路在安全利用率下满足它，且 P→D placement 不产生热点。还应分别测：P queue TTFT、compute TTFT、KV transfer、D queue、TPOT；总平均吞吐相等并不代表 P99 SLO 成立。

常见追问

• 为什么不能只按 prefill FLOPs/decode FLOPs 配比？ decode 主要读权重/KV，prefill 主要做大 GEMM；两者受不同屋顶、batch 和 SLO 控制。
• P 与 D 是否必须整数固定配比？ 不必。可建立共享弹性池或 mixed workers，但切换模型/缓存状态有成本。
• B 越大 TPS 一定越高吗？ 通常先升后饱和；context KV 读、A2A、调度和显存会让收益递减，且 TPOT 可能先违反 SLO。

结论

MoE overhead 不只是两次 All-to-All。完整路径是 gate/TopK → 计数与 prefix-sum → permute/pack → dispatch A2A → grouped expert GEMM → combine A2A → unpermute/加权归并；还包括专家负载不均、padding/capacity、同步、通信占用 SM、跨节点拥塞和小 GEMM 低利用率。

若每 rank 有 B 个原始 token，hidden size H，TopK=k，EP size=P，通信元素 b Byte，路由到 rank 均匀，则标准 token-expert-pair dispatch 的远端 payload 为：

整个 EP group 单向注入网络的远端 payload：

combine 若以相同精度返回每个 assignment 的 hidden output，大小相同。因此 dispatch+combine：

先声明“通信量”的统计口径

• 上式 fabric 是所有 rank 的远端 send payload 之和，local assignment 不上网络。
• 若监控把 send 和 receive NIC bytes 都相加，数字再乘 2。
• 一些库报告 logical bandwidth，会把 local traffic 也算进去；DeepEP README 明确提示其部分带宽数字包含 local rank traffic。
• padding、alignment、量化 scale、token index、routing weight 和协议 header 未包含在理想公式里。

DeepSeek-V3、EP32、每卡 B=16 的数值例子

假设：

• P=32
• B=16
• k=8
• H=7168
• dispatch 与 combine 都用 BF16，即 b=2
• expert 在 ranks 间均匀，忽略 shared/redundant expert。

每 rank 共 128 个 token-expert pairs，包含 local 的逻辑 payload：

乘远端比例 31/32：

• dispatch 每 rank 远端 send 约 1.695 MiB；
• dispatch 全 EP group 注入网络约 54.25 MiB；
• combine 同样约 54.25 MiB；
• 两阶段合计约 108.5 MiB/层/step 的远端 send payload。

如果 dispatch 用 FP8、combine 用 BF16，则约为：

若按 send+receive 两方向计 NIC 总字节，分别约 217 MiB 和 162.8 MiB。metadata 本身很小：例如每 assignment 一个 int32 expert/index 和 FP32 weight，单 rank约 128×8=1 KiB，但小 batch 时控制延迟仍不可忽略。

主要 overhead 分解

1. Router：gate GEMM、sigmoid/softmax、group-limited TopK、归一化。
2. Token layout：统计每 expert/rank token 数，prefix sum，permute、量化和 pack；大量不规则访存。
3. Dispatch A2A：跨 rank 发 token activation；跨节点时受 NIC、RTT、拥塞、拓扑影响。
4. 负载不均：collective/层延迟由最慢 rank 决定；热门 expert、capacity padding、空 expert 都降低利用率。
5. Grouped GEMM：expert M 小且形状不同，Tensor Core tile 浪费；权重是否命中 L2 也影响巨大。
6. Combine A2A：结果回源 rank，再 unpermute，按 TopK routing weight 加权求和。
7. 资源干扰：通信 kernel 使用 SM、shared memory、L2 和 HBM；与 expert GEMM/attention overlap 不当会互相抢资源。
8. 冗余/EPLB：复制 hot experts 可平衡负载，但增加权重显存、路由表更新和一致性管理。

优化方向

• 分层 A2A：先 IB 跨节点聚合，再 NVLink 节点内转发，匹配拓扑；或使用直接低延迟 P2P/RDMA。
• dispatch 使用 FP8，combine 按精度需求保留 BF16；scale 与量化融合到 pack/kernel。
• permute、通信、grouped GEMM、combine 流水化；双 microbatch 隐藏通信。DeepSeek-V3 prefill 官方就同时处理两个 workload 相近的 microbatch 来重叠 attention/MoE 与 dispatch/combine。
• redundant experts、周期性 EPLB 和动态路由解决热点；buffer 按 P99 token count，而不是最坏值无限预留。
• decode 小 batch 仍优先压 RTT 与同步，prefill 大 batch 更看重吞吐和通信—计算重叠；但不要把当前 DeepEP 描述成两套彼此独立、固定不变的 API。截至 2026-07-15，DeepEP V2 主线文档已在 ElasticBuffer 下统一高吞吐与低延迟接口，并按分析模型配置通信 SM / QP；其 V2 说明还明确写出不再支持旧版“0 SM RDMA low-latency EP”。面试时应同时说清 workload 差异与版本边界。

常见追问

• 为什么是 TopK 倍通信？ 标准 dispatch 对每个 token-expert assignment 发送一份 activation。若同一 token 的多个 experts 恰好在同一 rank，专门实现可先去重传输再在目标 rank 复制，但不能把这种优化当作通用公式。
• combine 能否只返回一个 token 向量？ 只有目标端已拥有该 token 的全部 TopK expert 结果并能先聚合时才行；一般 TopK experts 分散在不同 ranks，需逐 assignment 回源再合并。
• A2A bytes 不大为什么还慢？ decode 是大量小消息，常由 RTT、同步和不均衡主导；带宽公式只给下界，不给实际延迟。

结论

4 台全连接服务器最自然是排成 2×2 逻辑网格，使用 2D block matrix multiplication/SUMMA。每台负责一个 N/2×N/2 的 C block，计算量完全均衡；A panel 沿行广播、B panel 沿列广播。

每台容量为 N² 元素时，总容量 4N² 足够保存 A、B、C 共 3N² 元素，但每台只剩 N²/4 workspace，不能天真地一次收两个完整 N²/4 block。应把内积维再切成宽度不超过 N/4 的 panels，流式广播。

若“缓存只有 N²/8”是每台唯一可用存储，则 4 台合计只有 N²/2，连一个输入矩阵都放不下，问题在没有外部存储/流输入的前提下不可行。若 cache 指快存，而 A/B/C 在慢存或远端，则使用 out-of-core tiled SUMMA，选择 tile 使 A、B、C tile 和双缓冲装得下。

情况一：每台 memory = N²

将矩阵切成 2×2 blocks：

每个 block 为 N/2×N/2=N²/4 元素。服务器 Pij 初始保存 Aij、Bij，负责输出 Cij：

若不考虑临时 buffer，每台常驻：

剩余 N²/4。

Panelized SUMMA 调度

将 inner dimension k=0...N-1 按宽度 w 切 panel。一次 microstep：

• P(i, owner_k) 沿逻辑行广播 A[i, k:k+w]，大小 (N/2)×w；
• P(owner_k, j) 沿逻辑列广播 B[k:k+w, j]，大小 w×(N/2)；
• 每个 Pij 执行：

最坏情况下一个 rank 同时需要两个远端 panels，workspace：

要求 Nw≤N²/4，即 w≤N/4。选 w=N/4 时恰好使用剩余空间，需要 4 个 microsteps 覆盖整个 inner dimension；若要 A/B 双缓冲以覆盖通信，选 w≤N/8 更安全。

计算与通信量

每台负责 C 的四分之一：

• 标量乘加数 N³/4；若一次乘加算 2 FLOP，则 N³/2 FLOP；
• 每台需要远端的 A 半块 N²/4 和 B 半块 N²/4，理想接收总量约 N²/2 元素；发送量对称；
• panelization 改变消息个数和重叠，不改变理想总 payload。

全连接物理网络不代表应做 4 份完整矩阵复制。2D SUMMA 利用行/列子通信组；实际应把逻辑网格映射到带宽更高的拓扑，并用 nonblocking broadcast + GEMM overlap。

情况二：每台快缓存 = N²/8

先说可行性

若没有慢存、远端内存或输入流，4 台总容量：

连 A 都无法同时存在，因此任何“只在这 4 块 cache 内完成”的方案都不成立。面试中应先指出这一信息论层面的约束，再说明加入 backing store 后的方案。

有 backing store 时：out-of-core tiled SUMMA

选择方形 tile 边长 t。一次最基本计算需驻留一个 A tile、一个 B tile、一个 C accumulator：

若 A/B 都做双缓冲，则需 C+2A+2B：

工程上可选整齐的 t=N/8：双缓冲共 5N²/64 < N²/8，还留有 metadata/通信空间。

此时每维 8 个 tiles。将 C tiles 按 (tile_i mod 2, tile_j mod 2) cyclically 分给 4 台服务器。每台对自己的 C tile：

load/zero one C tile
for k_tile = 0..7:
    prefetch/broadcast A[i,k] and B[k,j]
    C[i,j] += A[i,k] × B[k,j]
write C tile back to slow memory

4 台各负责 64 个 C tiles 中的 16 个，每个 C tile 做 8 次 tile GEMM，因此计算仍均衡。双缓冲使第 k 次 GEMM 与第 k+1 次读/通信重叠。

该简单调度一次只保留一个 C tile，可能重复加载 A/B。可把剩余 cache 用于一个 C tile strip 或矩形 blocking，增加 A/B 复用；原则是 tile 尺寸取 Θ(√M)，接近经典矩阵乘通信下界：对 fast memory 容量 M，数据移动至少为 Ω(FLOPs/√M) 量级。

常见追问

• 为什么不用把 A 按行、B 按列各复制两份？ 可以换通信换内存，但这里每台内存有限；2D SUMMA 的每台存储、计算和通信更均衡。
• C 的 partial sum 要 AllReduce 吗？ 上述 2D 算法让每个 Cij 的所有 k panels 都在同一服务器累加，不需要最终 AllReduce；若沿 K 维做 1D split，则需要对完整/分块 C 做 Reduce-Scatter。
• Strassen 是否更好？ 工程中的大 dense GEMM 通常先优化数据分块、通信和 BLAS-3 效率；Strassen 常数、额外 workspace、通信和数值稳定性使它不是这道 Infra 题的默认答案。

一句话结论

MHA、GQA、MQA 的主要区别是 多少个 Query Head 共享一组 K/V Head。它们的 Query 数量和 Attention 主体计算量基本不变，但依次减少 K/V 投影、KV Cache 容量和 Decode 时读取 KV Cache 的 HBM 流量，因此 Decode 的计算访存比依次升高。MLA 不只是“共享 K/V Head”，而是把每个 token 的 K、V 联合压缩为一个低维 latent；Decode 时通过矩阵吸收直接在 latent 空间计算，用更多计算换取数量级更小的 KV Cache。

原始 MHA、MQA、GQA 和 MLA 分别见 Attention Is All You Need、Fast Transformer Decoding: One Write-Head is All You Need、GQA: Training Generalized Multi-Query Transformer Models from Multi-Head Checkpoints 和 DeepSeek-V2。

先统一符号

• B：batch 中的序列数；
• S：当前已缓存的上下文长度；
• L：Transformer 层数；
• H：Query Head 数；
• H_kv：K/V Head 数；
• d：每个 Attention Head 的维度，先假设 K、V 维度相同；
• e：每个 KV 元素的字节数。

四种结构可以写成：

因此，普通 MHA/GQA/MQA 的完整 KV Cache 大小是：

M_KV = 2 × B × S × L × H_kv × d × e  bytes

从 MHA 改为 GQA-G，KV Cache 缩小 H/G 倍；改为 MQA，缩小 H 倍。

对算子图的影响

MHA

典型算子链是：

X
├─ GEMM: XWq → Q
├─ GEMM: XWk → K ─→ 写 KV Cache
└─ GEMM: XWv → V ─→ 写 KV Cache

QKᵀ → scale/mask → softmax → PV → output projection

每个 Query Head 都读取自己对应的 K/V。Decode 时 QKᵀ 和 PV 通常是长向量点积/矩阵向量型计算，天然容易受 KV Cache 带宽限制。

GQA / MQA

算子种类不变，但形状发生变化：

• Q 投影输出仍为 H×d；
• K、V 投影输出从各 H×d 变为各 G×d 或 1×d，投影权重和投影 FLOPs 同步减少；
• Attention kernel 需要做 Query Head 到 KV Head 的映射，例如 GQA 中 kv_head = q_head / (H/G)；
• 不能真的把共享的 K/V repeat 成 H 份再写回 HBM，否则会把节省的带宽重新浪费掉。kernel 应在寄存器、Shared Memory 或 Cache 中复用一份 K/V；
• 如果 grid 只沿 KV Head 划分，MQA/GQA 会因为 H_kv 太小而损失并行度。常见做法是同时沿 batch、Query Head Group、sequence tile 划分，并让同一 CTA 内的多个 Query Head 复用 KV tile。

对长度为 T 的一批输入，Q/K/V 线性投影的主要 FLOPs 近似为：

F_qkv ≈ 2T × d_model × (H + 2H_kv) × d

MHA 对应 H_kv=H，即括号内是 3H；MQA 在 H 很大时括号内接近 H，所以 Q/K/V 三个投影的合计 FLOPs 最多可接近缩小到原来的三分之一。单看 K/V 两个投影，它们从 2H 个 Head 输出降为 2 个，约缩小 H 倍；但 output projection 不变，Attention 中每个 Query Head 的点积也仍要计算，不能说整个 Attention FLOPs 缩小了 H 倍。

MLA

MLA 先联合压缩 K/V：

c_KV = W_DKV × x                 # c_KV 维度 d_c
K     = W_UK  × c_KV
V     = W_UV  × c_KV

为兼容 RoPE，DeepSeek-V2 还缓存一个解耦的共享 RoPE Key，维度记为 d_r。每 token、每层只需缓存：

[c_KV, k_R]，共 d_c + d_r 个元素

关键是 Decode 的“矩阵吸收”：

• 把 W_UK 吸收到 Query 侧，先把每个 Query Head 变换到 d_c 维，然后直接与缓存的 c_KV 点积；
• 把 W_UV 吸收到 output projection，先用 Attention 权重对 c_KV 做加权和，再走合并后的输出投影；
• 因而不必在每一步为全部历史 token 重建完整 K/V。

这使 Decode Attention 的形态从普通的 d 维点积变成以 d_c 为主的点积/归约。MLA 的单次计算量更大、寄存器和 Shared Memory tile 更大，但读取的 KV Cache 小得多。

工程上通常为 MLA 准备至少两条路径：

• Prefill 路径：Q 长度较大，GEMM 和 tiled Attention 易获得较高利用率，可将 latent 批量上投影后使用类似 FlashAttention 的 kernel，或使用专门的 MLA Prefill kernel；
• Decode 路径：Q 长度通常为 1，采用矩阵吸收后的 latent Attention，避免为整个历史上下文重建 K/V；
• 两条路径的交叉点由 Query 长度、上下文长度、batch、GPU Tensor Core 吞吐和 HBM/L2 带宽共同决定，不能只凭模型结构硬编码。

Decode 的计算访存比推导

先只计算 Attention core 对历史 KV 的读取，不计权重读取、Q/输出张量、softmax 和 kernel 额外流量。

普通 MHA/GQA/MQA 在一步 Decode 中：

QKᵀ FLOPs ≈ 2BSHd
PV   FLOPs ≈ 2BSHd
总 FLOPs   ≈ 4BSHd

KV 读取字节 ≈ 2BSH_kv d e

所以理想的 KV 算术强度是：

I_KV ≈ (4BSHd) / (2BSH_kv d e)
     = (2/e) × (H/H_kv)  FLOP/Byte

对 FP16/BF16，e=2，公式进一步变成：

MHA: I_KV ≈ 1
GQA: I_KV ≈ H/G
MQA: I_KV ≈ H

这正是 MQA/GQA 对 Decode 有效的硬件原因：计算量大体不变，HBM 字节数下降。注意这只是 KV 流量的理想值；batch 很小时，模型权重读取也可能占主导，paged layout、Cache miss、重复加载和非合并访存也会降低实测强度。

MLA 经矩阵吸收后，每个历史 token 的 Attention core 主要计算量近似为：

score 的 latent 部分: 2H d_c
score 的 RoPE 部分:   2H d_r
latent value 加权和:   2H d_c

F_MLA ≈ 2SH(2d_c + d_r)

若 [c_KV, k_R] 从 HBM 理想地只加载一次：

I_MLA,ideal ≈ 2H(2d_c+d_r) / ((d_c+d_r)e)

实际 kernel 往往因为 Head Group 分块而把 latent 流加载 ρ 次，或由 L2 命中部分承接复用，此时更现实的表达是：

I_MLA ≈ 2H(2d_c+d_r) / (ρ(d_c+d_r)e)，1 ≤ ρ ≤ H

因此“MLA 的计算访存比是多少”不能只给一个常数：必须说明是 最小 HBM 流量模型，还是某个具体 tiling 下测得的 DRAM bytes。

DeepSeek-V2 参数的定量算例

DeepSeek-V2 原论文与官方配置给出：H=128、qk_nope_head_dim=128、qk_rope_head_dim=64、v_head_dim=128、d_c=512、L=60。这里必须注意：实际 Query/Key Head 是 128+64=192 维，Value Head 是 128 维，不能把二者都写成 128 维后再称为 DeepSeek-V2 的同形状对比。

按 FP16/BF16 缓存，一个保留相同 K/V 维度的朴素多头路径与 absorbed MLA 路径分别是：

同形状朴素 MHA：
  每 token 每层 KV = 128×(192+128)×2
                    = 81,920 B = 80 KiB

MLA：
  每 token 每层 KV = (512+64)×2
                    = 1,152 B = 1.125 KiB

本算例缩小倍数 = 81,920 / 1,152 ≈ 71.11×

DeepSeek-V2 论文另有“相对 DeepSeek 67B，KV Cache 减少 93.3%”的模型级报告；它比较的是两个不同模型基线，不是上面固定层数、Head 数与维度后的理论压缩比，二者不能混用。

把这组 H=128,d_k=192,d_v=128,e=2 参数放在一起比较，并令 GQA 有 G=8 个 KV Group：

这张表比较的是 Attention core 的历史 KV 路径。它没有计入 QKV/output projection、模型权重、softmax、分页元数据和重复加载，因此不能直接当成端到端加速比。

若 S=8192、单序列：

MHA：每层 640 MiB，60 层约 37.5 GiB
MLA：每层 9 MiB，60 层约 540 MiB

对应每个历史 token 的 Attention core：

MHA FLOPs = 2×128×(192+128) = 81,920
MLA FLOPs = 2×128×(2×512+64) = 278,528

MLA / MHA 计算量 = 3.4×
MLA 理想 I_KV = 278,528 / 1,152 ≈ 241.78 FLOP/Byte
MHA 理想 I_KV = 81,920 / 81,920 = 1 FLOP/Byte

这个同形状算例直观说明 MLA 的取舍：Attention core 多做约 3.4 倍计算，换取约 71.1 倍 KV Cache 压缩；在理想复用下，KV 相关算术强度从约 1 提高到约 242 FLOP/Byte。 端到端速度不会等于这些倍率，因为还包括模型权重、投影、softmax、调度和通信。

Prefill 为什么不能直接套上面的公式

Prefill 有很多 Query Token，Q/K/V tile 能在片上反复复用；其 Attention FLOPs 随序列长度近似二次增长，而必要输入输出流量近似一次增长，所以更容易走向 compute-bound。MQA/GQA 仍能降低 K/V 投影和 KV Cache 写入，但 Attention 中 H 个 Query Head 的配对计算仍然存在。上面的 I_KV 公式专门描述 单步 Decode 读取历史 KV，不能拿来声称 Prefill 也有同样的 H 倍加速。

常见追问

追问 1：GQA/MQA 为什么不一定按 KV 缩小倍数加速？

因为端到端还要读模型权重、做 Q/output projection、softmax 和调度；共享 K/V 后可并行的 KV Head 也减少。若 kernel 把 K/V 展开、重复加载或 grid 太小，理论带宽收益兑现不了。

追问 2：MQA 一定优于 GQA 吗？

不是。MQA 的 KV 最小，但表达能力和模型质量可能受影响；GQA 用多个 KV Group 在质量、带宽和并行度之间折中。它们是模型架构选择，不是只改一个推理 kernel 就能无损切换。

追问 3：MLA 一定比 GQA 快吗？

不是。MLA 显著省 KV，但 absorbed Attention 的维度更大，可能转为 Tensor Core/计算受限；小上下文、小 batch 或 Prefill 场景未必占优。应分别对 Prefill、Decode 和真实请求分布测量。

追问 4：如何判断实现是否真的吃到了 MQA/GQA/MLA 的收益？

用 profiler 同时检查 DRAM bytes、L2 hit rate、Tensor Core 利用率、kernel grid、occupancy 和 achieved bandwidth；再对照上述理论最小流量。只看 kernel latency，无法区分是算法少读了数据，还是测试形状变了。

一句话结论

算子优化不是“把 CUDA 写得更复杂”，而是先用 Roofline 判断瓶颈，再围绕四件事做取舍：减少总工作量、减少慢层级数据搬运、提高硬件执行单元利用率、减少调度与同步开销。最后必须在真实 shape 分布上验证，而不是只测一个漂亮的 microbenchmark。

Roofline 的核心上界是：

P_attainable ≤ min(P_peak, BW_sustained × I)

其中 I=FLOPs/Bytes。该模型来自 Roofline 原始论文。

1）算法级：先少算、少存，再谈指令优化

• 使用更低复杂度或更 IO-aware 的算法，例如 FlashAttention 采用在线 softmax 和 tiling，避免把完整 S×S Attention 矩阵写回 HBM；见 FlashAttention 原始论文。
• 利用稀疏性、结构化稀疏、剪枝或 Top-K，避免无效计算；但要计入索引、路由、负载不均和通信成本。
• 推理使用 FP16/BF16/FP8/INT8/INT4 等更低精度，减少字节并使用对应 Tensor Core；量化/反量化最好融合进主 kernel。
• 去掉重复变换和不必要的 layout conversion、padding、copy、全量中间张量。
• 对固定系数、shape 和小矩阵做常量折叠、预计算或 weight packing。

2）HBM 访存：让每个字节更有用

• 合并访存：一个 warp 访问连续、对齐地址，减少内存事务；SoA/AoS、leading dimension 和向量化 load/store 都要围绕实际访问模式设计。
• 提高复用：把会重复使用的 tile 从 Global Memory 搬到 Shared Memory，再从 Shared Memory 搬到寄存器；同一份数据尽量被多个 FMA 使用。
• 处理 bank conflict：Shared Memory 通过 padding、swizzle 或改变 tile layout，使同一 warp 的地址分布到不同 bank。
• 异步流水：使用异步 Global→Shared 拷贝、双缓冲或多 stage pipeline，让加载下一 tile 与计算当前 tile 重叠。
• 缓存友好：调整遍历顺序和工作分配，提高 L2/L1 命中；对反复使用的小工作集可考虑 L2 persistence/access policy，但必须测收益。
• Host↔Device 路径使用 pinned memory 和异步 copy，在有独立 copy engine 时与 kernel 重叠；不要为很小的 copy 支付大量 launch/同步开销。

CUDA C++ Best Practices Guide 对合并访存、Shared Memory、bank conflict、异步 copy 和 Host/Device 传输有官方说明。

3）片上计算：Tiling、Tensor Core 与寄存器

以一个 tile 的 GEMM C[m,n] = A[m,k]B[k,n] 为例，忽略缓存重复加载并假设 C 不需要预读时：

FLOPs = 2mnk
最少字节 ≈ e(mk + kn + mn)
I_tile ≈ 2mnk / (e(mk + kn + mn))

若语义是 C += AB，还要再计一次旧 C 的读取，即分母中的输出项由 mn 变为 2mn。

增大 m,n,k tile 通常能增加复用和算术强度，但也会消耗更多 Shared Memory 与寄存器。因此 tile 不是越大越好：寄存器溢出会落到 Local Memory，Shared Memory 太大则减少每个 SM 的并发 block 数。

常见手段包括：

• 选择匹配 Tensor Core 指令 shape 和数据类型的 tile，保证地址与维度满足对齐要求；
• 使用寄存器 blocking，让 accumulator 长时间留在寄存器；
• 通过 warp specialization，让部分 warp 负责搬运、部分 warp 负责计算；
• double/triple buffering，隐藏 HBM 或 Shared Memory pipeline 延迟；
• 使用向量化指令和合适的 layout，减少地址计算与指令条数；
• 控制 register pressure，避免 spill，同时保留足够 ILP（Instruction-Level Parallelism）。

4）并行划分：既要足够多的工作，也要减少归约代价

• 大矩阵使用二维/三维 tile；小 M/N 的 GEMM 可用 split-K 或 Stream-K 增加 CTA 数，但会引入 partial result 归约。
• Reduction 先 warp 内 shuffle，再 block 内 Shared Memory 归约，最后必要时做第二阶段 kernel；避免每个元素直接争抢一个 Global Atomic。
• 让每个 CTA 的工作量足够大以摊薄 launch 和索引开销，但不能大到只有少数 CTA，造成 SM 吃不满。
• 避免 warp divergence；不可避免的分支可按数据重排，把执行相同路径的线程放在一起。
• 对长序列 Attention 沿 Query、KV sequence 和 head 共同分块，避免只在 head 维提供并行度。

5）融合：省中间张量与 launch

常见融合包括：

• GEMM epilogue 融合 bias、activation、scale、quantize；
• Residual + Dropout + LayerNorm/RMSNorm；
• QKV projection 后的 reshape、RoPE、写 KV Cache；
• Attention 中 scale + mask + online softmax；
• MoE 中 permutation、quantize、通信前 packing 和通信后 unpermute。

假设两个 kernel 之间有一个 N 元素中间张量，单独执行至少多一次 N×e 写和一次 N×e 读；融合理论上可少 2Ne 字节，再省一次 launch。但融合也会增加寄存器、Shared Memory 和代码尺寸，可能降低 occupancy，所以必须检查资源账本。

6）调度与框架开销

• 大量短 kernel 可用 CUDA Graph 降低 CPU launch 开销；
• 使用 stream 做无依赖计算、通信和 memcpy 的重叠，但避免无意义的全局同步；
• 对极小算子可采用 persistent kernel 或批处理，把多请求/多任务合在一次 launch 中；
• 复用 workspace 和内存池，减少频繁 cudaMalloc/cudaFree；
• 多 GPU 场景让通信分块并与计算重叠，优先减少跨 PCIe/NVLink/网络的总字节。

7）正确的优化闭环

1. 记录真实 workload 的 shape、dtype、stride、batch 和长尾分布；
2. 建立 FLOPs、理论必要 bytes、workspace 和精度基线；
3. 用 Nsight Systems 定位 launch/同步/通信，用 Nsight Compute 看 DRAM、L2、stall、occupancy、Tensor Core；
4. 用 Roofline 判断下一步该减 bytes 还是提高计算利用率；
5. 每次只改变一个主要因素，并做数值误差与边界 shape 验证；
6. 与成熟库比较，不仅比较平均 latency，也看 P95/P99、吞吐和端到端收益。

常见追问

追问 1：occupancy 越高越好吗？

不是。occupancy 足以隐藏延迟后，继续牺牲寄存器复用去追求 100% occupancy 可能更慢。高性能 GEMM 常用大量寄存器保存 accumulator，occupancy 不一定很高，但 ILP 和数据复用很好。

追问 2：split-K 什么时候有用？

当 M/N 很小、K 很大，普通二维切分产生的 CTA 数不足以占满 GPU 时。代价是 partial output 的额外写回、归约和同步；K 不够大或原本并行度充足时可能更慢。

追问 3：融合为什么可能负优化？

融合后 live variable 更多，引发寄存器 spill；Shared Memory 增大使 resident CTA 减少；不同子算子的最佳 tile 也可能冲突。要比较省下的 HBM/launch 与新增资源压力。

追问 4：memory-bound 算子怎么优化？

优先减少 bytes、合并访存、提高缓存/片上复用、融合和降精度。单纯增加线程或 FMA 指令通常不能突破 BW×I 的 Roofline 上界。

一句话结论

硬件适配的核心不是按 GPU 型号写一堆 if，而是把 算子形状与算术强度 映射到 该硬件的计算吞吐、内存层级容量/带宽、并发资源和互联拓扑，然后做多版本 kernel dispatch 与真实负载 autotune。

1）先建立硬件画像

至少收集：

• GPU compute capability、SM 数、warp size、每 SM 最大线程/CTA；
• 各 dtype 的可用计算路径，例如 SIMT Core、Tensor Core 以及支持的累加精度；
• HBM/GDDR 容量与实测带宽、L2 容量、每 SM 的 L1/Shared Memory、register file；
• 异步 copy、Tensor Memory Accelerator、warp-group matrix 指令等架构能力；
• PCIe、NVLink/NVSwitch、GPU Direct RDMA 和网卡拓扑；
• CPU NUMA、Pinned Memory、GPU 与 NIC/CPU socket 的亲和性；
• 功率、频率、MIG/MPS 等运行约束。

这些值应通过 runtime/driver API、官方 tuning guide 和实测获得，而不是只抄产品峰值。CUDA 提供 cudaGetDeviceProperties、occupancy API 等查询；架构能力可参考 CUDA Programming Guide、Ampere Tuning Guide 和 Hopper Tuning Guide。

2）用 Ridge Point 决定优化方向

设实际可持续计算吞吐为 P_peak，实测 HBM 带宽为 BW：

I_ridge = P_peak / BW

• I < I_ridge：更可能 memory-bound，应优先减少/复用数据；
• I > I_ridge：更可能 compute-bound，应优先提高 Tensor Core、指令与调度利用率；
• 接近 ridge 时，数据布局、cache 命中和 pipeline overlap 都可能决定最终瓶颈。

例如一块假想 GPU 对某 dtype 的有效吞吐是 1000 TFLOP/s、实测 HBM 是 3 TB/s，ridge point 约为 333 FLOP/Byte。Decode MHA 的 FP16 KV 强度约 1 FLOP/Byte，无论怎样堆 FMA 都很难 compute-bound；应减 KV bytes、增 batch 复用权重或改 Attention 架构。

3）根据片上资源选 tile，而不是照搬另一代 GPU

一个 block 能否驻留，粗略受以下最小值约束：

resident_blocks_per_SM ≤ min(
  架构最大 blocks/SM,
  floor(max_threads_per_SM / threads_per_block),
  floor(registers_per_SM / registers_per_block),
  floor(shared_mem_per_SM / shared_mem_per_block)
)

因此同一个 tile 在不同架构上可能有完全不同结果：

• Shared Memory 增大后可以扩大 tile 或 pipeline stage；
• register file/分配粒度不同，可能让原 kernel 从 2 CTA/SM 降到 1 CTA/SM；
• Tensor Core 指令 shape 和支持 dtype 不同，需要重新安排 layout、warp 数与 accumulator；
• Ampere 上可用异步 Global→Shared copy；Hopper 可进一步用 TMA 和 warp-group 级矩阵计算。应按 capability 选择实现，而不是在不支持的架构上模拟同一条路径。

4）针对不同 shape 做多版本 dispatch

一个“GEMM kernel”至少会因以下维度需要不同策略：

• 大 M/N/K：标准大 tile，追求 Tensor Core 吞吐；
• 小 M、大 N/K（典型 Decode GEMM）：更多 split-K/Stream-K 或 persistent 调度；
• batched 小矩阵：合并 batch、grouped GEMM，摊薄 launch；
• skinny/irregular shape：避免 padding 浪费，调整 warp 划分；
• contiguous 与 strided/broadcast layout：采用不同 load 路径；
• FP16/BF16/FP8/INT8：不同转换、缩放、累加和对齐约束。

实际系统通常采用：预编译一组候选 kernel → 按 GPU capability、shape、dtype、stride 过滤 → 离线或在线 autotune → 缓存最佳配置。线上 autotune 必须限制探测成本并保证候选的数值语义一致。

5）针对内存层级适配数据布局

• HBM 带宽紧张：量化、压缩 KV、融合、避免 transpose/materialize；
• L2 较大：调整 batch/调度，让共享权重、KV block 或专家权重具有时间局部性；
• Shared Memory 足够：扩大复用 tile，但检查 bank conflict 和 occupancy；
• register 紧张：减小 accumulator、拆 epilogue 或减少 pipeline stage，避免 Local Memory spill；
• 支持异步搬运：让 copy 和 compute 形成 steady-state pipeline；
• Host↔Device 较慢：把预处理下沉 GPU、使用 pinned memory 和批量异步传输。

6）根据互联拓扑选择并行策略

• NVLink/NVSwitch 内优先放高频 All-Reduce/All-to-All 的并行组；
• 跨 PCIe 或跨节点通信更贵，尽量减少 Tensor Parallel 的高频小消息；
• MoE Expert Parallel 要同时看总通信量、专家负载不均和 NIC/GPU 亲和性；
• Pipeline Parallel 通信频率较低但有 bubble，应结合 microbatch；
• 通信 kernel 和计算 kernel 使用独立 stream/资源时，要验证是否真正 overlap，不能仅从时间线“有重叠颜色”就下结论。

7）把精度能力纳入硬件适配

• 支持某种低精度指令，不代表模型可直接切换；要定义量化粒度、scale 格式、累加精度和溢出策略；
• TensorFloat-32、FP16、BF16、FP8 的指数/尾数范围不同，要结合数据分布；
• 关键 reduction、normalization、logit 或累加可以保留更高精度；
• 确定性要求可能禁止某些最快的 atomic/reduction 算法，这也是“硬件适配”的约束，不只是正确性开关。

常见追问

追问 1：为什么不能只看标称 TFLOPS？

标称值依赖 dtype、指令模式、频率和理想利用率。真实 kernel 还受 shape、padding、指令混合、数据供给和 occupancy 限制；memory-bound 算子更与 TFLOPS 几乎无关。

追问 2：新 GPU 上旧 kernel 为什么可能更慢？

资源分配粒度、cache/shared 配置、Tensor Core 指令、调度和最优 tile 都可能变化。旧 tile 可能产生 wave quantization、寄存器瓶颈或无法使用新异步路径。

追问 3：如何证明是硬件适配而不是过拟合 benchmark？

在真实 shape 分布、并发和功率设置上报告端到端吞吐与 P95/P99；同时展示候选覆盖率、回退路径、数值测试和多架构结果。只展示一个固定 shape 的峰值不够。

一句话结论

同样输入可能得到 非 bitwise 一致 的输出，但这通常不是数学意义上的随机，而是并行调度、归约顺序或算法选择不固定，再叠加浮点运算不满足结合律。atomicAdd 只保证一次 read-modify-write 不被打断，并不规定不同线程的全局执行顺序；对浮点数，顺序不同会导致舍入结果不同。

CUDA 官方对 atomicAdd 的定义和作用域见 CUDA Programming Guide：Atomic Functions。这里必须区分 API：传统 atomicAdd 一类 legacy atomic 只提供 memory_order_relaxed，默认是 device scope，_block / _system 后缀分别改变作用域；较新的 cuda::atomic、cuda::atomic_ref 与 __nv_atomic_* 才支持显式选择 C++ memory order 与 thread scope。无论哪一种，单次原子 RMW 都不能自动替代完整的同步协议。

为什么浮点顺序会影响结果

IEEE 754 浮点加法每一步都可能舍入，因此：

(a + b) + c  不一定等于  a + (b + c)

例如 FP32 中三个贡献：

a = 1e20
b = -1e20
c = 3.14

(a+b)+c → 0+3.14 → 3.14
a+(b+c) → a+(-1e20) → 0

若数千个线程对同一地址做 atomicAdd(float)，硬件可能以不同合法次序串行化这些原子操作。每个操作都没有“丢写”，但最后几 bit 甚至在病态数据上更大的误差可能变化。

哪些场景会不确定

• Global Atomic 累加、scatter-add、embedding gradient、histogram；
• 并行 reduction 的 block 完成顺序或 reduction tree 改变；
• 同一库在多 stream 下根据 workspace/并发选择不同内部算法；
• 动态 shape、不同 SM 数、不同 toolkit/driver 导致 kernel 或 tiling 改变；
• FMA contraction、不同精度路径、TF32/混合精度造成不同舍入；
• 真正的数据竞争：非原子读写同一位置且缺少 happens-before。这不只是末位误差，而是程序错误，结果可能任意。

cuBLAS Results Reproducibility 说明：在给定 toolkit 版本、相同 GPU 架构和相同 SM 数等条件下，单 stream 路径通常保证 bitwise reproducibility；多并发 stream、toolkit 版本变化或某些模式会破坏该保证，并给出了独立 workspace/handle 等处理方式。

会造成什么问题

• 单测若使用 == 会偶发失败；
• Top-K、argmax 或阈值附近的微小变化可能放大为离散选择变化；
• 训练中一步的微差会经长期迭代放大，使 loss 曲线和最终 checkpoint 不一致；
• 分布式训练中不同 rank 若做出不一致控制流选择，严重时可能死锁；
• Debug、回归定位和故障复现变困难；
• 若使用低精度且累加项动态范围很大，误差可能不止 1 ULP。

如何实现确定性

1. 固定输入、seed、shape、GPU 型号、SM/MIG 配置、toolkit、driver 和库算法；
2. 使用固定的 reduction tree，例如 block 内固定树归约，再用第二个 kernel 按固定顺序归约 block partials；
3. 避免多个 block 直接对同一个 FP32 输出做无序 atomic，或先按 key 排序/分组；
4. 为每个 stream 提供独立 workspace/handle，使用库的 deterministic 配置；
5. 必要时提高累加精度、使用 pairwise/Kahan 类求和，降低误差，但注意“误差更小”不等于“执行顺序固定”；
6. 测试使用经过误差预算的 atol/rtol；若业务必须 bitwise 一致，则单独维护 deterministic path 并接受性能成本。

必须区分的三个概念

• Atomicity：一次 RMW 不被别的线程拆开；
• Ordering / synchronization：不同操作先后关系和可见性；
• Numerical reproducibility：多次运行是否得到相同 bit pattern。

atomicAdd 解决第一项，不自动解决后两项。

常见追问

追问 1：整数 atomicAdd 也会不确定吗？

如果只有加法、没有溢出或语义定义为模加，最终和通常与顺序无关；但返回给各线程的旧值仍取决于顺序。浮点加法因不满足结合律，最终值也可能变化。

追问 2：加一个 __syncthreads() 能让 Global Atomic 顺序固定吗？

只能同步同一 block 的线程，并不能规定不同 block 的执行顺序；即便分阶段同步，同一阶段内原子操作的顺序仍未必固定。

追问 3：确定性和跨平台可复现是一回事吗？

不是。一个 kernel 在同卡同版本上可以 deterministic，但换架构、编译器或库版本后因算法/FMA/近似函数变化而不同。跨平台 bitwise reproducibility 的约束更强。

一句话结论

不会溢出。 在约 9.9×10^9 的数量级，FP32 相邻数间距（1 ULP）是 1024，而 10^-9 只有约 9.77×10^-13 ULP，所以减法中的小数会被完全“吞掉”，结果与减法前的大数 bitwise 相同。

但“减法前的大数是多少”取决于题目要求怎样进行 FP32 求值：

• 若写成一个 FP32 字面量 9.9e9f，结果是 9,900,000,256；减去 1e-9f 后仍是该值；
• 若严格执行两次 FP32 运算 9.9f * 1e9f，乘积先舍入为 9,899,999,232；再减 1e-9f 后仍是该值。

面试时应先说明这个求值路径差异，再给主答案。

为什么 ULP 是 1024

FP32（binary32）有效精度为 24 个二进制位（含隐藏位）。9.9×10^9 位于 binade [2^33, 2^34) 中，因此该区间的间距是：

ULP = 2^(33-(24-1)) = 2^(33-23) = 1024

round-to-nearest 下，只有变化超过约半个 ULP，也就是 512，才可能舍入到相邻 FP32。这里：

1e-9 / 1024 = 9.765625e-13 ULP

远小于半个 ULP，所以减法结果不变。

本机 C 程序实测

为避免编译器只做常量折叠，测试使用 volatile float，并打印位型：

#include <math.h>
#include <stdint.h>
#include <stdio.h>
#include <string.h>

static uint32_t bits(float x) {
    uint32_t u;
    memcpy(&u, &x, sizeof u);
    return u;
}

int main(void) {
    volatile float small = 1.0e-9f;

    volatile float literal = 9.9e9f;
    volatile float r1 = literal - small;

    volatile float a = 9.9f;
    volatile float billion = 1.0e9f;
    volatile float product = a * billion;
    volatile float r2 = product - small;

    printf("literal %.0f 0x%08x, result %.0f 0x%08x\n",
           literal, bits(literal), r1, bits(r1));
    printf("product %.0f 0x%08x, result %.0f 0x%08x\n",
           product, bits(product), r2, bits(r2));
    printf("ULP(literal)=%.0f\n",
           nextafterf(literal, INFINITY) - literal);
}

实测输出：

9.9e9f                = 9,900,000,256  bits 0x50138581
9.9e9f - 1e-9f        = 9,900,000,256  bits 0x50138581

9.9f * 1e9f           = 9,899,999,232  bits 0x50138580
(9.9f*1e9f) - 1e-9f   = 9,899,999,232  bits 0x50138580

ULP                   = 1024
结果与减法前相差      = 0 ULP

差异来自 9.9f 本身先舍入为约 9.8999996185，再参与 FP32 乘法；而 9.9e9f 是把精确十进制常量 9.9×10^9 直接舍入一次到 binary32。两者恰好落在相邻的两个 FP32 数上。

是否发生 overflow / underflow

• 没有 overflow：FP32 最大有限值约为 3.402823466×10^38，9.9×10^9 离它非常远；
• 1e-9f 本身也没有 underflow：它远大于 FP32 最小正规数约 1.17549435×10^-38；
• 小数消失属于大数与小数相加减时的 rounding absorption，不是 overflow，也不是 underflow。

IEEE binary32 格式和舍入规则的权威定义见 IEEE 754-2019。

C/C++ 表达式陷阱

题面中的 10^9 是数学排版。如果原样写进 C/C++，^ 表示按位异或而不是乘方，并且与浮点操作数组合会编译失败。代码中应写 1e9f、powf(10.0f, 9.0f) 或明确的乘法；本题最自然的是 9.9e9f - 1.0e-9f。

常见追问

追问 1：为什么不是“精确的 9,900,000,000”？

因为该数量级的 FP32 只能表示 1024 的整数倍，9,900,000,000 不是可表示点；它会舍入到相邻的某个倍数。

追问 2：用 FP64 呢？

在约 9.9×10^9 处，binary64 的 ULP 约为 2^(33-52)≈1.907×10^-6，1e-9 仍小于半 ULP，所以即便 FP64，减法结果也通常与大数相同；要保留这个差值需要更高精度或改变数值表达方式。

追问 3：这是 catastrophic cancellation 吗？

不是典型的“大致相等两数相减导致有效位丢失”。这里是数量级差太大，小操作数在对齐尾数和舍入时被吸收。

一句话结论

CUDA 中要区分 编程可见的 memory space 与 物理存储/缓存层级。Global、Local、Shared、Constant、Texture 是地址空间或编程抽象；HBM/GDDR、register file、片上 SRAM 是物理实现；L1/L2 是缓存，不是需要 malloc 的独立地址空间；Unified Memory 是虚拟内存与迁移机制，也不是另一块物理显存。

CUDA 官方层级和作用域见 CUDA Programming Guide：GPU Memory 与 Writing CUDA SIMT Kernels：Memory Spaces。

主要 memory space

Cache 层级

• L1 / Unified Data Cache：每个 SM 独有，很多架构上与 Shared Memory 共享物理容量或可配置比例；
• L2：全 GPU 的 SM 共享，是 Global/Local 等访问的重要缓存，也是不同 SM 之间数据交汇的层级；
• Constant Cache：每 SM 独立，服务 constant/参数访问；
• Texture Cache、Instruction Cache：服务对应访问类型。

Cache 一般由硬件管理，Shared Memory 由程序显式分配和组织。说“把数据放到 L2”通常只是通过访问策略提高驻留概率，而不像 Shared Memory 那样精确控制每个字节。

Blackwell 的 Tensor Memory 不是通用“第六级内存”

PTX ISA 对第五代 Tensor Core 的说明新增了专用片上 Tensor Memory（TMEM）。使用 .cta_group::1 时，由当前 CTA 的一个 warp 执行 tcgen05.alloc / tcgen05.dealloc；使用 .cta_group::2 的相关指令时，则由一对 peer CTA 各出一个 warp 共同分配或释放，并可访问这两个 CTA 的 TMEM。所有已分配 TMEM 都必须在 kernel 退出前显式释放。它与 register file、Shared Memory 都不同，也不能外推到 H100 等不具备该指令族的 GPU；更准确的说法是“Blackwell 上服务第五代 Tensor Core、具有单 CTA 与 CTA-pair 粒度的架构特定存储”。

三个相似名字不能混用：

• Tensor Memory（TMEM） 是上面的片上存储；
• Tensor Memory Accelerator（TMA） 是搬运多维 tensor 的异步数据传输引擎，不是一块存储；
• 截至 CUDA 13.3，CUDA C++ 的 __tile__ 变量 被分配在 device global memory；这里说的是变量的 memory-space specifier，不是同名的函数 execution-space specifier，也不能把它解释成 TMEM。

精确定义见 PTX ISA：Fifth Generation TensorCore / Tensor Memory。

容易混淆的两类内存

Local Memory 为什么可能很慢

局部数组虽然写在 kernel 函数体内，但如果无法放进寄存器，会进入 Local Memory。其物理 backing 在 Device Memory，延迟/带宽特征接近 Global Memory，只是 CUDA 按线程 ID 布局，某些相同相对偏移访问可合并。

可通过编译器的 register/local 使用报告和 profiler 的 local load/store 指标检查 spill。

Unified Memory 是什么

cudaMallocManaged 或系统分配场景提供 CPU/GPU 可访问的统一内存映像，由 runtime/驱动按需迁移或建立访问。它解决可编程性、过量订阅和一致地址空间问题，但物理页某一时刻仍位于 GPU Memory 或 System Memory；Page Fault 和迁移可能成为性能瓶颈。

Host 侧相关内存

• Pageable Host Memory：普通 CPU 内存，异步传输通常需要 runtime 先做 staging；
• Pinned Host Memory：页锁定内存，可获得更高、更稳定的 DMA 带宽，并支持真正的异步 H2D/D2H；
• Mapped / Zero-copy Host Memory：GPU 可直接访问 System Memory，省显式 copy，但经 PCIe/互联访问，适合少量、一次性或集成式系统场景，不适合反复随机访问替代 HBM。

常见追问

追问 1：速度排序是不是 Register > Shared > L1 > L2 > HBM？

可作为粗略直觉，但不能当严格延迟表。命中、bank conflict、合并程度、并发、访问依赖和架构都会改变结果。Shared 的 bank conflict 访问甚至可能被严重串行化。

追问 2：const 变量一定在 Constant Memory 吗？

不一定。C++ 的 const 是类型/可修改性限定，不是 CUDA 地址空间限定；只有显式 __constant__、kernel parameter path 或编译器判定才可能走 constant memory/cache。

追问 3：L1 和 Shared 谁更快？

不能脱离模式比较。Shared 可显式复用、可预测，但需避免 bank conflict；L1 自动缓存，hit 时方便但驻留不可精确保证。高性能 kernel 常两者一起使用。

一句话结论

GM 有虚拟地址。 cudaMalloc 返回的是当前进程 CUDA 地址空间中的 Device Virtual Address，GPU MMU 再把它映射到物理显存。两个进程默认拥有独立 CUDA context/地址空间，不能把进程 A 的裸 device pointer 直接交给进程 B；但可以通过 CUDA IPC 或 VMM 导出/导入同一个物理 allocation，各进程获得自己的有效 device pointer 后并发访问。

CUDA 官方明确说明普通 device pointer 在创建它的进程之外无效；跨进程需要 IPC/VMM handle，见 CUDA Programming Guide：Interprocess Communication。

“同一个地址”要拆成三层

1. 同一个数值形式的 pointer：两个进程里都看到 0x...，不代表同一物理页；进程页表不同；
2. 同一个物理 GM allocation：通过 IPC/VMM share handle 映射，可以做到；
3. 同一个逻辑对象且并发正确：除了共享 allocation，还必须有生命周期、访问权限和同步协议。

Legacy CUDA IPC 的典型流程

进程 A：cudaMalloc
      → cudaIpcGetMemHandle
      → 通过 Unix socket / shared memory 把 handle 传给进程 B

进程 B：cudaIpcOpenMemHandle
      → 得到本进程内有效的 device pointer
      → kernel 访问同一底层 allocation

结束：B cudaIpcCloseMemHandle
     → A 确认无人使用后 cudaFree

不同进程对同一 IPC handle 调用 open，不保证得到数值相同的虚拟地址。程序不应把 pointer 数值相等作为共享成功的判断条件。

Legacy CUDA IPC 主要用于 Linux；Managed Memory 不能直接用这套 legacy IPC handle 分享。具体平台和设备限制以运行时文档为准。

还要注意 allocation 边界：cudaMalloc() 可能从更大的底层内存块做子分配，导出 IPC handle 时有机会把同一底层块中的其他子分配一并暴露给接收进程。CUDA 官方因此建议只共享大小按 2 MiB 对齐的 allocation；安全敏感场景更适合用 VMM 显式创建、授权和映射可共享 allocation。这是跨租户隔离问题，不只是性能细节。

VMM 如何说明“GM 有虚拟地址”

CUDA Driver VMM API 把过程显式拆为：

cuMemAddressReserve   # 只保留一段虚拟地址
cuMemCreate           # 创建物理 allocation handle，尚未映射
cuMemMap              # 把物理 allocation 映射到 VA
cuMemSetAccess        # 为设备设置访问权限

跨进程时，导出方用 cuMemExportToShareableHandle，接收方用 cuMemImportFromShareableHandle，再把同一 allocation 映射进自己的 VA 范围。官方流程见 CUDA Programming Guide：Virtual Memory Management 与 CUDA Driver API：Virtual Memory Management。

VMM 允许每个进程选择自己的 VA；即便刻意尝试保留相同数值的 VA，也不应让协议依赖这一点。

两个进程能否“同时写”

可以同时发 kernel 访问同一 allocation，但正确性分情况：

• 写不同、不重叠的区域：通常可以并行；
• 一个写、一个读同一区域：需要跨进程事件/外部 semaphore/host IPC 建立先后和可见性；
• 同时写同一位置：普通 load/store 是数据竞争；需要支持相应 scope 的 atomic 或更高层同步，并验证目标设备和 memory type 的原子语义；
• 普通 CUDA stream 顺序只约束本 stream，不能自动约束另一个进程的 stream。

Legacy IPC 可以共享以 cudaEventInterprocess 创建的 event handle，接收进程可在 stream 中等待；VMM 只解决 memory mapping，本身不提供完整生产者/消费者同步。

UVA、Unified Memory 和 IPC 的区别

• UVA（Unified Virtual Addressing）：在一个进程的统一虚拟地址空间里区分/管理 Host 与多个 GPU allocation；
• Unified Memory：提供统一内存映像、按需迁移/映射和一致性机制；
• IPC/VMM sharing：把 allocation capability 安全地传给另一个进程，并在对方地址空间建立映射。

“有 UVA”不等于另一个进程能使用我的裸 pointer。

常见追问

追问 1：MPS 会让两个进程天然共享 pointer 吗？

不会。MPS 主要共享 GPU 调度/执行服务，不是把所有客户端地址空间和 allocation 自动变成共享内存。显式共享仍应使用受支持的 IPC/VMM 机制。

追问 2：两个进程映射同一 allocation 后，L2 是否一致？

GPU 对受支持的共享映射提供其内存模型定义的可见性，但必须通过正确同步建立 happens-before；不能把“共享 L2”当成同步原语。

追问 3：为什么不能只传 64 位地址？

地址只在创建它的进程/context 映射中有意义，且没有携带 allocation 权限与生命周期。IPC/VMM handle 才是可移植的共享能力描述。

一句话结论

如果题目指 kernel launch 时传入的标量/小结构参数，它们位于 kernel parameter space，编译器通常通过 constant-memory/constant-cache 路径访问；如果是显式声明的全局只读常量，应使用 __constant__。但普通 C++ const 并不决定物理位置：字面量可能编码进指令，局部常量可能进寄存器，const T* 指向的数据仍可能在 Global Memory。

CUDA Programming Guide：GPU Memory 明确说明编译器可把 kernel parameters 放入 constant memory；__constant__ Memory 描述了其作用域、生命周期和典型容量。

按常量类型分类回答

1）Kernel 参数

__global__ void scale_kernel(float* y, const float* x,
                             int n, float alpha) {
    // n、alpha 以及两个 pointer 值属于 kernel 参数
}

• n、alpha、x 和 y 的 pointer 值本身 由 launch 参数传入，通常经 parameter/constant cache 路径读取；
• x 指向的数组仍在它实际分配的 memory 中，通常是 Global Memory；const float* 只表示 kernel 不应通过该指针修改数据；
• 较大的只读按值参数可使用 CUDA 的 __grid_constant__ 注解，避免编译器为每线程创建副本，具体语义见 CUDA __grid_constant__ Parameters。

2）显式 Device 常量

__constant__ float coeff[64];

// host 侧初始化
cudaMemcpyToSymbol(coeff, host_coeff, sizeof(host_coeff));

它具有 CUDA context 生命周期，每个 device 有自己的对象，kernel 中只读。适合所有线程反复读取相同或少量地址的小表、scale、卷积系数等。

constant cache 的最佳模式是一个 warp 的线程读取同一地址，此时可广播；若 32 个线程读取许多不同 constant 地址，请求会按不同地址拆分，可能比合并的 Global/L2 读取更差。

3）编译期字面量和局部常量

constexpr int TILE = 128;
const float inv_sqrt_d = 0.0883883476f;

编译器可能：

• 把数编码为 instruction immediate；
• 在使用前 materialize 到寄存器；
• 常量折叠后完全消失；
• 若局部对象较大或寄存器压力高，放到 Local Memory。

不能仅凭源码中的 const 推断它在 Constant Memory。

4）大型只读权重或查找表

Constant Memory 容量小，不适合模型大权重。大型只读张量通常仍放 Global Memory，依赖 L2/L1/read-only/texture 路径和 tiling 复用；若每线程索引不同，Global Memory 的合并访问往往比 constant cache 的分歧访问更合适。

常见追问

追问 1：Constant Memory 是片上的吗？

逻辑 constant memory 驻留在 device memory，SM 上有较小的 constant cache。命中时很快，但不能把整个典型 64 KiB constant address space 都理解成每个 SM 各有一份片上存储。

追问 2：为什么同地址广播快？

一个 warp 对同一 constant 地址的访问可以由一次 cache 请求广播给所有线程；访问 N 个不同地址通常需要拆为 N 组请求。

追问 3：kernel 参数过多怎么办？

可以打包小结构，但仍要关注 parameter limit 与访问方式；大型数据应传 pointer，不应把整个 tensor 按值作为参数。pointer 在 parameter space，pointee 在 Global/Shared 等实际地址空间。

一句话结论

如果“单核”指一个支持 SMT/Hyper-Threading 的物理 CPU Core，两个硬件线程在 单线程无法填满执行资源、经常因 cache miss、依赖或分支停顿 时，可以用另一个线程的就绪指令填补空槽，从而提高聚合吞吐、缩短完成两份工作的总时间。若单线程已经饱和关键执行端口、内存带宽或 cache，两线程通常无收益，甚至因资源竞争更慢。

SMT 的原始机制是让多个线程在同一周期向 superscalar core 的多个执行单元发射指令，见 Tullsen、Eggers、Levy 的 Simultaneous Multithreading: Maximizing On-Chip Parallelism（ISCA 1995）。

先明确“更快”比较的是什么

• 两个独立任务各跑一遍：比较一个线程串行完成两份任务，与两个 SMT 线程并行完成两份任务的 wall time；
• 一个任务拆成两个线程：除了 SMT 是否能填泡，还要求任务可并行，并计入拆分、同步和 cache coherence；
• 单请求 latency：SMT 开第二线程后，单个请求可能反而变慢，但单位时间完成的请求总数变多；
• 若 CPU 没有 SMT，同一时刻物理核只能执行一个线程，两个 runnable software thread 只会时间片切换，一般不会让纯计算更快。

两线程可能更快的条件

1）单线程存在长延迟 stall

• LLC/DRAM miss、TLB miss，且单线程自身 memory-level parallelism 不足；
• 长指令依赖链，后续指令无法发射；
• 分支错误预测造成 pipeline bubble；
• 某类高延迟指令使部分执行资源空闲。

当线程 A 等数据时，线程 B 可以发射就绪指令，隐藏部分延迟。

2）两个线程使用互补的执行资源

例如线程 A 偏整数/分支，线程 B 偏向量或 load/store，且没有共同打满同一个 bottleneck port。SMT 可以提高 issue width、execution port、ROB/queue 等资源的整体利用率。

3）还有带宽和缓存余量

两个线程增加 outstanding memory request，可能提高 memory-level parallelism；前提是内存带宽尚未饱和，两个 working set 不会把彼此挤出 cache，也没有严重 TLB/cache contention。

4）一个线程会阻塞或睡眠

若任务包含 I/O、锁等待或 page fault，另一个软件线程可在同一物理核被调度运行。即使没有 SMT，这也能减少物理核空闲时间；但这属于调度隐藏阻塞，不是两个纯计算线程同时执行。

一个简单的吞吐上界直觉

设单线程只利用了关键 issue/执行资源的比例 u，忽略额外争用时，两个相同线程的理想加速上界可粗略写为：

Speedup ≤ min(2, 1/u)

• 若 u≈0.5，理论上有接近 2× 的填泡空间；
• 若 u≈0.9，即使其他条件理想，上界也只有约 1.11×；
• 实际还要扣除共享 cache、队列、带宽、同步和频率下降，因此通常低于该上界。

这不是严格性能模型，但适合面试中说明为什么“一个核有两个逻辑线程”不等于“性能翻倍”。

两线程通常更慢的情况

• 单线程已经打满向量/FMA、load/store port 或 retirement bandwidth；
• 已经达到内存带宽上限，第二线程只增加排队；
• 两个 working set 争抢 L1/L2/TLB，cache miss 增多；
• false sharing、锁竞争、barrier 和原子操作很多；
• 两线程导致每线程可用 ROB、queue 或其他共享资源减少；
• 功率/温度限制使 all-thread 频率下降；
• 任务太短，线程创建、调度和同步开销大于节省时间。

如何验证

1. 把两个线程 pin 到同一个物理核的两个 SMT sibling，而不是两个不同物理核；
2. 分别测单线程 latency、双线程单任务 latency、aggregate throughput；
3. 观察 IPC、执行端口利用率、cache/TLB miss、memory bandwidth、上下文切换和频率；
4. 与“两个物理核各一个线程”对照，避免把 SMT 收益误认为多核收益；
5. 使用足够长且稳定的 steady-state workload，排除线程创建和 turbo 波动。

常见追问

追问 1：GPU 的一个 CUDA Core 上两个线程也是这个逻辑吗？

GPU 更常以 warp 为调度单位，通过大量 resident warp 隐藏延迟；它与 CPU SMT 的共享 superscalar core 概念相似但不等同。回答前应先确认面试官说的是 CPU 物理核、GPU SM，还是 CUDA Core。

追问 2：为什么两个线程的单任务 latency 变长但总耗时变短？

两个线程共享资源，每个任务可能获得更少执行带宽，所以 individual latency 上升；但原本空闲的执行槽被利用，aggregate throughput 上升，完成两份任务的总 wall time下降。

追问 3：memory-bound 就一定适合 SMT 吗？

不是。若是 latency-bound 且带宽有余量，SMT 可用更多 outstanding request 隐藏延迟；若已经 bandwidth-bound，第二线程通常无益。

结论：没有区别。 两者都是“指向只读 int 的指针”，const 修饰的是 int，不是指针本身。

const int* A = nullptr;
int const* B = nullptr;

它们都满足：

• 可以改变指针，让它指向另一个 int；
• 不能通过该指针修改所指对象；
• 所指对象不一定真的是常量，只是经由这个指针不能修改。

下面的代码可以编译：

#include <iostream>

int main() {
    int x = 1;
    int y = 2;

    const int* A = &x;
    int const* B = &y;

    A = &y;       // 可以修改 A 自身
    B = &x;       // 可以修改 B 自身
    // *A = 3;    // 编译错误：不能通过 A 修改所指对象

    x = 3;        // 可以直接修改 x，因为 x 本身并不是 const
    std::cout << *A << ' ' << *B << '\n'; // 2 3
}

容易混淆的是下面两种声明：

int* const p = nullptr;        // p 是常量指针：p 不能改，*p 可以改
const int* const q = nullptr;  // 指针和所指内容经由 q 都不能改

读复杂声明时可以从变量名向外读：int * const p 中，紧挨 p 的是 const，所以 p 自身是常量；const int *p 中，const 修饰的是 int。

结论：交换的是两个指针保存的地址，不是修改常量 1 和 2。 因此应把 A、B 声明为“可修改的、指向常量的指针”，然后使用 std::swap。

注意，不能对字面量 1、2 直接取地址。应先创建两个有存储位置的常量对象：

#include <cassert>
#include <utility>

int main() {
    const int one = 1;
    const int two = 2;

    const int* A = &one;
    const int* B = &two;

    std::swap(A, B);

    assert(*A == 2);
    assert(*B == 1);
    assert(one == 1 && two == 2); // 两个常量对象完全没有被修改
}

交换前后关系是：

交换前：A -> one(1)，B -> two(2)
交换后：A -> two(2)，B -> one(1)

边界与坑：

• 如果声明成 const int* const A = &one;，第二个 const 使 A 自身不可修改，此时不能交换。
• 用 const_cast 修改一个原本就被定义为 const 的对象是未定义行为，不是本题的解法。
• 函数参数如果写成 void swap(const int* a, const int* b)，参数是按值传入，函数内交换副本不会改变调用者；应使用引用 const int*& a、const int*& b，或直接在调用处使用 std::swap。

结论：拷贝构造函数用一个已有的同类型对象，初始化一个新的对象。 最典型的签名是：

T(const T& other);

“初始化新对象”和“给已有对象赋值”必须区分：

T a;
T b = a; // 拷贝构造：b 此前不存在

T c;
c = a;   // 拷贝赋值：c 已经存在，调用 operator=

一个可编译的例子：

#include <iostream>
#include <string>
#include <utility>

class Record {
public:
    explicit Record(std::string text) : text_(std::move(text)) {}

    Record(const Record& other) : text_(other.text_) {
        std::cout << "copy constructor\n";
    }

    Record& operator=(const Record& other) {
        std::cout << "copy assignment\n";
        text_ = other.text_;
        return *this;
    }

    const std::string& text() const { return text_; }

private:
    std::string text_;
};

int main() {
    Record a("AI Infra");
    Record b = a; // 拷贝构造

    Record c("old");
    c = a;        // 拷贝赋值

    std::cout << b.text() << ' ' << c.text() << '\n';
}

什么时候会调用？

常见场景包括：

• 用同类型左值初始化新对象；
• 按值传参、按值返回时，语义上需要复制对象；
• 容器扩容或插入时需要复制元素。

但从 C++17 开始，一些按值返回和纯右值初始化场景会发生强制拷贝消除，此时即使语义上像“返回一个对象”，拷贝构造函数也可能完全不执行。不要用“打印是否出现”来判断程序是否符合值语义。

编译器会不会自动生成？

如果用户没有声明某些特殊成员函数，编译器通常可以隐式声明拷贝构造函数，并按成员逐个复制。是否生成、是否被定义为 delete，会受成员类型和用户声明的移动构造、析构函数等影响。

工程上常用两条规则：

• Rule of Zero：优先让成员使用 std::string、std::vector、智能指针等 RAII 类型，从而不手写拷贝、移动和析构。
• 如果类直接管理裸资源，通常需要同时考虑析构、拷贝构造、拷贝赋值、移动构造、移动赋值，即 Rule of Five。

常见错误是写成 T(T other)：按值传入 other 本身就需要先拷贝，导致递归依赖；标准形式应是 T(const T&)。

结论：浅拷贝只复制资源句柄或地址，两个对象可能指向同一份资源；深拷贝会创建一份独立资源，并复制资源内容。

假设类中保存一个裸指针：

int* data_;

编译器默认的成员复制只会复制地址。复制后两个对象的 data_ 相同，会产生两个典型问题：

1. 修改一方会意外影响另一方；
2. 两个析构函数对同一地址执行 delete[]，造成 double free 和未定义行为。

下面给出一个完整的深拷贝实现。它不仅有拷贝构造，还实现了拷贝赋值和移动操作，使资源所有权完整：

#include <algorithm>
#include <cstddef>
#include <iostream>
#include <memory>
#include <utility>

class Buffer {
public:
    explicit Buffer(std::size_t size)
        : size_(size), data_(size ? new int[size]{} : nullptr) {}

    ~Buffer() {
        delete[] data_;
    }

    // 深拷贝构造：申请新空间，再复制内容。
    Buffer(const Buffer& other)
        : size_(other.size_),
          data_(other.size_ ? new int[other.size_] : nullptr) {
        if (size_ != 0) {
            std::copy_n(other.data_, size_, data_);
        }
    }

    // 拷贝赋值：copy-and-swap，兼顾自赋值和异常安全。
    Buffer& operator=(const Buffer& other) {
        if (this == &other) {
            return *this;
        }
        Buffer tmp(other);
        swap(tmp);
        return *this;
    }

    Buffer(Buffer&& other) noexcept
        : size_(std::exchange(other.size_, 0)),
          data_(std::exchange(other.data_, nullptr)) {}

    Buffer& operator=(Buffer&& other) noexcept {
        if (this != &other) {
            delete[] data_;
            size_ = std::exchange(other.size_, 0);
            data_ = std::exchange(other.data_, nullptr);
        }
        return *this;
    }

    void swap(Buffer& other) noexcept {
        std::swap(size_, other.size_);
        std::swap(data_, other.data_);
    }

    int& operator[](std::size_t i) { return data_[i]; }
    const int& operator[](std::size_t i) const { return data_[i]; }
    std::size_t size() const { return size_; }

private:
    std::size_t size_ = 0;
    int* data_ = nullptr;
};

int main() {
    Buffer a(2);
    a[0] = 10;
    a[1] = 20;

    Buffer b = a;
    b[0] = 99;

    std::cout << a[0] << ' ' << b[0] << '\n'; // 10 99
}

实际工程中更推荐直接写：

std::vector<int> data_;

std::vector 已经正确实现深拷贝、移动和析构，类可以遵循 Rule of Zero。深拷贝并不意味着所有指针成员都必须复制目标对象；例如观察者指针、共享资源和 std::shared_ptr 的语义本来就可能是共享，关键是先定义清楚所有权。

结论：sizeof 计算的是对象本身占用的字节数，包括非静态数据成员、基类子对象、对齐填充，以及实现为支持多态或虚继承加入的隐藏状态；它不包含成员函数代码，也不把整张虚函数表算进每个对象。

标准能保证什么？

• 非静态数据成员属于对象存储的一部分；静态数据成员不属于。
• 成员函数和虚函数的机器码不存放在每个对象中。
• 派生类对象包含其非虚基类子对象；虚基类在最派生对象中通常只有一份。
• 对象要满足对齐要求，因此成员之间和对象尾部可能有 padding。
• 空类的 sizeof 至少为 1，以保证不同完整对象可以有不同地址；但空基类可能触发 Empty Base Optimization，不额外占空间。

vptr / vtable 的常见实现

C++ 标准没有规定必须存在 vptr 或 vtable。在主流 Itanium ABI、MSVC ABI 中，含虚函数的多态类对象通常包含至少一个隐藏的虚表指针 vptr：

• 第一次让一个继承分支变成多态，通常会让对象多出一个指针大小，再叠加对齐；
• 再增加若干虚函数，通常不会继续增大每个对象，因为它们只是在类共享的 vtable 中增加表项；
• 多重继承中，如果存在多个需要独立多态分派的基类子对象，派生对象可能包含多个 vptr；
• 虚继承可能需要额外的指针或表项来定位虚基类，具体大小由 ABI 决定。

“虚函数表很大，所以每个对象也很大”是错误说法。vtable 通常由同一动态类型的所有对象共享，每个对象主要承担的是一个或多个指针的成本。

sizeof 看静态类型

sizeof 通常在编译期由表达式的静态类型决定，不会根据指针或引用所指对象的动态类型变化：

Base* p = new Derived;
sizeof(*p); // 等于 sizeof(Base)，不是 sizeof(Derived)

下面的程序可在目标编译器上观察实际 ABI 的布局结果；不要把某台机器输出的具体数字当成 C++ 标准保证：

#include <iostream>

struct Empty {};

struct Plain {
    int x;
};

struct Poly {
    virtual ~Poly() = default;
    virtual void f() {}
    int x;
};

struct MoreVirtual : Poly {
    void f() override {}
    virtual void g() {}
    // 新增虚函数通常只增加共享虚表的表项。
};

struct Left {
    virtual ~Left() = default;
    int l;
};

struct Right {
    virtual ~Right() = default;
    int r;
};

struct Multiple : Left, Right {
    int m;
};

struct VirtualBase {
    int v;
};

struct VirtualDerived : virtual VirtualBase {
    int d;
};

struct Ebo : Empty {
    int x;
};

int main() {
    std::cout << "Empty: " << sizeof(Empty) << '\n';
    std::cout << "Plain: " << sizeof(Plain) << '\n';
    std::cout << "Poly: " << sizeof(Poly) << '\n';
    std::cout << "MoreVirtual: " << sizeof(MoreVirtual) << '\n';
    std::cout << "Multiple: " << sizeof(Multiple) << '\n';
    std::cout << "VirtualDerived: " << sizeof(VirtualDerived) << '\n';
    std::cout << "Ebo: " << sizeof(Ebo) << '\n';
}

面试时最好按下面顺序分析一个具体类：

1. 展开所有基类子对象和非静态成员；
2. 判断常见 ABI 下需要几个 vptr / 虚基定位信息；
3. 按每个成员的对齐要求计算内部 padding；
4. 最后把总大小补齐到类的最大对齐倍数；
5. 强调最终精确值依赖 ABI、编译器、架构和编译选项。

结论：两种方法都能在特定前提下工作，但生产代码应使用 std::swap。 “不用额外变量”不代表更快；编译器可以把临时变量放在寄存器里。加减法与异或法都有同址别名陷阱，加减法还多一项有符号溢出风险。

加减法

void swapByArithmetic(int& a, int& b) {
    if (&a == &b) { // 两个引用绑定同一对象时无需交换
        return;
    }
    a = a + b;
    b = a - b;
    a = a - b;
}

代数上，执行过程为：

a = old_a + old_b
b = (old_a + old_b) - old_b = old_a
a = (old_a + old_b) - old_a = old_b

但在 C++ 中，有符号整数溢出属于未定义行为。即使 old_a、old_b 都是合法 int，old_a + old_b 仍可能溢出。它也不能省略同址检查：若 a、b 都绑定到 x，第一句把 x 改成 2x，第二句读取的两个操作数已经是同一个新值，结果会把 x 置为 0。因此这不是一个对所有引用和所有 int 都正确的通用函数。

异或法

void swapByXor(int& a, int& b) {
    if (&a == &b) { // 必须处理两个引用绑定同一对象
        return;
    }
    a ^= b;
    b ^= a;
    a ^= b;
}

因为 x ^ x = 0、x ^ 0 = x，所以在 a、b 是不同对象时可以完成交换。若不检查 &a == &b，调用 swapByXor(x, x) 的第一步就会把 x 清零。

推荐写法

#include <utility>

void swapInt(int& a, int& b) {
    std::swap(a, b);
}

int main() {
    int a = 1;
    int b = 2;
    swapInt(a, b);
}

异或交换还会建立串行数据依赖，常常不如普通交换利于优化。面试中可以先写出题目要求的两种方法，再主动指出：两者都要处理同址引用，加减法还要排除有符号溢出，工程上另有可读性与优化问题。

结论：令方向向量 $\mathbf d=B-A$，则 C 在直线 AB 上的正交投影为

推导

直线 AB 上任意一点都可写成：

投影点 $C’$ 满足误差向量 $C-C’$ 与直线方向 $\mathbf d$ 垂直：

代入 $C’=A+t\mathbf d$：

于是：

将 $t$ 代回即可得到开头公式。这个公式适用于二维、三维乃至任意欧氏空间。

如果问的是线段 AB 上的最近点

直线投影参数 $t$ 可能小于 0 或大于 1。若目标是线段而不是无限延伸的直线，需要截断：

• $0\le t\le1$：投影在线段内部；
• $t<0$：最近点是 A；
• $t>1$：最近点是 B。

边界条件是 $A=B$。此时 $|B-A|^2=0$，直线方向不存在；若把问题理解为“到退化线段的最近点”，答案就是 A（也等于 B）。浮点实现中应判断分母是否小于一个容差，而不是机械地和 0 精确比较。

结论：Sigmoid 把任意实数压缩到 $(0,1)$，定义为

它的导数可以复用函数值：

推导

把 $\sigma(x)=(1+e^{-x})^{-1}$，使用链式法则：

另一方面：

所以：

面试常见追问

• $\sigma(0)=0.5$。
• 导数最大值在 $x=0$，为 $1/4$。
• 当 $|x|$ 很大时，输出接近 0 或 1，导数接近 0，容易出现梯度消失。
• 直接计算很大的负数输入时，exp(-x) 可能溢出。稳定实现可分支计算：

在二分类损失中，通常直接使用框架提供的 BCEWithLogitsLoss 一类“logits + loss”融合实现，避免先算 Sigmoid 再取对数造成的数值精度问题。