本文针对GLM-5.2-FP8模型解码阶段进入Sparse FlashMLA前的计算流程进行优化，实验基于B300显卡，测试场景为同时处理两个请求，每个请求8个token，即M=16、KV=8192的配置。原本的计算流程需要约30个短kernel，总耗时80us，经过kernel fusion优化后降至61.631us。随后基于Megakernel思路实现的MegaRTP，将各个kernel拆分为tile级指令放入持久化kernel中执行，扣除固定框架开销后耗时约39.601us，完整运行耗时44.272us。后续通过CUDA Graph+Multistream+PDL（简称CMP）的组合优化，最终将耗时降至36.992us。

一、 融合后的显存利用率瓶颈

原本的实现包含约30个独立的CUDA kernel，包括Add+RMSNorm+Quant、各类GEMM、RoPE、缓存写回等操作。将相邻的计算合并到同一个kernel后，kernel数量降至11个，串行执行时间从80us缩短到61.631us，减少了多次kernel启动和中间数据读写的开销。

但这种优化并没有解决GPU的利用率问题：11个kernel在单条stream上依次执行，前一个kernel未完成时，后一个不会启动，同一时刻只有当前kernel的CTA驻留在SM上。部分kernel仅提交12、16或32个CTA，而B300拥有148个SM，例如12-CTA的B1 producer最多只能让12个SM同时运行，其余SM只能等待该kernel结束。

通过执行区间和CTA数量的记录可以看到，整段执行平均活跃SM占比仅为43.7%，kernel之间的bubble时间中位数仅为0.672us，占总时间的1.08%，说明优化空间并不在压缩kernel launch开销，而是在于执行顺序。原本的11个kernel并非完全串行依赖链：第一个kernel完成后，QKV-A、Head-Gate和Indexer-K三条路径互不依赖，QKV-A的结果又分别进入主attention路径和Indexer-Q路径，单stream执行将本可并行的工作强行排成了队列，因此需要将调度粒度从整个kernel下沉到tile级别。

二、 Megakernel的并行调度逻辑

Megakernel的核心是打破传统的kernel边界，不再将每个kernel作为整体启动，而是按照tile粒度拆分为多组任务。tile指的是单次CTA级的工作单元，例如RMSNorm按行切分为单个tile，GEMM按二维输出块切分为(blockM, blockN)的输出块。在MegaRTP中，Host将每个tile编码为一条指令，启动固定数量的worker CTA，每个worker按照CPU预先生成的指令队列逐条执行，而非在运行时重新推导计算图。

具体的并行机会分为两类：

• 无依赖路径的并行执行：以QKV-A、Head-Gate和Indexer-K为例，三者都依赖Add+RMSNorm+Quant，但彼此之间没有数据依赖。当基础操作完成对应tile并更新事件计数器后，等待该结果的指令就可以开始执行，三条路径的指令可以交错推进，共同使用空闲的SM，避免了单流执行时逐个完成整个kernel的情况。
• 有依赖任务的提前准备：对于存在数据依赖的前后任务，例如任务B依赖任务A的输出，任务B的矩阵乘法无法提前执行，但在第一次读取A的输出之前，任务B可以完成很多与A无关的工作，例如读取任务描述、构造调度状态、初始化barrier、准备TMA descriptor、加载自身的权重和scale。Megakernel可以让执行任务B的worker CTA在A仍在执行时就开始这些准备工作，仅让负责读取输入的warp等待A的结果。当A的对应输出满足依赖后，任务B的activation loader开始搬运输入，与已经准备好的权重在barrier处汇合，进入原本的kernel主循环，这种方式并没有绕过数据依赖，只是提前完成了无关的准备工作。
• tile级细粒度依赖：传统kernel依赖以整个grid完成为边界，即使A的第一个输出tile已经写完，B仍要等待A的所有CTA退出。而Megakernel将依赖细化到tile级别：A完成一个tile后更新对应的事件计数器，B的指令等待该计数器达到预设目标值后才读取并计算该tile，无需等待A的整个grid结束。例如1024行的RMSNorm后接blockM=128的GEMM，前128行完成并更新计数器后，GEMM就可以先计算对应tile，无需等待RMSNorm的其余行结束。不过当M较小时，例如M=16，上游kernel可能在一个wave内完成，单个tile的耗时接近整个kernel，此时下游没有空闲CTA可以提前执行，细化到tile级的依赖并不会缩短关键路径。

