[nvrohanv]

Motivation

At one to four concurrent requests, a GLM-5.3 NVFP4 decode step on 8 x B300 is bound by launch gaps and small-M GEMMs:
each MTP verify step runs 78 layers of attention, MoE and tensor-parallel all-reduces at 4 to 16 rows, and the NextN
draft step (3 draft-decode steps and one draft-extend step per decode step) runs SGLang's forward with a logits
all-gather over the full vocabulary and separate launches for the embedding, the eh norms and the eh projection. This
PR adds two GLM-5.3 specs to the plugin (they also serve GLM-5.2, below), through SGLang's megakernel layer
(--enable-nvidia-megakernel, sgl-project/sglang#44945), in the same layout as the DeepSeek-V4.1 spec: a target spec whose step-function executor
runs CuTe DSL kernels at the covered verify widths, each with a stock path that reproduces SGLang's verify step bit for
bit, and a draft spec for GlmMoeDsaForCausalLMNextN that runs the draft step in the draft-decode and draft-extend CUDA
graphs. Both serve MTP with 3 draft steps (4 draft tokens) and with 5 draft steps (6 draft tokens, 5/1/6), the draft
setting of SGLang's low-latency GLM-5.3 and GLM-5.2 NVFP4 recipes for B300.

Modifications

Target spec:

• sgl_nvidia_mega/glm5_3/ (new):
  • spec.py: the setups the spec serves, checked at argument resolution before any weight load: SM103, TP 8, BF16,
    full decode graphs, EAGLE / NEXTN with topk 1 and 4 or 6 draft tokens (3 or 5 draft steps), FP8 E4M3 KV with page
    size 64, the DSA backend
    with the TRT-LLM decode kernel, the FlashInfer TRT-LLM MoE runner with shared-experts fusion off, the checkpoint
    shapes and routing in HF_CONFIG (one expert group, sigmoid scores, the renormalized top-8), and the CuTe DSL
    version the kernels were validated on. Every other setup stops at startup,
    naming each reason; there is no silent fallback with the flag on. NUMERICS holds one tolerance per kernel
    family, each naming the test that measured it.
  • plan.py, runtime/: the build (buffers, derived weights and every kernel of every width compiled at load,
    before the KV cache is sized) and the executor. A step is covered when it is a target-verify forward of 1 to 4
    requests (4 / 8 / 12 / 16 rows) with 4 draft tokens, or of 1 or 2 requests (6 / 12 rows) with 6, recorded while
    SGLang captures a full decode graph; every other forward, and every larger verify step, runs SGLang's forward.
    The first capture of each width checks, on every rank, the
    facts the kernels need (among them an interleaved RoPE, which the fused MLA prologue assumes); a reason on any rank
    stops every rank.
  • At startup rank 0 logs the covered verify widths (1 to 4 requests, 4 / 8 / 12 / 16 rows, with 4 draft tokens; 1
    or 2 requests, 6 / 12 rows, with 6) and the derived weight copies' memory per GPU (about 20 GiB), and warns when
    --max-running-requests is set above 8: larger verify steps run SGLang's path while the copies still take
    KV-cache memory, so the megakernel helps little on high-concurrency servers.
  • sm103/hop/api.py builds the router-hoist weights of the o_proj hop under torch.no_grad(), so the build keeps
    no autograd state and its fp64 intermediates are freed before the KV cache is sized.
  • The hop mailboxes' quiesce barriers (sm103/hop/transport.py, passed in by boundary/) run on the TP group's CPU
    group, so the build does not create the TP device group's NCCL communicator before the KV cache is sized.
  • attention/, moe/, boundary/, logits/: the step functions and their stock bodies (adapted from SGLang's model
    code; each adapted file names its origin in its header).
  • sm103/: the CuTe DSL kernels: the small-M q_a / kv_a projection, the MLA prologue, the indexer query, the sparse
    MLA decode and the value up-projection; a small-M NVFP4 MoE engine for the routed experts and the BF16 shared
    expert (1, 2 or 3-4 groups of up to 4 rows: a request's rows with 4 draft tokens; 6 rows run as 4 + 2, 12 as
    4 + 4 + 4); the all-reduce + residual + RMSNorm hops over symmetric memory (o_proj with the
    router on MoE layers, the MoE finalize folded into the MLP-output all-reduce, the final norm). In the MoE
    engines every lane of a warp passes a per-warp named barrier before lane 0 releases a shared-memory buffer, so
    each hand-off is ordered by PTX-defined synchronization.
  • The MoE engines read the routed FC2 weights from compact per-(expert, quarter, CTA) records derived at load (216 MiB
    per MoE layer and GPU, 15.8 GiB per GPU).
• Slots: --enable-nvidia-megakernel runs every slot, with no slot list: attention, moe, hops,
  qkv_a_small_m_shared (the q_a / kv_a projection of the layers that reuse the previous layer's top-k, on the
  small-M kernel; one packed weight copy each), dnorm (the deferred input norm, which moves the rounding points
  within its tolerances) and draft (below). SGLANG_NVIDIA_MEGAKERNEL_SLOTS overrides that list for tests and
  bisecting only; a list without a slot another needs (dnorm needs qkv_a_small_m_shared and hops,
  qkv_a_small_m_shared needs attention) stops the server at startup, naming it.
• Reads of SGLang beyond the binding are declared in each module's SGLANG_READS, and every copied SGLang body is
  pinned in SGLANG_MIRRORS.
• Stated limits: the hops' cross-GPU wait bound is compiled into the kernels (20 s), so the spec refuses a
  --watchdog-timeout below 40 s (the default 300 passes); SGLang's attention branches are checked at the first capture
  of each width.

Draft spec:

• sgl_nvidia_mega/glm5_3/draft.py, draft_step.py, draft_attention.py (new): the draft spec (slot draft, on by
  default) and its step. In every captured draft graph:
  • the embedding gather and the two eh RMSNorms run as one kernel on the replicated embedding table (embedding
    replication, which the flag turns on), bitwise equal to SGLang's;
  • eh_proj runs as a column-sharded GEMV gathered by the target's all-reduce hop on its mailbox (1-4 rows, and 8,
    12 and 16 with 4 draft tokens or 6 and 12 with 6);
  • the top-1 is computed distributed, in place of the logits all-gather (up to 8 rows); the draft's forward returns
    the one-hot LogitsProcessorOutput SGLang's topk-1 draft path reads;
  • the draft layer's MLA prologue runs the target's attention kernels (up to 16 rows).
    Each kernel's other row counts run SGLang's code.
