一、架构断裂：当“单片思维”撞上“芯粒现实”

现代GPU架构已转向基于芯粒（Chiplet）的多芯片设计，例如AMD Instinct MI300X/MI350和NVIDIA Blackwell。然而，主流的CUDA/HIP执行模型并未完全适应这一根本性的架构变迁。一个核心痛点是：编程模型缺乏直接的方法来表达工作组群之间的数据亲和性，或将计算任务限定到特定芯粒的内存层次中。

AMD的研究团队在论文《Fleet: Hierarchical Task-based Abstraction for Megakernels on Multi-Die GPUs》中揭示了这一矛盾。当GPU硬件进入“联邦制”的多芯粒时代，软件编程模型却仍停留在“中央集权”的单片逻辑中，这导致了显著的性能瓶颈。

以AMD Instinct MI350为例，其包含8个XCD（加速计算芯粒），每个XCD拥有私有的4MB L2缓存，且这些L2缓存之间没有硬件维护的一致性。所有XCD共享末级缓存（MALL）和HBM。这种非一致性的分区L2缓存架构是传统编程模型的核心挑战：缺乏芯粒感知的调度会导致权重数据在不同XCD的L2中被重复加载，造成宝贵的缓存空间和HBM带宽的浪费。

研究数据量化了这一代价：在典型的LLM解码过程中，线性层操作贡献了95%的时延，但其L2缓存命中率仅有16.4%。问题的根源并非缓存容量不足，而在于缺乏一种能够“感知”并“利用”芯粒边界进行协同调度的编程抽象。

Fleet提出的解决方案是引入一个名为 Chiplet-task 的新任务层级。该层级精确映射到XCD的边界，并通过持久化内核（Persistent Kernel）运行时，实现芯片粒感知的协同任务执行。这标志着从“单片思维”到“芯粒现实”编程范式的关键转变。

二、架构与编程模型的断裂：芯粒化硬件与扁平化软件

表 1：从单片 GPU 到芯粒 GPU 的内存层次与编程模型映射。Fleet 通过引入 Chiplet-task 抽象，填补了芯粒级硬件资源与设备级软件视图之间的关键空白。

在传统的单片 GPU 架构中，硬件内存层次与编程模型抽象之间存在清晰、对称的映射关系：寄存器对应 SIMD 执行，共享内存（LDS）对应 CU/SM，而高带宽内存（HBM）则作为全局内存供整个设备访问。L2 缓存作为设备级资源，对所有计算单元透明，程序员无需关心数据在 L2 中的具体位置。

然而，多芯粒架构打破了这种对称性。以采用 Chiplet 设计的 AMD CDNA3/CDNA4 架构为例，其计算核心由多个 XCD（加速计算裸片）组成，每个 XCD 拥有私有的 L2 缓存。硬件上，L2 已从“设备级”资源转变为“芯粒级”资源。但主流编程模型（如 HIP）仍将设备视为一个扁平的、统一的内存空间。 这导致调度器可能将不同的工作组随机分配到不同的 XCD 上，使得同一份权重数据在不同 XCD 的 L2 中被重复缓存。这不仅浪费了宝贵的片上缓存容量，也产生了冗余的 HBM 访问流量。

AMD Instinct MI300X 多芯片加速器架构示意图。其包含 8 个 XCD 计算裸片，每个配备私有 L2 缓存，通过高速 Infinity Fabric 互联并与 HBM 连接。

要理解这一断裂带来的具体影响，需剖析芯片的访存路径。当某个 XCD 上的计算单元发起全局内存加载请求时，会首先查询其本地 L2。若未命中，请求需通过芯片间互联（Infinity Fabric）路由至 IO Die，进而访问 HBM 或持有数据的其他 XCD。AMD 的缓存一致性协议由内存指令中的范围位控制。在默认的“波前范围”设置下，L2 缓存行被视为 XCD 本地所有，不会主动发起跨 XCD 的一致性探测或失效操作。这意味着，如果两个位于不同 XCD 上的计算单元读取同一份权重数据，系统会从 HBM 执行两次独立的加载，并在两个 L2 分区中各保留一份副本。 更严重的是，在内核边界执行的设备级同步栅栏（如 buffer_wbl2）会强制刷新所有脏缓存行以确保跨 XCD 可见性，这导致 L2 中缓存的权重数据无法在内核间保留，每次算子调用都需从 HBM 重新加载。

这种“架构-编程模型”的断裂在大型语言模型（LLM）的解码阶段被急剧放大。以 Qwen3-8B 模型为例，其单层权重矩阵大小约为 368MB（bf16 格式），而批大小为 128 时，激活矩阵仅约 1MB——内存访问流量主要由权重加载主导。当前主流服务系统（如 vLLM）为每个算子启动独立的内核，单 token 解码跨越 36 层需要近 250 次内核调用。每次内核调用后 L2 内容因同步而被清空，迫使下一个算子必须从 HBM 重新加载全部所需权重。

表 2：在 MI350 架构上，无芯粒感知调度时 LLM 解码的性能特征分析。数据显示，线性层操作占据绝大部分计算时间，但 L2 缓存利用率极低。

表 2 的数据量化了问题的严重性：线性层操作占据了 95% 的解码时延，但在缺乏协同调度的情况下，XCD 上 4MB 的 L2 缓存仅能实现约 16.4% 的命中率。这揭示了一个关键瓶颈：即使拥有高达 5.3TB/s 的 HBM 带宽，低效的 L2 利用率仍使得实际性能远低于硬件潜力，计算单元大量时间在等待数据从 HBM 加载。

三、Fleet 任务模型：为内存层次“量体裁衣”

Fleet 编程范式的核心创新之一是提出了一个与多芯粒 GPU 硬件层次精确匹配的多级任务模型。其设计哲学在于：任务应在与其资源需求相匹配的硬件粒度上进行定义和调度。

表 3：Fleet 的四级任务层次结构，每一级都映射到特定的硬件资源边界及其内存作用域。

