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

framework 2604.15379 — Cross-paper Synthesis

L3 Relate · Fleet (2604.15379) #

1. 相关论文 #

#Entity关联维度关联强度
1TileRT (tilert-speed-scaling-law)persistent kernel + GPU-resident scheduling + warp/tile-level specialization — 同为 persistent megakernel 范式,TileRT 在 NV 端用 AOT compile 消除 kernel boundary,Fleet 在 AMD 端用 Chiplet-task 层级消除 L2 碎片化
2PrfaaS (2604.15039)PD 分离 + 跨 DC serving + 层间 prefill pipeline — PrfaaS 关注跨 DC KV 传输带宽,Fleet 关注单卡内部 L2 带宽;二者都在"暴露并利用被编程模型隐藏的硬件拓扑层级"
3DualPath (2602.21548)storage I/O disaggregation + agent serving弱-中 — DualPath 利用 DE SNIC 空闲带宽聚合 KV 加载,Fleet 利用 XCD 间 L2 分区做权重复用;二者互补:DualPath 解 storage→GPU 瓶颈,Fleet 解 HBM→L2 瓶颈
4PPD (2603.13358)多轮 PD 路由决策 + decode 干扰管控 — PPD 在 scheduling 层做 append-prefill 路由避免网络瓶颈,Fleet 在 kernel execution 层优化 decode tile 遍历。二者在 agent 场景下可叠加
5MFS (2603.17456)网络 flow 调度 + SLO-aware disaggregated serving — MFS 解决三阶段通信争用,Fleet 解决单 GPU 内 L2 争用。两者在不同网络层级(datacenter fabric vs on-die fabric)解决类似问题模式
6TensorHub (2604.09107)RDMA pipeline replication + RL training weight transfer — TensorHub 的 pipeline replication(部分完成即可服务)与 Fleet 的 M-major windowed traversal(部分权重列复用)有结构相似性:都是"让下游提前消费未完整数据"
7ZeRO-Prefill (2605.02960)MoE prefill + weight streaming + communication hiding — ZeRO-Prefill 反转 EP 数据流(weight 流向 activation 而非反之),Fleet 反转 tile 遍历方向(M-major 让 weight 被多个 worker 复用而非每个 worker 独立拉取)。核心 insight 共通:把 weight 当可调度资源而非固定放置
8KVServe (kvserve)KV-cache 压缩 + PD 通信优化 — KVServe 压缩 PD 间传输的 KV 数据,Fleet 压缩 HBM→L2 间的 weight 带宽需求。两者在不同 memory hierarchy 层级做"带宽压缩"

2. 本篇 vs 相关论文的 delta #

2.1 vs TileRT — 同为 persistent kernel,不同战场 #

TileRT 将整个 transformer 模型 AOT 编译为单个 Engine Kernel,以 warp/block/GPU 三级特化消除 inter-kernel idle [tilert-speed-scaling-law]。Fleet 同样使用 persistent megakernel 消除 kernel launch 开销,但核心贡献不在此——而在 Chiplet-task 抽象和 M-major 协作式 L2 复用 [2604.15379]

维度TileRTFleet
硬件NV H200 NVLAMD MI350 8-XCD
核心问题BS≈1 下 kernel boundary 主导延迟(~10× gap)chiplet L2 分区导致 weight 重复加载(11.5× overshoot)
Persistent kernel 角色核心贡献本身 — 消除 inter-kernel overhead前置条件 — 基于 Mirage MPK 已有能力
独特贡献层AOT compiler + warp specializationChiplet-task 层级 + 两级事件同步 + 三档 cache modifier
调度粒度tile(编译期静态)Chiplet-task → CU-task → wavefront-task(运行时软件调度)

Fleet 的 incremental delta:TileRT 假设 L2 是 monolithic(NV 端 L2 逻辑上相干),Fleet 是第一个把 persistent kernel 适配到 物理分区非相干 L2 架构的工作。Fleet 的两级事件同步(31× fence reduction)是 persistent kernel 在 MI350 上能成立的必要条件——没有这一步,248 个 worker 的 GPU-scope fence 就会打爆 Infinity Fabric [2604.15379]。TileRT 在 NV 端不需要这一步。

