持续执行范式：从 Kernel Boundary Elimination 到 Ultra-Low-Latency Engine #

1. 主题缘起 #

2024–2026 年，LLM 推理工作负载从大 batch 训练时代的规则化执行模式，全面转向延迟敏感的微 batch serving。Token-by-token decoding 每次 forward 仅涉及少量数据但触发上百个 μs 级小操作——attention QKV、softmax、activation、LayerNorm、KV cache update 等 [2604.17861]。GPU 的 kernel launch 机制仍为 ms 级大 kernel 设计，每次穿越 host-device 边界的固有开销在 3–7 μs [2604.17861]。当操作本身仅需 3–10 μs compute 时，launch 开销占 30–100%，GPU profiler 呈典型"梳齿形"——计算 burst 之间充斥 host 提交空白期 [2604.17861]。

以 8×H200 NVL 系统为例：聚合内存带宽 ~38 TB/s，GLM-5.1 每 decode token 激活参数量 ~42 GB，纯带宽约束下理论吞吐可达 ~1000 tok/s，而实际系统仅输出几十 tok/s——差距达一个数量级 [tilert-speed-scaling-law]。根因不是算力不足，而是 GPU 反复执行 launch→load→compute→store→synchronize 循环，每次 kernel boundary 打断数据流、销毁局部性、强制重新同步 [tilert-speed-scaling-law]。

与此同时，GPU 硬件从单 die 走向多 chiplet 架构——AMD MI300X/MI350 的 8 颗 XCD 各有 4 MB 私有 L2，跨 XCD 不相干 [2604.15379]。CUDA/HIP 编程模型仍把 L2 当设备级共享资源暴露——LLM decode 时 8 颗 XCD 独立从 HBM 拉同一份权重到各自 L2，带宽被浪费 [2604.15379]。FlashAttention2 的 round-robin 调度将同一 attention head 的 workgroup 打散到不同 XCD，导致 L2 命中率暴跌至 ~1% [2511.02132]。

TileRT 开源代码仓库首次以可审计形式暴露了"持续执行理想"在生产级推理引擎中的工程实态 [tile-ai-tilert]。代码揭示的执行模型与博客叙事存在关键偏差：博客描述 AOT 编译将整个模型展开为单个 Persistent Engine Kernel [tilert-speed-scaling-law]，而代码实际使用 CUDA graph 捕获-回放配合闭源 C++ tile 调度器实现 [tile-ai-tilert]。尽管如此，代码验证的性能数字——DeepSeek-V3.2 达 600 tok/s、GLM-5 达 500 tok/s [tile-ai-tilert]——证明 tile 级跨算子融合和整图执行确实可以大幅逼近理论带宽上限。

与此同时，TokenSpeed 开源推理引擎代表了一条不同于 persistent kernel 的 kernel boundary elimination 路线：C++ FSM 类型安全调度器将 KV cache 资源安全从运行时检查提升到编译期保证，Placement 编译器通过静态 SPMD 注解自动插入通信操作消除手写 AllReduce 的繁琐，分层内核注册表 + Blackwell MLA 特化内核在保持可扩展性的同时取得极致性能 [tokenspeed]。Kimi K2.5 on B200 上 Pareto 全面优于 TRT-LLM——min-latency 快 ~9%，100 TPS/User 吞吐高 ~11% [tokenspeed]。TokenSpeed 的出现表明：kernel boundary 消除不一定需要 persistent kernel 或 megakernel——结构化编译 + 类型安全调度 + CUDA graph 捕获同样可以逼近极致延迟 [framework]。

与前述聚焦 GPU 内部执行的路线不同，Blink 从另一个维度攻击同一问题：host CPU 本身常驻在每个 token 的关键路径上——负责调度、批处理、KV cache 管理与 CUDA kernel dispatch，即便叠加 CUDA graph 和 overlapped scheduling，调度器每个 decode step 仍必须返回 host [2604.07609]。这条脆弱的 CPU 路径不仅带来固有协调开销，还使推理对同机干扰高度敏感——pbzip2 干扰下 vLLM+H100 吞吐掉 3.8×、P99 TTFT 膨胀 139×，根因是跨地址空间的 TLB 失效与 LLC 污染把 LLC stall cycles 抬高 11.2× [2604.07609]。Blink 的处方是把整个 serving stack 委派出 host CPU：前端（HTTP、tokenization、请求管理、RDMA 传输）交给 SmartNIC/DPU（BlueField-3 的 16 个 ARM 核心），后端调度交给一个在 GPU 上永不退出的 persistent CUDA kernel（单 thread block、256 线程），host CPU 仅在启动时加载模型、捕获 CUDA graph，随后完全退出 [2604.07609]。

这些压力——kernel 边界开销、chiplet NUMA 效应、agentic workload 对超低延迟的刚性需求 [framework]、以及 host CPU 常驻关键路径带来的干扰脆弱性 [2604.07609]——独立驱动研究者走向同一方向：持续执行（persistent execution）或更广义地说 kernel boundary elimination，GPU 程序在操作之间不将控制权交还 host。2024–2026 年间五条独立路线同时涌现：AOT 编译的全模型 Engine Kernel [tilert-speed-scaling-law]（其代码实现为 20+ 跨算子 fused ops + CUDA graph 捕获 [tile-ai-tilert]）、JIT 动态 persistent dispatcher [2604.17861]、chiplet-aware persistent megakernel [2604.15379]、结构化编译 + 类型安全调度引擎 [tokenspeed]、以及 SmartNIC-delegated CPU-free serving stack + GPU-resident persistent scheduler [2604.07609]。这一主题横跨 kernel（算子级优化、硬件调度、通信重叠）和 framework（编译策略、运行时编排、部署架构），是当前推理基础设施中最活跃的跨领域交汇点 [kernel] [framework]。