三、 相关技术背景

Stanford Megakernels

面向单GPU、batch size为1的Llama解码场景，解决的是decode中每个算子都很短的问题，即使将kernel launch替换为CUDA Graph，算子之间的边界仍会阻止下一个算子提前准备权重，也会让本来可以交错的工作被整个grid隔开。其做法是将一轮forward的执行计划放入GPU内部的轻量指令执行器，Host侧先构造算子DAG，再将每个算子拆分为可由一个worker CTA执行的指令，并分配到静态队列。队列通常编码为[num_sms, max_queue_len, 32]的int32张量，每条指令为32个int32，较短的队列用NoOp补齐，保证所有worker可以按统一迭代次数推进。

Device侧启动持久化kernel，每个worker CTA负责一个SM的队列，CTA内使用warp specialization：controller warp读取指令，建立逻辑页到物理页的映射并初始化信号量；loader warp负责TMA权重和activation搬运；consumer warp执行矩阵乘和向量计算；storer warp负责结果写回并发布完成状态。不同角色通过指令到达、阶段信号量和页信号量协作，实现控制信息、权重搬运、Tensor Core计算和写回的并行。跨算子的正确性由全局计数器保证，生产者完成一个输出tile的异步写回后，通过atomicAdd发布完成数量，消费者仅在真正读取该activation前等待目标计数器，与输入无关的准备工作可以提前进行。SMEM被划分为可循环复用的页，当前指令释放页后，下一条指令可以使用该页预取权重，避免每次都从空流水开始。

MegaRTP继承了这套Host生成计划、GPU内解释执行、warp specialization、跨指令预取的思路，但针对GLM5的动态M/KV、tile级计数器依赖和Blackwell GEMM进行了调整，重新设计了指令ABI、worker角色和页生命周期。

Mirage和Event Tensor

Mirage Persistent Kernel将算子拆分为SM级task，并根据实际读写范围建立task之间的event，而非使用“整个算子完成”的粗粒度barrier，让MatMul的某个tile完成后，对应的通信或下游计算tile就可以开始，实现计算和通信的流水。其在生成tGraph时会做三类压缩：event fusion合并具有相同producer或consumer集合的event，减少event数量37-118倍；tGraph normalization仅在fork/join时插入辅助task/event，规整任务的触发和依赖关系；tGraph linearization重新排列task，让同一个event的后继task在队列中连续，缩小successor metadata的大小4.4-5.9倍，MoE模型上约15倍。

Event Tensor的亮点是将event本身表示为带维度和索引关系的对象，例如按batch、head、tile组织完成状态，使shape-dependent、data-dependent的依赖可以用同一套表示描述。实验表明，细粒度依赖并不会自动带来收益，静态调度的dense模型仅能获得6-8%的收益，而动态调度的dense模型性能甚至会降至基线的0.82-0.89x，说明动态调度的开销可能抵消重叠带来的收益。

TileRT

其CUDA kernel实现为闭源，一次forward在每个rank上有162个kernel，并非用一个kernel覆盖完整模型，而是每层attention和FFN分别启动kernel。attention拆分为两种GPU角色：rank0执行Indexer Q/K、扫描KI cache、计算score和Top-2048，然后将选中的位置发布给其余七张卡；rank1-7以TP7分片持有主MLA权重，先完成projection和cache准备，拿到Top-2048位置后再gather稀疏KV，继续完成QK、softmax、PV和输出投影。attention有两个跨rank同步点：第一个在稀疏gather前，rank1-7需要等待rank0发布的Top-2048索引才能读取KV；第二个在attention输出处，七个MLA rank产生的partial需要汇合，八张卡拿到一致的attention输出后才能进入FFN。

TileRT的MoE采用另一种分工：八个rank都执行相同的router和Top-8，都持有shared expert和256个routed expert，每张卡仅保存每个expert的1/8中间维，Up/Gate按输出维切分，Down按输入维切分，每张卡得到6144维partial，在FusedMoeExecutor内完成八卡归约和residual，属于expert内部的TP8，而非将不同expert分给不同GPU的EP8。基于8张B300显卡对GLM-5.1-FP8的本地实验，batch=1、关闭MTP、KV=8192时，单步延迟中位数为7.358ms，对应136.09 tok/s。各路径的有效带宽计算显示，Indexer、MLA和MoE分别仅达到单卡峰值的6.51%、11.10%和24.89%，带宽利用率低于预期。

