Fleet: Hierarchical Task-based Abstraction for Megakernels on Multi-Die GPUs #

Sangeeta Chowdhary, Ryan Swann, Sean Siddens, Muhammad Osama, Stephen Neuendorffer, Alexandru Dutu, Karthik Sangaiah, Sandeepa Bhuyan, Samuel Bayliss, Ganesh Dasika (AMD Research) | 2026-04 | https://arxiv.org/abs/2604.15379 Category: framework (primary) | Secondary tags: kernel, megakernel, persistent-kernel, chiplet, amd-mi350, cdna-isa, l2-locality, scheduler, runtime, llm-decode Read: 2026-04-21

Core Contribution #

Fleet 在 CUDA/HIP 的 wavefront → workgroup → grid 三层执行抽象之上补上了第四层 Chiplet-task（对应 AMD MI300/350 的 XCD 边界 + 4 MB 私有 L2），并以持久化 megakernel + 每 XCD 软件调度器 + 两级事件同步 + M-major 协作式 tile 遍历，把 8 颗 XCD 上本来互相踩踏的 L2 cache 变成协作复用的 L2，在 Qwen3-8B decode 上把 bs=1–16 比 vLLM 快 1.3–1.5×、bs=32 L2 hit rate 从 39% 提到 51%、HBM 读带宽降低 37%。

TL;DR #

现代 chiplet GPU（AMD MI300X/MI350 八颗 XCD 各自 4 MB 私有 L2；NV Blackwell 双 die）物理上把 L2 切碎了，CUDA/HIP 却还把 L2 当整块设备级资源暴露——结果 LLM decode 时 8 颗 XCD 各自独立把同一份权重从 HBM 拉进自己的 L2，把 40 MB L2 当成 4 MB 用。Fleet 做了三件事：(1) 在任务抽象里新增 Chiplet-task 绑定 "一颗 XCD + 其 L2 working set"；(2) 以 Mirage MPK 为基础实现持久化 megakernel，每颗 XCD 留一个 workgroup 做软件调度器，剩下 31 个做 worker，消除 250 次/token 的 kernel launch；(3) 用 M-major windowed 遍历 + cache-streaming 修饰符 让同 XCD 内多个 worker 共享同一个权重列，配合两级事件计数把 248 次 cross-chiplet fence 砍成 8 次。端到端数字：Qwen3-8B bs=1 decode 6.73 ms/token (vLLM 10.51 ms)，bs=32 L2 hit 38.9% → 51.0%、HBM 读降 18%，bs=64 HBM 读降 37%。

Q1 / Q2 / Q3 #

Q1 (痛点)：现代 chiplet GPU 的 L2 cache 物理分区 + 非相干（MI350 8 × 4 MB），但 CUDA/HIP 编程模型仍把 L2 当作设备级共享资源——LLM decode 时每个 XCD 独立从 HBM 拉同一份权重到各自 L2，L2 带宽被浪费；叠加每 token 250 次 kernel launch，小 batch 下调度开销主导延迟。

Q2 (方法)：提出四级任务抽象（wavefront / CU / Chiplet / device），新增的 Chiplet-task 把"一颗 XCD + 其 4 MB L2"作为显式作用域；以 Mirage MPK 为基础搭建持久化 megakernel + 每 XCD 软件调度器 + 两级事件同步（intra-XCD 无 fence / 仅"最后一人"出 fence）+ M-major 协作式 tile 遍历 + CDNA 三档 cache modifier（cache-streaming 权重 / NT bypass 激活 / volatile 调度通信）。区别前作：HazyResearch MegaKernel / FlashFormer / Mirage MPK 都假设 L2 是 monolithic；HipKittens 做了 XCD-aware 但只优化单 kernel。

Q3 (结果)：MI350 + Qwen3-8B dense bf16：bs=1 decode 6.73 ms/token（vLLM 10.51 ms，1.56× 加速），bs=1–16 普遍 1.3–1.5× 快于 vLLM；bs=32 L2 hit rate 38.9% → 51.0%，HBM 读流量 3,906 GB → 3,190 GB (−18%)；bs=64 L2 hit 61.4%，HBM 读 6,203 GB → 3,925 GB (−37%)。Ablation 证明：小 batch 的加速来自 Chiplet-task 调度（543 任务 vs 1,407）降低 dispatch overhead；大 batch 的加速来自 M-major 协作式 L2 复用；bs ≥ 64 Fleet 进入 compute-bound 后反而慢于 hipBLASLt（未实现 K-split）。

Summary #

动机：LLM decode 是 memory-bound（Qwen3-8B 每层权重 368 MB，激活仅 1 MB）且 kernel-launch-bound（36 层 × 近 7 op = 250 次 launch/token）。monolithic GPU 时代解决方案是 persistent megakernel（HazyResearch、Mirage），但 2024–2026 的 MI300X / MI350 / Blackwell 都走了 chiplet 设计——AMD 8 颗 XCD 各有 4 MB 私有 L2，跨 XCD 不相干；NV Blackwell 2 die 各有 L2 分区（逻辑上相干但物理 NUMA）。CUDA/HIP 语法里没有 chiplet 这一层概念，标准 GEMM dispatch 把 workgroup 均匀撒到 8 颗 XCD，每颗 XCD 独立把完整权重拉进自己 4 MB L2 → L2 被 reload 8 次，HBM 带宽被浪费。

方法：Fleet = 任务模型 + 运行时 + 代码生成三件套。（1）任务模型：四级任务 wavefront-task / CU-task / Chiplet-task / device-task，其中 Chiplet-task 是全新的一层，把"一颗 XCD 上所有 31 个 worker + 4 MB L2"当成原子单位。GEMM 按输出列 N-split 给 8 个 Chiplet-task，每个 XCD 的 31 个 worker 按 M-major windowed traversal 协作把 R=min(31, m_tiles) 个 M-tile 打包处理同一个权重列，从 HBM 只读一次权重就可以喂多个 M-tile → L2 hit 率解析式 (R-1)/R。（2）运行时：基于 Mirage MPK 改造，launch 一个 megakernel 占满所有 256 CU，每颗 XCD 1 个 workgroup 做 scheduler、31 个做 worker，scheduler 通过 HW_ID 寄存器发现自己在哪颗 XCD。两级事件同步：intra-XCD worker 对 L[xcd][e] 做 device-scope atomic（无 fence，因为 L2 内相干），只有 last-worker-per-XCD 出一次 buffer_wbl2 + GPU-scope atomic 更新 HBM 全局 G[e] → 248 次 fence 降到 8 次/事件（W=31 倍 reduction）。三档 cache modifier：权重用 sc1=1 nt=1（cache-streaming，短暂驻留）、激活 NT=1（bypass L2，不驱逐权重）、调度通信 volatile。（3）代码生成：在 Mirage 编译器里为 Chiplet-task 补上新的节点描述符，为每个 XCD 生成 base_ptr + k × (N/8) × K × sizeof(bf16) 的 pointer。

结果：AMD MI350X + Qwen3-8B dense bf16，与 vLLM v0.17.2 eager + Mirage MPK (chiplet-unaware) 对比。Fleet 在 bs=1 decode 6.73–6.82 ms/token（vLLM 10.51 ms, 1.56×；Mirage 7.83 ms, 1.16×）；bs=32 L2 hit 从 Mirage 的 38.9% 升至 51.0%、HBM 读减 18%；bs=64 L2 hit 61.4%、HBM 读减 37%。核心 ablation：M-split 变体（有 Chiplet-task 但 XCD 间不共享权重）在 bs=1–16 与 M-tile 打平（证明此区间加速来自 dispatch 减少），bs ≥ 32 差距打开（证明 L2 协作式复用独立于 dispatch）。bs ≥ 64 Fleet 被 vLLM 反超，因为 Fleet 的 GEMM 没实现 wave-level K-split，compute-bound 时 hipBLASLt 占优。

Key Findings #

• Chiplet-task 是 CUDA/HIP 缺失的那一层抽象：memory hierarchy 有 Register / LDS / L2 / HBM 四级，但 CUDA/HIP 任务层次只有 wavefront / workgroup / grid 三级——多出来的 L2 scope 在 chiplet 架构下成为 NUMA 性能悬崖的主因。
• 两种加速机制在不同 batch 区间起作用：bs=1–16 加速来自 "543 任务 vs 1,407 任务" 的调度开销降低（M-tile 与 M-split 打平）；bs ≥ 32 加速来自 L2 协作式复用（M-tile 领先 M-split 15%+）。M-split ablation 是整篇论文最干净的单因子分离。
• L2 hit 率解析式 (R-1)/R 在 bs=32 精准命中：预测 50%、实测 51%；bs=64 预测 75% 实测 61.4% 的 13 pp 差距暴露 cache-streaming 修饰符的副作用——R=4 时第 4 个 worker 触达之前权重已被标记逐出。
• 两级同步是 chiplet 架构下 persistent kernel 能 work 的必要条件：248 个 worker 都出 buffer_wbl2 会把 Infinity Fabric 打爆；Fleet 把 fence 次数压到每事件 8 次（仅 last-worker-per-XCD），31× 减少是 persistent megakernel 在 MI350 成立的前置条件。
• Persistent megakernel 扛掉 kernel launch 成本是 chiplet 感知的前提：vLLM 每 token 近 250 次 kernel launch，且每次 launch 都会 flush L2 以保证 device-sync 语义，任何 L2 优化都无法跨 kernel 复用；Fleet 付一次 launch 成本后 L2 状态整个 decode 周期都保留。
• cache modifier 三档策略是 AMD 独有优势：CDNA3/CDNA4 的 SC0/SC1/NT 组合允许按数据类型精细控制 L2 分配策略；NVIDIA Hopper 没有等价 per-instruction 粒度的 cache policy control，所以 Fleet 方案在 AMD 上更干净。
• Register pressure 让 persistent megakernel 只能跑 1 wave/SIMD：所有 task 类型编译进同一函数，寄存器用量取联合最大值 → 放弃 wave-switching latency hiding → 每次 L2 miss 直接 stall MFMA pipeline → L2 hit rate 从"锦上添花"升格为"生死攸关"。