2. 覆盖的 category 分布 #

Kernel 类别贡献底层执行原语：persistent kernel 调度机制、chiplet NUMA 感知调度、tile-level 编程抽象、计算-通信分离策略 [kernel]。Framework 类别贡献系统级架构：全模型编译流水线、运行时异构 worker 编排、跨 GPU 通信集成、类型安全调度与资源管理 [framework]。TileRT 代码仓库的加入为 framework 类别提供了实现级证据——从源码审计可见 20+ fused op 如何跨越 RMSNorm/GEMM/AllReduce 算子边界 [tile-ai-tilert]，以及 CUDA graph 整图捕获如何消除逐 kernel launch 开销 [tile-ai-tilert]。TokenSpeed 的加入提供了一种结构化替代方案——通过 Placement 编译器在编译期静态规划通信、通过 C++ FSM 在调度期保证资源安全，不依赖 persistent kernel 也能逼近极致延迟 [tokenspeed]。Blink 的加入把 framework 类别的边界推向了一个新维度——不再是"如何在 GPU 内消除 kernel 边界"，而是"如何把整个 serving stack 的控制权从 host CPU 迁出"：前端委派给 SmartNIC/DPU、后端调度交给 GPU-resident persistent kernel，host CPU 在稳态推理中零参与 [2604.07609]。两个类别的交汇产生了核心张力——kernel 的硬件亲和性与 framework 的模型泛化性之间的 trade-off [kernel]。

3. 时间线 #

从 ConCCL 到 TokenSpeed 的 18 个月间，优化粒度从"隔离通信与计算的资源竞争"逐步扩大到"将整个推理管线融合为单一持续执行体"或"用编译期类型安全替代运行时协调" [framework]。TileRT 代码仓库的版本演进记录了融合方向的关键节点：从初始 baseline（v0.1.0）到 tile 调度优化获得 3–4× 加速（v0.1.1），到 MTP 投机解码的 590 tok/s（v0.1.2），再到双模型 500/600 tok/s 的生产性能（v0.1.3），平均每 2–3 个月一次里程碑式跃升 [tile-ai-tilert]。TokenSpeed 代表了并行发展的编译器路线：3 个月内从零搭建到 Pareto 优于 TRT-LLM，MLA 内核被上游社区（vLLM）采纳 [tokenspeed]。而 Blink 则把优化对象从"GPU 内的 launch 开销"扩展到"整条 host CPU 控制路径"——用一个 GPU-resident persistent kernel 取代 host 在 decode loop 中的角色、用 SmartNIC 承接前端，验证了 host CPU 可以完全退出稳态推理 [2604.07609]。

4. 技术谱系 #

谱系包含六个关键分支点：

编译策略分叉。GPUOS 选择 JIT 保持动态性——NVRTC 模板编译 + dual-slot aliasing 零停机热更新 [2604.17861]。Fleet 选择静态 task graph 预编译，获取跨 operator 寄存器分配等编译器优化空间 [2604.15379]。TileRT 博客声称 AOT 全模型静态展开，获取最大优化边界 [tilert-speed-scaling-law]。TokenSpeed 选择 Placement 注解 + 静态 SPMD 编译，在编译期规划通信但保持 Python 模型定义的灵活性 [tokenspeed]。四者是灵活性-性能 Pareto 曲线上的不同采样点。

博客叙事 vs 代码实现的分叉。TileRT 的技术谱系存在叙事层与实现层的显著分裂。博客描绘的路径从 Persistent Kernel 范式出发，走向 AOT 全模型静态展开为单个 Engine Kernel [tilert-speed-scaling-law]。但代码审计揭示的实际实现有双重谱系来源：一是从 Operator Fusion 传统沿袭的激进跨算子融合——20+ fused ops 将 RMSNorm+Projection+Quantization+AllReduce 四类操作跨边界融合 [tile-ai-tilert]；二是从 CUDA Graphs 发展而来的整图捕获-回放机制——prepare_money() 一次性捕获全部 61 层 forward pass [tile-ai-tilert]。CUDA graph 是录制-回放模型（kernel 在 replay 时仍由硬件 scheduler 逐个调度），而 persistent kernel 是 GPU 端常驻执行（永不退出，device-side polling 接收新任务），两者的 host-device 交互模型根本不同 [2604.17861]。

Persistent kernel vs 结构化编译的分叉。TokenSpeed 验证了一条不经过 persistent kernel 的路径：C++ FSM 用 std::variant<13 states> + RAII move 语义在编译期保证 KV cache 资源安全 [tokenspeed]，Placement 编译器自动插入通信操作 [tokenspeed]，5-band 内核注册表选择最优 kernel 实现 [tokenspeed]。这套方案不需要 GPU 端常驻 kernel——调度安全通过类型系统而非 device-side polling 保证，launch overhead 通过 CUDA graph + 优化 kernel 而非 persistent kernel 消除。它代表了从"消除 launch"到"让 launch 的代价相对于优化后的 kernel 可忽略"的思路转变。