Fleet 的任务模型打破了传统 GPU 编程中单一的网格调度模式，将控制权交还给程序员，使其能够显式地将计算映射到特定的内存范围。从细到粗，其四级任务分别为：

• 波前任务：在 64 个线程内执行，利用寄存器和 LDS 进行细粒度逐元素操作。
• CU 任务：占据单个计算单元，通过 LDS 在块内进行线程间通信。
• 芯粒任务：这是 Fleet 的核心创新抽象。 它将一个 XCD 上的所有计算单元（如 31 个 CU）编组为一个统一的执行单元，并显式地以该 XCD 的私有 L2 缓存（如 4MB）作为其共享的、可管理的数据工作集。
• 设备任务：协调多个芯粒任务（如 8 个），在全局 HBM 层面进行结果的聚合与同步。

这种分层设计不仅还原了被扁平编程模型所掩盖的硬件内存拓扑信息，更为后续的协同数据布局与调度优化提供了基础框架。

3.1 Chiplet-task：芯粒级抽象的深层意义

Chiplet-task 的价值远不止于增加了一个任务层级。它所解锁的关键能力是：允许程序员精确指定一个 XCD 内所有工作线程在 L2 缓存中的协同工作模式。

以一个 [M,K] × [K,N] 的矩阵乘法（GEMM）为例，Fleet 可以将输出矩阵按列划分为 8 个独立的子矩阵。每个芯粒任务负责计算其中一个 [M,K] × [K, N/8] = [M, N/8] 的切片。这种划分策略实现了多重优化：

1. 计算隔离：每个芯粒处理独立的权重列分区，无需跨 XCD 同步即可完成计算。
2. 缓存协同：芯粒内的所有工作线程围绕同一份权重分片进行访问。首个线程将权重块从 HBM 加载至 L2 后，后续线程可直接从 L2 命中，极大提升了缓存复用率。
3. 消除干扰：不同芯粒间的 L2 缓存完全隔离，避免了虚假共享和不必要的一致性协议开销。

这一设计同时解决了三个关键问题：
* 通过输出维度的列划分，天然消除了跨芯粒的通信依赖。
* 通过芯粒内线程协同访问，将多个线程的“缓存容量需求”聚合，有效放大了 L2 的可用性。
* 将调度单元从细粒度的 CU 提升至粗粒度的 XCD，显著减少了任务分发与管理的开销。

此外，Fleet 使用事件表达任务依赖关系，一个芯粒任务仅需一个事件即可代表整个 XCD 上所有工作线程的完成状态。这与传统模型中需要为每个工作线程或工作组管理事件相比，将跨芯粒的同步事件数量减少了数十倍（例如 31 倍），这是带来显著性能提升的重要因素之一。

3.2 协同权重分块：将 L2 命中率从 16.4% 提升至 54% 的核心策略

Fleet 提升 L2 利用率的核心是 M-major 窗口化遍历策略。

• 在传统的调度方式中，同一 XCD 上的不同工作线程可能被分配处理输出矩阵的不同区块，因此它们会加载不同的权重列。这导致多个不相关的权重数据流同时冲刷有限的 L2 缓存，命中率低下。
• Fleet 的 M-major 策略改变了遍历顺序：芯粒内的每个工作线程会先完整计算一个输出块 [m, n]，遍历完该块所需的所有 K 维度权重 [0:K, n] 和激活值后，再共同移至下一个权重列。 这意味着，在给定的时间窗口内，该芯粒上的所有线程都在访问相同的权重列。这种“步调一致”的访问模式，使得权重数据在 L2 中被加载一次后，即可被所有线程重复利用，形成了高效的“流水线式”L2 复用，从而将 L2 命中率从不足 20% 提升至超过 50%。

图 3：单个芯粒任务内的分块遍历顺序（4×6 输出分块网格，图示 4 个工作线程）。单元格色调表示访问的权重列，颜色越深表示时间步越晚。(a) N 主导遍历：连续分块沿权重列推进，并发工作线程加载不同权重数据，增加 L2 缓存压力；(b) M 主导遍历：连续分块沿激活行向下推进并共享同一权重列，提升 L2 时间局部性。该图对比了两种遍历方式的缓存效率差异。N 主导遍历导致线程访问不同权重，L2 无法复用；M 主导遍历使同芯粒线程按批处理维度遍历，共享权重列，首个线程加载后，后续线程可直接命中 L2。结合缓存流策略，该设计既保证了数据的短时复用，又避免了无用数据长期驻留，是 Fleet 提升内存效率的关键算法。

这一策略的数学本质可以用一个简洁的模型刻画。设 $w$ 为共享同一权重块的芯粒内工作线程数量，其中 $w = W times B$，$W$ 是每个 XCD 的工作线程数（MI350 上为 31），$B$ 是沿批处理维度的分块数。

在 M 主导遍历下，预期 L2 命中率为：

$$
text{Hit Rate} = frac{w – 1}{w}
$$

这个公式揭示了两个关键洞察：
* 当批大小较小时（如 bs=1，$B=1$），$w=31$，模型预测权重复用率为 0，实际测得的约 17% 命中率源于融合算子（如 SiLU 与 GEMM 融合）带来的激活数据复用，而非权重共享。
* 当批大小增至 32（$B=32$），$w=992$，理论命中率跳升至约 99.9%，实测为 51.0%。此差距源于 L2 容量限制，无法容纳所有协同访问的权重块。
* 当批大小继续增至 64（$B=64$），理论命中率趋近 100%，实测为 61.4%——差距进一步揭示了 L2 容量限制导致的权重块相互驱逐。

这一分析表明，Fleet 的优化效果是可预测、可建模的系统行为。其本质是一种“批大小驱动的时间局部性放大”：通过强制多个工作线程在时间上对齐权重访问模式，将原本平摊到各线程的独立访存流，合并为一条能被 L2 缓存有效截获的复用流。

2.3 缓存修饰符策略：对 L2 缓存的精细化控制

Fleet 对 AMD 缓存修饰符的策略运用实现了精细化的流量管理。CDNA3/CDNA4 架构允许每条内存指令携带范围位（SC1, SC0）和非时序位（NT），以精确控制数据在 L2 缓存中的分配和逐出行为。Fleet 将访存流量分为三类并配置差异化策略：