Limitations #

• 评测仅一个模型（Qwen3-8B dense bf16）一个硬件（MI350X），无 MoE / 无量化 / 无多 GPU TP / 无 prefill。
• bs ≥ 64 compute-bound 区间 Fleet 反而比 vLLM 慢 ~5–10%，因为 persistent kernel 的 GEMM 实现里没有 wave-level K-splitting；这是实现缺陷而非架构缺陷，但论文承认尚未修复。
• Chiplet-task 数量硬编码为 8（= MI350 的 XCD 数），任务图与调度器代码通过查询运行时获取 X/W/C 参数，但未评测其它 chiplet 拓扑（MI300X 有 8 XCD 但 38 CU/XCD、NV Blackwell 仅 2 die）。
• 8 个 CU（3.1% on MI350）被 scheduler 占掉；当 XCD 个数增加到 16/32 时这部分开销会线性增长且不易回收。
• M-major 协作需要 m_tiles ≥ 2（即 B ≥ 2 × T_M）；当 batch 小于 T_M 时协作失效，只剩 dispatch-overhead 这一个加速源。
• Mirage 编译器的 super-optimizer 尚未理解 Chiplet-task 的 cost model，当前只能手动提供输入 task graph。
• 未与 NVIDIA thread block cluster / DSMEM 做头对头比较（作者无 Blackwell GPU，论文明确说明）。

Infrastructure Impact #

• Algorithm: 几乎无影响。Fleet 是推理运行时优化，不改变训练算法，也不改变模型权重分布。
• Kernel: 高度相关。把 GEMM / Attention / RMSNorm / SiLU 的 tile 遍历顺序从 N-major 切到 M-major，把 cache-all 改为 cache-streaming，为 chiplet-aware kernel 定义了新的模板。
• Framework: 核心影响层。提出 Chiplet-task 概念等于给 vLLM / SGLang / TRT-LLM 提了一个新的调度粒度：serving framework 未来需要在 PagedAttention / RadixAttention 之外，加一层 "block-to-chiplet affinity"。
• LLM: 仅在 dense transformer decode 验证，MoE 的 expert-route-per-token 会破坏 M-major 协作前提；对 long-context、speculative decoding 的影响未评估。
• Agent: 间接利好——agent serving（多轮、小 batch、强延迟敏感）恰好落在 Fleet 最优 bs=1–16 区间，1.3–1.5× decode 延迟降低直接翻译成 time-to-first-tool-call 的同比例降低。
• Code: 开源路径尚未披露（arXiv 预印本）。Mirage MPK 是开源的，Fleet 作为 Mirage 扩展原则上可 upstream；HipKittens 的 XCD-grouping 与 Fleet 互补。

Key Figures #

Figure 1: AMD Instinct MI350 memory hierarchy #

What it shows: MI350 的物理封装——8 颗 XCD 各自 32 CU + 私有 4 MB L2，经 Infinity Fabric Crossbar 连到共享的 256 MB Infinity Cache (LLC)，再连到 8 组 HBM3。

Why it matters: 论文所有论证的起点。CUDA/HIP 把 L2 视为单块设备级缓存，但物理上 MI350 的 L2 是 8 个不相干的 4 MB 分区——这是 chiplet-unaware scheduling 产生 L2 miss storm 的根本原因。

Detailed description: 顶部 8 个方框代表 8 个 XCD（仅画 XCD 0/1/6/7），每个方框里有 32 个 CU 小方格 + 一个 4 MB L2 绿色标签。所有 XCD 通过中间灰色 "Infinity Fabric Crossbar" 连接到下方的 256 MB Infinity Cache (LLC) 黄色条，LLC 再往下连到 8 个 HBM stack。跨 XCD 的数据访问必须穿过 Crossbar → LLC → HBM，延迟显著高于 L2-local。

Key Tables #

Table 1: Memory hierarchy → programming model mapping #

Takeaway: 一张表把论文的全部新意说清楚——CUDA/HIP 在 chiplet 这一层是空白的，Fleet 用 Chiplet-task 填进去。

Table 2: Chiplet-unaware LLM decode 特征 #

Takeaway: Linear/GEMM 占 95% 时间；每 XCD 分到 46 MB 权重而 L2 只有 4 MB（11.5× overshoot），没有协作时 L2 hit 仅 16.4%——Fleet 的所有优化都围绕这 95% 的时间展开。

Table 4: L2 Hit% 与 HBM 流量（相对 Mirage 归一） #

Takeaway: bs=32/64 M-tile 与 M-split 的分叉是全篇最关键的单点证据——证明 L2 协作式复用（而非仅 Chiplet-task 调度）才是大 batch 下的加速源。bs=64 HBM 读从 1.00× 降到 0.63×（−37%）。

Deep Analysis (framework) #

1. System Scope #

• 目标：chiplet-aware 持久化 megakernel 推理运行时，专攻 LLM decode 延迟敏感场景。
• 规模：单卡（MI350X，8 XCD），论文明确说未评测多卡 TP。
• 场景：Online / latency-sensitive。评测 TPOT（Time Per Output Token），decode-only。
• workload：dense transformer LLM decode（Qwen3-8B: d=4096, d_ff=12288, 36 layers, GQA 32q/8kv heads），bf16，bs=1–64。

2. Architecture & Data Flow #

Fleet 运行时 = 1 次 kernel launch → 8 颗 XCD × (1 scheduler + 31 worker) 持续运行直到整个 decode sequence 结束。

2a. 数据流：每 decode token 走一圈 #

控制面 vs 数据面：scheduler 是控制面（只做 pointer 运算、计数器读写），全都在 L2 local 或 HBM non-temporal；worker 是数据面（MFMA + HBM weight fetch）。分离得很干净，没有共享的同步热点。

state：task descriptor 表在 HBM 中只读（launch 前填好）；per-worker queue 和 event counters 是运行时可变 state。

失败恢复：论文不处理（单 GPU 推理、persistent kernel 作用域 = 一次 generate 调用）；如果 worker hang 或 XCD power event 导致 kernel abort，整个 generate 失败，由上层 serving framework 重试。

2b. 数据移动热点 top 3 #

1. 每 XCD 权重加载：每层 46 MB / XCD，每 token 36 层 → 每 XCD 每 token 1.65 GB。M-tile 模式下 R=4 (bs=64) 共享 → 降至 413 MB，与 HBM 5.3 TB/s 的预算吻合。与 MFMA 不重叠（persistent kernel 只有 1 wave/SIMD，没有 wave switching）——论文明确说 "Every L2 miss directly stalls the MFMA pipeline"。
2. buffer_wbl2 引发的 L2 → HBM writeback：每事件 8 次 fence × 若干 MB dirty 行。随 task 数线性增长，是 Fleet 的核心控制开销。通过两级计数从 248→8 次压缩。
3. scheduler ↔ global event counter 的轮询流量：8 scheduler 周期性读 G[e]（HBM non-temporal load），读节律由 scheduler busy-wait 循环决定，实测论文说"workers rarely stall waiting"，说明该热点未成为瓶颈。
flowchart TB title["Hierarchical 2-level Event Synchronization — turns 248 global fences into 8 (§5.2)"] flatTitle["(a) Flat / naïve: every worker issues GPU-scope atomic + fence"] flatWorkers["248 workers (31 per XCD × 8) each execute: __threadfence + atomic_add(G[e], 1) with sc0|sc1 → 248 cross-chiplet fences per event → L2 dirty-line writebacks scale with per-worker traffic → long stalls on Infinit..."] flatG["Global G[e] on HBM (contended: 248 writers)"] flatCost["Total fences / event = 248 Expected stall ~ 248 × t_wbl2 + contention"] hierTitle["(b) Fleet hierarchical: 2-level counter, fence only on last-worker-per-XCD"] hierL1["Level 1: intra-XCD (L2-local, no fence) Each worker: atomic_add(L[xcd_id][e], 1) → resolves in local L2 only, no Infinity Fabric traffic"] hierLast["Last-worker test: if L[xcd_id][e] == T then issue buffer_wbl2 (1 L2→HBM writeback) flat_atomic_add(G[e], 1) with sc0|sc1"] hierG["Global G[e] on HBM (only 8 writers, one per XCD)"] hierCost["Total fences / event = 8 → 31× reduction per event (W = 31 workers/XCD)"] decomposeTitle["Decomposition rationale (paper §5.2)"] decompose1["① Task queue — read-only. No sync needed (populated pre-launch)."] decompose2["② Scheduler → worker dispatch — device-scope atomic, but writer and readers all on same XCD → resolves in local L2 without fabric traffic."] decompose3["③ Worker → worker intra-Chiplet-task — device-scope atomic on L[xcd][e] counter; all participants share L2 partition, so no cross-XCD coherence and no fence required."] decompose4["④ XCD → global — only one worker per XCD makes the expensive trip: 1 buffer_wbl2 + 1 GPU-scope atomic. Amortizes fabric cost across the whole Chiplet-task."] decompose5["For wavefront-task and CU-task (1 worker), two-level counting is skipped — the single worker directly does the global atomic. Fleet's linear events (8 Chiplet-tasks/event for a dense GEMM) → 8 fences/event exactly."]