Persistent kernel 与 CUDA graph 的合流 + 控制权外迁。在 TileRT 的叙事里，persistent kernel（device-side 常驻）和 CUDA graph（capture-replay）被描述为根本不同的两种 host-device 交互模型 [2604.17861]。Blink 恰恰把两者合流——它的 persistent GPU kernel 本身在 device 端 fire-and-forget 地 launch 预编译好的 CUDA graph（≈2 µs，比 host launch 快 5–8×），并用 window-based tail-launch 机制绕过 120 次 fire-and-forget 的未记录硬限制 [2604.07609] [2604.07609]。与此同时，Blink 把 GPUOS/Fleet 都保留在 host 侧的"谁拥有推理控制循环"这一问题彻底外迁：前端（HTTP/tokenization/请求管理/RDMA）交给 SmartNIC/DPU，后端 decode loop 交给 GPU-resident persistent kernel，两者通过 GPU 显存中的 lock-free ring buffer + per-slot 状态机 + atomic CAS 协调，host CPU 在稳态中完全不参与 [2604.07609] [2604.07609]。这是本主题中第一条把"持续执行"从 GPU 内部扩展到"跨 DPU-GPU 的 CPU-free serving stack"的路线。

Chiplet 感知与持续执行的汇流。NUMA-aware Attention 首先识别了 chiplet 架构下 L2 分区的性能悬崖 [2511.02132]。HipKittens 提供了 tile-level 的 XCD 调度算法（联合优化 L2+LLC，GEMM 加速 19%）[2511.08083]。Fleet 将两者统合进 persistent megakernel：kernel 边界消除后 L2 状态在整个 decode 周期得以保留 [2604.15379]。

通信策略从外部隔离到内部融合。ConCCL 通过 DMA offload 实现计算与通信的物理隔离 [2412.14335]。TileRT 代码走向融合——AllReduce 被嵌入每个 fused op 内部（UnProjOAllReduce、ExpertDownAllReduce 等），通信与计算在 tile 粒度交织 [tile-ai-tilert]。TokenSpeed 走结构化路线——Placement 编译器在组件边界自动插入 AllReduce/AllGather，支持融合优化如 FusedReduceNormOp [tokenspeed]。

5. 技术线交错 #

5.1 AOT 刚性 vs 运行时灵活性 vs 编译期类型安全 #

TileRT 博客描述的 AOT 编译在编译期静态规划 GPU 间通信的 tile 对齐和数据依赖，运行时无法动态调整 GPU 角色分配或通信拓扑 [tilert-speed-scaling-law]。代码仓库揭示了一层额外的刚性：变更 sampling 参数（temperature/top_p/top_k）需要完全 teardown 并重新捕获整个 CUDA graph [tile-ai-tilert]，因为 sampling 参数被 bake 到 graph 指令中。

GPUOS 走向对立面——NVRTC JIT + dual-slot aliasing + version counter 实现零停机 operator 热更新，新 operator 从编译到全设备可调用仅需 ms 级 [2604.17861]。但 template-based compilation 限制了表达能力，无法注入需要精细控制 shared memory layout 和 warp-level primitive 的自定义 kernel [2604.17861]。

Fleet 介于两者之间——task graph 编译时确定，但运行时由 per-XCD 软件 scheduler 动态分派 [2604.15379]。Mirage 超优化器尚未理解 Chiplet-task cost model，当前只能手动提供输入 task graph [2604.15379]。

TokenSpeed 提出了第四种方案——Placement 注解 + 静态 SPMD 编译。模型作者用轻量级 dataclass 声明并行策略（Replicate(ATTN_TP) / Partial(MOE_TP_EP)），编译器自动在组件边界插入通信操作 [tokenspeed]。这保持了模型定义的 Python 灵活性（新模型上线仅需注解而非 C++ 改写），同时在推理时实现零 dispatch 开销——比 DTensor 更轻量，比手写 AllReduce 更安全 [tokenspeed]。三个独立 ParallelGroup（ATTN_TP, DENSE_TP, MOE_TP_EP）支持不对称 TP——这是 TileRT 和 GPUOS 均未实现的能力 [tokenspeed]。

这四条路线映射到 ML 生态的不同阶段：GPUOS 的 JIT 方案适合 exploration-heavy 的研发阶段；Fleet 的 static+runtime-scheduler 适合 model-fixed 的部署阶段；TileRT 的全 AOT 适合 everything-fixed 的极致生产环境；TokenSpeed 的 Placement 编译器适合 fast-iteration 的 agent 推理场景（每 2 周新模型上线）[2604.17861]。

5.2 NVLink 域约束 vs Chiplet NUMA 约束 #