• Coverage: the draft spec claims every captured draft-decode and draft-extend graph, so the draft kernels run in every
  captured draft graph at the row counts each supports (the fused draft-decode attention at batch 1 to 16). The target
  spec's limit of 4 requests (16 verify rows) applies to the target verify step only: above it, the target verify step
  runs SGLang's forward.
• sgl_nvidia_mega/glm5_3/sm103/draft/ (new): the CuTe DSL draft kernels (embedding + eh norm front, eh projection
  shard, distributed argmax).
• With slot draft, the target spec also requires embedding replication (SGLANG_ENABLE_EMBED_REPLICATION, which
  SGLang sets for the flag when it is unset, spec.FLAG_SETS_ENV; an explicit 0 is refused) and greedy draft
  acceptance (no --speculative-use-rejection-sampling), and the draft spec refuses an SGLang whose EAGLE draft graph
  runners do not run a draft model's capture check (spec.EAGLE_DRAFT_CAPTURE_CHECK). The draft finds the target's
  hop kernels through the embedding weight the two share; a requested slot it cannot find stops every rank at startup.

5 draft steps (6 draft tokens, 5/1/6):

• Served: EAGLE / NEXTN with topk 1 and --speculative-num-draft-tokens = --speculative-num-steps + 1, at 4 (3 / 1 / 4,
  as before) or 6 (5 / 1 / 6). Any other draft length stops the server at argument resolution, before any weight load,
  naming the served values (plan.SERVED_TOKENS_PER_REQUEST), and the build refuses it again.
• Widths follow the verify rows a request T (plan.supported_rows: every whole number of requests up to the kernels' 16
  rows). At T 6, verify steps of 1 or 2 requests (6 / 12 rows) run the kernels; steps of 3 or 4 requests (18 / 24 rows)
  run SGLang's forward, as steps above 4 requests do at T 4, and rank 0's coverage line says so.
• MoE: moe_split(rows) groups a step's rows by 4 with the remainder last: 6 rows run on the two-group engine as (4, 2)
  and 12 rows on the four-group engine as (4, 4, 4). The engine's groups are token sets over one expert list, not
  requests, and every map of the same rows gives the same bytes (kernel tests below).
• Draft: the eh_proj shard's rows are derived from the widths (eh_rows): 1-4, 8, 12 and 16 at T 4, as before, and
  1-4, 6 and 12 at T 6.
• The one-request forms at 6 rows: a one-request verify step at T 6 has 6 rows, past the 4-row limit of the faster forms
  that serve one request at T 4. The split MoE fold hop (SPLIT_MAX_ROWS 6; still taken only where the occupancy probe
  at build admits it), the plain all-reduce hops on the dense layers (DENSE_MAX_ROWS 6) and the 16-warp sparse MLA
  decode now serve up to 6 rows, so one-request steps at 6 rows run them, as one-request steps at 4 rows do. Besides
  faster one-request decode, this makes 5 draft steps give the tokens of 3 draft steps (up to 2,048 tokens of context)
  in the one-request test under Accuracy Tests. No kernel source changed.
• At T 4 every width, MoE row map and engine, eh_proj row set and kernel form is the one served before (no T 4 width
  has 5 or 6 rows).

Both:

• plugin.py: two rows in SPECS (GlmMoeDsaForCausalLM, GlmMoeDsaForCausalLMNextN). README.md: the model and
  draft rows, the slot table, requirements and the launch command. NOTICE: one paragraph for the SGLang-adapted stock
  step. .github/CODEOWNERS: one line.
• Launch: the README's GLM-5.3 / GLM-5.2 command, the checkpoint's serving flags plus --enable-nvidia-megakernel only;
  no slot list, no environment variable and no other plugin setting (the SGLang PR below turns on the embedding replication
  slot draft needs).