四、 MegaRTP的实现框架

MegaRTP的核心是用一个持久化kernel运行不同的模型图，区别于Stanford Megakernels将固定模型的执行路径固化到kernel中，MegaRTP需要支持模型扩展和大M下的高性能GEMM。

首先，模型图必须可扩展：同一个runtime既要能执行不同模型的完整图，也要能执行任意选定的子图；当M和KV length改变时，仅需要重新生成Host侧的计划，无需修改常驻kernel的执行循环。其次，大M下需要高性能GEMM：GLM5的batch或MTP增大后，需要高吞吐的Blackwell GEMM，不能沿用仅适合M=1的matvec。为此，保留常驻worker的执行方式，但将模型图、tile、地址和event依赖从kernel代码中移到runtime plan，并接入可按M特化的DeepGEMM。

一次forward开始前，Host将模型图拆分为tile，为每个tile生成一条指令，再将这些指令排进148个worker CTA各自的队列。GPU仅启动一个持久化kernel，每个worker CTA从自己的队列读取下一条指令，根据opcode调用已经编译好的Add、RMSNorm或GEMM实现；输入未就绪时等待指令指定的event，执行和写回完成后更新event，通知依赖它的后续指令。这个持久化kernel不再仅执行固定的算子，而是执行Host为不同模型图生成的指令序列。

Host侧流程：捕获OpGraph与展开TaskGraph

前端像调用普通算子一样编排模型，记录算子和tensor的def-use关系，例如构建attention_prefix的图，包含x、residual作为运行时参数，norm_w、proj_w作为绑定权重，经过add、rmsnorm、blackwell_gemm后得到输出。OpGraph仅捕获模型的DAG，并未回答每个CTA处理哪部分数据、需要等待多少结果。

执行前，runtime根据当前shape将图展开为具体计划：模型图定义→OpGraph→每个Op按tile展开→TaskGraph→KernelPlan→指令张量+事件计数器+操作数指针表。每个Op按自身的输出tile展开：A0的RMSNorm/quant每行是一个tile；GEMM按(blockM, blockN)的输出块切分，再按Split-K展开；E1按request、Q-block和KV split展开。以M=16、KV=8192为例，计划的tile数量为：A0有16行，B0为1*21*4=84，C0为1*12*4=48，B1为1*8=8，D1为1*32*4=128，D0为1*64*2=128，D2为1*64*4=256，E1为4*16=64，合计732条tile task。

TaskGraph：tile到counter event的转换

OpGraph仅记录两个Op之间存在tensor依赖，lowering到TaskGraph时，Host根据(producer Op, consumer Op)查找配对策略。已知组合可以注册细粒度策略，根据producer写入区域和consumer读取区域生成event；未注册的组合自动退化为whole策略，让consumer等待整个producer完成。

配对策略分为四种：whole策略不分析tile对应关系，所有producer指令各递增一次counter，consumer等待完整计数；tile策略表示一块producer region恰好对应一块consumer region，每个region一个event，threshold为1；tile_cover策略表示多块互不重叠的producer region完整覆盖一块consumer region，producer分别递增同一counter，threshold为覆盖块数；tile_reduce策略表示多个producer partial对应同一块consumer region，producer分别递增同一counter，threshold为partial数量。从Device的event机制看，tile_cover和tile_reduce完全相同，区别仅在数据关系和Host检查：tile_cover是拼接，例如A0的16个row tile分别写入不同行，共同覆盖B0要读取的M16输入；tile_reduce是归约，例如C0的48个K-split都对同一个gate tile贡献partial，全部归约完成后D1才能读取。Counter仅记录完成的producer数量，不负责执行数值归约。

以M=16、KV=8192的计划为例，A0发布hidden FP8/scale和norm两类结果，各有16个row tile，因此event0>=16供B0、B1使用，event1>=16供C0使用；B0的16个最终Q/scale tile递增event2，C0的48个gate partial递增event3，D1的每条指令同时等待event2>=16和event3>=48，即B0+C0的join；D1最终产生32个Indexer-Q/head tile，递增event68，B1的cache写回完成后递增event69，E1等待event68>=32且event69>=1时才读取输入；D0的128个projection tile按输出区域两两归并为64个event，每个event的threshold为2，D2对应区域满足后即可开始。该计划共生成70个event，当M、batch或KV length改变时，仅改变tile数量、生产次数和threshold，event的等待/发布规则无需修改常驻kernel。