TileRT 在 8×B200 NVL 域内工作，核心约束是 TP synchronization 的通信开销——博客描述 GPU0 特化为 Sparse Indexer Worker、GPU1–7 为 MLA Workers [tilert-speed-scaling-law]，但代码仓库显示所有 8 个 GPU 构建完全相同的 Dsa 模型栈 [tile-ai-tilert]——博客描述的异构 Worker 角色分配在开源代码中未得到体现。TokenSpeed 同样在 B200 NVL 域内工作，但其 Placement 编译器通过不对称 TP 实现了更灵活的角色分配——Attention TP4 + MoE TP4EP 可以使用不同的并行度，编译器自动在边界 resharding [tokenspeed]。Fleet 在单块 MI350 的 8-XCD 域内工作，核心约束是 L2 不相干——M-major windowed traversal 让同 XCD workers 共享权重列 [2604.15379]。

异构特化在不同作用域展开：Fleet 是 XCD 粒度（同 GPU 内 8 个 XCD 分工），TileRT 是 GPU 粒度（8 块 GPU 分工），TokenSpeed 是组件粒度（attention 和 MoE 使用不同 TP 度）。HipKittens 的 XCD 调度算法（Algorithm 1）证明 chiplet 感知在 kernel 层面可带来 19% GEMM 加速 [2511.08083]，但只在单 kernel 级别——无法跨 kernel 维持 L2 affinity。Fleet 的 persistent megakernel 将此扩展到 decode 全路径 [2604.15379]。

5.3 Kernel 级 tile 抽象 vs Framework 级 tile pipeline vs 内核注册表 #

HipKittens 定义 tile 作为 kernel 内部的编程抽象——typed tiles 在 register/shared/global memory 间流转，PyTorch 风格算子操作这些 tiles [2511.08083]。TileRT 将 tile 提升为 framework 级调度抽象——compute、communication、async I/O 均为 tile-level task [tilert-speed-scaling-law]。代码审计揭示了 tile 抽象的实现细节：51 个 DsaTempVarIdx 固定索引构成静态 tensor 寄存器文件，所有 GPU 中间结果通过这些固定索引在 fused ops 间传递，零动态分配 [tile-ai-tilert]。

TokenSpeed 采用了不同的内核组织方案——5-band 优先级内核注册表 [tokenspeed]。select_kernel("attention", "decode", dtype=fp8, platform=SM100) 根据能力要求、数据类型、平台自动选择最优 kernel 实现。5 个优先级带（REFERENCE < PORTABLE < PERFORMANT < SPECIALIZED < PLUGIN）提供了"局部推理"能力——外部插件作者不需要审计所有内置注册就能确定覆盖优先级。这是 TileRT 的 monolithic fused op 和 GPUOS 的 device function pointer table 之外的第三种内核管理范式。

从 kernel-level tile 到 framework-level tile pipeline 再到 pluggable kernel registry，抽象边界的上移伴随灵活性-性能曲线的不同选择——这正是 kernel 和 framework 两个类别在 kernel boundary elimination 范式下的核心交错。

5.4 计算-通信分离 vs 融合 vs 编译器自动化 #

ConCCL 证明将通信 offload 到 DMA 引擎可完全消除 compute interference——通信和计算使用不同硬件执行器，互不干扰 [2412.14335]。TileRT 代码印证了"深度融合"策略——每个涉及跨 GPU 通信的 fused op 都以 AllReduce 后缀命名 [tile-ai-tilert]。TokenSpeed 的 Placement 编译器走了自动化路线——编译器在组件边界自动插入通信操作，并支持 FusedReduceNormOp（AllReduce + LayerNorm 合并）和 ResidualAllGatherOp（残差连接与 AllGather 重叠）等融合优化 [tokenspeed]。三种策略适用不同通信规模：ConCCL 在 ≥128 MB 时与 RCCL 持平 [2412.14335]，TileRT 的 pipeline 内通信适合 tile 粒度的细粒度 broadcast/reduce，TokenSpeed 的编译器自动化适合快速适配新模型架构。

5.5 博客叙事 vs 开源代码：持续执行的理想与现实 #

TileRT 是唯一同时拥有设计博客和开源代码仓库的持续执行系统，两者之间的裂缝构成了本主题中独特的技术张力。

Persistent Engine Kernel vs CUDA Graph。博客核心叙事是"AOT 编译将整个模型展开为单个 Persistent Engine Kernel，host 仅启动一次" [tilert-speed-scaling-law]。代码显示执行模型是 dsa_show_hands_prepare_money() 捕获 CUDA graph、dsa_show_hands() 单次回放 [tile-ai-tilert]。Persistent kernel（如 GPUOS/Fleet 采用的方案）是 launch 一次永不退出、通过 device-side polling 接收新任务 [2604.17861]；CUDA graph 则是将 kernel 序列录制为 DAG 后一次性 replay。TokenSpeed 同样使用 CUDA graph 捕获但不声称是 persistent kernel——其设计哲学是通过类型系统和编译器消除协调开销，而非通过 device-side scheduling [tokenspeed]。这一对比表明"CUDA graph + aggressive fusion"可能是比 true persistent kernel 更实用的 kernel boundary elimination 方案。

异构 Worker vs 对称 TP vs 不对称 TP。TileRT 博客详细描述 GPU0 = Sparse Indexer Worker、GPU1–7 = MLA Workers 的异构角色分配 [tilert-speed-scaling-law]。代码中 8 个 GPU 构建完全相同的模型栈 [tile-ai-tilert]。TokenSpeed 通过不对称 TP 实现了编译器级的异构——三个独立 ParallelGroup 允许 attention、dense MLP、MoE 使用不同并行度，编译器自动 resharding [tokenspeed]。这是一种比 Worker 级异构更细粒度的组件级异构。