1. 权重加载：缓存流式（sc1=1, nt=1）。该策略指示 L2 临时分配缓存行以服务短期复用，但在需要空间时立即标记为可逐出。这恰好匹配权重在 M 主导遍历中的访问模式——权重块被连续的工作线程短时间内密集访问，但不会跨计算层持久保留。流式策略既捕获了协同复用的时间窗口，又防止权重数据污染 L2 中更需持久驻留的数据。
2. 激活存储：非时序（NT=1）。GEMM 输出、RMSNorm 结果和 SiLU 激活等中间数据的写回操作绕过 L2 缓存，直接写回 HBM。这避免了瞬态数据驱逐正在被协同复用的权重块——在标准缓存策略下，写分配会为每个存储操作分配 L2 缓存行，从而快速冲刷掉宝贵的权重缓存。
3. 调度器通信：差异化范围。跨 XCD 的事件轮询使用非时序加载，绕过可能过时的 L2 副本，直接读取 HBM 中的最新值；XCD 内部通信则使用易失性加载，充分利用共享 L2 的低延迟优势。

这一策略组合的本质是对 L2 缓存替换策略的“软件定义”改造。在硬件 LRU 策略下，任何访问都会将数据推向最近使用位置，流式访问的权重和瞬态的中间结果很容易相互驱逐。Fleet 通过显式标注每个内存访问的语义意图（复用权重、一次性激活、控制消息），使 L2 缓存的行为更贴近程序员的访问模式，而非依赖硬件的通用启发式规则。 论文数据显示，仅将 SiLU 激活函数融合进 gate+up GEMM 芯粒任务这一项优化，就将 bs=1 时的 L2 命中率从 9.4% 提升至 17.4%——这消除了中间缓冲区的读写循环，将激活数据保留在寄存器或本地数据共享存储器中，而不流经 L2。

图 2：(a) 标准块调度将同一 XCD 上的工作线程分配到不同的 GEMM 任务，导致 L2 缓存颠簸；(b) Fleet 调度在 XCD 间划分权重列，协调各 XCD 上的所有工作线程读取同一权重分区，将 L2 缺失转化为命中。该图对比展示了两种调度策略下的 L2 利用差异。标准调度打乱了芯粒内的数据局部性，引发 L2 频繁颠簸；Fleet 按权重列进行分区，使同芯粒线程协作访问同一数据分块，最大化 L2 的时间局部性。该设计直接针对芯粒 GPU 的 L2 缓存分区特点，将高延迟的 HBM 访存转化为低延迟的 L2 缓存访问，是 Fleet 降低延迟、减少数据流量的核心调度策略。

三、运行时系统设计：一个微型的 GPU 操作系统

图 4 呈现了 Fleet 系统架构的全貌。(a) 部分展示了单个 Transformer 层在批大小为 1 时的任务图对比。

图 4(a)：单 Transformer 层任务图对比。Fleet 将任务数从 1,407 个压缩至 543 个，降幅达 2.6 倍。

可以看出：
* 标准调度将每个 GEMM 运算分解为 96 至 256 个独立的计算单元任务（总计 1,407 个）。
* Fleet 仅用 8 个芯粒任务来表示每个 GEMM 运算（总计 543 个，减少 2.6 倍），并将 SiLU 激活函数融合进了 gate+up 芯粒任务中。

而图 4(b) 则揭示了运行时调度的核心机制：每个 XCD 有一个调度器工作组，其余计算单元作为工作线程。调度器从设备内存中的任务描述符结构拉取任务，并按芯粒本地亲和性进行分发。

图 4(b)：运行时调度机制。每个 XCD 的调度器将 CU 任务轮询分配给个体工作线程，或将芯粒任务广播给所有工作线程进行协同执行。任务队列和事件计数器驻留在 HBM 中，但会自然缓存在 XCD 本地的 L2 中；XCD 内部通信通过共享 L2 隐式保持一致，无需全局的获取/释放语义；跨 XCD 的信令则需要显式的 buffer_wbl2 内存栅栏。这张图揭示了 Fleet 高效的架构根源：它将调度和同步的“控制面”流量压缩到极致，使“数据面”的访存带宽几乎完全服务于权重和激活数据的传输。

图 5 则详细刻画了分级同步协议的时序逻辑。

图 5：分级同步协议。工作线程在无内存栅栏的情况下递增 XCD 本地 L2 缓存中的计数器（绿色部分）。仅每个 XCD 上的最后一个工作线程会发出单个 buffer_wbl2 内存栅栏并更新全局 HBM 中的计数器（红色部分）。调度器轮询全局计数器以调度下游任务。相比每个工作线程都进行全局信令，这一协议大幅减少了跨 XCD 的内存栅栏数量。这是 Fleet 在芯粒架构上实现高效任务协调的关键机制设计。

• 工作线程在 XCD 本地 L2 计数器上累加完成状态，无需内存栅栏（图中绿色部分）。
• 仅每个 XCD 的最后一个工作线程发出单个 buffer_wbl2 内存栅栏并更新全局 HBM 计数器（图中红色部分）。
• 调度器轮询全局计数器以触发下游任务。

这一设计将跨 XCD 内存栅栏的数量从每工作线程一次降低为每芯粒每事件一次。

3.1 持久化内核运行时：一次性启动，持续运行

Fleet 的运行时延续并深化了持久化内核的理念，并针对芯粒架构进行了关键改造。

在传统 GPU 编程模型中，每个算子对应一次内核启动——启动开销包括主机到设备的命令推送、设备端调度器初始化以及内核完成后的同步。对于 LLM 解码这类需要数百次算子调用的场景，累积开销不容忽视。现有方案虽可通过 CUDA Graph 降低启动成本，但 Graph 需要为每个批大小预先录制，且跨内核的数据复用（尤其是 L2 缓存驻留）会因内核边界处的隐式同步而丧失。