KernelPlan：任务队列分配

KernelPlan负责将732条task放入固定的worker队列。当前版本按round-robin将732条指令分到148个worker队列，其中140个worker有5条有效指令，8个worker有4条有效指令，不足的用NoOp补齐。因此Device看到的张量形状固定为[148,5,32]，无需在GPU上维护需要不断扫描的ready queue。

每条指令是固定大小的任务描述，包含opcode、wait_count、signal_count、operand_row_id、wait(event_id, threshold)、signal(event_id, increment, publication_id)、payload(tile坐标和少量动态参数)、padding，共32个uint32，即128B。以E1的一条指令为例，仅保存request、Q-block、KV split范围、需要等待的event68>=32和event69>=1，以及operand_row_id=7，不保存绝对CUDA地址。opcode+spec_id选择编译期生成的CUDA Op family，动态部分仅包含tile坐标、event threshold和少量stride。

为了避免重复存储地址，设计了OperandPtrTable，将每个Op的地址保存为连续的CUDA张量，形状为[num_ops,16]，dtype为uint64，例如A0到E1的图形状为[8,16]，共1024B。每个Op占一行，同一个Op拆出的所有tile instruction共用这行，16个slot的含义由Op自己定义，未使用的slot填0。例如E1的所有指令都保存operand_row_id=7，Device端通过table+7*16取出第8行，保存Q、KV Cache、KV scale、head weight的TMA descriptor地址，context_lens地址，block_table地址，logits地址等。更换输入、权重、KV Cache或workspace时，只需重建或更新OperandPtrTable，通过TMA访问的tensor如果地址或shape变化，只需重建对应的TMA descriptor，无需修改指令和worker队列。仅当M/KV改变tile数或event threshold，或模型连接变化时，才需要重新生成task、event和worker队列。

Device侧执行逻辑

Host将模型图展开为指令并分到worker队列后，Device不再推导模型依赖或重新划分tile，启动148 CTA x 384 threads的持久化kernel，每个CTA长期循环执行自己队列中的指令。一条指令对应一个CTA负责的tile，同一个CTA在不同迭代中可以执行不同Op的tile。

普通kernel仅执行固定计算，状态随CTA结束释放，而MegaRTP的常驻CTA需要连续执行不同Op的指令，因此必须在SMEM中为每条正在执行的指令保存独立状态，每份状态占5120B，主要包括128B指令、13个逻辑页到物理页的映射、slot生命周期和Op内部使用的local mbarrier，以及4KiB instruction-local scratch。

为了让下一条指令的准备与当前指令的计算重叠，每个CTA配置两个slot，共占10240B：当前指令使用一个slot，Controller在另一个slot中写入下一条指令、建立页映射并初始化local mbarrier；旧指令完成后，两个slot交替复用。384个线程按warp specialization固定分工：warp0是Controller，warp1是MMAer，warp2是ALoader，warp3是WLoader，warp4-11是Epilogue。它们是五个同时推进的role loop，每个role仅等待自己需要的状态。

具体的交接流程为：1. Controller等待当前slot的上一代指令完成，将下一条指令拷入slot，更新页映射，初始化本条Op的local mbarrier，这些mbarrier用于CTA内部的数据阶段同步：WLoader/ALoader通知MMAer输入到达，MMAer通知Epilogue消费结果；2. Controller完成准备后，arrive instruction_arrived，这是MegaRTP专用的slot发布协议，保证其他role只有在描述、页映射和信号量都写好后，才能读取该slot；3. WLoader看到发布后可以预取与上游无关的权重和scale，ALoader直到真正读取上游activation的位置才等待global event，依赖满足后，ALoader搬运activation，并用local mbarrier交给MMAer；4. MMAer等待A/B两类local mbarrier后进入tcgen05/UMMA主循环，此时Controller可以在另一个slot准备i+1的指令，WLoader也可以预取i+1的权重，实现i+1的prologue与i的MMA重叠；5. i的Epilogue从TMEM取结果、完成转换和写回，发布下游global event，所有非Controller role完成收尾后arrive instruction_finished，Controller才能覆盖该slot。