闭源核心的可审计性困境。TileRT 的核心 tile 调度逻辑封装在闭源 C++ torch.ops.tilert.* 中 [tile-ai-tilert]。GPUOS 提供 3793 行全开源 [2604.17861]。TokenSpeed 以 MIT 开源 ~205K LOC [tokenspeed]，但 MLA 内核的最优路径依赖未开源的 .so binary。Fleet 代码未开源 [2604.15379]。Blink 至发表时无公开仓库（GPU backend ≈16K + DPU frontend ≈17K，共 ~33K 行 CUDA/C++）[2604.07609] [2604.07609]。可审计性从高到低：GPUOS > TokenSpeed > TileRT > Fleet ≈ Blink。

5.6 谁驱动持续执行循环：GPU-resident scheduler vs host CPU vs SmartNIC 委派 #

前面五篇工作虽然实现路线各异，却共享一个隐含前提：host CPU 仍是 serving stack 的所有者——GPUOS/Fleet 的 persistent kernel 在 GPU 上常驻，但请求接入、tokenization、批处理决策仍在 host [2604.17861]；TileRT/TokenSpeed 用 CUDA graph 消除 launch 开销，但 graph 的触发和结果回收仍经由 host 每步返回 [tokenspeed]。Blink 打破了这个前提，把"谁驱动持续执行循环"这个轴拉到了极限：host CPU 完全退出，控制权一分为二——SmartNIC/DPU 拥有前端（请求管理 + RDMA 传输 + tokenization），GPU-resident persistent kernel 拥有 decode loop（扫 ring buffer → CAS claim → 选/发 CUDA graph → 轮询完成 → 发布 token）[2604.07609]。

这条轴上的位置差异带来两个可量化的后果。其一是延迟结构：device-side fire-and-forget graph launch ≈2 µs，而 host launch 需 11–17 µs（快 5–8×），256 线程 1–5 µs 扫完 4096 个 ring buffer slot [2604.07609] [2604.07609]；同用 TensorRT 引擎的对照下，Blink 相对 TRT-LLM P99 TTFT 低 1.35–3.45×、吞吐最高高 37%，且 MoE 模型（compute-to-orchestration 比最低）收益最大 [2604.07609]。其二是干扰免疫：因为关键路径上没有 CPU，同机 pbzip2 + Ninja 干扰下 Blink 的 TTFT 膨胀仅 0.92–1.14×、吞吐保持 99–100%，而基线膨胀 1.54–18.84×、吞吐仅存 28–64% [2604.07609]。这与 GPUOS/Fleet 追求的"延迟下限"是正交的收益——它们仍暴露在 host 抖动之下，而 Blink 把 host 抖动这一维度整体移除。

代价同样清晰：Blink 依赖 BlueField-3 DPU + 200 Gbps RDMA 这一 NVIDIA 独家硬件组合，DPU 的采购与运维 TCO 未计入其能效核算 [2604.07609]；且其 120-launch 硬限制的规避依赖对 CUDA 未公开行为的逆向工程，无 API 合约保护 [2604.07609]。这让"控制权外迁"成为一条硬件耦合度极高、可移植性最低的持续执行分支。

6. 共识与分歧 #

6.1 共识 #

Inter-kernel overhead 是当前微 batch 推理的头号敌人。 GPUOS 量化了这一开销：null kernel launch 3–7 μs，100 ops/token 累计 300–700 μs 纯协调开销 [2604.17861]。TileRT 博客从系统角度验证：理论 ~1000 tok/s vs 实际几十 tok/s，差距一个数量级 [tilert-speed-scaling-law]。TileRT 代码以工程产出证明消除这一 overhead 的可行性——600/500/590 tok/s [tile-ai-tilert]。Fleet 从 chiplet 角度补充：每 token 近 250 次 launch，且每次 launch 都会 flush L2 以保证 device-sync 语义 [2604.15379]。TokenSpeed 从 agentic workload 角度确认：vLLM Python 调度器在高并发下成为瓶颈 [tokenspeed]。Blink 把这一开销的归因从"GPU 内 launch"上移到"host CPU 协调本身"——即便消除了 LLC 竞争，host orchestration 仍使 attention dispatch +104%、cudaLaunchKernel +115%、KV-cache dispatch +172%，而 GPU kernel 时间不变 [2604.07609] [2604.07609]。所有 9 个研究与实现都直接或间接地以消除 kernel 间/host-device 间协调开销为目标 [kernel] [framework]。

"消除 launch 的概念"是共识方向，但实现路径分化为 persistent execution 和 structured compilation 两大阵营。 GPUOS 进程 launch 一次后永不退出 [2604.17861]；Fleet megakernel 占满所有 CU 直到 decode 结束 [2604.15379]。TileRT 代码通过 CUDA graph 一次性捕获全部 61 层 forward [tile-ai-tilert]。TokenSpeed 通过 C++ FSM + CUDA graph 消除调度开销，不依赖 persistent kernel [tokenspeed]。Blink 用一个永不退出、在 device 端 fire-and-forget launch CUDA graph 的 persistent kernel 承接整个 decode loop，并把前端委派给 SmartNIC [2604.07609]。核心共识：问题不在"优化每次 launch"，而在"消除 launch 的概念" [2604.17861]——但达成方式可以是 device-side scheduling（GPUOS/Fleet/Blink）、compile-time safety（TokenSpeed），也可以是把控制权整体迁出 host CPU（Blink）。