3. Design Space & Constraint Analysis #

时代定位 (Era Positioning) #

2025 下半年起 LLM 推理优化的"低垂果实"——continuous batching / paged attention / 算子融合 / FlashAttn-3——基本采完，研究重心转向两个方向：(a) 跨 op 持久化内核（HazyResearch MegaKernel 2025-05、FlashFormer 2025-05、Mirage MPK 2025 CoRR）和 (b) chiplet-aware 细粒度调度（HipKittens 2025-11、Fleet 2026-04）。Fleet 是这两股流合一的产物，也是第一次把 AMD CDNA3/4 的 XCD 抽象塞进 persistent kernel 编程模型。时间上这也是在 MI300X 规模化出货（2024 中）+ MI350/MI355 即将发布的关键 1–2 年窗口出现。

3a. Alternative Approaches #

方案	思路	为何不被 Fleet 采用 / 为何失败
标准 kernel-per-op（vLLM 默认）	让 hipBLASLt / Tensile 自己解决 L2 locality	launch 开销 每 token 250 次 launch；kernel 边界强制 flush L2（保 device-sync 语义）→ 跨 op 复用不可能
CUDA/HIP graph capture	录制 kernel 序列消除 launch overhead	每个 batch size 要独立 graph；仍是 kernel 粒度，跨 kernel L2 不复用；mismatch fallback 成本巨大（vLLM 用的就是这条）
Thread block swizzling + CUTLASS persistent GEMM	单 kernel 内通过 block index 重排提升 L2 hit	只优化单 kernel；跨 op 无能；且假设统一 L2
Thread block cluster (NV Hopper/Blackwell)	GPC-scope 内 block 协作，DSMEM shared	GPC ≠ chiplet 边界，GPC 更细（不对应 L2）；AMD 无等价硬件
HipKittens XCD-grouping	kernel launch 时把 block 绑到 XCD	仍是 per-kernel 优化，不能把 decode 整条路径编织到一个 megakernel 里
K-split on multi-XCD	跨 XCD partition K 维度	要跨 XCD reduction → 每个 partial sum 都经 Infinity Fabric 做 atomic → 小 batch 下延迟主导
硬件做 XCD-aware dispatch	让 hw scheduler 自动感知	硬件已有但受 grid launch 粒度限制，无法跨 kernel 维持 affinity

3b. Constraint Derivation #

• 为何不用 K-split 做 chiplet 划分？ 因为 K-split 需要 cross-XCD reduction，每个 partial 都要跨 Infinity Fabric atomic，单位延迟 ~100s ns × K/k 次，在小 batch（compute 总量本就小）下完全被通信吃掉。论文 §4.1 明确说 "N-split dominates at small batches where persistent kernel scheduling is the bottleneck"。
• 为何不跨 XCD 共享 L2？ MI300/350 的 L2 物理就不相干；跨 XCD visibility 必须 buffer_wbl2 全局回写 + 跨 Fabric probe，比 L2-local 慢 10–100×。
• 为何不让 scheduler 跑在 CPU？ 每 task dispatch 延迟 = PCIe 往返 + CPU→GPU signaling 几 μs，而 CU-task 本身只 ~3.8K cycles (≈2 μs)，调度延迟会超过任务执行延迟 → 必须放 GPU 端。
• 为何每 XCD 留 1 个 workgroup 当 scheduler？ (1) 要本地轮询 L2 local 的 event counter，放远端 XCD 会引入 fabric 延迟；(2) workgroup 粒度最少是 1 CU，进一步细粒度（wavefront-level scheduler）会被 occupancy 限制。代价 = 8/256 = 3.1% CU。
• 为何不按 row-major workgroup index 就近分 XCD？ 因为 HW 默认 grid 分派策略是 round-robin 跨 XCD；要 override 必须拿到 HW_ID 再重新排队 → 这正是 Fleet 做的。

3c. Assumption Audit #

• A1 (Chain step 2-3 load-bearing)：LLM decode 是 memory-bound，linear 占 95% → Fleet 所有优化围绕 linear。若将来 attention 因 long-context 反超（seq_len > 32K, KV cache > weights），这一前提动摇。风险评估：当前 agent 场景 context 4K–32K 仍以 linear 为主；若走向 128K+ long-context，attention 成本相对上升，Fleet 需要为 attention 也做协作式 L2 策略。
• A2 (Chain step 4 load-bearing)：batch 大到 m_tiles ≥ 2（即 B ≥ 32 配 T_M=16）时 M-tile 才有 L2 复用优势。这个门槛卡在 serving 典型的 bs=8–32 临界点，不稳健。若某 agent serving 负载 bs 稳定 ≤ 16，Fleet 只能吃 dispatch overhead 这一部分（1.15× 左右），拿不到大加速。
• A3 (Chain step 5 load-bearing)：cache-streaming 修饰符 (sc1=1, nt=1) 按 LRU "短暂驻留" 的语义——但论文 bs=64 的 61.4% 实测 vs 75% 预测存在 13 pp 差距，说明假设在 R=4 时已有偏差。若后续 batch 进一步增大或 K-chunk tile 更大，实际 L2 hit 会进一步偏离 Eq.(1)。论文未给出 streaming 修饰符的 L2 residence time，这是被作者刻意省略的"可攻击面"。
• A4 (Chain step 1 load-bearing)：Fleet 把 L2 当"通过 per-instruction cache modifier 精细控制"的资源——这是 AMD CDNA3/CDNA4 独有的语义（SC0/SC1/NT 位）。若迁移 NV Hopper/Blackwell，硬件只提供粗粒度的 _ro/_cg/_cs load modifier 和 cluster-scope barrier，无法复现同样的 3 档策略。论文 §8 声称 "could implement with platform-specific runtime backend"，但未量化 NV 端性能。
• A5 (Chain step 5 decorative)：register pressure 导致 1 wave/SIMD（无 wave-switching）→ L2 miss 直接 stall MFMA。该前提若被动态寄存器分配（AMD RDNA 有、CDNA 还没）打破，persistent megakernel 的 register 负担可减轻 → Fleet 加速幅度会缩水（因为动态寄存器分配让 kernel-per-op 方案也能恢复部分 occupancy）。

3d. Core Technical Barrier #

两级事件同步 × intra-XCD L2-local counter。这是 "persistent megakernel in chiplet GPU" 能成立的唯一支点：

• 248 workers × 每事件 1 fence × 每层 ~6 事件 × 36 层 = 53k fence/token。按 fabric-scope fence 平均 ~200 ns 计，就是 10 ms/token——已经超过整个 decode 预算。
• Fleet 两级计数 + last-worker-gate 把 fence 压到 8 × 6 × 36 = 1.7k fence/token，约 340 μs/token。
• 实现门槛：(1) 必须知道 last-worker 是谁 → 依赖 L2-local atomic 返回值 == Tth 这条判定；(2) L[xcd][e] 计数器必须 不出 fence（否则就白做）→ 依赖同 XCD 的 L2 读写自动相干；(3) 全局 G[e] 必须用 flat_atomic_add with sc0|sc1 才能跨 XCD 可见——这三条里任何一条在硬件层面弱化，整个方案崩盘。

为什么这是"低层工程洞察"而非"高层思想"：高层思想（两级 counter）是标准 hierarchical reduction 的模板，任何 parallel computing 教科书都讲；真正难的是找到 CDNA3/4 ISA 里哪组 scope 位组合能让 intra-XCD 的 atomic 既有原子性又免 fence——论文把这条实现到 buffer_wbl2 vs flat_atomic_add sc0 sc1 的精确边界。

3e. Design Binding Critique #

• Binding 1 (attacks chain step 3)：强绑定 AMD CDNA3/CDNA4 的 per-instruction scope/NT 位。外部生态校验：NV Hopper/Blackwell 仅有粗粒度 _cs/_cg/_ca 修饰符，无 XCD-scope 概念；但 2026 年 NV/AMD chiplet GPU 各占约一半 AI 算力（MI300X + MI350 + B200），若 Fleet 不做 NV 后端，生态影响面被限制在 AMD 阵营。风险等级：medium-beta（AMD 阵营持续扩张，bet 安全）。
• Binding 2 (attacks chain step 4)：强绑定 Mirage MPK runtime。Mirage 本身是 open-source 但尚未成为主流 serving 框架；若 vLLM / SGLang 自己做 megakernel（参考 HazyResearch 风格），Fleet 需要迁移。外部校验：截至 2026-04，vLLM 和 SGLang 均未原生集成 megakernel；主流生产路径仍是 eager/CUDAGraph。Fleet 通过 Mirage 生效意味着短期只能作为 AMD research 实验室工具链。风险等级：medium。
• Binding 3 (attacks chain step 4)：强绑定 dense transformer。MoE 模型里每 token 的 expert-route 不同，Fleet 的 M-major 协作假设 "多个 M-tile 共享同一个权重列" 被直接破坏——不同 token 要走不同 expert 权重。外部校验：Mixtral / DeepSeek-V3 / Qwen3-MoE / Llama-4 全是 MoE 趋势；2025–2026 新一代 flagship 几乎一致转向 MoE。这是论文最严重的 binding，但 paper chain step 1 的"痛点"定义在 dense 上，所以 chain 内部自洽；真正受影响的是外部 adoption 范围——MoE 场景需要一个新的协作模型（可能要 Expert-task 而非 Chiplet-task）。风险等级：high-beta。
• Binding 4 (attacks chain step 5 decorative)：强绑定 memory-bound regime（bs ≤ ~32）。bs ≥ 64 进入 compute-bound 后 vLLM + hipBLASLt 反超 Fleet。外部校验：agent serving 和 chatbot（1-on-1 低 QPS）天然落在 bs ≤ 16；离线批处理则 bs 常达 64–256。Fleet 明确把自己定位在 latency-sensitive decode，这个 binding 对 agent/chatbot 场景不造成问题，但在离线批推理会让位。风险等级：low（定位明确）。