• Merge order: the SGLang PR that adds the DSA binding view (ForwardView.dsa) and the EAGLE draft graph capture check
  ([Feat] [NVIDIA] megakernels: DSA / MLA binding view, EAGLE draft capture checks, early draft embedding sharing sglang#46210) first; without the binding view this PR refuses at startup, naming each missing member, and without the
  capture check the draft spec refuses. That PR also has the flag turn on the embedding replication slot draft needs;
  on an SGLang without that step, a launch must set SGLANG_ENABLE_EMBED_REPLICATION=1 itself.
• SGLang's megakernel API version 2 ([Feat] [NVIDIA] megakernels: Qwen3.5 binding and EAGLE entry wiring for NVIDIA megakernel plugins sglang#45266, on SGLang main): a built object may also offer the EAGLE
  entries draft_loop, draft_extend and begin_draft_round and attention graph_variants, which SGLang reads from
  every built object. The GLM-5.3 / GLM-5.2 target and draft built objects offer none of them (None: SGLang's own
  path, as before), and both specs declare version 2, as main's DeepSeek-V4.1 and Qwen3.5 specs do. SGLang refuses a
  spec of another version, so this PR now needs an SGLang with version 2; on an older SGLang (version 1, such as SGLang
  main dba22e9b68) the specs are refused at startup.
• SGLang's megakernel interface: [Refactor] kernels: Add the plugin kernel backend and split plugin contracts from their srt integration sglang#45389 moved it from sglang.srt.megakernels to
  sglang.kernels.plugin_kernels.nvidia.spec (the spec contract) and sglang.srt.kernel_plugins.nvidia.binding (the
  binding, guards and runtime). The package imports each name from the new location and falls back to
  sglang.srt.megakernels on an older SGLang, as deepseek_v41 does, and each SGLANG_READS key names both files
  (new|old). With API version 2 (above), only an SGLang with the new location serves the specs; the fallback keeps
  the package importable on an older SGLang, which then refuses the specs, naming the version.
• test/unit/test_contract.py (existing): checks the DSA binding view's members and its TOLERATED_DSA_* lists, and
  each SGLANG_MIRRORS pin against the installed SGLang; against an SGLang without the DSA binding view, those
  checks are skipped, naming it, and the read of the draft capture check marker is skipped, naming it, on an SGLang
  without it. It reads each SGLang file at the first of its locations that exists, and it counts the public
  module-level functions of the binding file (SGLang's megakernel API, such as megakernel_agree, which #45389 moved
  there) as API rather than binding view callables. Its SGLANG_MIRRORS check is main's ([Perf] [GB300] qwen3_5: attention, GDN front and out_proj, shared expert and the MTP draft in the plugin; TP4, and two and four requests at TP8 behind an opt-in #20:
  mirror_digest and tuple pins); the 18 GLM-5.3 pins are written in that digest, and the four whose SGLang body
  changed on SGLang main 286e018594 (below) are tuples of the digests before and after the change, so they pass on
  that main and on dba22e9b68. On SGLang main 286e018594, and again on SGLang main 00394f9f35, each with the
  SGLang PR above and the EAGLE step-graph PR merged with that main, none of pytest test/unit's GLM-5.3 / GLM-5.2
  checks fails or is skipped; the ids that fail there also fail with sgl-nvidia-mega main alone (other models' rows on
  that SGLang).
• Tests in this PR: the contract test update above and the GLM-5.2 config test (below). The GLM-5.3 kernel tests (each
  kernel family against SGLang's or FlashInfer's op, eager and in CUDA graph replay, with 6- and 12-row cases for 6
  draft tokens) and the CPU tests of the 5-draft-step support (13 tests: the served and refused draft lengths at
  argument resolution and at build, the widths, the T 4 maps, eh rows and forms, the 6- and 12-row maps, the build and
  the startup line at T 6) follow in a separate PR; the kernel-test rows below ran them.

GLM-5.2: GLM-5.2 has GLM-5.3's architecture and the same value of every config field the two specs pin, so the same
specs serve GLM-5.2-NVFP4. The GLM-5.2 support changes no code path in the package: it names both models where the
plugin names the model (the spec names GLM-5.3 / GLM-5.2 and GLM-5.3 / GLM-5.2 NextN draft, and the log, refusal
and rank-agreement texts), lists GLM-5.2 in the README, and adds test/unit/test_glm5_3_glm5_2_config.py (CPU), which
runs the published nvidia/GLM-5.2-NVFP4 config through SGLang's requirement check for both specs and checks that a
config that differs in any pinned value, or lacks a pinned field (index_topk), is refused. On GLM-5.2-NVFP4 (8 x B300, TP 8, MTP), on the full
stack (see Accuracy Tests) with this PR at its commit before the rebase (the same kernels), the megakernel engages,
the test-only parity mode finds the plugin's stock target verify and draft steps equal to SGLang's (max abs difference
0), and GSM8K-200 scores 0.975 with the megakernel and 0.960 without it. GLM-5.2 results with 5 draft steps are below.