Chiplet 架构下 L2 cache 不可再被视为统一资源。 NUMA-aware Attention 证明 round-robin 调度在 MI300X 上导致 L2 命中率 ~1% [2511.02132]；Fleet 证明无 chiplet-aware 调度时 decode 性能被 L2 复用率限死 [2604.15379]；HipKittens 证明联合优化 L2+LLC 可带来 19% GEMM 加速 [2511.08083]。

6.2 分歧 #

Persistent kernel vs 结构化编译 vs CUDA Graph + Fused Ops。 GPUOS 通过 NVRTC JIT 保持完全动态 shape/operator 支持 [2604.17861]。TileRT 代码揭示的实际执行模型是"手写 fused ops + CUDA graph 捕获" [tile-ai-tilert]。TokenSpeed 通过 Placement 编译器 + FSM 调度器 + CUDA graph 实现极致延迟 [tokenspeed]。三者根本假设不同——GPUOS 假设 workload 是"大量异构操作的动态混合流"，TileRT 假设是"固定模型在固定硬件上的确定性执行"，TokenSpeed 假设是"模型快速迭代但运行时模式可预测" [2604.17861]。GPUOS 的 dual-slot aliasing 支持零停机热更新，TileRT 的 sampling 参数变更需要完整 CUDA graph re-capture [tile-ai-tilert]，TokenSpeed 的 FSM Retraction 机制支持 memory pressure 下的优雅降级但不支持 hot operator injection [tokenspeed]——这是三条路线的根本性分叉。

Persistent kernel 的资源分配哲学。 GPUOS 极其保守——每 SM 仅 1 个 thread block（2–4% 总线程），定位"填谷"，大 GEMM 仍走传统路径 [2604.17861]。Fleet 走向另一极端——persistent megakernel 占满所有 256 CU [2604.15379]。Blink 与 GPUOS 同属保守派但更极致——其 persistent scheduler 仅占用单个 thread block（256 线程），只负责调度而不占用计算资源，实际推理仍由 device 端 launch 的 CUDA graph 承担 [2604.07609]；这使得调度与计算在资源上解耦，256 线程只做 ring buffer 扫描与 CAS claim。TileRT 代码通过 CUDA graph 整图捕获消除 launch 开销，但个别 kernel 仍是标准执行模型 [tile-ai-tilert]。TokenSpeed 不使用 persistent kernel——调度器在 C++ 中微秒级完成，执行在 Python 面通过 CUDA graph 完成 [tokenspeed]。技术根源在于 register pressure：Fleet 将所有 task 类型编进同一函数，寄存器用量取联合最大值，被迫只跑 1 wave/SIMD [2604.15379]；GPUOS 因只处理微操作避免了这一约束；TileRT 和 TokenSpeed 的 fused op / kernel 分别编译，寄存器用量独立优化 [tile-ai-tilert]。

Wave/Warp specialization 的有效性是 architecture-dependent。 TileRT 博客描述 warp specialization 用于 Engine Kernel 内部的三级特化 [tilert-speed-scaling-law]。HipKittens 在 AMD CDNA4 上明确拒绝 wave specialization（仅达 80% peak），转而采用 8-wave ping-pong——因为 AMD 的 SIMD 寄存器分配是编译时静态的，producer waves 空占寄存器却不贡献计算 [2511.08083]。而 NVIDIA Blackwell 上 warp specialization 仍然有效——AVO 使用 warp-specialized pipeline（MMA/Softmax/Correction/Load warps）达到 1668 TFLOPS [kernel]。这一分歧的根源是硬件寄存器分配机制的差异，不可跨 vendor 泛化。

硬件阵营分裂。 GPUOS 深度绑定 NVIDIA（NVRTC、CUDA Driver API、PTX）[2604.17861]。Fleet、NUMA-aware Attention、HipKittens 深度绑定 AMD CDNA [2604.15379] [2511.02132] [2511.08083]。TileRT 运行在 NVIDIA B200 上，代码中 PTX ISA 级 swizzle 硬编码 [tile-ai-tilert]。TokenSpeed 以 Blackwell 为主要目标，Hopper/MI350 标注 ongoing [tokenspeed]。HipKittens 和 ThunderKittens 已在 tile 接口层实现跨厂商统一 [2511.08083]，但底层调度必须根本性重写。当前 kernel boundary elimination 范式完全沿厂商线分裂，无跨平台方案存在。

7. 根本性困难 #

7.1 动态模型架构 vs 静态编译 #