3f. Solution Justification — Throughput / Cache-Hit Model 完整走读 #

Notation Table（from paper §4.1, §6.4） #

符号	含义	单位
$B$	batch size (M 维度总长度)	tokens
$W$	workers per XCD	workers（MI350: 31）
$X$	chiplet 数	chiplets（MI350: 8）
$C$	L2 capacity / chiplet	bytes（MI350: 4 MB）
$T_M$	GEMM output tile height along M	elements (16)
$T_N$	GEMM output tile width along N	elements (64 或 256)
$m\_tiles$	$\lceil B / T_M \rceil$	—
$R$	$\min(W, m\_tiles)$（共享同一权重列的 worker 数）	workers
$N$	全 GEMM 输出列数	elements
$N/8$	每 XCD 分到的输出列	elements
AI	arithmetic intensity = FLOP / HBM byte	FLOP/byte

核心方程：Eq.(1) L2 Hit Rate #

$$\text{L2 Hit}_{\text{weight}} = \frac{R-1}{R} = 1 - \frac{1}{\min(W, m\_tiles)}$$

形式推导：M-major 遍历下，worker $i$ 在 $n=n_k$ 列先处理 M-tile $i$，再处理 M-tile $i+W$...；当 $m\_tiles \geq W$ 时，一个权重列被 $W$ 个 worker 近乎同时消费，第 1 个 worker 从 HBM 加载权重进 L2，后续 $W-1$ 个 worker 全部 L2 命中 → hit rate = $(W-1)/W$。当 $m\_tiles < W$ 时，只有 $m\_tiles$ 个 M-tile 需要处理，余下 $W - m\_tiles$ 个 worker 去做别的列，仍有 $m\_tiles$ 个 worker 共享一列权重 → hit rate = $(m\_tiles-1)/m\_tiles$。合起来就是 $R = \min(W, m\_tiles)$。

为什么这个形式而不是别的：

• 为什么用 min？ 因为 R 受 supply (W worker) 和 demand (m_tiles 个 M-row) 双边约束；其中较小者是 bottleneck。
• 为什么是 (R-1)/R 而不是 1−1/W？ 因为 bs 很小时 W 被 m_tiles 压住，实际共享参与者数小于 W。
• 为什么没出现 N/8 或 K？ 因为 M-major 遍历保证了权重列在"一个完整的 M 扫描"时间窗口内持续驻留 L2，N 维度只影响轮数、不影响单列的 hit 率；K 维度在 tile 内串行消费，不影响列间复用率。

单调性：对 $B$ 单调非减（因为 $m\_tiles$ 非减），对 $W$ 也单调非减。当 $B \to \infty$，$R \to W$，上限 hit rate = $(W-1)/W = 30/31 \approx 96.8\%$。实测 bs=64 只到 61.4%，说明实际执行中有大量其他因素蚕食了预测极限——这就是 attack surface。

辅助方程：AI_eff 推导（Figure 7） #

Standard scheduling 下 AI = B（batch size in FLOP/byte，因为每 weight byte 被 B 个 activation 复用）。Fleet 的 L2 reuse 等效于把 HBM 字节数降到 $(1 - \text{hit}) \times $ 原值：

$$\text{AI}_{\text{eff}} = \frac{B}{1 - \text{L2 hit rate}}$$

物理含义：HBM 作为主存来看，有效 $B$ 字节被"压缩"成 $B \cdot (1-\text{hit})$ 字节。

形式推导：AI 定义是 FLOPs / 外部 memory bytes；L2 hit 不产生 HBM 请求，只有 L2 miss 的部分贡献到分母。

为什么这个形式：把 roofline 图里的点水平右移 $\tfrac{1}{1-\text{hit}}$ 倍——这是 roofline 分析的标准用法。

与 ridge point 交互：MI350 MFMA ridge $\approx 245$ FLOP/byte。bs=32 原 $\text{AI}=32$，Fleet 51% hit → $\text{AI}_{\text{eff}} = 32/0.49 \approx 65$，仍在 bandwidth-bound 区（$65 < 245$），但离 ridge 缩短 33%。

Monotonicity / Convexity / Uniqueness（论文未明写，补） #

• Eq.(1) 对 $B$ 严格单调非减；一阶差分 $\Delta \text{hit}/\Delta B = O(1/R^2)$，随 $B$ 增大回报递减 → M-major 的边际收益在 $B > W \cdot T_M \approx 496$ 后饱和（对 Qwen3-8B 这规模 batch 极少到达，所以 practically 不会饱和）。
• $R = \min(W, m\_tiles)$ 是凹函数（两条斜率递减曲线取 min 永远是 concave）；Eq.(1) 又是 $1 - \tfrac{1}{R}$，对 $R$ 单增且凹 → hit rate 对 $B$ 是凹的，收益前高后低。
• 唯一性：给定 $(B, W, T_M)$，Eq.(1) 无任何可调参数，L2 hit rate 是 closed-form，无需 sweep 任何 knob。这是 作者证明 最强的地方。

Model → Reported Numbers 的一阶映射验证 #

用 MI350 参数 $W=31, T_M=16$ 代入：

B	$m\_tiles$	$R = \min(31, m)$	Eq.(1) 预测	实测 (Table 4, Fleet M-tile)	残差	评价
1	1	1	0%	16.9%	+16.9%	来自 fused SiLU 共享 activation，论文自陈为基线偏移
16	1	1	0%	23.7%	+23.7%	同上，fused SiLU 受益越大
32	2	2	50%	51.0%	+1.0%	★ 模型精准命中 ±1 pp
64	4	4	75%	61.4%	−13.6%	streaming 修饰符提前驱逐权重，预测偏高

第一个结论：在 bs=32 这个关键点上，模型误差只有 1 pp，这是 作者证明 是"load-bearing" 的直接证据——case study 的 51% 数字 真的是 从 Eq.(1) 一阶解出来的，而不是 sweep 出来再回头套模型。

第二个结论：bs=1/16 处的 16.9%/23.7% 完全不是 Eq.(1) 的预测；论文坦白是 fused SiLU 带来的"基线偏移"。这部分应该加一个独立项：$\text{hit}_{\text{baseline}} = f(\text{fusion})$，而 Eq.(1) 只描述 M-major 协作这一块增量。论文没显式拆开，是 作者证明 的一个表达瑕疵。

第三个结论：bs=64 的 −13.6% 残差直接暴露了 作者证明 的 attack surface（下文）。

"No Case Study" Defense（只靠 Eq.(1) 能否说服读者？） #

可以。单看 Eq.(1)：bs=32 + MI350 配置，模型预测 L2 hit 率 50%，相对 chiplet-unaware 的理论 0%（因为每 XCD 独立读权重，无论何时都是 miss）提升 50 pp → 按 AI_eff = B/(1−hit) 推算，effective AI 从 32 提升到 64，HBM 流量降低 50% → decode 延迟中 memory-bound 部分降低 50%。考虑 linear 占 95%，整体预估降低 47.5% → 从 Mirage 15.62 ms 应降至 8.2 ms。实测 12.35 ms，说明"到 50%"这个假设在 streaming 修饰符下兑现了 ~50%，predicted-over-actual ratio ≈ 63%，在一阶量级对齐。

作者证明 Attack Surface（真正的 §3c 钩子） #

1. Streaming 修饰符行为假设：Eq.(1) 隐式假定"在 R 个 worker 全部访问完前，cache line 不被逐出"。实测 bs=64 (R=4) 已偏差 13.6 pp，说明 cache-streaming 的 "短暂驻留" 时间窗比 R=4 需要的时间窗短——这是论文没给出定量表征的隐藏参数。若 batch 再增大 (bs=128, R=8)，预测值会离真实值越来越远。
2. Fused SiLU 基线偏移未纳入模型：bs=1/16 的 16–24% hit 完全来自 fused activation 的 L2 驻留，Eq.(1) 对此无预测能力；论文在 §6.4 用 "Without fused SiLU, L2 hit drops to ~9%" 一句话说明但未建模。
3. MALL (256 MB Infinity Cache) 的贡献被忽略：Table 5 per-layer 总权重 368 MB > 256 MB MALL (×1.4)，所以 MALL 不能 在一层内全放下；跨层是否受益取决于下一层是否复用本层权重——显然不复用，所以 MALL 只当 L2 victim 用，不在 Eq.(1) 的主预测路径。但若把 MALL 当二级共享缓存建模，理论上能再挖 5–10 pp hit rate。
4. Activation L2 bypass 在 register pressure 下是否成立：Eq.(1) 假设 activation NT=1 不驱逐权重。若某 op 的 intermediate 太大无法塞进 LDS+register → fallback 到 L2 → 破坏假设。论文未给寄存器预算表（与 LLaMA 参数规模的耦合），这是 attack surface 最大的一块。
5. Fence 开销未进入延迟模型：Fleet 不只省 HBM 流量，也省 fence，但 fence 成本只能在 Fleet 内部看（Table 4 不拆 fence 时间）。若 baseline 给 fence 时间 vs Fleet 的 fence 时间对比表，SJM 会更完整。

核心技术壁垒（Core Technical Barrier）- summary #