Base and revisions: this branch is this PR's commit ad0bf9103d on sgl-nvidia-mega main 6fd7ac3037, a README
correction of the GPU test commands (f6135dbba1), a merge of sgl-nvidia-mega main 9a21cb037a, six commits on top
(the 5-draft-step support (two commits, above) and its README rows, the 18 GLM-5.3 mirror pins written in main's mirror
digest (the same SGLang code), four mirror pins updated for SGLang main 286e018594, and the port to SGLang's megakernel
API version 2 (above)), a merge of sgl-nvidia-mega main 54b4289217, a docstring correction in
glm5_3/boundary/__init__.py (the dense layers' plain hops serve up to 6 rows), and two mirror pins updated for
SGLang main 9e9296fdf3 (below).

• Merging main 9a21cb037a, NOTICE, README.md and test/unit/test_contract.py conflicted with the Qwen3.5 specs of
  [Perf] [GB300] qwen3_5: attention, GDN front and out_proj, shared expert and the MTP draft in the plugin; TP4, and two and four requests at TP8 behind an opt-in #20; each keeps both sides, and test_contract.py keeps one SGLANG_MIRRORS check,
  main's. .github/CODEOWNERS and sgl_nvidia_mega/plugin.py merged without a conflict.
• Merging main 54b4289217 ([Fix] deepseek_v41: fix engram hosttable no attribute layout problem. #29, [Fix] deepseek_v41: Read MoeDeferredFinalize from its new module first #28 and [Fix] deepseek_v41: Build FlashMLA schedules ahead only where SGLang would build them #27, all in the DeepSeek-V4.1 package),
  test/unit/test_contract.py conflicted with [Fix] deepseek_v41: fix engram hosttable no attribute layout problem. #29, which lets a read name accept either of two fields
  (replication|layout). It keeps both: this PR's skips of the DSA view and draft capture check reads on an SGLang
  without them, then the reads check in [Fix] deepseek_v41: fix engram hosttable no attribute layout problem. #29's form. No GLM-5.3 file changed in this merge.
• For SGLang main 286e018594, four mirrored bodies changed; each pin is now a tuple of the digests before and after,
  and the copied steps stay as they are: forward_absorb_prepare ([Feat] dcp: Replicate a block-FP8 q_b_proj under --dcp-replicate-q-proj sglang#46359: every changed line is on
  the --dcp-replicate-q-proj path, which runs only with decode context parallelism, which the specs refuse),
  DeepseekSparseAttnBackend._forward_trtllm ([Refactor] attention: Use typed KV views and shared layer routing sglang#46990: the KV store and read go through the pool's
  layer router, which on the DSA pool makes the same set_mla_kv_buffer and get_key_buffer calls with the same
  arguments), DeepseekSparseAttnBackend.forward_decode ([Refactor] attention: Use typed KV views and shared layer routing sglang#46990 as above, and
  [Feat] [AMD] dsa: Enable DCP on the triton sparse-MLA backend sglang#44322, a branch taken only on HIP with decode context parallelism) and LogitsProcessor.forward
  ([Feat] spec: Capture aux hidden states on the target without a speculative draft sglang#43449: the new lines run only with --enable-return-aux-hidden-states, which turns speculative
  decoding off). The other 14 mirrored bodies are unchanged on that main.
• For SGLang main 9e9296fdf3, [Feat] sail: support sail fa3 and flashmla [3/N] sglang#45946 changed forward_absorb_prepare and forward_absorb_core:
  fp8 conditions that differ only on SAIL and, in forward_absorb_prepare, a llama_4_scaling step taken only on the
  FlashMLA backend. The DSA path these specs serve is unchanged, so both pins now also accept the new digests.
• The GLM-5.3 package (sgl_nvidia_mega/glm5_3/) differs from the build measured for the 5-draft-step results below
  only by the mirror pins, the API-version port and that docstring; its kernels (sm103/) are byte-identical.
• On SGLang main 286e018594 with [Perf] hicache: Enable async ready-count sync by default sglang#45795 and with the hooks PR and the EAGLE step-graph PR as merged
  with that main, this branch was checked on 8 x B300 (end of Accuracy Tests). After both SGLang PRs merged SGLang
  main 00394f9f35, the CPU tests above ran again with them. After both then merged SGLang main 9e9296fdf3, this
  branch was checked with them and [Perf] hicache: Enable async ready-count sync by default sglang#45795 on 8 x B300 with GLM-5.3 NVFP4, launched with
  --enable-nvidia-megakernel --enable-speculative-step-cuda-graph --hicache-async-ready-sync, at 3 and at 5 draft
  steps: the target and draft specs were active on 8 / 8 ranks, and GSM8K-200 (5-shot, greedy) at 4 concurrent
  requests scored 0.995 with 3 draft steps and 0.990 with 5 (one run each).

Earlier revisions:

• The accuracy results and the per-user decode table below were measured before the rebase, on sgl-nvidia-mega main
  d47deb3390 and SGLang main af856717c2, with this PR's commit of that time or the earlier revisions each one
  names. Rebasing onto 6fd7ac3037, .github/CODEOWNERS, README.md, sgl_nvidia_mega/plugin.py and
  test/unit/test_contract.py conflicted with [Feat] [GB300] kimi_k3: Add NVIDIA megakernel support for Kimi K3 TP8 #17 and [Fix] plugin: Import SGLang's megakernel interface from its new location first #21 on main, and each keeps both sides.
• For SGLang main 74c08f3dc5, four SGLANG_MIRRORS pins were updated (forward_absorb_prepare,
  forward_absorb_core, LogitsProcessor._get_logits and the indexer's _fused_q_prepare_and_store): main changed
  these functions in parts that GLM-5.3 does not copy or does not use, except a new stream order in the indexer, which
  the stock _index_k_and_q now follows (same values; stock index path only).
• The agentic-trace results below were measured on SGLang main 74c08f3dc5 with this PR's previous revision
  ef1ead3681 and the SGLang PRs' previous revisions.
• Since then, ad0bf9103d added the GLM-5.2 support (above: string constants, docstrings and comments in the package,
  the README and tests), SGLang's new interface location with the fallback (above), and, for SGLang main dba22e9b68,
  four more mirror pins (forward_absorb_prepare, forward_absorb_core, Indexer.forward_cuda and
  DeepseekSparseAttnBackend._forward_trtllm). Main changed these for attention context parallelism
  ([Perf] dsa: Overlap CP KV gather with indexer computation sglang#46179, and [Feat] glm5.3flash: Support padded CP BCG and prefix replay with PD/DFlash2 sglang#43774 for _forward_trtllm); the new code is reached only with
  attention context parallelism, which the specs refuse, and #43774's slice of the sequence lengths leaves
  them unchanged without it. No file under sm103/ (the kernels) changed.
• On SGLang main dba22e9b68, with the hooks PR, the EAGLE step-graph PR and [Perf] hicache: Enable async ready-count sync by default sglang#45795 ported there,
  ad0bf9103d was checked on one 8-GPU host with dummy weights: the package loads SGLang's interface from the new
  location; with --enable-nvidia-megakernel alone the target and draft specs are active on 8 / 8 ranks; and in the
  test-only parity mode the stock target verify and draft steps equal SGLang's (max abs difference 0). On that base,
  ad0bf9103d was measured for speed only together with the hooks PR, the EAGLE step-graph PR and
  [Perf] hicache: Enable async ready-count sync by default sglang#45795 (see the step-graph PR).

Accuracy Tests

GLM-5.3 NVFP4, 8 x B300, TP 8, MTP (EAGLE, 3 steps, topk 1, 4 draft tokens), on SGLang main af856717c2. The stack
column names what each run had on top of SGLang main: "the hooks PR" is the SGLang megakernel hooks PR (sgl-project/sglang#46210:
DSA binding view and EAGLE draft-graph capture checks); "the full stack" is the hooks PR, the SGLang EAGLE step-graph
PR (sgl-project/sglang#46211) and an earlier version of the deferred HiCache ready-count sync now proposed in sgl-project/sglang#45795,
with the step graph and the deferred sync present but switched off, plus this PR.

check	stack measured	result
stock target path vs SGLang (test-only parity mode: every eager verify step of a covered width runs SGLang's forward and then the plugin's stock step)	the full stack, slot draft only (no target kernel slot, so the covered verify step runs its stock code), CUDA graphs off	68 verify steps at 4 / 8 / 12 / 16 rows, each compared on all 8 ranks: max abs difference 0
stock draft step vs SGLang (test-only parity mode, every eager draft step)	the full stack, slot draft only (no target kernel slot, so the covered verify step runs its stock code), CUDA graphs off	204 draft steps (136 draft decode at 1-4 rows, 68 draft extend at 4-16 rows), each compared on all 8 ranks: max abs difference 0
KL of covered steps against the prefill path	the full stack, slots attention,moe,hops,draft	mean 0.0287 (bound 0.03)
passkey past the indexer's top-k span	the full stack, slots attention,moe,hops,draft	found
GSM8K-200 (5-shot, greedy, 4 concurrent requests, batches move across every width)	SGLang main + the DSA binding view part of the hooks PR + the target spec part of this PR before its last fixes, slots attention,moe,hops	0.990
startup and capture lines	SGLang main + the hooks PR + this PR before its last fixes	the draft spec active and its capture check run on 8 / 8 ranks; the eh gather on the target's mailbox found on 8 / 8
GSM8K-200 (5-shot, greedy, 4 concurrent requests), slots attention,moe,hops,draft	SGLang main + the hooks PR + this PR before its last fixes	0.985, mean acceptance length 3.18
GSM8K-200 by slot set	the full stack, this PR before its last fixes	attention,moe,hops 0.990; + qkv_a_small_m_shared 0.995; + dnorm 0.990; all six slots, the defaults (+ draft) 0.995, mean acceptance length 3.50
GSM8K-200 with the plugin installed and the megakernel off, same server setup	the full stack with an earlier revision of this PR installed, --enable-nvidia-megakernel not set	0.985
kernel tests (to follow in a separate PR: each family against SGLang's / FlashInfer's op, eager and graph replay)	the full stack	within spec.NUMERICS and draft.NUMERICS

On this PR's own stack (SGLang main + the hooks PR + this PR) the startup and capture lines and GSM8K were run; the
parity, KL and passkey rows come from the full stack. The rows above ran earlier revisions of this PR and of the
SGLang PRs, each with the slots its row names. "Before its last fixes" means this PR's code without its final review
changes: the routing pins in HF_CONFIG and the RoPE-style check, which only refuse unsupported checkpoints, and the
per-warp barriers in the MoE engines, whose outputs are bit-identical with and without them (168 MoE kernel-test cases
and 40 served prompts compared). Every run with slot draft in this section and in the speed table below set
SGLANG_ENABLE_EMBED_REPLICATION=1 by hand (it was required then); with the hooks PR, the flag sets that value itself,
as in the agentic-trace run below.

Sanitizers: compute-sanitizer memcheck, synccheck and initcheck pass on the MoE engines with the per-warp barriers, and
synccheck passes on this branch's engines at all five served row maps (4, 8, 12 and 16 rows in groups of 4, and 6 rows
as (4, 2)), with the shared expert on and off. racecheck cannot complete a launch of these persistent engines (with or
without the barriers): every warp arrives at a barrier across a group of CTAs at the start and some wait on it only at
exit, a split barrier that PTX allows and that completes on the GPU but that racecheck does not model (a minimal kernel
with only this barrier pattern hangs under racecheck the same way). A test-only copy of the engines whose only change
moves those waits earlier passes racecheck with 0 hazards at all five maps, with the shared expert on and off, and the
MoE kernel test file passes on it under racecheck with 0 hazards; a race seeded into the copy's epilogue (one named
barrier removed) is reported at every map. The served engines' memory ordering rests on review against the PTX memory
model and the bitwise comparisons above, with that copy's racecheck as support; what no racecheck run covers is stated
with the 5-draft-step results below.

5 draft steps (6 draft tokens, 5/1/6): 8 x B300, TP 8, MTP (EAGLE, 5 steps, topk 1, 6 draft tokens), real acceptance.
"All on" is SGLang main dba22e9b68 with the hooks PR (at 465a9dd5b7), the EAGLE step-graph PR (at 3ba3fd1d2e),
sgl-project/sglang#45795 (at 29e8125875) and this PR's 5-draft-step change on f6135dbba1, launched with
--enable-nvidia-megakernel --enable-speculative-step-cuda-graph --hicache-async-ready-sync; "stock" is SGLang main
dba22e9b68 with none of these PRs or flags. GSM8K servers: --max-running-requests 8, --cuda-graph-max-bs-decode 8,
HiCache write-back with a 270 GB host pool.

check	GLM-5.2 NVFP4	GLM-5.3 NVFP4
startup lines (all on)	target and draft specs active on 8 / 8 ranks; widths 6 / 12; routed batch sizes 1 and 2 captured; draft eh rows 1, 2, 3, 4, 6, 12; the coverage line names 1-2 requests (6 / 12 rows)	the same
GSM8K-200 (5-shot, greedy, chat completions), 2 concurrent requests (every verify step is 1 or 2 requests), all on / stock	0.965 / 0.960	0.990 / 0.990
GSM8K-200, 4 concurrent requests (in steps of 3 or 4 requests the target verify runs SGLang's path; the draft kernels still run), all on / stock	0.935 / 0.935	0.995 / 0.985
mean acceptance length in these runs, all on / stock	5.256 / 5.265	4.251 / 4.098
stock target verify and draft steps vs SGLang (test-only parity mode: every eager verify step of a covered width runs SGLang's forward and then the plugin's stock step; the all-on SGLang with --enable-nvidia-megakernel only, CUDA graphs and HiCache off, 1 to 4 concurrent greedy requests)	280 verify steps at 6 / 12 rows and 2,320 draft steps: max abs difference 0	336 verify steps at 6 / 12 rows and 2,600 draft steps: max abs difference 0
SGLang's GLM-5.2 MTP end-to-end test with 5 draft steps (test_dsa_glm52_nvfp4_tp_mtp_step_graph.py of the EAGLE step-graph PR, on SGLang's 5 / 1 / 6 fixture), run as a local subclass at TP 8 with this PR's launch arguments (as filed, it runs TP 4 with the CuTe DSL all-reduce backend, which the spec refuses)	with --enable-nvidia-megakernel: 3 tests OK, GSM8K 0.942, acceptance length 4.0201, batch-1 acceptance length 4.531 (thresholds 0.925 / 3.7 / 4.0); without it: 0.946 / 4.0352 / 4.377	-

With 3 draft steps and the 5-draft-step change (GLM-5.3 NVFP4): in the parity mode above, 544 verify steps at 4 / 8 / 12
/ 16 rows and 1,632 draft steps, max abs difference 0; the test branch's serving switch test
(test/e2e/test_glm5_3_switch.py) passes (3 tests), its parity check at 552 verify and 1,656 draft steps with max abs
difference 0.

Kernel tests (the GLM-5.3 test branch with its 6- and 12-row cases): all pass. The MoE engines match FlashInfer's
TRT-LLM MoE bit for bit at 6 rows as (4, 2), (3, 3) and (2, 4) and at 12 rows as (3, 3, 3, 3) and (4, 2, 4, 2), and
every map of the same rows gives the same bytes; the attention, hop and draft cases at 6 and 12 rows are within
spec.NUMERICS and draft.NUMERICS; the split fold is served up to 6 rows.

Sanitizers at 6 and 12 rows: compute-sanitizer memcheck, synccheck and initcheck pass on every new or changed form (the
MoE engines at the 6- and 12-row maps, the attention kernels, the hops and the draft eh projection), with no record in
this package's kernels. racecheck passes on the attention, draft and hop tests at these rows (the hops' two CUDA graph
replay tests in a reduced form: their 64 replays as two runs of 32 with a host synchronization between, since unchanged
they exceeded the memory limit set for the run under racecheck), and synccheck passes on those two unchanged. For the
MoE engines, as above, racecheck cannot complete a launch of the served engines and only the test-only copy passes it,
at all five maps, the 6- and 12-row ones included. No racecheck run covers the receive buffers through
which CTAs pass partial results to each other (stored into another CTA's shared memory and read after an mbarrier wait),
which the MoE engines, the sparse MLA decode and the o_proj hop use: in the engine copy above, a race seeded on such a
buffer (that wait removed) was not reported, although it changed the output. These buffers are covered by review against
the PTX memory model, by the kernel tests (bit for bit for the MoE engines) and by repeated launches (400 per kernel and
row count tested), whose outputs were bitwise identical.

The KL check of the GLM-5.3 server test (covered steps against the prefill path, 4 requests in flight) was also run at 5
draft steps. There most decode steps have 3 or 4 requests, so their target verify runs SGLang's path (with the flag, the
draft kernels still run), and the check's value moved across its bound from run to run both with and without
--enable-nvidia-megakernel, so it is not reported as a result here. On GLM-5.2 NVFP4 the same check (the test's own
launch with the GLM-5.2-NVFP4 checkpoint, four runs per server, on the SGLang of the "all on" runs above) gave a mean KL
of 0.0285 with this change and 0.0332 with this PR's filed revision ad0bf9103d at 3 draft steps, and at 5 draft steps
0.0325 with --enable-nvidia-megakernel and 0.0283 without it. Both differences are below 0.005, the run-to-run noise
set for this check before the runs; there, too, only 8-9 % of the sampled decode steps had 1 or 2 requests, so most
target verify steps ran SGLang's path with and without the flag.

On the up-to-date branch: SGLang main 286e018594 with the hooks PR (ba86b681b5), EAGLE step-graph PR
(0d39d7fdbb) and sgl-project/sglang#45795 (at 6ffe9cd46c), and this branch at b8a85c739b (the merge after it
changed no GLM-5.3 file, and the commit after that only one docstring line), launched as above (8 x B300, TP 8, real
acceptance; GLM-5.3 NVFP4 unless the row names GLM-5.2):

check	3 draft steps (4 draft tokens)	5 draft steps (6 draft tokens)
startup lines (all on)	target and draft specs active on 8 / 8 ranks; widths 4 / 8 / 12 / 16; the coverage line names 1-4 requests; the flag's embedding replication line logged once	the same on GLM-5.3 and GLM-5.2, with widths 6 / 12 and the coverage line naming 1-2 requests (6 / 12 rows)
stock target verify and draft steps vs SGLang (test-only parity mode, as above; logits and hidden states)	560 verify steps at 4 / 8 / 12 / 16 rows and 1,680 draft steps: max abs difference 0	336 verify steps at 6 / 12 rows and 2,760 draft steps: max abs difference 0
GSM8K-200 (5-shot, greedy), all on / stock	0.985 / 0.990 at 4 concurrent requests, 0.985 / 0.995 at 8	0.985 / 0.995 at 4 concurrent requests; mean acceptance length 3.710 / 4.162 (below)
GLM-5.2 NVFP4: GSM8K-200, 4 concurrent requests, all on / stock	-	0.945 / 0.960; mean acceptance length 5.273 / 5.182
kernel tests (the GLM-5.3 test branch on this branch; one GPU each, the hops on 8)	10 / 10 files pass, 0 failed; the split fold is served up to 6 rows	(the same files)

At 5 draft steps on GLM-5.3, the 200-question GSM8K runs had a lower mean acceptance length with
--enable-nvidia-megakernel: 3.710 against 4.162 in the run above, or 3.868 counting only the steps after the server
was ready (the step-graph PR's start-up warm-up, whose short requests the stock server does not send, is in the first
figure). A second pair of runs, with a third server that runs the all-on launch without --enable-nvidia-megakernel,
gave 3.989 all on, 4.418 stock and 4.387 without the flag after the server was ready (GSM8K-200 0.990, 0.995 and 0.995).
The all-on server had two answers at the 16,384-token output limit in the first run and one of them again in the second;
in the three runs without the flag that question stopped (at 11,860, 2,852 and 3,037 tokens). In these 200-question runs
the stack without the flag matched stock, so the lower acceptance length and the answers that ran to the limit came with
the flag. Most of the second run's gap is that one question: the decode log lines with one request in flight (81 of 140,
all from it, after every other question had finished) average 3.71, and with 3 or 4 requests in flight the three servers
are level (4.37 all on, 4.41 stock, 4.38 without the flag); in the first run the all-on server was 0.21 lower with 4
requests in flight (4.17 against 4.38).

With one request at a time (10 GSM8K questions, greedy, each run twice), 5 draft steps with the flag gave the tokens of
3 draft steps with the flag: 9 answers identical, with bitwise equal log-probabilities, and the tenth, the question that
ran to the limit in both all-on runs above and the only one whose context passed 2,048 tokens, identical until its
context passed 2,048 tokens (its log-probabilities differed at 19 positions, all past that point, by at most 0.026 nat).
Past 2,048 tokens of context, where the sparse attention starts selecting among the cached tokens, two runs of one
request on one server can differ, with and without the flag. Against the same server without the flag, the megakernel's
tokens first differ at the same 10 positions with 5 and with 3 draft steps, and at each of them the megakernel's token
is the first or second choice of the server without the flag (its first at 6 of the 10 when that server scores the
prefix in prefill). The mean per-request acceptance length was 4.336 with
the flag and 4.340 without it (3 draft steps: 3.388 and 3.393).

On the full GSM8K test set (1,319 questions, 4 concurrent requests, one run per server, the same stack with and without
the flag), 5 draft steps scored 97.27 % with the flag and 97.50 % without it (5 questions right only with the flag, 8
only without; exact McNemar p 0.58) and 3 draft steps 97.50 % and 97.57 % (4 and 5; p 1.0); 2 and 1 answers reached the
16,384-token limit at 5 draft steps, 3 and 4 at 3 draft steps. The question at index 423 of the test split reached it
with the flag at both settings and was answered correctly without it; the question above (index 119) reached it at 3
draft steps with and without the flag and was answered correctly in both 5-draft-step runs. Over the 200-question and
full-set runs and runs with two requests at a time, at 3 and 5 draft steps on SGLang main 286e018594, index 119 ran to
the limit or was answered wrongly in 10 of 14 runs with the flag and in 5 of 15 without it; in the one-request runs
above it stopped, with and without the flag. Acceptance-length parity with stock is not claimed for 5 draft steps on
GLM-5.3; GLM-5.2 on this branch (5.273 / 5.182) and GLM-5.3 on SGLang main dba22e9b68 (4.251 / 4.098, above) did not
show the gap.

Speed Tests and Profiling

Per-user decode throughput on one machine, benchmark/static_table.py protocol (fixed acceptance length 2.99,
--max-running-requests 5 and --cuda-graph-max-bs-decode 4, two launches per server in palindrome order, two passes
per launch; each cell is the mean of the two launches, and each launch's value the mean of its two passes' median
per-user throughput, embedding replication on in every arm). Every arm ran SGLang with earlier revisions of the full
stack's SGLang changes (step graph and deferred sync off, no HiCache): "off" with the plugin not installed, then an
earlier revision of this PR with slots attention,moe,hops (target only) and with slots attention,moe,hops,draft (+ draft; measured before qkv_a_small_m_shared and dnorm became
default slots, so no column runs all six):

concurrency	--max-running-requests	off: plugin not installed (tok/s per user)	attention,moe,hops (tok/s per user)	attention,moe,hops,draft (tok/s per user)	attention,moe,hops,draft / off
1	5	385.0	569.5	580.3	1.507
2	5	331.3	446.6	452.2	1.365
3	5	292.8	355.1	354.8	1.212
4	5	270.6	303.6	305.7	1.130

At concurrency 5, past the largest width, the megakernel steps aside and the server runs its usual path, which here is
eager decode in every arm (batch 5 is above --cuda-graph-max-bs-decode 4); attention,moe,hops,draft / off there
was 0.997.

On a multi-turn agentic coding trace, the public Hugging Face dataset semianalysisai/cc-traces-weka-062126 (split
train), replayed as recorded (about one hour per run): fixed acceptance length 2.99, HiCache write-back with a 270
GB host pool, --cuda-graph-max-bs-decode equal to --max-running-requests (1 at client concurrency 1), one run per
cell; the runs of each comparison were made on one machine in one session, each with SGLang's compiled Rust radix-tree
core. "Stock" is SGLang main 74c08f3dc5 with --enable-nvidia-megakernel not set. The megakernel run uses this PR's
previous revision (ef1ead3681) and the hooks PR's previous revision on SGLang main 74c08f3dc5 and launches with the
flag only (no slot list and no megakernel environment variable):

• "this PR + hooks": --enable-nvidia-megakernel; the EAGLE step-graph PR (its previous revision) and the deferred HiCache
  sync of [Perf] hicache: Enable async ready-count sync by default sglang#45795 (at its head 29e8125875) are applied but off.

client concurrency	metric	stock	this PR + hooks
1	p90 interactivity (tokens/s per user)	352.11	505.05 (1.43x)
1	p90 end-to-end normalized interactivity (tokens/s per user)	157.64	200.61 (1.27x)
1	p90 time to first token (s)	1.271	1.130 (0.89x)

At one user, this PR with the hooks PR gives 1.43x the stock p90 interactivity and 1.27x the stock p90 end-to-end
normalized interactivity, with a p90 time to first token of 0.89x stock, one run each. It composes with the EAGLE
step-graph PR (sgl-project/sglang#46211) and the deferred HiCache sync of
sgl-project/sglang#45795; see the step-graph PR for runs with all three on. The derived weight copies (about 20.3 GiB per GPU, logged at
startup) leave 3,241,024 KV-cache tokens, against 3,722,304 for stock (-12.9 %). p90 interactivity is 1 / the p90
inter-token latency and p90 end-to-end normalized interactivity is 1 / the p90 end-to-end time per output token, both
in tokens/s per user (higher is better); p90 time to first token is in seconds (lower is better).

With 5 draft steps (6 draft tokens, 5/1/6), on the same trace: GLM-5.2 NVFP4 on 8 x B300, TP 8, fixed acceptance length
3.61 with 5 draft steps and 2.99 with 3 draft steps, HiCache write-back with a 270 GB host pool,
--cuda-graph-max-bs-decode equal to --max-running-requests, SGLang's compiled Rust radix-tree core in every run.
"Stock" is SGLang main dba22e9b68 with none of these PRs or flags; "all on" is SGLang main dba22e9b68 with the hooks
PR (465a9dd5b7), the EAGLE step-graph PR (3ba3fd1d2e), sgl-project/sglang#45795 (29e8125875) and this PR's
5-draft-step change on f6135dbba1, launched with --enable-nvidia-megakernel --enable-speculative-step-cuda-graph --hicache-async-ready-sync. One run per cell; the four servers of each client concurrency ran on one host in one
session.

client concurrency	--max-running-requests	draft steps (fixed acceptance length)	p90 interactivity, stock / all on	p90 end-to-end normalized interactivity, stock / all on	p90 time to first token (s), stock / all on
1	1	5 (3.61)	370.37 / 529.10 (1.43x)	175.41 / 204.44 (1.17x)	1.158 / 1.035 (0.89x)
1	1	3 (2.99)	348.43 / 512.82 (1.47x)	162.46 / 198.22 (1.22x)	1.157 / 1.025 (0.89x)
2	4	5 (3.61)	313.48 / 378.79 (1.21x)	149.80 / 204.80 (1.37x)	0.749 / 0.736 (0.98x)
2	4	3 (2.99)	300.30 / 401.61 (1.34x)	171.47 / 204.82 (1.19x)	0.654 / 0.753 (1.15x)

With 5 draft steps, all on gives 1.43x the stock p90 interactivity at one user and 1.21x at two users. Against 3 draft
steps with all on, 5 draft steps gives 529.10 against 512.82 at one user (+3 %) and 378.79 against 401.61 at two users
(-6 %), with the same p90 end-to-end normalized interactivity at two users (204.80 against 204.82). So the 5-draft-step
support is for serving the 5-step recipes with the megakernel; at these fixed acceptance lengths it gives no speed gain
over 3 draft steps. Real acceptance was not measured in these runs, and neither were steps past the largest width at 5 draft
steps (3 or 4 requests, whose target verify runs SGLang's path). Usage counts on this SGLang main include a known
reasoning-token over-count at max_tokens; these numbers count each request's output tokens as its usage
completion_tokens.

With 3 draft steps the 5-draft-step change keeps every width, map and kernel form, and the startup and coverage lines
were identical with and without it. On the trace (GLM-5.3 NVFP4, fixed acceptance length 2.99, the all-on stack above
with --max-running-requests 1, 4, 8, 16, 32 and 64 at client concurrency 1, 2, 4, 8, 16 and 32; one run per server,
both servers of a pair on one host in one session), the package with the change against the package without it gave
p90 interactivity ratios of 0.995, 0.992, 0.963 and 1.141 (two pairs), 1.020, 1.019 and 1.033, and p90 end-to-end
normalized interactivity ratios of 1.006, 1.046, 0.987 and 1.087, 0.991, 0.984 and 0.975.

Checklist

• pre-commit run --all-files passes.
• Tests under test/ cover the change: this PR updates the contract test (test/unit/test_contract.py) and adds the GLM-5.2 config test (test/unit/test_glm5_3_glm5_2_config.py); the GLM-5.3 kernel tests and the 5-draft-step unit tests follow in a separate PR.
• The SGLang commit or PR this was tested against is named above.
• Accuracy and speed results are attached when kernel or forward code changes.
• A new model adds one line to .github/CODEOWNERS that names its owner.

[YAMY1234]

Rechecked the static table on 8×B300 TP8, GLM-5.3 NVFP4, with ISL8192/OSL1024 and simulated acceptance 2.99. Measured snapshots: mega ad0bf9103d, hooks 465a9dd5b7 + step 3ba3fd1d2e on SGLang base dba22e9b68.

Concurrency	Off	Target	Target + draft	Target + draft / off
1	391.9	587.0	595.8	1.520×
2	348.1	472.5	477.6	1.372×
3	318.1	382.6	387.9	1.219×
4	298.1	332.1	337.5	1.132×

Throughput is tok/s/user. Same static_table.py protocol and slot sets as the original table: attention,moe,hops, then +draft; embedding replication on, step CUDA graphs and HiCache off. Two fresh launches per configuration in palindrome order, two passes per launch; each cell averages four pass medians.

Relative gains closely match the original table. Absolute throughput is higher in both the baseline and enabled arms, by roughly 8–10% at c3/c4. Target-only and target+draft are effectively tied at c3/c4 within the observed pass variation.