MoE 模型中每 token 的 expert routing 是 input-dependent 的，从根本上挑战了 AOT 全模型静态编译。TileRT 代码揭示了其 MoE 处理方案：ExpertSelectUpGateSiLU（729 LOC）将 top-8 选择 + gate/up projection + SiLU 融合为单 kernel [tile-ai-tilert]——实用但限制了跨 expert 边界的优化空间 [tile-ai-tilert]。Fleet 的 M-major 协作假设在 MoE 下被直接破坏——不同 token 走不同 expert 权重 [2604.15379]。但对立视角同样成立：MoE 的 expert 稀疏激活与 chiplet 私有 L2 天然亲和——8 颗 XCD 恰好对应典型 MoE 的 expert 组数，将 Chiplet-task 从"XCD-local weight partition"扩展到"XCD-local expert affinity"是未探索但价值很高的方向 [2604.15379]。TokenSpeed 的不对称 TP 编译器通过三个独立 ParallelGroup 自动处理 MoE 的并行策略差异 [tokenspeed]，但编译器目前不理解 expert routing 的动态性。

GPUOS 的 template-based JIT 无法注入需要精细 shared memory 控制的 custom kernel [2604.17861]。预编译所有可能路径导致编译产物爆炸，在 kernel 内实现动态分发则本质回到运行时调度。需要 kernel 级的动态 expert 分发和 framework 级的 AOT/JIT 混合编译共同演进。预计难度：2–3 年。

7.2 Chiplet 执行的跨厂商统一 #

Fleet 的两级事件同步 + CDNA 三档 cache modifier 深度绑定 AMD ISA——buffer_wbl2 / flat_atomic_add sc0 sc1 / sc1=1 nt=1 的精确语义是 persistent megakernel 在 8-die GPU 上成立的唯一支点 [2604.15379]。NVIDIA Blackwell 仅有粗粒度 _cs/_cg/_ca 修饰符，无 XCD-scope 概念 [2604.15379]。HipKittens 和 ThunderKittens 已在 tile 接口层实现跨厂商统一 [2511.08083]，但底层调度必须根本性重写。

跨厂商统一需要统一 cache scope 语义和同步原语——这是 ISA 级别分歧，非软件层可消弭。预计难度：3–5 年，取决于硬件厂商标准化意愿。

7.3 硬件极限下的执行稳定性 #

TileRT 在生产流量下面临的核心挑战不是峰值性能而是执行稳定性。代码仓库揭示了若干脆弱性来源：整图 CUDA graph 捕获意味着 sampling 参数变更需要完全 teardown + re-capture [tile-ai-tilert]；batch=1 硬编码使得无法支持连续 batching 或 prefix caching [tile-ai-tilert]。TokenSpeed 的 Retraction FSM 提供了一种优雅降级方案——memory pressure 下将请求从 Decoding 状态转移到 Retracted 状态、KV 写回 host 后重新调度 [tokenspeed]——但这需要 13 个 FSM 状态的复杂性。Fleet 面临类似约束——register pressure 导致每次 L2 miss 直接 stall MFMA pipeline，L2 命中率从"锦上添花"升格为"生死攸关" [2604.15379]。

接近硬件极限的持续执行系统没有安全裕度。传统 kernel-per-op 模型中单个 kernel 失败只影响一个操作；persistent kernel 中任何异常都可能导致整个 decode 失败。预计难度：持续性挑战。

7.4 Persistent kernel 与集体通信的多 GPU 共存 #

GPUOS 的全部评估在单 GPU 完成 [2604.17861]，Fleet 同样仅评估单 GPU [2604.15379]。Blink 也是单 GPU、单节点系统，模型必须放入 GPU 显存，multi-GPU 被明确列为 future work（方向为 GPU-native collectives 或 GPU-initiated RDMA）[2604.07609] [2604.07609]——其 persistent scheduler 与外部 DPU 通过 RDMA-visible ring buffer 协调的模型若要跨 GPU 扩展，需解决跨 GPU persistent kernel 间同步（替代 NCCL 的 CPU-initiated collective）与 device-side KV-cache 管理 [2604.07609]。TileRT 代码是当前唯一在 multi-GPU（8×B200）上验证的持续执行实现 [tile-ai-tilert]，其方案是将 AllReduce 嵌入每个 fused op 内部 [tile-ai-tilert]——但核心调度逻辑闭源。TokenSpeed 通过 Placement 编译器实现 multi-GPU 通信的编译期规划 [tokenspeed]，但 PD disaggregation 功能尚未合并（preview 状态）[tokenspeed]。ConCCL 的 DMA offload 解决了通信侧——不再占 CU，GEMM 全速运行 [2412.14335]——但其 CPU 编排在小传输 (<32 MB) 时慢 4× [2412.14335]。

Production LLM serving 几乎必然是 multi-GPU（≥7B 模型需 TP≥2），persistent kernel 的 SM 常驻与 collective 的全 SM 需求之间缺乏硬件级隔离机制 [2604.17861]。结构化编译方案（TokenSpeed）通过编译期静态规划通信回避了这一问题，但牺牲了运行时灵活性；Blink 式的 CPU-free 路线则把难题进一步复合——一旦跨 GPU，就同时需要 device-side collective 与 multi-DPU 协调（哪个 DPU 做 master scheduler），且 device-side CUDA graph launch 目前只能作用于同 context 捕获的 graph [2604.07609] [2604.07609]。预计难度：1–2 年。

7.5 Persistent kernel + 量化 Attention 的融合空白 #