Fleet 的解决方案是，一次性启动一个占据全部计算资源的持久化内核，该内核在整个推理生命周期内驻留于 GPU。 具体组织方式为：每个 XCD 指定一个工作组（例如，32 个计算单元中占用 1 个）作为调度器，其余 31 个作为工作线程。调度器通过读取硬件 HW_ID 寄存器获知其所在的 XCD 身份，并在全局内存中维护每个工作线程的任务队列。这 8 个调度器并行工作，各自负责本芯粒内的任务分发，形成一种“联邦式调度”拓扑。

一个常被忽视的细节是：调度器占用的 CU（8/256=3.1%）不参与计算。 由于调度器仅执行轻量级控制操作（如读取计数器、写入队列、指针运算），工作线程极少因等待任务分配而停顿。但在任务执行时间极短的场景下，调度开销可能成为瓶颈，这为后续优化留下了空间。

3.2 分级同步：将跨芯粒栅栏减少两个数量级

分级同步是 Fleet 运行时设计中技术深度最高的部分。要理解其价值，需先了解跨芯粒同步的代价。在 MI350X 上，一个设备范围的原子操作必须对所有 8 个 XCD 可见，这需要跨芯粒互联结构的一致性流量，其延迟远高于 L2 本地操作。更严重的是，设备范围的栅栏（如 buffer_wbl2）会刷新本地 L2 中所有脏缓存行以确保跨 XCD 可见性——该操作的开销与脏行数量成正比，并会阻塞后续内存访问直至刷新完成。

Fleet 的分级同步协议通过将通信模式匹配到最窄的充分内存范围，系统性地规避了这些开销：

1. 任务队列（只读）：任务描述符的全局队列在内核启动前填充完毕，执行期间不可变。工作线程和调度器读取任务元数据无需任何同步。
2. 调度器 → 工作线程分派（L2 本地）：调度器使用设备范围的存储和原子操作，写入本芯粒内工作线程的每工作线程队列。由于调度器及其工作线程位于同一 XCD，这些访问在本地 L2 中解析，无跨 XCD 一致性流量。
3. 工作线程 ↔ 工作线程协调（L2 本地）：芯粒任务内的工作线程使用设备范围原子操作，在每 XCD 计数器上累加子任务完成数。这些操作同样在本地 L2 解析。无需栅栏，因为所有参与工作线程共享同一 L2 分区。
4. XCD → 全局信令（GPU 范围）：仅当某个 XCD 上的最后一个工作线程完成时，才发出单个 threadfence 并更新全局事件计数器（使用 GPU 范围原子操作 flat_atomic_add with sc0 sc1）。这把昂贵的跨 XCD 一致性成本均摊到整个 XCD 的所有工作线程上。一旦全局计数器达到阈值，完成的工作线程将事件入队到负责调度器的队列，从而触发下游任务。

这一协议的效果体现在数量级上：对于芯粒任务，两级计数为每个事件最多触发每个 XCD 一次栅栏，在无任务的 XCD 上则为零次栅栏。Fleet 的线性事件精确对应 8 个任务（每 XCD 一个），因此每个事件仅触发恰好 8 次栅栏。 相比之下，若每个工作线程都进行全局信令，将产生 256 次跨 XCD 栅栏——分级同步将其减少了 31 倍。这一减少直接转化为更低的任务切换延迟和更高的调度吞吐。

3.3 任务注册与代码生成：编译器支持

Fleet 的代码生成流程基于 Mirage 编译器扩展。Mirage 在模型加载时将抽象任务图编译为运行时描述，其中包含每个节点的库代码。为生成基于芯粒任务的图，Fleet 向 Mirage 编译器提供不同的输入代码，但暂未启用编译器的超优化搜索——论文明确指出这一方向留待未来工作，仅需为芯粒任务设计合适的成本模型即可。

当算子被分解为芯粒任务时，编译器在代码生成阶段计算每个分块的指针。例如，对于跨 8 个 XCD 的权重分区，任务 k 接收的指针为 base_ptr + k × (N/8) × K × sizeof(bf16)。这意味着 GPU 内核代码无需感知正在执行哪个芯粒任务——它仅读取指向当前数据块的指针。在每个芯粒任务内，所有 CU 接收相同的指针，调度器也无需在任务分派时进行指针运算。 取而代之的是，从 Mirage 编译器发出的调用被修改为包含 XCD 内的分块索引，生成对 _execute_xcd_task(tile_idx) 的调用，库实现负责根据分块索引为每个 CU 执行不同行为。

这种设计体现了编译器与运行时之间清晰的职责分离：编译器负责静态的数据分区和指针预计算，运行时负责动态的任务调度和同步协调。两者通过简洁的接口（分块索引）耦合，既保持了灵活性，又避免了运行时的复杂指针运算开销。

四、端到端评估：数据背后的技术叙事

图 6 展示了四条性能曲线的对比：vLLM（外部基线，使用 hipBLASLt GEMM，每个算子独立内核）、Mirage MPK（内部基线，持久化内核但对芯粒无感知）、Fleet M-split（芯粒任务调度但无协同权重共享）、Fleet M-tile（芯粒任务调度 + 协同权重共享）。

图 6：MI350 上 Qwen3-8B（稠密，bf16）的纯解码每输出 token 耗时（ms/token）。Fleet（M-tile）采用 M 主导窗口遍历，XCD 内连续工作线程在 L2 中共享权重分块；Fleet（M-split）为每个 XCD 分配独立 M 分块，无工作线程间权重共享。两者均使用 16×64×256 CK GEMM 分块与缓存流权重加载。该图量化全批次性能：小批次下，Fleet 两种模式均比 vLLM 快约 1.5 倍，提速来自持久化内核与调度优化；大批次下，M-tile 因协作分块使 L2 复用最大化，性能远超 M-split 与 Mirage。这验证了小批次的瓶颈是调度开销，大批次的瓶颈是 L2 利用率，Fleet 的双策略覆盖了全场景，在 LLM 推理延迟优化上优势显著。