两级事件同步 + CDNA 三档 cache modifier 的组合。看懂是一回事，复现是另一回事：必须在 GPU assembly / ISA 级别理解 AMD buffer_wbl2 / flat_atomic_add sc0 sc1 / sc1=1 nt=1 load 修饰符的相对语义，并把它们编织进 Mirage 的 IR 生成路径。对复现者要求：有 AMD 内部硬件文档 + 熟悉 Mirage 编译器 + 能做 HW_ID 寄存器级调度——三个条件合起来基本只有 AMD Research 团队具备。

论证链重构（Argument Chain） #

Step	Premise	Conclusion	Evidence
1	Modern chiplet GPUs (MI300/MI350/Blackwell) 用私有 L2 分区	CUDA/HIP 的 3 层抽象 (wavefront/workgroup/grid) 对应不上 L2 层的边界；编程模型留下 "Chiplet scope" 空白	§1 Intro, §2 Table 1, Figure 1
2	LLM decode 95% 时间在 linear GEMM，368 MB 权重 vs 4 MB L2 per XCD 11.5× overshoot	chiplet-unaware dispatch 让 8 颗 XCD 各自独立拉完整权重 → L2 bandwidth 被浪费，decode 卡在 HBM	§2.1, §2.2, Table 2 (95% + 16.4% L2 hit)
3	CUDA/HIP graph capture 把 launch overhead 减少但不消除、且每 batch 要独立 graph、跨 kernel 不复用 L2	必须换到 persistent megakernel 才能同时解 launch + 跨 op L2 复用	§2.3, §7 Related Work
4	monolithic 时代 persistent kernel (HazyResearch/FlashFormer/Mirage) 可行；chiplet 时代 naïve 实现会被 248× fence 打爆	必须引入两级事件同步 + Chiplet-task 抽象 + M-major 协作 tiling	§5.1, §5.2, Figure 4/5
5	Eq.(1) L2 hit = (R−1)/R 预测 bs=32 达 50%；实测 51%（误差 1 pp）	Chiplet-task 方案在 load-bearing 的 bs 区间兑现预测 → 1.56× decode 加速	Table 4, §6.4, Figure 6

Load-bearing: step 2（痛点定义）、step 4（方法必要性）、step 5（SJM + 实证）。

Decorative: step 3（文献综述性质，去掉不影响主 claim）。

作者预先应对的反对意见：§7 Related Work 里把 HazyResearch/FlashFormer/Mirage/HipKittens/TB cluster 一一对齐，预先回应 "这不就是又一个 megakernel 吗？" 的反问——答案是 "他们假设 L2 是 monolithic"。

4. Key Innovations #

flowchart TB title["Fleet Persistent Megakernel Runtime Layout on MI350 (§5.1, §5.2)"] hbm["Global HBM (288 GB): • Task descriptor queues (populated before kernel launch, read-only during execution) • Global event counters G[e] (updated only by last-worker-per-XCD with sc0|sc1) • Model weights, KV cache..."] llc["256 MB Infinity Cache (shared LLC) — victim cache for all 8 L2 partitions"] fabric["Infinity Fabric Crossbar — cross-XCD coherence traffic flows here (expensive: buffer_wbl2 + GPU-scope atomics)"] xcd0[" "] xcd0_label["XCD 0 (L2 coherent domain, 32 CUs)"] sched0["Scheduler (1 workgroup on 1 CU) • reads HW_ID reg • polls global G[e] • dispatches tasks to per-worker queues • broadcasts Chiplet-task to all 31 workers"] workers0["Workers W₀ … W₃₀ (31 CUs) Each worker holds per-worker task queue in HBM (naturally cached in local L2 via wave-scope loads). Executes: Wavefront-task | CU-task | Chiplet-task body (MFMA GEMM tiles, attentio..."] counter0["L2-local event counter L[0][e] (device-scope atomic_add; resolves inside L2; no fence)"] l20["Private 4 MB L2 (TCC) — cache modifier policy: • Weights: sc1=1 nt=1 (cache-streaming, reused briefly then evicted) • Activations: NT=1 (L2 bypass — don't evict weight tiles) • Intra-XCD sched↔worker: volatile loads (..."] last0["Last-worker check → single buffer_wbl2 fence (L2 → HBM writeback) → flat_atomic_add G[e] with sc0|sc1"] xcd7[" "] xcd7_label["XCD 7 (symmetric; XCD 1..6 omitted for space)"] sched7["Scheduler (XCD_id=7) polls G[e] ≥ need"] workers7["Workers W₂₁₇ … W₂₄₇ (31 CUs) Consume their assigned Chiplet-task slice: output cols [7 × N/8, 8 × N/8) M-major windowed traversal → R = min(31, m_tiles) workers share each weight tile"] counter7["L2-local event counter L[7][e]"] l27["Private 4 MB L2 — same cache modifier policy as XCD 0"] last7["Single buffer_wbl2 + GPU-scope atomic_add G[e]"] last0 -->|" "| fabric last7 -->|" "| fabric fabric -->|" "| llc llc -->|" "| hbm

Innovation	Mechanism	Benefit	Cost/Tradeoff
Chiplet-task abstraction	把 "一个 XCD + 其 L2" 作为新的任务层级，GEMM 按输出列分成 8 个 Chiplet-task	任务数 1,407 → 543 (2.6×)；intra-chiplet L2 coherence 免 fence	只绑 AMD XCD 粒度；MoE 不适用
M-major windowed traversal	worker 先沿 M（batch）方向遍历再进下一个 N 列，让 $R=\min(W, m\_tiles)$ 个 worker 共享同一权重列	L2 hit 38.9% → 51.0% (bs=32)；HBM 读 −18%	要求 $B \geq 2 T_M$（bs ≥ 32），小 batch 无增益
Three-tier cache modifier	权重 sc1=1 nt=1 streaming / 激活 NT=1 bypass / 调度 volatile	激活不驱逐权重；权重生命周期对齐共享窗口	仅 CDNA3/4 有 per-instruction scope 位；NV 不可复现
Hierarchical 2-level sync	intra-XCD worker 只做 L2-local atomic，仅 last-worker 出 buffer_wbl2 + flat_atomic_add sc0	sc1	每事件 fence 从 248 降到 8（31×）	需 HW_ID 发现自身 XCD、需正确识别 "last worker"
Per-XCD scheduler + HW_ID-based dispatch	每颗 XCD 1 个 workgroup 做 scheduler，通过读 HW_ID 寄存器发现自己在哪颗 XCD	workers 自主调度，launch 成本摊到整个 generate 周期	3.1% CU 被 scheduler 占用

flowchart TB title["Cooperative M-major Weight Tiling = L2 Hit Rate Model (§4.1, §6.4)"] setup["Setup: One Chiplet-task processes [M, K] × [K, N/8] = [M, N/8] sub-GEMM on one XCD. Output tile shape is [T_M, T_N] (e.g. 16 × 256). Along M (batch) → m_tiles = ⌈B / T_M⌉ tiles. Along N/8 → n_tiles = (N/8)/T_N...."] eq1Box["L2 hit model (Eq.(1) §6.4): R = min(W, m_tiles) — # workers that concurrently touch the same weight column L2 Hit_weight = (R − 1) / R = 1 − 1 / min(W, m_tiles)"] thH0["Batch B"] thH1["m_tiles (T_M=16)"] thH2["R = min(31, m_tiles)"] thH3["Predicted Hit% = 1 − 1/R"] thH4["Measured M-tile Hit%"] thH5["Residual (meas − pred)"] thH6["Comment"] r1_0["1"] r1_1["1"] r1_2["1"] r1_3["0%"] r1_4["16.9%"] r1_5["+16.9%"] r1_6["fused SiLU baseline"] r16_0["16"] r16_1["1"] r16_2["1"] r16_3["0%"] r16_4["23.7%"] r16_5["+23.7%"] r16_6["pure fusion, no coop"] r32_0["32"] r32_1["2"] r32_2["2"] r32_3["50.0%"] r32_4["51.0%"] r32_5["+1.0%"] r32_6["★ model load-bearing"] r64_0["64"] r64_1["4"] r64_2["4"] r64_3["75.0%"] r64_4["61.4%"] r64_5["−13.6%"] r64_6["streaming evict shortfall"] aisft["Roofline shift (Figure 7): AI_eff = B / (1 − L2 hit rate) — Fleet's cooperative tiling raises effective arithmetic intensity. At bs=32, 51% hit → AI_eff = 32 / 0.49 ≈ 65 (2.0× rightward shift toward MFMA ridge at AI=245)"] attacks["SJM attack surface: (i) Eq.(1) assumes perfect LRU/streaming: in practice the cache-streaming hint (sc1=1, nt=1) evicts weights before the R-th worker touches them — residual gap at bs=64 (−13.6%) is the tell. ..."]

5. Scheduling & Resource Management #

• Batch formation：不负责——由上层 serving framework 决定 batch，Fleet 只负责单个 decode batch 内的 task 调度。
• Memory management：权重在 launch 前加载完成；KV cache 由上层管理；Fleet 不做 paged attention（decode-only，KV cache 访问模式固定）。
• GPU utilization：persistent megakernel 占 100% CU（减 scheduler 3.1%）；靠 scheduler 控制 worker 的任务流保证计算一直有活。register pressure 导致 1 wave/SIMD，没有 wave-switching latency hiding → 必须靠 L2 hit 率保证 MFMA pipeline 不 stall。
• Multi-tenancy：不支持——megakernel 期间整块 GPU 不可抢占。serving 场景的 multi-tenancy 要靠上层切分 request 流。
• SLO-aware：不支持优先级；decode 内部 FIFO。
• Scheduler idle/busy balance：论文说 "workers rarely stall waiting for task assignment"——scheduler 足够快，瓶颈在 worker 的 MFMA / HBM 读。