五种状态分别保护不同对象：instruction_arrived表示slot的指令已准备好，其他role可以读取；Op-local mbarrier用于CTA内loader→MMAer→Epilogue的阶段交接；global event counter用于跨CTA的tile数据就绪计数；page_finished表示某个物理SMEM/TMEM page的最后一个使用者释放它，下一条指令才能覆盖该页；instruction_finished表示当前slot的所有worker role都已收尾，Controller可以复用该slot。

共享内存分页机制

双slot仅复制指令状态，并未为两条指令各准备一整套数据缓冲区，两条指令仍共享同一份用于存放权重、activation和输出tile的SMEM。该SMEM被切分为13个16KiB的physical page。当前指令尚未用完部分page时，下一条指令可能已经开始加载数据，因此需要解决两个问题：下一条指令使用哪个physical page，以及何时可以覆盖该page。

第一个问题由pid_order解决，每个Op的实现使用logical page id，每个instruction slot保存13项的pid_order表，将logical page映射到本轮实际使用的physical page，Op通过该表取得地址，无需将physical page编号写死。Controller准备下一条指令B时，为B的每个logical page生成映射，例如B的logical page0，调用前一条指令A的release_lid(0)，获取应该继承的A的logical page，假设返回2，而A的映射是pid_order_A[2]=7，则pid_order_B[0]=7，即B仍将该缓冲区称为logical page0，但实际使用physical page7。release_lid为13个logical page分别给出继承关系，形成B的整张映射。各个Op根据自身的page使用顺序安排映射，尽量让下一条指令的权重加载落到上一条指令较早释放或未使用的page上，第一条指令没有前驱，直接采用logical page i→physical page i。

第二个问题由每个physical page的page_finished mbarrier解决：B的loader在覆盖page7前先等待page_finished[7]，A对page7的最后一次读写完成后，由最后一个使用者arrive该mbarrier，B才能开始加载。其他page各自交接，因此B不必等待A的整条指令结束，只需等待自己要覆盖的page。pid_order保存在双slot中，仅保留相邻两代，instruction i+1读取slot i的映射并写入slot i+1，准备instruction i+2时读取slot i+1，slot i仅在instruction i的所有role完成后才会被覆盖。

大M下的GEMM实现

共享page解决了不同Op复用片上内存的问题，但GEMM还有另一个问题：同一个持久化kernel中会出现多种矩阵乘。以attention-pre为例，Q投影、Indexer-K和Main-Q的权重矩阵K/N都不同，batch或MTP改变后，输入行数M也会改变，M=16时可能只有一两个M tile，M=128时需要处理多个M tile，适合的blockM、stage和epilogue组合也可能不同。

如果将某一套GEMM模板和配置写死在kernel中，只能支持一种权重shape和一种M，换GEMM要么复用不合适的tile，要么退回仅适合小M的GEMV，提前准备固定配置也无法覆盖任意模型、batch和MTP。将所有可能的模板编译进去会让执行类越来越大，且无法覆盖新的shape。

MegaRTP将“本轮需要哪些GEMM实现”交给Host在OpGraph→TaskGraph时决定，前端遍历本轮模型图中的GEMM，根据权重shape、输入M和选定的tile划分生成配置记录，并去重。例如本轮模型图中，GEMM-A：K=6144, N=2624→spec0，GEMM-B：K=6144, N=32→spec1，GEMM-C：K=6144, N=2624→复用spec0。随后codegen仅为spec0和spec1实例化对应的C++ GEMM类型、pipeline和epilogue，并编入同一个执行类，生成的是类和模板特化，而非每个spec一个新的kernel，常驻kernel仍然只有一个。

Runtime指令仅保存opcode、spec_id和当前tile的坐标，CTA在Device端读取spec_id，通过编译时展开的spec分支选择对应实现，真正被调用的每个spec都已经编译好，TMA shape、Tensor Core主循环、SMEM/TMEM布局和epilogue无需在Device端重新生成，因此模型图、权重shape、batch和MTP可以变化，大M GEMM的路径仍保持编译期特化。

五、 优化效果与实现成本

在M=16、KV=8192的测试中，扣除固定执行框架的开销后，MegaRTP的计算与调度部分耗时约39.601us，相对于11个kernel串行执行的61.631us baseline，减少了22.030us，约35.8%。收益来自三个方面：独立分支的tile可以交错执行；producer完成一个tile后通过global event让consumer开始，无需等待producer的所有tile完成；同一个CTA中，下一条指令的Controller/WLoader prologue可以与当前指令的MMA/Epilogue重叠。