2.2 vs ZeRO-Prefill — "weight 作为可调度资源"的两种实现 #

ZeRO-Prefill 把 MoE expert 从"activation 路由到 expert 所在 GPU"反转为"expert weight 流向 activation 所在 GPU" [2605.02960]。Fleet 把 GEMM 权重从"每 XCD 独立从 HBM 拉取"变为"M-major 遍历让 R 个 worker 共享同一权重列" [2604.15379]

共通 insight:weight 不是静态参数,而是可调度的带宽资源。两者都利用 compute window 来 overlap/amortize weight 传输:

关键差异:ZeRO-Prefill 需要 NVLink 级带宽(D2D AllGather 在后台),Fleet 只需 L2-local 一致性(无需跨 XCD 传输权重——权重从 HBM 到 L2,不跨 chiplet)。Fleet 的数据移动方向是垂直的(HBM→L2),ZeRO-Prefill 是水平的(GPU→GPU)

2.3 vs PrfaaS — 暴露隐藏拓扑的两个层级 #

PrfaaS 把"跨 DC Ethernet 带宽"从 PD 分离的隐式约束变为显式可调度资源(长度阈值 $t$ 路由 + 双时间尺度调度)[2604.15039]。Fleet 把"XCD 私有 L2 分区"从 CUDA/HIP 的隐式 NUMA 变为显式 Chiplet-task 作用域 [2604.15379]

二者属于同一范式演进主线(详见 Fleet L2 §E):

论文隐式约束显式资源
PrfaaS (2604.15039)跨 DC Ethernet BW长度阈值路由 + 选择性 offload
Fleet (2604.15379)L2 物理分区Chiplet-task binding

矛盾/互补:PrfaaS 的 layer-wise prefill pipelining 使 KV 传输与计算重叠 [2604.15039]——这与 Fleet 的 M-major 遍历让 weight load 与 MFMA 重叠有结构相似性。但 Fleet 明确说 persistent kernel 的 1 wave/SIMD 让 L2 miss 直接 stall MFMA [2604.15379]——即 Fleet 的场景下 load 和 compute 不能 重叠。这是 persistent kernel register pressure 的代价:TileRT 通过 warp specialization 解决了这个问题(不同 warp group 分别做 load 和 compute),Fleet 在 AMD 端没有等价方案。

2.4 vs DualPath — agent serving 的两层优化 #

DualPath 在 storage→GPU 层面通过聚合 DE SNIC 带宽解决 agentic 推理的 I/O 瓶颈(1.87× 离线 JCT 提升)[2602.21548]。Fleet 在 HBM→L2 层面通过 M-major 遍历减少 37% HBM 读流量 [2604.15379]

叠加性分析:Agent serving(bs=1–8,多轮短 generation)恰好落在两者的最优区间。DualPath 解决"KV cache 从 storage 到 GPU 太慢",Fleet 解决"weight 从 HBM 到 compute 太慢"。二者在 memory hierarchy 的不同层级各自消除一个瓶颈,理论上可完全叠加——前提是 Fleet 的 AMD-only 限制和 DualPath 的 NVIDIA-only 限制不冲突(当前互斥,不能同时部署)。

3. 可攻击面 #

Attack 1:bs=1–16 的加速不是 Fleet 的真正贡献 #

Fleet 声称 bs=1 达 1.56× vLLM 加速 [2604.15379],但 L2 hit 分析显示此区间 R=1、Eq.(1) 预测 0% L2 复用 [2604.15379]。实测 16.9% hit 来自 fused SiLU 的基线偏移,与 Chiplet-task 无关。Fleet 在 bs≤16 的加速全部来自 dispatch overhead 消减(1,407→543 任务数)——这本质上是 Mirage MPK 的已有能力(megakernel 消除 launch),加上 HW_ID routing 减少 scheduler 冗余。