6. Target Scenarios & Workload Characterization #

场景	Workload	SLO	Fleet 相对 vLLM 优势	瓶颈
Interactive chatbot decode	bs=1–4, seq 长短混合	TPOT < 10 ms	1.3–1.5× TPOT	kernel launch + L2 miss
Agent tool-call decode	bs=1–8, 多轮短 gen	TPOT < 10 ms	1.3–1.5× TPOT	kernel launch + L2 miss
Voice agent / real-time	bs=1, 每 token 预算极紧	TPOT < 5 ms	1.5×（最大收益区间）	与 chatbot 同
Batch summarization	bs=32–64, 长 context	总吞吐	bs=32: 1.27×；bs=64 持平 / 略慢	MFMA compute
MoE decode	每 token route 不同 expert	—	不适用（binding 3）	expert 不共享

primary bottleneck：在 Fleet 目标区间（bs ≤ 16）是 scheduling-bound + memory-bound 各占一半；bs=32 变成 memory-bound 主导；bs ≥ 64 切换到 compute-bound（Fleet 失效的转折点）。

7. Performance Evaluation & Before-After Comparison #

7a. Metrics #

Metric	Definition	Unit
TPOT	Decode-only time per output token, 64-input-1024-output	ms/token
L2 Hit Rate	rocprofiler 测 L2 访问命中率	%
HBM Rd/Wr	测到的 HBM 读/写字节数	GB/token sequence
Effective AI	B / (1 − L2 hit rate)	FLOP/byte

7b. Before-After #

Figure 6: Decode-only TPOT (ms/token) across batch sizes #



What it shows: 4 条曲线对比 TPOT：vLLM (蓝)、Mirage MPK (灰)、Fleet M-tile (红实线)、Fleet M-split (橙)。

Why it matters: 论文的"头号实验图"。在 bs ≤ 16 Fleet 两个变体与 vLLM 拉开差距（约 1.5× 快）；bs=32 M-tile 开始领先 M-split 和 Mirage；bs=64 M-tile 仍领先 Mirage 1.30× 但被 vLLM（compute-bound 区借助 hipBLASLt K-split）反超。

Detailed description: 横轴 log2 batch size 1–64，纵轴 ms/token 0–30。vLLM 曲线在 bs=1–16 横在 10–12 ms；Mirage 从 7.8 涨到 10.8；Fleet M-tile/M-split 从 6.7 涨到 ~9（bs=16）。bs=32 处 4 条线分叉：vLLM ~12，Mirage 15.6，M-split 13.4，M-tile 12.4。bs=64 vLLM 保持 ~12，Mirage 24.1，M-split 23.4，M-tile 18.6。

Optimization	Metric	Baseline (Mirage)	Fleet M-tile	Improvement	Conditions
Chiplet-task dispatch	TPOT @ bs=1	7.83 ms	6.82 ms	1.16×	MI350, Qwen3-8B
Chiplet-task + M-major	TPOT @ bs=32	15.62 ms	12.35 ms	1.27×
Chiplet-task + M-major	TPOT @ bs=64	24.10 ms	18.61 ms	1.30×
(vs external vLLM)	TPOT @ bs=1	10.51 ms (vLLM)	6.73 ms (M-split)	1.56×
M-major L2 reuse	L2 hit @ bs=32	38.9%	51.0%	+12.1 pp
M-major L2 reuse	HBM read @ bs=64	6,203 GB	3,925 GB	−37%
Fused SiLU alone	L2 hit @ bs=1	9.4%	17.4%	+8 pp	Fleet config

7c. Bottleneck Shift Analysis #

vLLM: launch-bound (250×) + memory-bound → Mirage MPK: memory-bound (16% L2 hit) → Fleet M-split: memory-bound (same hit, less dispatch) → Fleet M-tile: memory-bound (51% hit @ bs=32) → @ bs=64: compute-bound (MFMA pipeline) → Fleet 失效, vLLM 反超

决定性 shift 有三个：

1. 去掉 launch overhead（vLLM → Mirage）：bs=1 有 1.34× 加速。
2. 降低 dispatch 开销（Mirage → Fleet M-split）：bs=1 再加 1.16×。
3. 把 L2 利用率拉起来（Fleet M-split → M-tile）：bs=32+ 再加 1.10–1.27×。

7d. Baselines & Fairness #

• 公平性：vLLM 0.17.2 在 ROCm，明确 --enforce-eager 禁用 CUDA Graph（否则对 Fleet 不公）；Mirage 是 Fleet 作者移植到 AMD 的内部 baseline，区别只在是否 chiplet-aware——最干净的内部 ablation。
• 遗漏基线：HipKittens (2511.08083) 做了 XCD-aware GEMM 但只 kernel-level；论文引用但未做端到端对比。如果 HipKittens + vLLM 集成，bs=32 以下可能逼近 Fleet。
• overhead：8 scheduler CU = 3.1% 计算资源永久占用；未独立量化 scheduler 对总功耗的影响。
• Fleet 会输的场景：bs ≥ 64（compute-bound，K-split 未实现）、MoE、long-context attention-bound workload、多 GPU TP。