固定框架的开销为4.671us，来自持久化kernel启动、CTA状态和page barrier初始化、指令读取与发布、role握手、TMEM生命周期以及退出同步，对于几十微秒的M16子图，该固定成本已经足以影响最终方案选择。

迁移普通CUDA kernel到MegaRTP需要额外成本：不能仅将kernel函数搬进常驻CTA，需要重新按tile定义输入输出区域和event，拆出Controller、ALoader、WLoader、MMAer、Epilogue的职责，并为每个page标出最后一个使用者。原kernel中隐含的线程协作、shared-memory地址、barrier和写回顺序，都需要改写成指令、slot、page和publication的约定。

部分kernel的执行资源与当前worker模型不匹配，例如2-SM的tcgen05.mma要求一对CTA同时到达并共同执行GEMM，而MegaRTP的CTA从各自的worker队列独立取指令，无法保证同时领取配对的task，需要额外的成对预留和同步调度。TopK也是一个问题：当前高性能实现通常让单个CTA使用1024个线程，而MegaRTP的常驻CTA固定为384个线程，强行使用384线程会导致性能回退。

六、 收益来源的深入分析

61.631us到44.272us的优化证明了打破kernel串行顺序可以缩短路径，但kernel内的trace让我们需要进一步分析这17.359us的收益来自Megakernel特有的tile级依赖，还是更容易实现的kernel并行和提前加载。

Trace中最明显的两类overlap：第一类是DAG中独立分支的并行执行，A0之后的B0、C0和B1分属三条路径，它们的输入分别在8.6-9.5us就绪，随后同时计算并在20-22us完成；B0之后，D0与等待B0+C0的D1也在20.5us同时进入主计算，原本被单流排开的GEMM现在共同使用同一时间段的SM。第二类是下游指令提前进入并在等待输入时完成准备工作，以D1为例，128条指令最晚在11.232us已经Visible，但主输入到20.576-20.768us才Ready，中间约9us的窗口中，Controller已经准备好指令，WLoader可以加载不依赖B0/C0输出的权重和scale，B0、C0和B1同样在A0输出就绪前就已经Visible。不过trace只能证明准备窗口存在，不能将Visible→MainInputReady的整段都算作有效预取，真正能隐藏的时间需要通过matched PDL实验测量。

累计接入实验从A0+C0开始逐步加入B1、B0、D0、D1、D2和E1，记录子图总耗时的增加量和新增Op单独执行的时间。实验结果显示，B1和B0单独执行分别需要9.088us和10.463us，但接入后仅让子图增加4.528us和4.064us，E1也仅将关键路径拉长3.823us，低于单独执行的5.343us，说明新增Op的相当一部分执行时间与已有子图重叠。不过接入D0和D2后，子图增加的时间略高于新增Op单独执行的时间，原因是Op并发后会争用SM、Tensor Core和显存带宽，无法直接通过逐项相减得到耗时。

回到TaskGraph和实际生成的event，tile级依赖允许producer每完成一个tile就更新event，让consumer不必等待整个上游Op，但在本例的关键路径中并没有这样的窗口。本例中的主要计算都是GEMM，而M=16恰好等于这些GEMM的blockM=16，因此每个GEMM在M方向只有一个tile。更关键的是，上游GEMM沿N方向输出的多个tile，往往会成为下游GEMM的完整K维输入，下游只有等这些tile全部到齐才能计算自己的唯一一个M tile，实际生成的event阈值接近“上游kernel的相关输出全部完成”。

例如A0→B0/B1需要event0>=16，即A0的16行hidden/scale全部就绪；A0→C0需要event1>=16，即A0的16行norm输出全部就绪；B0→D0需要event2>=16，即组成K=2048输入的16个Q/scale tile全部就绪；B0+C0→D1需要event2>=16且event3>=48，即B0的完整Q/scale和C0的全部48个partial都已就绪；D1+B1→E1需要event68>=32且event69>=1，即32个Indexer-Q/head tile和整次K Cache写回都已完成。这里没有第二个M tile可以提前交给下游，第一块数据就是全部M=16，因此critical path上的consumer仍要等到上游GEMM的主体计算结束，tile级event并没有比kernel级数据依赖更早地释放主计算。