反驳 path:如果对手复现"Mirage MPK on MI350 + XCD-aware HW_ID dispatch"(不含 M-major 协作),在 bs≤16 应能获得 ~90% 的 Fleet 加速。Fleet 的独门贡献(M-major L2 复用)仅在 bs≥32 生效。论文 Table 4 的 M-tile vs M-split 对比其实承认了这一点 [2604.15379],但 abstract 和 headline 用 bs=1 的 1.56× 做标题是 cherry-pick 最大数字

Attack 2:cache-streaming 修饰符的行为未建模——bs=64 的 13.6 pp 残差是 iceberg tip #

Eq.(1) 假设"R 个 worker 访问完前,cache line 不被逐出"。bs=64 时预测 75% 实测 61.4%(−13.6 pp)[2604.15379]。论文未给出 streaming 修饰符的 L2 residence time 参数。如果 batch 进一步增大(bs=128),理论预测的 ~97% hit 与实际可能偏差 30+ pp——整个 closed-form model 在大 batch 下崩溃

与 ZeRO-Prefill 的饱和阈值 T = t_EP × F_GPU × γ 对比 [2605.02960]:ZeRO-Prefill 的 T 是可测量的物理量(profile run 直接标定),而 Fleet 的 Eq.(1) 中 streaming eviction 行为是不可观测的 cache 微架构行为。Fleet 的模型在 load-bearing 区间(bs=32)偶然精准(±1 pp),但这更可能是参数空间碰巧对齐而非模型鲁棒性的证据。

Attack 3:Persistent megakernel 的 register pressure 注定 1 wave/SIMD — Fleet 的加速天花板被 compute-bound 锁死 #

所有 task 类型编译进同一 persistent kernel → register 取联合最大值 → 只能跑 1 wave/SIMD → 无 wave-switching latency hiding [2604.15379]。这直接导致:

TileRT 通过 warp specialization(不同 warp group 承担不同职责)部分缓解此问题 [tilert-speed-scaling-law],但 Fleet 在 AMD CDNA 端没有等价的细粒度 warp 特化机制——CDNA 的 wave64 粒度比 NV 的 warp32 更粗,特化空间更小。

Attack 4:MoE 场景的致命盲区 #

Fleet 的 M-major 协作假设"多个 M-tile 共享同一个权重列"——MoE 模型中每 token 路由不同 expert,权重列不再共享 [2604.15379]。ZeRO-Prefill 的 AsyncEP 在 MoE 上实现 1.35× 提升 [2605.02960],证明MoE 优化是 real workload demand。Fleet 在 2026 年最热门的 MoE 架构(DeepSeek-V3/V4、Qwen3-MoE、Llama-4)上完全失效

Fleet L2 §C 的"对立观点"提出了 Expert-Chiplet-task 方向(expert↔XCD affinity routing)[2604.15379],但这只是猜想——ZeRO-Prefill 已经用实验证明了"把 expert 流向 compute"比"把 compute 路由到 expert"更优。

Attack 5:单模型单硬件评测 + 代码未公开 #

评测仅 Qwen3-8B dense bf16 一个模型、MI350X 一个硬件。无量化(FP8/INT4)、无 MoE、无多 GPU TP、无 prefill、无 long-context attention [2604.15379]。代码未公开,Mirage upstream 完全没有 AMD 后端 [2604.15379]

对比:ZeRO-Prefill 评测 4 种硬件/精度组合 × 多种并行度 × 真实 workload [2605.02960];DualPath 评测 3 种模型 × 1152 GPU 扩展 [2602.21548]。Fleet 的 evaluation breadth 远不及同期 framework papers。

4. 生态位 #

范式定位 #

Fleet 处于 "硬件拓扑显式化" 范式演进的 chiplet 层级节点:


FlashAttention (2022): HBM→SRAM 显式化
  → PagedAttention (2023): KV memory→page 显式化
    → Fleet (2026): L2 分区→Chiplet-task 显式化
      → (future): Infinity Cache/NVLink→Island-task 显式化