图 7 的 Roofline 分析为这个故事提供了数学注脚：Fleet 的 L2 复用将有效算术强度从名义值（批大小）推高至 B / (1 - L2 命中率)。在 bs=32 时，51% 的命中率使有效强度从 32 翻倍至 65 FLOP/字节，将操作点从带宽瓶颈区推向计算瓶颈区的边缘。

图 7：MI350 上 GEMM 运算的峰值性能分析（bf16 MFMA）。标准调度以标称算术强度 AI=B（批次大小，FLOP/byte）运行。Fleet 的 L2 复用通过减少 HBM 流量提升有效算术强度：AIeff = B / (1 – L2 命中率)。批次 32、L2 命中率 51% 时，Fleet 将有效 AI 从 32 提升至 65，向峰值点（245）右移 2.0 倍。该峰值分析证明 Fleet 通过 L2 复用提升有效计算强度，将内存绑定任务向计算绑定转移。标准调度受限于 HBM 带宽，有效强度低；Fleet 协作分块提升 L2 命中率，减少 HBM 依赖，有效强度翻倍。从架构层面验证了芯粒感知缓存复用是突破多芯粒 GPU 带宽瓶颈的核心路径。

4.1 实验设置与基线选择

评估在单卡 AMD Instinct MI350X 系统上进行。该设备拥有 256 个 CU（分布于 8 个 XCD，每 XCD 32 CU）、每 XCD 4MB L2、256MB 共享 MALL（Infinity Cache）、288GB HBM3。对比两条基线：

• Mirage MPK（移植至 MI350X）：芯粒无感知的持久化内核，使用标准 2D GEMM 分块，无协同调度。作为内部基线，用于隔离 Fleet 芯粒感知优化的效果。
• vLLM v0.17.2 on ROCm（–enforce-eager）：标准的每算子内核启动，无持久化内核。作为外部基线，代表当前服务框架的 SOTA 水平。

性能指标通过 HIP 级计时 API、GPU 片上周期计数器、性能计数器和底层追踪的组合采集。所有端到端延迟结果使用 64 个输入 token 和 1024 个输出 token，报告仅解码 TPOT（每输出 token 时间），通过 GPU 侧每迭代时间戳（s_memrealtime）测量，排除了预填充阶段。

4.2 延迟分解：调度开销 vs. L2 复用

在批大小为1时，vLLM的TPOT为10.51ms，Mirage MPK为7.83ms（实现了1.34倍加速），而Fleet M-tile与M-split方案则分别达到6.82ms和6.73ms（相较于vLLM分别获得1.54倍和1.56倍加速）。此阶段的性能提升主要源于持久化内核消除了内核启动开销，以及芯粒任务调度减少了调度器与工作单元之间的通信开销。

表4的数据进一步证实了上述分析：
* 在批大小为1至8时，三种配置的L2缓存命中率与HBM读取流量高度一致（命中率约16-24%，HBM读取约2400-2800 GB）。这表明该阶段的加速完全来自调度优化：芯粒任务调度将每次GEMM运算的调度分派，从96-256次CU级别调度压缩为8次XCD级别调度。
* 当批大小增至32时，情况发生变化。M-tile方案通过协同权重共享激活，L2命中率达到51.0%，HBM读取流量降至3,190GB（相较于Mirage的3,906GB减少18%），TPOT为12.35ms（相对Mirage的15.62ms提速1.27倍）。而M-split方案的L2命中率为39.5%，HBM读取流量反而增至4,295GB，TPOT为13.37ms（加速比降至1.17倍）。两者约1ms的延迟差距完全归因于L2权重复用。
* 在批大小为64时，协同复用的主导作用更加显著。
* M-tile方案的L2命中率攀升至61.4%，TPOT降至18.61ms（相对Mirage的24.10ms提速1.30倍），HBM读取流量从基线6,203GB锐减至3,925GB（降幅达37%）。
* M-split方案的性能则跌落至几乎与Mirage持平（TPOT 23.40ms，仅1.03倍加速），HBM读取流量增至7,426GB。

表 4：Qwen3-8B在MI350上的L2缓存命中率与HBM流量（解码1024个令牌）。HBM读/写流量以Mirage方案为基准进行归一化（1.00x）。数据通过rocprofiler-sdk设备计数模式测量。此表揭示了三个关键信息：(1) 批大小1-8时，三者的L2命中率与HBM读流量高度一致，印证了此阶段加速完全源于调度优化；(2) 批大小32-64时，M-tile的命中率显著高于其他两者，且HBM读流量大幅降低，量化了协同复用的独立贡献；(3) M-split在批大小≥32时HBM读流量不降反升，说明仅有调度优化而无协同复用，可能因缓存干扰而导致性能劣化。

值得注意的是，当批大小≥32时，vLLM的延迟曲线趋于平坦（约11-12ms），这是因为其底层hipBLASLt的GEMM内核采用了波前级K-split优化，在计算密集区域表现优异。Fleet M-tile方案在批大小32时仍比vLLM快1.06倍，但在批大小64时被反超——论文指出，其当前持久化内核的GEMM实现尚未集成K-split优化与注意力优化。这一局限性将在后续章节深入探讨。

4.3 Roofline模型视角下的优化本质

Roofline模型为理解Fleet的优化提供了宏观框架。名义算术强度（AI）为批大小（FLOP/字节），即每从HBM加载一字节权重所执行的浮点操作数。在标准调度下，批大小32时的AI为32，远低于MI350的Roofline脊点（约245），计算受限于内存带宽。

Fleet的L2缓存复用改变了有效算术强度的计算方式。设L2命中率为 ( h )，则实际需要从HBM加载的权重比例降为 ( 1-h )。有效算术强度 ( AI_{eff} ) 变为：
[ AI_{eff} = frac{B}{1-h} quad text{(FLOP/字节)} ]
当批大小 ( B=32 ) 且命中率 ( h=51% ) 时，( AI_{eff} approx 65 )，是名义值（32）的两倍。这使得运算点在Roofline图上向右平移，更接近计算瓶颈区。图7中的实测数据点（空心标记）与理论预测（实心标记）高度吻合，验证了该分析框架的有效性。