唯一的例外是D0→D2，D0为每个head生成独立event，threshold为2，对应head的两个N tile完成后，D2的四个tile就能开始，不必等待另外63个head。trace中最早的D2在29.25us输入就绪，而D0最晚到32.90us才完成，说明这里确实发生了约3.6us的逐head overlap，但本次trace最晚结束的是E1，时间为41.120us，D0/D2这条支路没有决定最终完成时间，因此该细粒度overlap并未转化为本例关键路径的额外收益。

综上，MegaRTP在本例中真正决定整体时间的机会为两类：独立kernel并行执行，以及dependent kernel在输入就绪前完成prologue和权重预取。前者可以由Multistream实现，后者可以由PDL实现，同时每个CUDA kernel还能保留最合适的线程数、cluster、SMEM/TMEM布局和矩阵乘法配置。

七、 Multistream与PDL的优化方案

Multistream：无依赖分支并行执行

CUDA stream仅保证同一条stream内的kernel按顺序执行，不同stream之间如果没有event依赖，kernel就可以并发执行。将QKV-A、Head-Gate和Indexer-K三条无依赖分支放入不同的stream后，一条分支未占满的SM可以执行其他分支的CTA，该过程不改变任何kernel内部实现，解决了独立分支被单流强制串行的问题。

PDL：启动许可与数据就绪分离

PDL将“consumer可以进入GPU”和“consumer可以读取producer输出”拆分为两个时刻。Producer的每个CTA到达选定位置后，由一个线程调用cudaTriggerProgrammaticLaunchCompletion()，等所有CTA都trigger或退出后，consumer才有机会被调度，该信号仅允许launch，并不表示producer的输出已经写完。因此trigger可以放在producer较早的安全位置，例如A0发出与输入无关的norm-weight TMA后就允许后续kernel进入，B0和D1的producer也可以先允许各自的finalizer预取静态参数。

Consumer在读取producer输出之前需要调用cudaGridDependencySynchronize()（简称GDC），该调用会等待直接依赖的producer grid完成，并保证后续读取能够看到producer的写回。GDC仅阻塞执行它的线程，并非整个CTA的barrier，因此不读取producer输出的线程仍可继续工作。Trigger提前并不会破坏正确性，因为真正的数据读取仍受GDC约束，如果将GDC放在kernel入口，consumer的初始化和权重预取也会一起被挡住，如果放到第一次activation load之后，结果可能错误。

Programmatic Event：跨stream依赖

PDL不仅适用于同一条stream或CUDA Graph，CUDA还允许将cudaEvent_t绑定到producer kernel的programmatic launch completion，再让另一条stream等待该event。普通event需要等待producer stream中此前的工作全部完成，而Programmatic Event在producer的所有CTA都trigger或退出后就能唤醒consumer stream，此时producer kernel可以仍在执行。

Host侧的关键是在启动producer时通过cudaLaunchAttributeProgrammaticEvent绑定event，具体流程为创建cudaEvent_t，设置launch属性，配置launch参数，启动producer kernel，然后让consumer stream等待该event，最后启动consumer kernel。Producer的每个CTA在选定位置调用cudaTriggerProgrammaticLaunchCompletion()，event被触发后，consumer可以跨stream提前启动。Consumer可以先做prologue、权重预取等不依赖producer输出的工作，任何读取producer输出的线程都必须先执行GDC，因此Programmatic Event传递的是“可以启动”，而非“数据已经就绪”。后文的Graph rewrite方案使用Graph programmatic edge，手动Multistream使用此处的Programmatic Event，两种接口表达的都是同一套PDL关系：producer trigger，consumer提前进入，真正读取数据前再GDC。

DeepGEMM的PDL改造

原始DeepGEMM在入口等待上游，随后由同一个TMA loader搬运activation A、权重B以及两者的scale，这种方式在A尚未就绪时，完全独立的B也无法提前加载。为支持PDL，为GEMM增加split-loader路径：warp0仅负责activation A及其scale SFA，warp3仅负责静态权重B及其scale SFB。

Kernel完成公共的scheduler、descriptor和barrier初始化后，只有warp3跳过GDC，先把权重流水线填起来；其他warp仍执行GDC，避免任何依赖本轮activation的工作越过等待点。拆分后，每个stage的full barrier等待ALoader和WLoader两次arrival，WLoader最多提前填满kNumStages个stage，再往前会阻塞在empty barrier，因此不会覆盖MMA尚未消费的权重。Activation ready后，ALoader补上A/SFA，两路数据在full barrier汇合，后续的矩阵乘法逻辑不变。