Fleet 不是这条主线上影响力最大的节点(FlashAttn/PagedAttn 远超),但是 第一个把范式外推到 chiplet 层级的严肃实现 [2604.15379]

采纳证据与壁垒 #

vs peer systems 的部署场景分化 #

场景最优系统Fleet 角色
Agent decode bs=1–8, AMD MI350Fleet主力(1.3–1.5× decode 加速)
Agent decode bs=1–8, NV H200TileRT不适用
Agentic serving, storage I/O boundDualPath互补但平台互斥
MoE prefill-onlyZeRO-Prefill不适用
Multi-turn PD routingPPD正交可叠加
Disaggregated network schedulingMFS正交(不同网络层级)
Cross-DC prefill offloadPrfaaS互补
KV transfer compressionKVServe互补

5. 未探索方向 #

Direction 1:Fleet + ZeRO-Prefill 的 Hybrid — MoE on Chiplet GPU #

ZeRO-Prefill 的 AsyncEP 把 expert weight 后台 AllGather 到本地 GPU [2605.02960]。Fleet 的 Chiplet-task 把 weight 按输出列分配到 XCD [2604.15379]组合方向:在 MI350 上,先用 AsyncEP 把 routed expert 从远端 GPU AllGather 到本地,再用 Fleet 的 M-major 协作把 expert weight 在 8 颗 XCD 间做 L2 复用。这解决了 Fleet 的 MoE 盲区 + ZeRO-Prefill 的 chiplet-unaware 问题。

需要解决:expert 数量(64–128)远超 XCD 数量(8),需要多轮 weight streaming + Chiplet-task 重绑定的动态调度。

Direction 2:Adaptive Fusion Granularity Switch #

Fleet L2 §B 识别了 persistent megakernel vs hipBLASLt 的 structural trade-off 轴 [2604.15379]。未探索的 hybrid:megakernel 内部维护性能预测器,当 B·T_M ≥ L2_capacity_threshold 时走 Fleet 路径,否则切换到 hipBLASLt K-split GEMM。实现方式:预留 "bypass task" 类型,scheduler 看到此类型就委托给 hipBLASLt handle。

价值:消除 Fleet 在 bs≥64 的劣势,实现 memory-bound→compute-bound 全覆盖的 Pareto。

Direction 3:TileRT 的 warp specialization 移植到 CDNA #

Fleet 的 1 wave/SIMD 限制(register pressure)是 performance ceiling。TileRT 用 warp specialization 让 load/compute/comm 在不同 warp group 中并行 [tilert-speed-scaling-law]。如果 AMD CDNA5 提供更细粒度的寄存器分区(类似 NV 的 setmaxnreg),可以实现"load warp 用少寄存器 + compute warp 用多寄存器"的异构配置,打破 Fleet 的 occupancy=1 限制。

Direction 4:PrfaaS 的层间 pipeline + Fleet 的 Chiplet-task = 跨 DC chiplet-aware serving #

PrfaaS 依赖层间 prefill pipeline 使 KV 传输与计算重叠(Eq.3 的 min 形式)[2604.15039]。Fleet 在单 GPU 内不能 overlap(1 wave/SIMD)。但如果将 Fleet 部署在 PD 分离架构的 decode 端:decode 的 weight load 由 Fleet 的 L2 协作优化,KV load 由 PrfaaS 的跨 DC pipeline 优化——两条路径在不同 memory hierarchy 上并行。

Direction 5:Hardware-assisted Chiplet-task #

Fleet L2 §D 推演了硬件可以简化 Fleet 的 4 项特性 [2604.15379]。最有价值的是 硬件管理的 XCD-local 任务队列——消除 3.1% scheduler CU 开销。结合 MFS 的 DSCP 标记思想 [2603.17456]:如果 GPU 硬件提供类似 DSCP 的 per-tile priority tagging(标注 weight tile 的 L2 生命周期),cache-streaming 行为可以从"hint"升级为"guarantee",消除 Eq.(1) 的 streaming eviction 残差。