图 7：MI350上GEMM运算的Roofline分析（bf16 MFMA）。标准调度以名义算术强度 ( AI = B )（批大小，FLOP/字节）运行。Fleet的L2缓存复用通过减少HBM流量提升了有效算术强度：( AI_{eff} = B / (1 – text{L2命中率}) )。在批大小32、L2命中率51%的情况下，Fleet将有效AI从32提升至65，向峰值点（245）右移了2.0倍。该分析证明，Fleet通过L2缓存复用提升有效计算强度，将内存受限任务向计算受限方向转移。

从Roofline视角看，Fleet优化的本质是“通过缓存复用，将访存密集型操作伪装成计算密集型操作”。它并未改变硬件的峰值带宽或算力，而是通过减少有效的HBM访问流量，使运算在Roofline图上的投影位置发生横向移动。这一洞察的价值超越了LLM推理本身——任何具有数据复用潜力但在扁平调度下受限于带宽瓶颈的工作负载，都可能从类似的芯粒感知调度中获益。

4.4 逐GEMM权重分析与L2命中率模型验证

表5列出了Qwen3-8B模型每层中四个关键GEMM运算的权重尺寸。

表 5：Qwen3-8B各GEMM权重尺寸（bf16）。L2窗口指每个工作单元（worker）的活跃工作集（单个K维分块）。此表从权重尺寸角度解释了协同分块为何能在超出L2总容量的情况下仍取得高命中率：关键在于活跃工作集（每个worker约32KB × 31个worker ≈ 1MB）远小于L2容量（4MB），而M-major遍历顺序确保了这1MB的工作集在时间窗口内被充分复用。表中对MALL（256MB Infinity Cache）的标注“否（1.4×）”表明，即使末级缓存也无法完全容纳所有权重，凸显了L2级优化的必要性。

qkv_proj、o_proj、gate_up_proj、down_proj四个GEMM运算在每个XCD上的分区大小分别为6MB、4MB、24MB和12MB。这些尺寸均无法完整放入4MB的L2缓存。然而，协同分块的关键在于：实际驻留在L2中的是“活跃工作集”——即每个工作单元当前正在处理的K维分块（约32KB）。31个工作单元合计约1MB，远小于L2总容量。

表5同时回答了一个关键问题：MALL（256MB Infinity Cache）能否容纳所有权重？ 所有四层GEMM在每个XCD上的权重总和为46MB，8个XCD合计368MB，超出MALL容量（256MB）约1.4倍。这意味着MALL无法作为权重的完全驻留层，L2级别的协同复用仍然至关重要。

前文提出的命中率模型（公式1）在实验数据中获得了良好验证：
* 批大小 ( B=1 ) 时，模型预测复用率为零，实测命中率17%主要来自融合算子带来的激活复用（未融合时约9%）。
* 批大小 ( B=32 ) 时，模型预测命中率50%，实测为51%。
* 批大小 ( B=64 ) 时，模型预测命中率75%，实测为61%。

实测值与预测值的差距主要源于L2容量限制导致的权重块相互驱逐，这是简化模型未考虑的二级效应。

总结而言，Fleet通过芯粒任务（Chiplet-task）抽象和M-major协同权重分块，在MI350上实现了显著优化：将Qwen3-8B的解码延迟降低至vLLM的1.3-1.5倍（批大小1-8）；在批大小32时，将L2命中率从38.9%提升至51.0%，HBM读取流量减少18%，相对芯粒无感知基线获得1.27倍加速。

Roofline分析进一步揭示了优化的本质：通过缓存复用，将有效算术强度翻倍，使运算从内存带宽瓶颈区向计算瓶颈区平移。

五、相关工作：处于“图计算”与“芯粒感知”的交叉点

Fleet的研究贡献需置于三条技术脉络的交叉点上来审视：LLM推理的Megakernel化、GPU缓存感知调度、以及芯粒架构下的系统软件适配。本节通过对比分析，厘清Fleet与现有工作的继承、差异与突破。

5.1 Megakernel 路线：从 FlashAttention 到全模型融合

Megakernel 范式在 LLM 推理中的演进路径清晰可循。
* FlashAttention 开创了 IO 感知的算子融合先河，将注意力机制的多个步骤融合为单一内核，显著减少了高带宽内存的读写。
* HazyResearch 的 Megakernel 将这一思路推向极致，将整个 Llama 解码器融合进一个持久化内核，通过片上 GPU 解释器编排算子执行，在批大小为 1 时达到了 H100 内存带宽的 78%。
* FlashFormer 采用了类似的全模型融合方案，使用流水线共享缓冲区。
* Mirage MPK 引入了编译器生成的任务图和事件驱动同步机制，在 A100/H100 上相对于 SGLang/vLLM 取得了 1.0-1.7 倍的加速。

这些工作的共同前提是“单片 GPU 假设”——它们都面向具有统一二级缓存的 NVIDIA GPU，无需考虑芯粒边界。Fleet 的价值不在于发明了 Megakernel 或持久化内核，而在于识别出当这一范式迁移至芯粒架构时，扁平的任务调度将导致严重的二级缓存碎片化和利用率塌陷。Fleet 的芯粒任务调度是对 Megakernel 范式的“架构适配层”，使其能够跨越从单片到芯粒的硬件断层。

5.2 缓存感知与芯粒感知调度：从启发式到显式控制

线程块交错调度和 CUTLASS 持久化 GEMM 通过调整工作组访问模式来改善二级缓存复用，但它们假设统一的二级缓存。
* HipKittens 将 ThunderKittens 移植至 AMD CDNA3/CDNA4，引入了 XCD 分组机制，在独立 GEMM 上取得了最高 19% 的性能提升。

然而，这些方法都局限于“每内核”优化——它们优化的是单个内核内的二级缓存行为，无法跨算子保留缓存状态。

Fleet 的突破在于将芯粒感知从“单内核优化”提升至“全模型持久化运行时”层面。它通过分级同步协议，使得权重数据能够在跨越多个算子的持久化内核生命周期中驻留于二级缓存，而不被内核边界处的栅栏冲刷。这一能力是独立内核优化无法企及的——因为内核边界本身就是状态破坏性的。

