第一次看到 Weave 这篇论文的标题我停下来看了两遍。MoE 大内核、细粒度动态 SM 调度、H100这几个词放在一起基本把最近一年 GPU kernel 优化圈子里最挠头的痛点点破了。做 MoE 推理的人应该都有同感模型架构、路由策略、显存管理这些层面已经被挖得很深但真正把一个大内核丢到 H100 上跑SM 利用率常常低得让人不想看 profiling 结果。Weave 这篇论文做了一件听起来简单、做起来极难的事——把大内核内部的 SM 调度从静态分配改成运行时动态调度在 4×H100 上拿到了 2.89× 的层加速。我先把结论放在前面如果你做的事是 MoE 模型的高性能推理、kernel 融合或者 GPU 调度策略设计这篇论文值得精读而且值得对照自己的 profiling 数据去判断收益边界。我读完之后最大的感受是Weave 并没有发明什么全新的硬件魔法它只是把一个很多人在工程里隐隐约约想过、但不敢下决心实现的事情做完整了——让 SM 自己去找活干而不是开工前就被绑死在某个任务上。这篇文章我会按自己的理解把 Weave 为什么盯着大内核、动态调度具体怎么落地、2.89× 这个数字怎么拆解、以及复现时容易踩的坑一层层讲清楚。里面会掺不少我做 GPU kernel 优化时的个人经验和判断论文里没写全的部分我会尽量说明这是我的推测方便你对照原文分辨。1. 先看结论MoE 大内核为什么在 H100 上跑不满1.1 被忽略的瓶颈大内核是一堆任务的组合不是单个计算先对齐一个概念什么是大内核。在 MoE 架构里一个 decoder 层通常包含 attention、router、多个专家的 FFN、以及 token 的 permute/combine。以前的做法是把这些拆成几十个小 kernel 顺序发射每次 launch 都有开销中间结果还得反复写全局内存。后来大家开始做 kernel 融合把多个 expert 的 GEMM、甚至是整个 MoE 层的计算塞进一个大 kernel 里一次启动尽量在寄存器、共享内存和 L2 里把数据流转消化掉。这种融合出来的大内核内部结构远比一个普通 GEMM 复杂它可能同时存在 attention 阶段、多个专家并行计算、以及专家之间的依赖关系。问题在于大内核内部的任务形状高度不均。不同 expert 拿到的 token 数不一样不同阶段的计算量不一样有些分支是矩阵乘有些分支是 elementwise还有些分支要等前面的结果做完才能开始。这就好比一个餐厅里菜单上的菜式五花八门厨师们能力相近但如果开工前就把每道菜固定分配给某个厨师结果必然是有人忙到冒烟、有人提前收工。大内核跑在 H100 上SM 就是那些厨师任务就是各种计算分片。传统 GPU kernel 的硬件调度器能做的只是在 kernel 启动时把线程块按顺序分发到 SM 上它感知不到任务内部的依赖关系更没法在某个 SM 提前空闲时把别人手里的活匀过来。这也是我会把大内核称为“一堆任务的组合”而不是“单个计算”的原因。如果你只把它当成一个普通 kernel 去优化看到 profiling 里 SM 利用率只有 50% 多半会一头雾水但如果你把它拆成“几十个并行任务 依赖边”问题一下子就清楚了浪费的根本原因是任务粒度和 SM 资源之间错配。1.2 三个看不见的浪费波次取整、尾延迟、专家负载偏差大内核在 H100 上的 SM 浪费我习惯分成三类来看对应三种不同的 profiling 特征。第一个是波次取整。H100 SXM 版有 132 个 SMGPU kernel 的线程块按 wave 分发每 wave 最多同时跑 132 个块。如果一个 kernel 总共产生 200 个线程块第一 wave 132 个块跑满第二 wave 只剩 68 个块相当于有 64 个 SM 在第二 wave 里干坐着。这类浪费在普通 GEMM 里已经被研究得很透stream-K 之类的调度就是为了解决它但大内核里的 wave 问题更隐蔽因为任务不是均匀的 200 块而是各种尺寸的混合体。第二个是尾延迟。大内核里的多阶段依赖很容易造成“全 wave 等最慢的一个”。比如一个 expert 的 up 投影做完数据才能交给 down 投影如果某个 SM 负责的 up 分片特别大其他 SM 就算把活干完也得等这个慢的 SM 落盘才能进入下一阶段。GPU 上的全局同步是极其昂贵的操作而大内核里隐式的阶段同步到处都是。这种等待在 Nsight Compute 里通常表现为 barrier stall或者 SM busy 率不低但实际有效计算时间很少。第三个是专家负载偏差。MoE 路由的结果由输入 token 决定每个 expert 实际拿到的 token 数天然不均匀。有的 expert 分到一小撮 token有的分到一大把即使 token 数接近不同 expert 的 FFN 维度也可能不一样。静态地把 4 个 SM 分给 expert A、8 个 SM 分给 expert B几乎必然导致资源错配。我在实际 profiling 里经常看到某个专家已经结束对应 SM 空出来但旁边专家的计算还在排队——这就是最典型的“看起来都在忙、实际在空转”。这三个浪费经常叠加出现。Weave 的做法说穿了就是把这三个浪费合并到一个统一的机制里去解决让 SM 不再固定属于某个任务而是在运行时动态认领。1.3 为什么静态方案救不了编译期算不准运行时有人可能会问能不能在 kernel 启动前根据模型配置和 batch size 精确算好每个 SM 的任务分配答案是很难。问题在于 MoE 的路由结果要到运行时才知道token 在哪个 expert 上、每个 expert 分到多少 token只有实际跑一遍 router 才能确定。这就意味着任何基于编译期信息的静态分配本质上都是拿昨天的天气预判明天的雨。模型变得越细粒度这个问题越严重。现在的 MoE 架构流行细粒度专家设计专家数量从几十个涨到几百个单个专家变小计算量更碎片化。碎片化带来的直接后果是SM 之间的负载偏差被放大一个专家少算几千 token在粗粒度模型里可能只是一个小波动在细粒度模型里就足以让一个 SM 整个 wave 空转。静态方案还有一个天然短板它无法应对“任务实际上不需要那么多 SM”的情况。一个专家如果只分配到一个很小的 GEMM你给它 4 个 SM它跑不完时其他 SM 也不能来帮忙跑完了这 4 个 SM 就闲着。真正的解法只有一个——把“谁干哪块活”的决定权从编译期移到运行期让空闲的 SM 主动去取新的任务。Weave 就是沿着这个思路把动态调度完整地做进了大内核内部。2. Weave 的核心机制细粒度动态 SM 调度怎么运转2.1 第一步把大内核按依赖关系拆成调度单元动态调度的前提是你得有“可以动态分配的最小单位”。Weave 的做法不是把 kernel 拆到指令级那太细了而是按计算边界拆成一组“调度单元”。我的理解是每个 expert 的 up 投影是一个调度单元down 投影是另一个调度单元如果某个 GEMM 本身太大还会按 split-K 的方式切成多个分片每个分片也是调度单元。attention 部分同样可以按阶段切QK 计算和 PV 计算是不同阶段各自成为独立调度单元。拆的时候必须带着依赖关系一起拆。up 投影没算完down 投影不能开始QK 的结果没出来PV 也动不了。所有调度单元之间形成一个有向无环图SM 执行完一个单元后只能去领取那些“前置依赖已完成”的单元。这个图就是动态调度正确性的基石。如果只拆不管依赖就会拿到一个看似能并行、实际数据还没准备好的单元跑出来全是错的。这里有一个值得注意的细节依赖关系不只在单个 expert 内部还可能跨 expert。MoE 里的 all-to-all 通信、token 的 permute/combine 阶段会把不同专家之间的数据流串起来。所以 Weave 要维护的 DAG 是整个 MoE 层级别的而不是单个 expert 级别的。这也意味着调度器必须对层结构有全局视野不能只看自己手上那一段。2.2 第二步运行时任务队列 SM 空闲状态感知拆好调度单元之后核心机制就变成了一个全局任务队列。所有调度单元进入队列每个 SM更准确地说是每个 persistent 线程块或 warp group在执行完当前单元后从队列里领取下一个可执行单元。这正是 persistent kernel 的经典套路kernel 启动后线程块常驻循环处理任务直到所有任务耗尽。为了直观我写一个最简化的示意模型注意这不代表论文源码只是帮你理解机制轮廓__device__ int next_task 0; // 全局任务计数器 __global__ void weave_skeleton(...) { while (true) { int my_task atomicAdd(next_task, 1); if (my_task total_tasks) break; execute_task(my_task); // 执行对应调度单元 __threadfence(); // 发布当前 SM 对全局内存的写入 } }这个简化版本的问题很明显所有线程块争抢同一个原子计数器调度开销全花在 atomicAdd 上了而且只做了任务分发没考虑依赖关系和局部性。Weave 真正的实现显然要复杂得多但我希望这个骨架能让你建立基本心象——动态调度本质上就是“用一个全局可原子访问的状态驱动多个常驻 SM 循环取任务”。任务领取策略其实还有讲究不一定每次只领一个单元也可以一次领一批减少原子访问频率。领队的顺序、是否优先领与本 SM 数据亲缘性高的任务都影响性能。论文里真正的贡献恰恰是把这些决策机制系统化了。2.3 第三步阶段间的轻量通知避免全局同步动态调度最容易做砸的地方是同步。多阶段依赖摆在那里最笨的办法是每结束一个阶段就来一次全局 barrier所有 SM 到齐了再放行。问题是全局 barrier 的代价极高而且它天然造成“全 wave 等最慢者”。我在自己做过的大内核优化里吃过不少全局同步的亏最后基本都是靠把同步粒度打碎来解决。Weave 采用的应该是细粒度依赖通知每个调度单元执行完更新一个就绪计数器只有真正依赖它的后续单元被触发时相关 SM 才进入等待。这样不同阶段的任务可以重叠执行等待只发生在真实的依赖缺口中。它本质上把“全量同步”换成了“按需同步”。你可以想象成流水线每个工位只需要等上游那一个零件而不是等整条产线所有零件都到位。这个设计是 Weave 最值得读的部分。调度开销是 1% 还是 20%往往就取决于依赖通知的实现方式。我猜测论文里会用原子计数器加 acquire/release 语义来维护每个调度单元的 ready 状态配合 warp specialization让专门的调度 warp 来推进状态而不是让所有计算 warp 都去轮询抢锁。下面这个小节继续展开这些工程决策。读论文时可以先看它怎么处理依赖通知再看它怎么切调度单元。这两个地方是决定动态调度能不能真正落地的关键。3. 设计取舍为了动态调度Weave 付出了什么3.1 调度粒度不是越小越好要跟计算强度对齐动态调度的直觉是粒度越细负载均衡越精细但工程现实是调度本身有成本。每个调度单元都要经过队列操作、原子更新、依赖检查如果任务本身只有几千次浮点运算调度开销可能比计算还贵那还不如静态分配。所以调度粒度必须和计算强度对齐一个调度单元至少要能覆盖掉调度开销再加上数据搬运开销才值得被当作独立任务。我给一个经验比例供参考在我接触过的 kernel 里调度相关开销加起来超过内核总耗时 5% 的时候就要警惕了。如果发现动态调度只带来了负载均衡收益却被调度开销吃掉一大半就应该把两个小单元合并、减少原子操作次数。Weave 的调度粒度大概率不是固定的而是按层结构自适应调整——专家大就多切几块专家小就整块作为单元。这就像快递分拣你把每个包裹单独派给一个分拣员看起来最公平但交接成本会拖垮效率把几个包裹合成一箱再派发整体吞吐反而更高。粒度的艺术就在这里。3.2 数据局部性任务迁走之后共享内存和 L1 怎么办动态调度最大的隐形代价是数据局部性。GPU 的每个 SM 有独立的一级缓存和共享内存数据在 SM A 上算完如果能直接在 SM A 上继续下一阶段就能享受共享内存的低延迟但如果下一阶段被调度到 SM B数据必须从 SM A 的共享内存写回全局内存再由 SM B 从 L2 或全局内存重新读进来。这一来一回可能把动态调度省下的等待时间全赔进去还倒贴。所以好的动态调度必须考虑任务亲缘性。我的理解是Weave 会让同一个 expert 的连续阶段尽量落在同一个 SM 上只有当前一个阶段的计算量实在不均衡时才允许任务跨 SM 迁移。跨 SM 迁移也不是不能做但要走显式路径先保证写入对全局可见再让目标 SM 读取必要时用异步拷贝把数据搬到共享内存避免直接读全局内存的长延迟。这一点在复现时要格外小心。如果只是抄了“动态取任务”的机制没有处理数据局部性性能大概率比静态分配还差。我在类似系统上踩过这个坑看起来 SM 利用率涨了 20%但因为任务迁移导致 L2 访问量翻倍端到端时间反而倒退了。3.3 原子操作与内存可见性正确性比性能更难动态调度把调度器从编译期移到了运行时随之而来的是正确性压力。多个 SM 同时抢任务必须保证同一个调度单元不会被两个 SM 同时领取一个 SM 完成阶段写入的数据必须对领取下一阶段任务的 SM 可见。GPU 的内存模型里这对应的是原子操作的 release/acquire 语义稍微写错就会出现数据竞争而且这种竞争往往不是必现的只在特定负载和时序下冒出来。从工程角度看动态调度 kernel 的 bug 排查是出名的痛苦。我印象最深的一次是任务队列在收尾阶段出现空窗口某个 SM 判断队列为空准备退出但另一个 SM 还在提交新任务结果整个 kernel 提前结束输出全是错的。这种问题用普通单测很难复现压力大了才现形。所以运行这类代码时我建议从第一天就把 Compute Sanitizer 的 race 检测跑起来不要等性能调完再查正确性。3.4 这套设计对硬件特性的依赖动态调度不是在任何 GPU 上都能拿到同样的收益。H100 的 Hopper 架构给了几个重要支撑更强的全局原子操作吞吐、更大的 L2 缓存、以及 thread block cluster 这类新特性。比如如果原子操作吞吐不足调度队列会变成全局瓶颈如果 L2 不够大任务迁移的数据都要挤在内存带宽上延迟会非常难看。Weave 选 H100 做实验很合理因为它是当前数据中心主力卡特性齐、资源足动态调度能发挥出来。如果你是冲着复现去的我建议想清楚目标硬件。A100 上跑同一套机制收益大概率打折不是因为算法错了而是硬件的原子延迟、缓存带宽和 SM 数量都在拖后腿。理解 Weave 的设计时把它当成“Hopper 特性之上的调度系统”来看会比当成通用方案更准确。4. 实测结果解读4×H100 上 2.89× 是怎么拆出来的4.1 实验环境与对比基线先明确一个容易混淆的点标题里的“2.89× 层加速”指的是单个 MoE 层的计算加速不是整个模型端到端的加速。做推理系统的人都知道一次完整请求要经过路由、通信、attention、多个 expert、输出投影还有很多额外开销。Weave 优化的是其中“大内核”这一核心部分所以论文很严谨地强调了层加速。实验环境是 4×H100这个配置对应的是张量并行下常见的部署形态。4 张卡通过 NVLink 互联每张卡处理一部分层和一部分专家。测层加速时比较对象是未采用细粒度动态调度的大内核基线通常是一个高效的静态版本 fused kernel。为什么用 4 卡而不是单卡因为贴近真实部署同时也能看出在跨卡通信存在的前提下纯计算部分的优化空间还有多大。对比基线非常关键。如果基线很弱2.89× 就说明不了问题如果基线已经很强2.89× 才真正有价值。所以读论文时一定要看清楚基线长什么样是 CUTLASS 的 kernel是官方实现还是他们自己写的静态 fused 版本。我在我的优化实践中见过太多“对比对象写得很差所以新方法大比分获胜”的案例这类论文可信度要大打折扣。Weave 的 2.89× 是真是假很大程度上取决于基线的质量。4.2 加速比构成拆解负载均衡、波次整形与调度开销2.89× 这个数字不能只看成一个整体得拆开看它的构成。我按常见的 kernel 时间模型做个示意拆解帮你建立一个理解框架。假设静态版本的大内核里真正花在计算上的时间只占 35%剩下 65% 都在各种等待波次取整、尾延迟、专家负载偏差。Weave 动态调度后等待被压缩到 15%同时调度本身引入了 5% 的额外开销那么新内核的时间占比约 0.35 0.15 0.05 0.55加速比就是 1 / 0.55 ≈ 1.8 倍。要凑到 2.89 倍说明静态版本的浪费可能比 65% 还要高或者调度开销没到 5%。这个拆解是示意性的但逻辑方向是对的2.89× 不是把 GEMM 本身算得更快而是把 SM 填得更满把等待时间转化成了有效计算。时间组成静态版本Weave 动态版本有效计算35%35%等待/空闲65%15%调度开销0%5%总时间归一化1.00.55~0.35要验证这个拆解你在复现时应该盯两个指标SM busy rate 和 wave 利用率。Nsight Compute 里可以直接看到 SM 的吞吐百分比和 warp stall 原因。如果 Weave 跑了之后SM busy 从 55% 涨到 85%同时“等待依赖”“barrier”这两类 stall 显著下降那 2.89× 就有了依据。如果 SM busy 原本就很高那动态调度的收益空间就值得怀疑。4.3 加速比的边界换个模型、换张卡结果会怎么变任何加速数字都有边界条件2.89× 也不例外。第一个边界是模型结构专家越碎、负载越倾斜Weave 的收益越大如果模型只有 8 个胖专家每个专家的 GEMM 又大又均匀静态调度已经能跑得不错动态调度的发挥空间就很小。第二个边界是 batch size小 batch 时任务数量本身不够填满所有 SM调度器能做的事有限极大 batch 时任务足够多硬件调度本身已经接近满载收益也可能变小。收益最大的往往是中间区域——任务数量适中但专家负载偏差明显。第三个边界是硬件。H100 上能拿到 2.89×换到 A100 大概率要缩水。这不一定是 Weave 实现的问题而是 Atom 吞吐、L2 容量、SM 数量都在影响动态调度的天花板。第四个边界是数据分布。MoE 路由的分布随数据集变化如果测试集路由非常均匀加速比自然会向 1 收敛如果路由特别偏动态调度就特别有存在感。所以看到 2.89× 时最好追问一句在什么模型、什么 batch、什么数据集下测出来的这些条件一变数字就会大幅浮动。5. 对工程实践的启示做 MoE 推理优化的人可以带走什么5.1 给推理框架集成者的建议从 CUDA Graph 里解耦调度现在的主流推理框架大量使用 CUDA Graph 来消除 kernel launch 开销但 CUDA Graph 是静态捕获的图一旦建好每个 kernel 的 block 分配就固定了。这意味着一个很尴尬的局面你想做动态调度但外层框架把执行路径静态化。Weave 的思路给了一个出路——动态调度发生在 kernel 内部而不是 kernel 之间。CUDA Graph 仍然是静态图但图里的某个 kernel 内部是一个 persistent 循环循环里通过原子队列动态拿任务。这样既保住了 Graph 的低 launch 开销又获得了 kernel 内部的动态性。集成时有个重要的工程约束动态调度需要的原子队列、任务状态都必须放在 kernel 内不能用 device-side launch也不能在 capture 期间做 malloc 或 stream 同步。这要求框架层把“调度状态”显式地当作 kernel 参数传入而不是依赖 launch 边界上的隐式同步。我见过不少团队在集成这类 kernel 时被 CUDA Graph 的 capture 限制折磨其实提前把这些状态设计成“常驻全局内存 kernel 参数传递”就能绕开。另外如果框架里有算子融合抽象可以考虑把多个 expert 的 FFN 融合成一个 mega kernel再把 Weave 式调度嵌进去。这一步的实际收益是双重的既减少了 kernel launch 次数又让 SM 有了跨 expert 动态调度的空间。5.2 给 Kernel 开发者的三个落地技巧把 Weave 的思想落地到自己项目里我有三个实际建议都是踩过坑之后的经验。第一先量化再动手。不要因为论文标题漂亮就立刻重写 kernel。先用 Nsight Compute 量一次基线 kernel 的 SM 利用率、warp stall 分布、wave 利用情况。如果 SM busy 低于 70%并且排在前面的 stall 原因是等待依赖或 barrier那动态调度大概率有效如果 SM busy 已经超过 90%瓶颈在访存或计算本身动态调度带来的收益会非常有限。第二用 warp specialization 隔离调度逻辑。不要让每个计算 warp 都去抢任务队列那样原子竞争会把性能拖垮。专门留一个 warp 做任务分发和依赖状态更新其他 warp 只从本地拿“已经分配好”的任务。这个做法还有额外好处调度 warp 可以做预取和依赖预判把计算 warp 要用的下一批任务提前准备好进一步藏住调度延迟。第三调优参数时先扫“任务领取批量”和“任务粒度”再动调度算法。领取批量太小原子访问太频繁领取批量太大负载均衡又变差。任务粒度同理。我通常把这两个参数做成可配置的启动前用一个小规模校准测试扫一遍找到当前模型和 batch 下的最优值再固化成配置。这比每次手工调快得多也稳得多。5.3 Weave 与 stream-K、persistent kernel 的关系辨析如果你熟悉 GPU kernel 优化看到 Weave 的第一反应大概率是这不就是 persistent kernel 加 stream-K 吗我跟你说关系确实很近但范围不同。persistent kernel 是基础骨架kernel 常驻循环取任务避免反复 launch。stream-K 则是把单个 GEMM 的归约任务切得更细用动态调度解决 wave quantization让 SM 在归约段之间动态接活。Weave 做的事情是把 stream-K 的思想从单个 GEMM 推广到整个 MoE 大内核这里的任务不仅有 GEMM 分片还有专家间、阶段间的多种调度单元不仅要解决波次取整还要解决依赖和数据局部性。复杂度比单 GEMM 高一个量级。用一句话概括stream-K 在解决一道菜怎么分给多个厨师Weave 在解决一桌菜怎么在厨师之间动态派单。前者不需要管菜与菜之间的依赖后者必须维护一张完整的依赖图。这就是为什么 Weave 会比单纯应用 stream-K 更复杂也更贴近真实 MoE 推理负载。6. 阅读论文与复现建议别只盯着 2.89×6.1 读论文时重点看哪三类数据看这种系统论文我习惯带着三个问题去读数据部分比只看摘要有效得多。第一调度开销占比。论文有没有明确给出动态调度带来的额外开销是百分之几如果调度开销超过 10%那这套机制换到低配 GPU 上很可能会反噬如果只有 1~3%说明实现相当干净。这个数字直接决定了机制的可移植性。第二消融实验。论文有没有做“关掉依赖感知”“关掉亲缘性调度”“减少调度单元粒度”这类对照消融实验做得越细越能说明是哪个机制真正带来了 2.89×。如果只有“开/关动态调度”两组实验那说服力会弱很多因为你无法判断收益来自任务拆分、依赖处理还是局部性优化。第三负载分布特征。论文有没有画出 expert 负载不均的分布或者不同阶段的时间占比如果数据集中负载已经很均匀而 Weave 仍然拿到高加速那说明它除了负载均衡还解决了别的问题比如波次取整如果负载极偏那加速多半来自均衡化适用场景反而更窄。6.2 复现实验的软硬件准备与典型坑如果你打算在自己的环境里复现或借鉴 Weave我建议按下面的清单准备。硬件方面至少需要一块 H100最好有 4 卡 NVLink 环境贴近论文设定。A100 也能跑但我前面说过收益很可能打折。软件方面CUDA 12.x、CUTLASS 作为 baseline 参考、Nsight Compute 做 profiling 分析基本是标配。代码库本身如果不是论文公开源码可以按自己的 kernel 框架实现关键组件就是原子任务队列、依赖状态表、调度 warp 三块。典型的坑有三个。第一CUDA Graph capture 的限制。如果外层框架用 CUDA Graph 包裹整个模型动态调度代码必须完全在 kernel 内完成不能在 capture 期间做动态内存分配或 device 端 launch。第二测量时要隔离通信。4 卡环境下all-to-all 通信会占用时间测 kernel 层加速必须把通信阶段从计时区间里剔除否则 2.89× 会被通信时间稀释成 1.2×看起来好像失效了。第三正确性验证一定不要省。动态调度代码建议从一开始就开启 race 检测不要等性能调优完成后再补正确性检查否则你会被凭空出现的偶发数据错误折磨疯。我个人的建议是从小型复现开始单卡、单层、固定 router 分布先验证动态调度机制是否符合预期再逐步扩展到 4 卡和真实模型。这样即使遇到问题定位范围也小很多。6.3 我认为值得跟进的两个后续方向Weave 打开了不只一个口子。我自己最关注的是两个方向。第一个是把动态调度和 Hopper 架构的异步能力进一步结合。H100 上有 Tensor Memory Accelerator 和异步拷贝如果把数据搬运和计算重叠起来调度器甚至可以在一个 SM 等待数据期间就把它分配给下一个任务进一步藏延迟。Weave 目前的机制解决了“SM 空等任务”下一步自然要解决“任务等数据”。第二个是端到端框架集成。Weave 目前是层加速但如果把它放进 vLLM、SGLang 这类系统里和 paged attention、chunked prefill 共存还能保留多少收益这才是工程界最关心的问题。层内调度和层间调度之间的协同比如让上一层专家计算和下一层路由通信重叠很可能是下一个主要突破点。我自己的体会是读这类论文最容易犯的错是把层加速直接当端到端收益去汇报。2.89× 确实漂亮但真正值钱的是它把“大内核怎么调度”这个原本靠手感和试错的命题变成了一个有明确机制的系统设计。如果你手头正好在调 MoE 的 kernel我建议先别急着上 Weave 的全套方案先用 profiler 量化你自己的 SM 空闲和 wave 利用率确认瓶颈确实是调度而不是 memory-bound。如果 profiling 显示 SM 大量时间在等待依赖或者干脆空闲那 Weave 的思路值得你认真抄一遍作业。