{{drawio:2604.15379_arch.drawio#page=4&height=720}}

Figure 7: Roofline analysis — Fleet shifts effective AI rightward #



What it shows: log-log roofline 图，x 轴 operational intensity (FLOP/byte)、y 轴 attainable performance (TFLOPS)。HBM 5.3 TB/s 斜线 + MFMA 1299 TFLOPS ridge @ AI=245。每个 bs 对应两组点：实心蓝 = Standard 理论、实心红 = Fleet 理论、空心蓝 = Mirage 实测、空心红 = Fleet 实测。

Why it matters: 直观解释 Fleet 的核心收益——把 memory-bound 的点沿 bandwidth 斜线水平右移；bs=32 的 Fleet 红三角从 AI=32 向右推到 AI≈65，离 ridge 缩短一半距离。

Detailed description: 从 bs=1 的 6 TFLOPS (AI≈1) 到 bs=32 的 120 TFLOPS (AI≈65) 呈对角线分布；Fleet 实测点都在 Mirage 实测点的右侧且更高。

8. API & Usability #

• API 未披露完整。论文说 "provide different input code to Mirage compiler"——当前仍是手写 task graph，不是自动编译出来。
• 模型支持：仅验证 Qwen3-8B dense；任何 dense transformer 理论可复用，MoE 不可。
• 部署：未讨论容器化或 K8s；当前是 AMD research 阶段原型。
• 可配参数：T_M、T_N、cache modifier 组合——多数硬编码到 Mirage 生成的 library 代码里。

9. Infrastructure Impact #

Layer	Impact
Algorithm	无直接影响（推理运行时优化）
Kernel	定义了新的 tile 遍历模板（M-major windowed），推动 kernel DSL 支持 per-instruction cache modifier
LLM	验证于 dense transformer；MoE / long-context 未验证
Agent	agent serving 多是 bs=1–8 短 gen，完美落在 Fleet 最优区间
Ops	仍是单 GPU；未讨论 autoscaling / 监控

10. Comparison Matrix #

Feature	Fleet	vLLM	SGLang	HazyResearch MegaKernel	FlashFormer	Mirage MPK	HipKittens
Persistent megakernel	✓	✗	✗	✓	✓	✓	✗ (per-kernel)
Chiplet-aware L2	✓	✗	✗	✗	✗	✗	部分 (XCD grouping)
Continuous batching	上层	✓	✓	上层	上层	上层	—
Paged attention	上层	✓	✓	上层	上层	上层	—
AMD target	✓ (MI350)	✓	✓	✗ (H100)	✗ (H100)	✗ (A100/H100)	✓ (MI350)
NV target	可移植 (未实现)	✓	✓	✓	✓	✓	✗
Open source	未披露	✓	✓	blog post	✓	✓	✓

11a. Code Cross-reference — Mirage MPK upstream (AMD 端实现不在 upstream) #

Shallow-cloned 到 knowledge-meta/cache/repos/mirage-project-mirage，版本 2026-04 HEAD。

走查 include/mirage/persistent_kernel/persistent_kernel.cuh (1,496 行) +

runtime_header.h (335 行) + python/mirage/mpk/persistent_kernel.py (2,225 行)。

Finding 1 — Upstream Mirage 完全没有 AMD / HIP / CDNA / XCD 代码 #

ls include/mirage/persistent_kernel/tasks/
# ampere blackwell common cute deprecated hopper speculative_decoding
grep -rn "XCD\|chiplet\|MI300\|MI350\|CDNA\|gfx9\|HIP_" include/mirage/persistent_kernel/
# (no matches)

含义：Fleet 论文里 "Mirage MPK ported to AMD MI350X, serving as internal baseline"

一句话轻描淡写 —— 实际上 upstream Mirage 只有 Ampere/Hopper/Blackwell 后端，

没有任何 AMD path。Fleet 团队除了论文主贡献（Chiplet-task 抽象），**还隐式完成了

一次完整的 Mirage HIP/CDNA 后端移植**，这部分工作量在论文中完全未量化。对读者

的启示：Fleet 的工程投入 ≠ 论文 §3–§5 描述的内容，真实投入远大于此。

Finding 2 — Mirage upstream 已有 两层 scheduler/worker 架构，但假设 device scope #

// persistent_kernel.cuh L454-L525 核心分派逻辑
int const num_schedulers =
 config.num_local_schedulers + config.num_remote_schedulers;
int const num_schedulers_per_sm = std::min((int)blockDim.x / 32, 4);
int const sched_id = blockIdx.x * num_schedulers_per_sm + warp_id + offset;

• Mirage 的 scheduler 是 warp 级别 (threadIdx.x % 32 == 0，每 SM 最多 4 个 scheduler warp)
• Fleet 的 scheduler 是 workgroup 级别（1 workgroup/XCD = 1 CU 被占，3.1% on MI350）
• 粒度差别的原因：Mirage NV 端假设 L2 monolithic，worker/sched 共享 L2 无损；Fleet 必须把 sched 绑在 XCD-local L2，最小单位是 workgroup

// event 分派：Mirage 用 random scheduler id
int sched_id = get_rand_sched_id(event_index, worker_id,
 config.num_workers,
 config.num_local_schedulers);

• Mirage 随机路由 event → scheduler
• Fleet 按 HW_ID（XCD id）亲和路由，让 event 保持在 L2-local 边界内
• Fleet 的 "per-chiplet scheduling" 本质是把这行 get_rand_sched_id 替换成 get_xcd_local_sched_id(HW_ID)

Finding 3 — Event counter + trigger 机制在 upstream 已存在 #

// runtime_header.h
EventCounter *all_event_counters;
int *all_event_num_triggers; // 每 event 的触发阈值

// persistent_kernel.cuh: 典型 trigger 逻辑
atom_cas_release_gpu_u64(&config.sched_queue_last_ready_event_id[sched_id],
 last_event_pos, last_event_pos + 1);

• Mirage 已有 num_triggers 阈值 + "最后一人触发"模式
• Fleet 的 "两级事件计数"复用了这个骨架，但把 outer counter 从 device-scope atomic 换成 L2-local atomic（intra-XCD 免 fence），inner counter 维持 device-scope
• 论文 §5.2 描述的 "last worker issues buffer_wbl2 + flat_atomic_add sc0|sc1" 实际是把

Mirage 的 atom_cas_release_gpu_u64（NV release 语义）改写为 AMD 等价语义

等价性验证：Mirage NV 端 atom_cas_release_gpu_u64 = atomicCAS_block + membar.sys；

Fleet AMD 端需要 s_waitcnt lgkmcnt(0) + buffer_wbl2 + flat_atomic_add sc0|sc1。

粒度不同但语义对齐。

Finding 4 — per-worker queue 结构 upstream 已存在，Fleet 未改动 #

// persistent_kernel.cuh prepare_kernel
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
 i < 2 * config.num_workers; i += ...)
 config.worker_queue_last_ready_task_id[i] = 0;
// ^^^^^ 2 × num_workers = local + remote queue pair per worker

• 每 worker 一对 queue（local + remote）已是 Mirage 标配
• Fleet 只是把 local queue 的读写限制在同 XCD 内（通过 scheduler 的 HW_ID 路由）
• 这解释了为什么 Fleet §5.1 只有 "extends Mirage Persistent Kernel system with chiplet-aware concepts" 一句话 — 数据结构没变，变化全在分派策略

Summary — Fleet 相对 upstream Mirage 的 delta（从代码角度） #

维度	Upstream Mirage	Fleet (paper)	Fleet 工程投入
后端	ampere / hopper / blackwell	+ AMD CDNA (MI350)	完整新后端移植（论文未量化）
Scheduler 粒度	warp-level (4/SM)	workgroup-level (1/XCD)	架构微调
Event routing	get_rand_sched_id	get_xcd_local_sched_id(HW_ID)	1 函数改写
Event atomic	atom_cas_release_gpu_u64 (NV release)	`buffer_wbl2 + flat_atomic_add sc0\	sc1` (AMD)	ISA 层等价重写
Task 抽象	wavefront / CU / device (3 层)	+ Chiplet (新增 1 层)	任务模型扩展
Per-worker queue	local + remote pair	完全复用	0
Task graph 生成	Mirage muGraph 超优化	Fleet 手写 input task graph	退步（作者承认 §5.3 future work）

诚实承认：本次 cross-reference 的局限 #

• Fleet 的 AMD 后端源码完全未公开（论文未给 repo link）。我们只能看到 Mirage upstream (NV 端) 来推断差异。
• 无法验证 Fleet 论文里 "8 schedulers × 1 CU" 的精确 SM/CU 占用数据——需要 AMD backend 源码才能确认。
• Eq.(1) (R-1)/R 的实际实现在 HIP kernel 中如何保证 worker 按 M-major 顺序执行，无法从 upstream 代码推断（Mirage upstream 的 traversal pattern 是通过 task graph 编码，不走 kernel-level loop）。
• MI350 CDNA3/4 的 buffer_wbl2 和 sc0|sc1 atomic 语义只能依赖 AMD ISA 手册（AMD ISA reference）交叉验证。

结论：Fleet 的实际工程体量比论文暗示的大得多（完整 AMD 后端移植被省略没提），

paper 的 "extends Mirage with chiplet-aware concepts" 不能字面理解 —— 正确理解是

"把 Mirage 的 device-scope 假设在 AMD 8-die 架构上彻底解构重写"。这是代码

cross-reference 最有价值的输出：paper 的轻描淡写 vs 代码揭示的真实工程量。

11. Adoption, Maturity & Ecosystem Influence #

• 开源状态：预印本 2026-04-15 submitted；代码未公开链接；可能会通过 Mirage MPK 下游以 PR 形式 upstream。
• AMD 内部路径：论文第一作者 + 全部合作者来自 AMD Research → 大概率进入 ROCm 生态工具链（但不一定走 vLLM ROCm backend）。
• 下游影响预判：
• HipKittens (2511.08083) 已做了 XCD grouping；Fleet 可看作它的 persistent-kernel 兄弟篇——可能二者合并为 AMD 端到端 kernel DSL。
• vLLM ROCm backend 与 SGLang ROCm 如果要采纳，需要从 eager / CUDAGraph 路径切到 megakernel 路径，工程改动大；短期不太可能。
• Mirage 超优化器（Wu et al. 2025 OSDI） 未来扩展 cost model 支持 Chiplet-task 是自然方向。
• 采纳门槛：需要 Mirage + AMD HIP + CDNA4 ISA 知识 + 对 megakernel persistent 运行时能熟练 debug；估计前 1 年仅 AMD 内部使用。
• 生态影响追踪：paper 发表 <1 周，尚无外部引用；关注 2026-Q3 后是否出现在 SGLang / vLLM ROCm 路线图。

Open Questions #

1. NV Blackwell 端（2 × L2 partition）是否能 1:1 实现 Chiplet-task？cluster_scope barrier + DSMEM 是否足够替代 buffer_wbl2？性能增益有多大？
2. MoE 场景下 Chiplet-task 如何适配——每 token 只激活 2 个 expert，M-major 协作前提崩溃；需要 Expert-task 抽象？
3. 多 GPU TP + Chiplet-task 组合时，N-split 维度会被 TP 先吃掉一部分（N → N/TP），导致 per-XCD 分区过小跌破 L2 容量阈值——在什么 TP 度下 Fleet 仍然有收益？
4. K-split wave-level reduction 如何补到持久化 megakernel 里？bs ≥ 64 的 compute-bound 区 Fleet 能否翻身？
5. scheduler 3.1% CU 开销在 MI400 这类 16 XCD 架构下会线性放大（bigger chiplet count），能否用 wavefront-level scheduler 降开销？
6. long-context（seq ≥ 32K）时 attention 成本相对上升到 40–50%，Fleet 的 linear-only 优化是否需要延伸到 attention 的协作式 L2 策略？
7. Fleet 的 SJM（Eq.(1)）在 streaming 修饰符下 bs=64 residual 13.6 pp，能否给出 streaming eviction 的解析模型（e.g. LRU window × allocation rate）？

第二轮批判审视（外部 review 交叉校验） #

本节由 2026-04 NeuralTalk Fleet 解读（微信公众号）触发。对比该独立 解读与本笔记，surface 了若干原始笔记没覆盖的批判视角。这些视角已经 沉淀为 的新规则（根本性 vs 缓解性 / Trade-off 轴 / 对立观点 / 软件→硬件反推 / 范式演进 / multi-tenancy）。按 Skill- Modification Protocol 回溯补入本笔记。

A. 根本性 vs 缓解性审视 #

Fleet 对"内存墙"的根本性突破仅存在于批处理场景 (bs ≥ 32)。

场景	加速来源	是否根本性
bs = 1 (agent / chatbot 延迟敏感)	dispatch 压缩 (543 vs 1407 task)	❌ 非根本性。L2 hit 16.9% ≈ Mirage 的 16.4%，权重复用完全未发生。Fleet 在此 regime 下只是"launch overhead eliminator"，和 CUDA Graph + hipBLASLt 的组合效果差距有限
bs = 8-16	dispatch 压缩	❌ 非根本性。同上。M-tile 和 M-split 变体打平（Table 4 L2 hit 20–24% 几乎一致）印证没有 L2 复用发生
bs = 32 (m_tiles=2, R=2)	dispatch 压缩 + L2 协作 (预测 50% / 实测 51%)	✅ 根本性。L2 复用是 Fleet 独门贡献
bs = 64 (m_tiles=4, R=4)	L2 协作 (预测 75% / 实测 61.4%)	✅ 根本性但 compute-bound 区让位 hipBLASLt

前提 (condition): Fleet 的"根本性"L2 突破要求 m_tiles ≥ 2 ⇔ B ≥ 2 · T_M = 32（T_M=16）。

前提失效时的效果归零：bs < 32 时 Eq.(1) 的 R=1 → 预测 0% L2 reuse，完全不产生权重复用加速；仅剩 dispatch 开销压缩（约 1.15× vs Mirage）。

论证的结构性结论：Fleet 对 "内存墙"的根本性解决仅在 batched serving。对 chatbot / agent / voice agent 这类 bs=1 延迟敏感场景，"内存墙"本身需要硬件层面变革（跨 chiplet L2 共享、一致性协议改写），软件方案的天花板已经被 Fleet 打到了。外部 review 的原话值得引用："真正根本性的解决方案可能需要硬件层面的变革"。

B. Trade-off 轴识别 + 未探索 Hybrid #

bs=64 Fleet 被 vLLM 反超不是 "K-split 未实现" 那么简单 —— 论文把它归因为"实现未完善"是一种规避真相的叙事。真相是存在一根结构性 trade-off 轴：

$$\text{Trade-off 轴:}\quad \underbrace{\text{跨算子持久化复用}}_{\text{Fleet 方向}} \ \longleftrightarrow\ \underbrace{\text{单 kernel 极致调优}}_{\text{hipBLASLt 方向}}$$

• Fleet 方向 必须: 所有 task 类型编译进单函数 → register budget 取并集 → 只能跑 1 wave/SIMD → 放弃 wave-switching latency hiding → compute-bound 区必然劣于 kernel-per-op baseline
• hipBLASLt 方向 必须: 每 kernel 独立 register 预算 + 自动调优 tile + 支持 K-split/stream-K/tensor-core 动态组合 → compute-bound 区饱和 MFMA → 但每个 kernel 边界都会 flush L2，memory-bound 区被 HBM 带宽限死

两个方向在 compute-bound vs memory-bound 谱上天然互斥，不是"Fleet 补上 K-split 就能两全"。

未探索的 Hybrid 方向（Fleet 论文完全没讨论）：

Adaptive Fusion Granularity Switch — persistent megakernel 内部携带一个轻量级性能预测器，每次迭代开始时根据 (batch_size, sequence_length, hit_rate_last_step) 决定： - 当 B · T_M ≥ L2_capacity_threshold (≈ bs ≥ 32 on MI350) → 走 Fleet megakernel 路径 - 否则 → 切换到 hipBLASLt K-split GEMM + 独立 kernel 具体实现：megakernel 留出一个 "bypass task" 类型，任务描述符里指向 hipBLASLt 的 handle；scheduler 看到 bypass 类型就把对应 GEMM 委托出去。代价是每次切换约 1 次 kernel launch，但在 bs ≥ 64 能吃到 hipBLASLt 的 compute-bound 极值。

这条 hybrid 方向是一个完整的独立研究问题，值得放进 Open Questions。

C. 对立观点主动搜索 —— MoE 其实可能比 dense 更契合 Fleet #

原本 §3e 的 Binding 3 写的是 "MoE 不适用（最严重）"，因为"每 token 路由不同 expert，M-major 协作假设破坏"。但对立观点审视发现：

我们原本视角	对立视角 (NeuralTalk review)
视角起点：Fleet 现行设计的 M-major 协作	视角起点：重新设计 Fleet 给 MoE 用
结论：MoE binding 3 最严重	结论：MoE 天然契合 Fleet
理由：token 之间不共享 expert 权重 → 协作失效	理由：不同 expert 驻留不同 XCD 的 L2，gating network 做 expert→XCD affinity routing，expert 权重从此 XCD-local，比 dense 的 N-split 更自然

重要观察：MI350 有 8 颗 XCD —— 这个数字恰好对应典型 MoE 的 num_routed_experts（Mixtral 8×7B 精确对齐；DeepSeek-V2 的 64 routed + 2 shared 可以做 8 个 XCD 各 8 expert；Qwen3-Next 80B-MoE 的 128 routed 可 16 expert / XCD）。Expert ↔ XCD 映射是 MoE 专属的 chiplet-affinity opportunity，Fleet 的 Chiplet-task 抽象原则上可以直接表达（需要把 Chiplet-task 类型扩展出 Expert-Chiplet-task 变体）。

结论：我们原本的 Binding 3 critique 在 Fleet 的 M-major 现行设计下成立，但作为研究方向是误导性的。正确表述应该是：

Binding 3（精修版）: Fleet 的 M-major GEMM 协作设计不直接适用 MoE，但 MoE 的 expert 稀疏激活与 chiplet 私有 L2 天然亲和——将 Chiplet-task 从"XCD-local weight partition"语义扩展到"XCD-local expert affinity"是论文未探索但价值很高的方向。chain step 3 的 premise 需从 "dense-only" 放宽到 "partition 可沿 expert 或 token 维度"。

D. 软件 → 硬件反向推演 (framework §12 conditional — hardware-proximal 触发) #

Fleet 的纯软件方案给以下未来硬件特性提供了存在性证明 + 量化依据：

未来硬件特性	Fleet 证明其价值的数据	硬件支持后软件开销预计降低
硬件管理的 XCD-local 任务队列	Fleet 用 8 workgroup（3.1% CU）做 scheduler，workers rarely stall waiting for task assignment (§5.1)	scheduler CU 开销从 3.1% → 0%；MI400 类 16 XCD 架构尤其显著（当前 6.25% → 0%）
跨 chiplet 轻量级 fence (SC + atomic 的单周期版本)	Fleet 两级计数把每事件 fence 从 248 降到 8，量化了 fabric fence 的代价主导	每事件 fence 继续降到 0（硬件直通 counter），持久化 megakernel scheduling throughput 再涨 20-30%
可编程 cache scope 分区 (per-stream LRU class)	Fleet 用 3 档 cache modifier 隔离 weight/activation/sched 流量	动态 partition 无需 paper §4 描述的静态 modifier hint；能适应 MoE expert 切换的突发访问
HW-accelerated HW_ID routing	Fleet scheduler 通过 HW_ID 寄存器读 + software routing 决策	硬件实现 scheduler-local queue 会让"每 XCD 一个 scheduler"变成免费

推论：AMD MI400 / MI500 代际的 CDNA 5/6 架构应当考虑将上述 4 项从 ISA 扩展到 microarchitecture；这是 Fleet 论文给 AMD 架构路线图贡献的隐式 design wishlist。NVIDIA Blackwell Ultra / Rubin 代也可借鉴——Blackwell 的 thread-block cluster 已经向 "硬件 chiplet 感知同步原语" 迈了一步但粒度不对 L2 层级，Fleet 的证明给 Rubin 代做到 "chiplet-scope cluster" 提供了量化论据。

E. 范式演进定位 #

Fleet 的真实贡献不只是 "1.5× vLLM decode 加速"，而是把一个"隐式硬件约束 → 显式软件资源"的抽象模板补完了：

Paper	隐式硬件约束	变成显式软件资源
vLLM PagedAttention (2023)	KV cache 必须连续分配	page 显式资源（paged memory manager）
FlashAttention (2022)	HBM 访问是瓶颈	tile in SRAM 显式资源（IO-aware scheduling）
Fleet (2026)	L2 物理分区（chiplet）	Chiplet-task 显式资源
(hypothetical next)	MALL Infinity Cache 分层	LLC-task ?
(hypothetical next)	跨 GPU NVLink bandwidth	Island-task ?

这是一条可以推广的范式演进主线：每当硬件引入一种新的隐式拓扑约束，软件抽象就需要把它变成显式一等公民。Fleet 在这条主线上的贡献不是最大的（PagedAttn / FlashAttn 更有影响力），但是这条主线被外推到芯粒层级的第一个严肃实现。下一篇类似工作很可能处理 Infinity Cache 或跨 GPU L2 —— 这是值得关注的方向。

F. Multi-tenancy & 侧信道审计 #

Fleet 论文完全未讨论云多租户场景。但持久化 megakernel + 分区 L2 引入了至少 3 类侧信道攻击面：

1. L2 contention fingerprinting：co-tenant 的 Fleet 实例占用同一 XCD 时，L2 cache 替换率泄露对方 Chiplet-task 的 working set size → 模型架构指纹（通过周期性 L2 miss 模式识别对方是 MoE 还是 dense）
2. Scheduler queue depth side-channel：全局 G[e] event counter 的 atomic 竞争延迟可被对方 poll 到 → batch 到达率/并发度泄露
3. buffer_wbl2 timing channel：每次跨 XCD flush 的延迟和 dirty-line 数量成正比 → 对方 working set 变化率泄露，进而推出 sequence length / layer depth 特征

Fleet 当前适用范围（基于安全视角）：

• ✅ 单租户 inference（个人 / 企业专用 serving）
• ✅ 受信任 multi-tenancy（同一 organization 内部）
• ❌ 不受信任 multi-tenancy（公有云跨租户直接共享单 GPU）—— 需要加上 XCD-level hardware isolation + L2 scrubbing between requests 才能进入此场景

结论：Fleet 走向生产需要一整块 multi-tenancy security 补课。这是论文的"盲区 limitation"，论文未列、我们原笔记也没列。

设计哲学摘要 #

一句话总结：Fleet 是"给 persistent megakernel 补上 chiplet 这一层的 OS kernel"——把 CUDA 的 hw dispatcher 换成软件 per-XCD scheduler，把 L2 分区从编程模型 bug 变成显式资源。

技术细节彩蛋：

• fused SiLU 的副作用意外给出 bs=1 基线 L2 hit 16.9%：gate+up 拼成单个 GEMM 后共享一份 activation，驱动 L2 hit 从 9% 涨到 17%——Eq.(1) 完全不 cover 这部分，是论文未建模的"免费午餐"。
• volatile load for intra-XCD scheduler comm：不用 L2 因为 scheduler 要看到最新写入；这是一个很少被外界讨论的 AMD CDNA3/4 语义 corner。