5.3 线程块集群与拓扑感知执行：硬件支持的可能性

NVIDIA 在 Hopper/Blackwell 中引入的线程块集群提供了一定的硬件支持：协同调度同一图形处理集群上的块，提供集群范围栅栏和分布式共享内存，以及张量内存访问多播。然而，图形处理集群是比芯粒更细的粒度，且缺乏与二级缓存层级的直接对应。Fleet 的芯粒任务模型允许线程跨整个二级缓存范围协调，且完全以软件实现，不依赖特定硬件特性。

两者的关系是互补而非竞争：线程块集群为更细粒度的协作提供硬件加速，Fleet 为更粗粒度的芯粒级协调提供软件抽象。未来的 GPU 架构若能将 Fleet 式的芯粒感知原语固化为硬件支持，将可能进一步降低同步开销。

5.4 LLM 服务系统：vLLM 与 SGLang 的局限

vLLM 和 SGLang 作为主流开源服务框架，在内存管理和请求调度层面贡献卓著。然而，也存在如下局限：
* 它们在内核执行层面依赖 CUDA/HIP 图捕获或直接调用 rocBLAS/hipBLASLt。
* 图捕获虽然减少了启动开销，但仍需为每个批大小单独录制，且内核边界处的二级缓存刷新不可避免。
* 底层库虽可实现某些二级缓存优化，但难以表达应用级的数据局部性意图。

Fleet 与这些服务系统的关系是“互补的层次”：vLLM/SGLang 管理“哪些请求一起处理”（批形成、KV 缓存管理），Fleet 管理“如何在芯片上执行这些批次”（芯粒级任务调度、二级缓存协同复用）。两者可以集成——服务系统将批处理好的请求交给 Fleet 持久化内核执行，享受其芯粒感知优化，同时保留上层调度灵活性。

5.5 Mirage 超优化器：融合什么 vs. 调度何处

Fleet 直接构建于 Mirage 的 MPK 运行时之上，但添加了芯粒感知维度。两者的分工清晰：Mirage 的 muGraph 超优化决定“哪些算子应该融合”，Fleet 决定“融合后的任务应该调度到哪个芯粒、以何种协同模式执行”。论文明确指出两者的互补性，并提到编译器驱动的超优化（仅需为芯粒任务设计成本模型）是未来工作方向。

这一分工折射出 AI 编译器领域的一个深层趋势：优化决策正从“单一层级”向“跨层级协同”演进。传统编译器关注指令级和寄存器级优化，Mirage 将优化提升至算子图级，Fleet 进一步将硬件拓扑信息（芯粒边界、二级缓存分区）纳入优化空间。这种“从逻辑到物理”的优化粒度下沉，是弥合软件抽象与硬件实现之间日益扩大的鸿沟的必然路径。

六、结论与展望

6.1 结论总结

Fleet 提出并验证了一套面向多芯粒 GPU 的层次化任务模型与运行时系统。其核心贡献体现在三个层面：

抽象层面：引入 Chiplet-task 作为编程模型的新原语，填补了计算单元/工作组与设备/网格之间的语义空缺。这一抽象将芯粒边界从“隐式的性能影响因素”提升为“显式的可编程资源”，使程序员能够精确控制工作组在二级缓存中的协作模式。

机制层面：设计了分级同步协议，将跨芯粒栅栏数量从每工作项一次压缩至每芯粒每事件一次（减少 31 倍）。通过将同步范围精确匹配到通信需求（二级缓存本地用于芯粒内协调，GPU 全局仅用于跨芯粒汇聚），系统性规避了芯粒架构下昂贵的一致性流量。

算法层面：提出 M-major 窗口化权重遍历策略，使同一芯粒内的多个工作项在时间窗口内对齐权重访问，将二级缓存命中率从 12-39% 提升至 51-61%。Roofline 分析揭示该策略的本质是将有效算术强度翻倍，使操作从带宽瓶颈区向计算瓶颈区平移。

在 AMD Instinct MI350X 上，Fleet 将 Qwen3-8B 的解码延迟降低至 vLLM 的 1.3-1.5 倍（批大小=1-8），相对芯粒无感知基线提速 1.27-1.30 倍（批大小=32-64）。这些数据证明，软件管理的调度能够在无需硬件改动的情况下，有效恢复被芯粒架构碎片化的二级缓存局部性。

6.2 进阶分析

抛开论文作者的自然倾向，从问题解决的底层逻辑进行冷峻审视，Fleet 的工作存在若干需要坦诚面对的边界条件与潜在问题。

根本性 vs. 缓解性

Fleet 是否从根本上解决了芯粒 GPU 的二级缓存割裂问题，还是在特定假设下的阶段性缓解？

答案是后者。Fleet 的核心策略——M-major 遍历与协同分块——依赖于一个关键前提：批大小足够大，使得协同复用方能激活。当批大小为 1 时（这是 LLM 服务中常见的交互式场景），权重复用完全无法发生，加速仅来自调度开销的降低。

论文数据显示批大小为 1 时 Fleet 的二级缓存命中率为 16.9%，与 Mirage 的 16.4% 几乎持平。这意味着 Fleet 对“内存墙”的根本性突破仅存在于批处理场景，而在对延迟最敏感的单请求场景中，其优化效果并未触及二级缓存复用这一核心机制。真正根本性的解决方案可能需要硬件层面的变革——例如支持跨芯粒的二级缓存共享或一致性协议优化。

方法论边界与隐性成本

论文实验设计存在若干未充分讨论的边界条件。

局限性分析

论文展示的评估结果主要基于单GPU、单模型（Qwen3-8B密集架构）的场景。其芯粒分区策略在更广泛场景下的有效性仍有待验证：
* 对于需要张量并行的大模型、MoE架构的稀疏模型或混合精度（FP8/INT8）推理，目前缺乏数据支撑。
* 论文承认，持久化内核将所有任务编译进单一函数，导致寄存器压力增大，将占用率压制为每SIMD单波，丧失了波切换的延迟隐藏能力。“每一L2未命中直接阻塞MFMA流水线”这一代价缺乏量化分析，在L2命中率更差的场景下，这种占用率损失是否会反噬性能尚不明确。
* 调度器占用的8个CU（约3.1%算力）在算子执行时间极短的情况下（如逐元素激活），其开销可能从“可忽略”变为“主导性”。