GPUOS 消除了微操作的 launch overhead [2604.17861]，SageAttention 系列加速了 attention 计算本身（INT8/FP8/FP4 量化达 2.1×–5× over FA2）[kernel]。两者融合的场景——在 persistent kernel 中直接 dispatch 量化 attention——尚未被探索。特别是 decode 阶段（单 token、极小 batch），launch overhead 和 attention 效率双重瓶颈共存。TokenSpeed 的 MLA q-head 折叠 [tokenspeed] 和 TileRT 的 FlashSparseMLA [tile-ai-tilert] 代表了 attention kernel 优化的不同方向，但均未与 persistent kernel 调度结合 [kernel]。预计难度：1–2 年。

8. 成熟度判断 #

整体判断：Kernel boundary elimination 范式从 research frontier 快速迈向 early production。两条并行路线正在竞赛：一是 TileRT 代表的"极致 fused ops + CUDA graph + 闭源调度器"路线，已有生产部署但通用性受限（batch=1 硬编码、单一硬件）[tile-ai-tilert]；二是 TokenSpeed 代表的"结构化编译 + 类型安全调度 + 开源内核"路线，尚处 preview 但生态影响力已现（MLA 内核被 vLLM 采纳）[tokenspeed]。GPUOS、Fleet 与 Blink 的 persistent kernel 方案提供了重要的概念验证，但向生产系统的跨越仍需解决 multi-GPU、dynamic batching、跨平台等根本性问题 [framework]。Blink 尤其把范式的边界从"GPU 内消除 launch"推进到"host CPU 整体退出稳态推理"，并首次给出干扰免疫这一正交收益的完整实证（CPU 干扰下吞吐保持 99–100% vs 基线 28–64%）[2604.07609]，但其对 BlueField-3 DPU 的硬件耦合与单 GPU 限制使其暂处研究原型阶段 [2604.07609]。趋势方向：加速中，且正从"persistent execution"这一狭义定义扩展为更广义的"kernel boundary elimination"乃至"CPU-free serving"。

9. 邻接 topic #

cluster-llm-deployment。持续执行当前局限于单卡或单 NVL 域。TileRT 代码虽然工作在 8×B200 上，但硬编码 num_devices=8 且仅支持 batch=1 [tile-ai-tilert]。TokenSpeed 的 PD disaggregation 功能尚未合并 [tokenspeed]。Fleet 在 bs ≥ 64 被 vLLM 反超 [2604.15379]，暗示集群级调度需要在 persistent 和 traditional 路径之间动态切换。Blink 同样是单 GPU 系统，且其 CPU-free 架构的集群化是最紧迫的开放方向——需要跨 GPU persistent kernel 间的 device-side collective 与 multi-DPU 协调 [2604.07609]；同时 Blink 把 RDMA ring buffer 从机间通信引入 intra-node DPU-GPU 协调，与 node 间 RDMA 通信方案分层组合后可构成"全链路 CPU-bypass 分布式推理"栈 [2604.07609]。Framework 类别的 DynaServe 用 micro-request 统一 colocation 和 disaggregation [framework]，为 kernel boundary elimination 系统的集群集成提供了调度框架参考。Blink 的干扰免疫（CPU 完全移出关键路径）还为 colocation 部署提供了新论据——高密度混部下持续执行系统不再受同机 CPU 干扰拖累 [2604.07609]。

agent-system。Agent serving（多轮、短追加、长上下文、DAG 依赖）落在持续执行最优的 bs=1–16 区间 [2604.15379]。TileRT 的"更快推理 = 更多 rollout = 更深推理链"叙事直接映射到 agent 场景的 test-time scaling 需求 [tilert-speed-scaling-law]。TokenSpeed 将 agentic workload 作为设计核心——C++ FSM 的 Retraction 机制专门处理 agent 长 context 的 memory pressure [tokenspeed]。Framework 类别中 ThunderAgent、Helium、Halo 等 agentic serving 系统 [framework] 与超低延迟引擎的集成是自然延伸——两个 topic 的边界在 agent workload 的 KV cache 管理和 tool call 延迟处模糊化。

attention-optimization。SageAttention 系列的量化 attention 加速（2.1×–5× over FA2）[kernel] 与 persistent kernel 调度的融合是未探索的组合空间 [kernel]。TokenSpeed 的 MLA q-head 折叠 [tokenspeed] 和 TileRT 的 FlashSparseMLA [tile-ai-tilert] 已在 attention kernel 维度做了优化，但与 persistent execution 的交叉仍有大量未覆盖区域。

潜在演化：如果硬件层面解决了 launch overhead（如 NVIDIA GB10 已部分实现 [2604.17861]），本 topic 中 persistent kernel 方向可能收敛为纯 chiplet 优化方向。如果 CUDA graph + aggressive fusion + structured compilation 被证明是比 true persistent kernel 更实用的方案（TileRT 代码和 TokenSpeed 均提供了证据），本 topic 可能从"persistent execution"重定义为"kernel boundary elimination"——一个更广泛的范式定义，涵盖从 device-side scheduling 到 compile-time safety 的完整谱系。Blink 进一步提示了第三个演化维度：如果 SmartNIC/DPU 委派前端 + GPU-resident 调度被证明能在生产中稳定运行，"控制权从 host CPU 整体外迁"（CPU-free serving）将成为该谱系中与 device-side scheduling、compile-time safety 并列的第三根主轴 [2604.07609]。

10. 参考 #