最后需要检查编译后的指令顺序，曾观察到指向producer输出的参数带const __restrict__时，编译器会将activation load移到GDC之前，因此需要去掉危险限定，并用最终SASS确认acquire在load之前，再用污染输入和重复Graph replay验证没有提前读取。

PDL收益拆解

为了分析consumer提前进入GPU、prologue和权重预取各自的收益，固定一对从Add+RMSNorm+Quant到QKV-A GEMM的producer和consumer，保持GEMM的数学语义、tile、scheduler和epilogue不变，仅改变consumer的GDC wait位置和weight loader结构。

三种consumer的差异为：entry模式在kernel入口立即等待GDC，producer完成前不做任何kernel内部工作，通用prologue包括预取TMA descriptor、初始化shared memory中的mbarrier、分配TMEM、配置scheduler和本地流水状态、完成必要的CTA/cluster同步，之后进入携带数据依赖的加载和计算；deepgemm模式将GDC wait移到通用prologue之后，但A/B仍由同一个loader在wait后加载；prefetch模式将不读取producer输出的静态权重B/SFB分给独立WLoader，使其可以在GDC释放前填充weight ring，三种模式都在真正读取activation前等待producer，因此数据依赖没有放松。

比较三种实现的PDL-off和PDL-on，从producer开始到consumer完成的总span节省如下：M=16时，entry PDL收益0.608us（4.4%），deepgemm PDL收益1.344us（9.6%），prefetch PDL收益2.912us（20.7%）；M=32时分别为0.512us（3.7%）、1.375us（9.7%）、2.688us（19.2%）；M=64时为0.960us（6.4%）、0.896us（6.0%）、2.144us（14.5%）；M=128时为0.896us（5.8%）、1.216us（7.8%）、1.856us（12.0%）；M=256时为0.704us（4.1%）、1.408us（8.0%）、1.360us（7.9%）。

entry模式在producer完成前既不执行prologue也不加载权重，但仍能节省0.51-0.96us，这部分收益来自consumer提前进入GPU和producer完成后的更紧handoff。移动wait的两个实现使用unified A/B loader，仅将wait从kernel入口移到通用prologue之后，0.19-0.54us可以归因于提前执行descriptor prefetch、barrier/TMEM初始化等工作。拆出WLoader的两个实现保持相同GEMM config，仅将B/SFB拆给不等待producer的独立WLoader，0.21-1.50us是split-loader和weight prefetch的净关键路径价值。

权重预取的收益不仅取决于weight ring的深度，还取决于consumer多早开始执行，当M=256时，consumer开始时producer已经完成约77%，可用的overlap窗口仅约1.25us，即使weight ring可容纳9个stage，权重预取的增量收益也仅剩0.208us。提前进入GPU也不是免费的，M=16/32时，prefetch版consumer提前搬运权重，使producer分别变慢0.160/0.352us，但consumer提前完成的工作更多，总span仍相对deepgemm版缩短1.504/1.440us。因此PDL的收益不能仅看consumer duration，也不能将时间线上的整段overlap都算作有效工作，最终必须比较producer开始到consumer完成的总span。

另一个反例是Attention输出投影到MoE RMSNorm的依赖边，consumer在GDC前只能发出一条12KiB的norm-weight TMA，没有可以持续填充的多stage权重流水，当M=16时，consumer开始时producer完成比例为6.9%，timeline overlap为27.81us，span收益为0.54us；M=64时为6.6%、29.54us、0.48us；M=256时为5.3%、33.95us、-0.86us，此时consumer提前参与资源调度的代价已经超过一条norm-weight TMA能隐藏的工作，总span反而回退。

八、 总结

本文针对GLM-5.2-FP8解码阶段的计算流程进行了多维度优化，从最初的80us到最终通过CMP方案降至36.992us。分析表明，Megakernel的核心收益并非tile级依赖，而是独立分支的并行执行和依赖任务的提前准备，这两点可以通过Multistream和PDL高效实现。在M=16的场景下，tile级依赖并未带来显著的额外收益，因为关键路径上的consumer仍需等待上游GEMM的全部输出完成。同时，Megakernel的实现存在固定框架开销和算子迁移成本，在该场景下相比CMP方案并无明显优势。该研究为大模型解码阶段的计算优化提供了新的思路，即在不同的场景下选择最合适的优化方案，平衡实现复杂度和性能收益。