性能反超的场景分析

图6中一个值得关注但未充分讨论的现象是：在批大小（bs）为64时，vLLM的延迟（约11.5ms）反超了Fleet M-tile方案（18.61ms）。

图6：MI350上Qwen3-8B（稠密，bf16）的纯解码每输出token耗时（ms/token）。Fleet（M分块）采用M主导窗口遍历，XCD内连续工作线程在L2中共享权重分块；Fleet（M分割）为每个XCD分配独立M分块，无工作线程间权重共享。两者均使用16×64×256 CK GEMM分块与缓存流权重加载。

作者将此归因于“持久化内核的GEMM实现尚未集成K-split优化”。这揭示了一个深层问题：Fleet的软件调度栈在追求跨算子L2数据驻留的同时，可能牺牲了单个GEMM内核内部的极致优化能力（如波级K-split、自动调优的分块尺寸）。这构成了 “跨算子复用”与“单算子极致优化”之间的张力——当批大小增大到使单算子成为计算瓶颈时，算子内部的优化优先级上升，Fleet的持久化融合策略可能不再是帕累托最优。

由此推论，最优策略可能是混合式的：在小批大小下使用Fleet式持久化融合以最大化数据复用，在批大小突破阈值后切换至hipBLASLt式独立优化内核以追求峰值算力。这种自适应调度机制在论文中未被探讨。

可移植性的“软硬边界”

Fleet对AMD特定缓存修饰符（sc1/nt位）的依赖，意味着其L2策略高度绑定CDNA3/CDNA4架构。尽管作者声称“任何具有分区L2缓存的芯粒GPU都可实现Chiplet-task”，这一论断在抽象层面成立，但在具体实现层面掩盖了大量平台相关的适配工作。例如，NVIDIA Blackwell虽有多Die设计，但其缓存一致性模型和内存指令修饰符与AMD存在本质差异，同步原语、缓存控制指令、调度器硬件ID获取方式等均需逐平台重新实现。

未来工作展望

论文指明的方向

论文在多个位置明确了未来的研究方向：
1. 编译器驱动的超优化：当前Fleet依赖手工编写输入代码生成芯粒任务图，未来可集成Mirage的muGraph超优化搜索，为芯粒任务设计成本模型，将优化扩展至芯粒感知维度。
2. 张量并行的多GPU评估：当前实验限于单GPU。张量并行（TP）先将权重矩阵跨GPU分片，Fleet再在每个GPU内部进行芯粒级分区，两者可自然组合，但缺乏实测验证。
3. 预填充阶段性能评估：论文仅评估了解码阶段（内存瓶颈），未涉及预填充阶段（计算瓶颈）。芯粒感知调度在计算密集场景下的效果是未解问题。
4. 寄存器压力缓解：论文承认持久化内核的寄存器占用限制了波级并行度。潜在解决方案包括每任务独立编译（会牺牲单次启动特性），或依赖硬件支持动态寄存器分配。
5. MLIR集成：将Fleet的芯粒感知调度与MLIR内核编译器集成，在统一IR中共同表达芯粒感知调度和分块级代码生成。

更宏观的研究路径

从领域趋势与技术瓶颈出发，Fleet的成果可能催生以下新的研究路径：
* “可编程内存层次”编程模型：Fleet的核心贡献之一是将L2缓存从“透明的硬件自动管理资源”转变为“可显式寻址和控制的软件管理资源”。这一思路若推广至整个内存层次（如MALL缓存、HBM），可能催生一种新的编程范式：程序员以“内存范围”为第一公民组织计算，运行时系统负责在范围间编排数据移动与任务调度。
* 自适应融合策略：“小批融合占优，大批独立占优”的现象指向一个通用问题：如何根据运行时特征（批大小、序列长度、硬件拓扑）动态选择融合粒度？一个可能方向是在持久化内核中内置轻量级性能模型，在每次迭代前动态决策是“保持融合”还是“切换到独立优化内核”，即将静态任务图扩展为“动态可重配置任务图”。
* 芯粒架构的软硬件协同设计：Fleet以纯软件方式实现芯粒感知，证明了软件的灵活性。这也引发反向思考：未来的GPU架构是否应提供更直接的硬件支持？例如硬件管理的芯粒本地任务队列、跨芯粒轻量级栅栏、可编程缓存分区策略。Fleet的设计为这类硬件特性的价值提供了“存在性证明”。
* 稀疏模型与MoE的芯粒感知：当前评估限于密集模型。在MoE架构中，专家的稀疏激活天然具有“数据并行”特性，不同专家可驻留在不同芯粒的L2中，门控网络将token路由至对应芯粒。这种模式与Fleet的芯粒任务抽象高度契合，可能实现更显著的加速。
* 持久化内核的安全性与隔离性：在云环境等多租户场景下，持久化内核的长时间驻留特性带来了新的安全挑战，例如如何防止通过L2缓存竞争的侧信道信息泄露。这是Fleet走向生产环境必须面对的问题。

Fleet的价值不仅在于为LLM推理带来加速，更在于揭示了GPU编程模型演进的一个方向：当硬件进入多芯粒时代，软件抽象必须从“扁平化”走向“层次化”，从“硬件透明”走向“拓扑感知”。Chiplet-task是这一方向上的首个显式原语，但绝非最后一个。

• Event Tensor：通过动态 MegaKernel 消除重编译，实现 GPU 核间通信与计算重叠
• Mirage Persistent Kernel：首个自动巨核化多 GPU LLM 推理的编译器-运行时系统，实现细粒度计算-通信重叠
• Luminal：自动编译生成 MegaKernels，解决 GPU SM 负载不均，消除内核启动开销与内存气泡

关注“鲸栖”小程序，掌握最新AI资讯