Skip to content

[Refactor] kernels/diffusion/tests: Prune new dead code and make silently skipped tests run - #48366

Merged
BBuf merged 70 commits into
sgl-project:mainfrom
BBuf:bbuf/cleanup-round2
Nov 6, 2026
Merged

BBuf merged 70 commits into
sgl-project:mainfrom
BBuf:bbuf/cleanup-round2

Conversation

@BBuf

@BBuf BBuf commented Nov 4, 2026 •

Copy link
Copy Markdown
Collaborator

Motivation

This is the second pass of #44690. Three weeks after that PR, python/sglang/kernels had gained another 66K lines, python/sglang/multimodal_gen 43K and test/registered 162K. This PR removes what in that code is dead or redundant. More importantly, it fixes the places where CI is green without testing anything, and the runtime bugs found along the way.

433 files, +5,551 / −13,052.

  • About 1.4K of the additions are tests moved to test/manual/, because no CI runner has the hardware or package they need.
  • About 0.8K are shared test helpers that replace copies.

Bugs fixed (please review these first)

Area Problem Fix
Kernel CI base-b-kernel-unit-test-8-gpu-h200 has no registered test. On every kernel PR and three times a day it books an 8×H200 runner, installs everything, and logs No tests found. Matrix row and suite removed.
Kernel CI Tests registered on lanes whose hardware never satisfies their skip gate. The production kernels behind them had no CI at all. Registered on hardware that runs them, and each was run there:
• SM120 FlashMLA fallback (Ran 15; skipped=15 on H100)
• Blackwell FP4 indexer quant (36 of 39 cases skipped)
• Q8KV8 sparse prefill (25 cases)
• FlashInfer KDA prefill graph and input pack
• SM100 Mamba stochastic rounding
• the FA3 extend-init regression
• KDA MTP slot stride
• Gemma4 Blackwell sampling split
• gfx950-only kernels on the MI300 lane
• diffusion FA4 forward_into
• SubBlock sm100/sm120 kernels (one case skipped on sm_100, see below)
• 4-rank and 2-rank diffusion cases on 1-GPU lanes
Q8KV8 sparse prefill Once the Blackwell cases ran, 15 of them failed on main. #45909 packed each request's compressed history contiguously and updated the H100 cases, but nobody ran the Blackwell cases, so their expected offsets stayed stale. Their output-vs-paged-attention assertion passes. Registration and expectation fix land together; 37 passed on sm_103.
Kernel tests test_kpool_topk_transform_fused.py runs under unittest.main(), which drops the module-level pytest functions. 4 of its cases (ties, overfull bins, page-table mapping under graph replay) never ran. pytest.main entry; scanned the kernel tests for the same pattern.
Kernel tests Assertions that cannot fail:
• a tensor compared with itself (test_rmsnorm);
• an AMD all-reduce "determinism" test with no assert, which only prints ✓/✗ and tests RCCL rather than SGLang;
• *_jit_vs_aot comparisons against AOT ops that do not exist (always skipped);
• a MoE wna16 GEMM test that always skipped (it now checks against a dequant reference and runs its 24 cases).
Fixed, or deleted where nothing SGLang-owned was tested.
Unit tests Triton kernels tested only from CPU-registered files, so their CUDA cases never ran: row_logsumexp[_topk], free_unbind_inplace, gather_page64_kv_latent, the token-id translation Triton path. GPU registrations added (5090).
Unit tests Tests of copies of the code instead of the code:
• the decode CUDA-graph profile options (a test-local copy);
• four MoE shared-expert fusion gates asserting only assertTrue(reason), which the CPU early-return satisfies;
• test_h3_request.py, which failed 8 of 12 cases when run alone because its ComfyUI stubs had rotted. It passed only if another test happened to import the real module first.
The tests now call the real method and assert each branch's message. The stubs are fixed.
CI registration test_gated_launch.py used nightly=True on a CUDA registration, so no job ever ran it. Flag removed; it now runs in nightly. check_registered_tests.py rejects nightly=True on CUDA/CPU.
CI tooling update_est_time.py matched registrations with a single-line regex. It missed 150 of the 2,459 CUDA/CPU registrations (wrapped, positional, float), so their est_time never refreshed. The literal is found through the AST: 0 missed.
Diffusion --srt-encoder-connection-timeout was stored under a name that is not a field, so the documented flag (and the one GLM-Image's timeout error recommends) was silently dropped. type=int also rejected the documented default 3.05. dest= the field, type=float.
Diffusion VMoBA The runtime imported kernel.attn.vmoba_attn.vmoba, which exists nowhere. The chunk-selection fix in #44286 went into the in-tree copy that nothing imported, and its tests sat outside CI discovery. The runtime and resolver import the in-tree module. Its tests moved into multimodal_gen/test/unit.
Diffusion exact tier Wan and LingBot fuse Q/K/V into one GEMM unconditionally. At the model shapes it is bit-identical on H200 and GB300. At M ≤ 256, cuBLAS picks split-K and the result differs, so exact requests were not bit-identical. Wan's fused RMSNorm+RoPE had no check either. The QKV fusion is gated at lossless (same math, rounding moves), the default tier, so default behavior is unchanged. The norm+RoPE fusion goes through BitExactFusionGate.run.
MiniMax-H3 The pipelined Ulysses attention never ran its first-call equality check. Enabling the check exposed that the sequential reference left non-zero pad rows that #47566 zeroes, which would have turned the pipeline off. The reference zeroes the alignment rows (never read downstream); the check runs and passes (same=1).
Registry gemm.dsv3_fused_a_gemm pointed at sgl_kernel:dsv3_fused_a_gemm, which does not exist. The namespace test only checked that the module existed. Spec removed. test_inventory_targets_resolve now resolves every target's attribute statically (incl. the AOT package) on the CPU lane.
CuTe DSL GDN/KDA decode The wrappers accepted scale, use_qk_l2norm_in_kernel, softplus_beta and softplus_threshold, then compiled with fixed values. Today's callers pass exactly those values. Non-matching values raise ValueError (+ test).
ROCm block-FP8 use_w8a8_block_fp8_matmul_unrolledx4 evaluated its condition without return, so production never selected the unrolledx4 kernel. Only the tuner ran it. Kernel, selector and tuner branch removed. Runtime behavior is unchanged.
Unlimited-OCR / QK-norm gate Their PDL probes checked capability >= 9, which is true on gfx942/950, so they emitted PDL on ROCm. current_arch().supports_pdl.

Modifications

Kernels (python/sglang/kernels, 108 files, −3.1K net)

  • Dead code (each re-checked by whole-repo grep, incl. string targets):
    • 865 lines of jit/csrc/moe/expert_specialization/*.cuh that nothing loads;
    • ngram decode fast paths that were never wired in, and their benchmarks;
    • the MuseSpark router_gemm JIT kernel (the model only uses its token limit);
    • the test-only TextSeal standalone selector;
    • two unused LoRA MoE custom ops;
    • LPLB shmem helpers, row_argmax and create_argmax_workspace;
    • the f32x2 helpers left by a reverted PR;
    • finish_flat_indexer_topk;
    • the small_field_copy descriptor path;
    • an unloadable LoRA plans table;
    • losing mla_qk_prep variants (the default path is byte-identical on H200 and GB300 over 64 cases);
    • unused use_tma/topk/config parameters;
    • the MuseSpark softmax routing branches;
    • aliases, constants and re-exports nothing uses.
  • Classification ([Refactor] kernels: Regroup operators by computation and dissolve ops/muse_spark #45112 rules):
    • engram_gather_peer (peer-memory read) → communication.
    • MSA-v1 compact_dequant (KV-cache gather + dequant) → kvcache.
    • The two diffusion op ids that targeted quantization/activation modules are renamed to their real groups; their facade re-exports are gone.
    • qwen_moe_topk_output moves to its only caller in qwen2_moe.py.
    • The gfx1250 per-token FP8 quant → quantization.
  • Arch queries: about 15 sites that still probed the device by hand now use sglang.kernels.arch ([Refactor] kernels: Centralize GPU architecture queries in sglang.kernels.arch #47782) with the same semantics. argmax/softmax JIT adopt max_vec_bytes. Import-time device-0 probes moved to call time.
  • Dedups:
    • the gfx950 gluon MoE helpers (identical AMDGCN for all 68 cases);
    • the hipcc compile-and-load block;
    • the SM90 FP8 CuTeDSL LoRA class, now a subclass of its BF16 sibling (bit-exact on H200);
    • the gfx950 vattn HIP binding.
  • Env: raw os.environ reads in kernels/models are declared in environ.py with the same names and defaults. Two settings are removed:
    • SGLANG_HIP_MHC_BOUNDARY_WIDE_TILE, a kill switch for a bitwise-identical tile;
    • SGLANG_XPU_LN_MOD_NUM_WARPS, a layout-only setting.
  • HiCache: the page-unified layout is admitted only if both the write-back and the load-back JIT probes pass; until now the load-back probe was never called.
  • Docs: cookbook lines that told users to set env vars nothing reads are removed.

Diffusion (python/sglang/multimodal_gen, 118 files, −1.5K net)

  • Dead code:
    • 7 disaggregation flags nothing read;
    • --mask-strategy-file-path (sliding-tile attention reads --attention-backend-config mask_strategy_file_path=..., and the error now says so);
    • the /stats timeout counter that was always 0;
    • vestigial PortArgs fields that leaked three temp files per construction;
    • about 25 unused helpers, constants and config fields;
    • test oracles moved out of runtime modules into the tests.
  • Shared helpers:
    • four more hand-written verify-once gates use BitExactFusionGate.run / can_attempt (flux3 video VAE, FLUX.2 XPU, KL-VAE, Qwen-Image-2.1);
    • Wan and LingBot share one row-packing helper;
    • the duplicated Nomas helpers, OpenAI image/video/mesh API helpers and flux3's copy of the launcher conditions are merged;
    • the two disagreeing "unrecoverable device error" classifiers are now one.
  • Docs:
    • --enable-attention-backend-autotune was documented but not a CLI flag; it is now one.
    • ring_sp_performance.mdx documented a flag that does not exist; it now uses the real --kv-gather-degree.

Tests and CI

  • Deleted:
    • tautologies (argparse pass-through, default and constant echoes, inspect.getsource/AST source tests, mock-the-SUT, torch-math);
    • duplicate cases and files;
    • benchmarks of test-local copies;
    • always-skipped tests;
    • #47603 leftovers (dead model lists, a test that returns in CI);
    • tests with a hard-coded developer path.
  • Merged copy-pasted fakes, fixtures and oracles into python/sglang/test/ helpers. Two test classes were collected twice through imports (45 extra runs per CPU job).
  • Trimmed grids with no coverage gain: page-unified dtype sweep (252 → 168 cases), MiniMax decode top-k (144 → 48).
  • Faster CPU gate:
    • test_runtime_context.py: 132 s → 13 s;
    • test_flashinfer_autotune_sync.py: 30 s → 10 s;
    • three one-purpose files merged.
  • Moved to test/manual/: tests that need NCCL ≥ 2.31.2, a private checkpoint, an unset env var, or 4 GPUs of a manual-only class; plus a benchmark CLI that lived in a registered test.
  • CI:
    • the empty 8-GPU XPU nightly job is removed;
    • stale suite names, globs and rerun_test_groups entries are removed;
    • the Muse Spark HF token follows the runner, not the stage;
    • comfy-kitchen (no deps) is installed for the diffusion unit/B200 jobs, so its three tests stop skipping.

Behavior changes

  • Diffusion CLI:
    • the 7 unused disaggregation flags and --mask-strategy-file-path are now rejected;
    • --srt-encoder-connection-timeout takes effect and accepts floats.
  • Diffusion exact requests no longer use the fused QKV GEMM (Wan, LingBot). lossless and high outputs are byte-identical to main.
  • HiDream-O1 rejects --flow-shift/--scheduler-class-override, as it already does for its other fixed settings. HunyuanImage-3 honors a user --flow-shift instead of overwriting it.
  • /v1/images/edits returns 400 when mask is given; it was silently ignored.
  • The diffusion client treats CUDA error:, HIP error: and torch.AcceleratorError as unrecoverable, like the worker does.
  • ROCm: no PDL in the Unlimited-OCR DASH / QK-norm gate kernels. The gfx1250 kpool indexer uses Triton act_quant, like the other two DSA paths.
  • CuTe DSL LoRA prepare: SM120 now gets a clear NotImplementedError instead of compiling the SM100 class and failing.
  • Removed env vars: SGLANG_HIP_MHC_BOUNDARY_WIDE_TILE, SGLANG_XPU_LN_MOD_NUM_WARPS. Newly declared boolean envs (SGLANG_ROCM_SMALLM_FP8_BS, SGLANG_DISABLE_LFM_FUSED_CONV) parse like every other EnvBool.

Left alone (owner decisions, not in this PR)

Accuracy Tests

Full regression sweep, base (4c7c775c) vs this branch, same box. Each test file runs in its own pytest process. Failing test ids are compared:

Box Files Test ids failing only on this branch
H200 (sm_90) test/registered/kernels/ops (448), test/registered/unit (1153), multimodal_gen/test/unit (395) 0
GB300 (sm_103) test/registered/kernels/ops (448) 0 (15 fewer: the Q8KV8 Blackwell cases)
RTX 5090 (sm_120) test/registered/kernels/ops (448) 0 (the same 833 pre-existing ids)

The remaining failures are identical on both trees: multi-GPU tests on a 1-GPU box, optional packages missing from the image, and kernels that need hardware the box lacks.

End to end, fixed seed, base vs this branch. An instrumented run confirmed the fused paths were active on both sides.

Model Box Result
Wan2.1-1.3B H200 lossless, exact and high mp4 byte-identical; fused QKV GEMM active (M=14040); norm+RoPE gate verified
Wan2.1-14B GB300 lossless and exact byte-identical
Qwen-Image-2.1 H200 byte-identical; 10 DiT gates verified (incl. the converted ones)
LingBot-World realtime (CI case) H200 exact and lossless raw frames identical; fused GEMM active under lossless
MiniMax-H3 T2VA, Ulysses 2 GB300 lossless and exact byte-identical; pipelined attention first-call check same=1

Kernel dedups, bit-for-bit (torch.equal on dumped outputs, base vs branch):

  • the CuTeDSL FP8 LoRA class merge, gelu_and_mul_triton, MuseSpark routing, argmax and the SM10x JIT softmax: 88 tensors on H200 and 112 on GB300;
  • mla_qk_prep: 64 cases on H200 and GB300;
  • the gfx950 gluon helper hoist: identical AMDGCN for all 68 cases (372 compiles via Triton's AMD backend).

Lane changes were run on the hardware of their new lane:

  • SM120 cases on an RTX 5090;
  • SM100/SM103 cases on GB300;
  • H100-lane cases on H200;
  • the new 4-GPU diffusion cases on 4×GB300.

CI tooling:

  • pre-commit run --from-ref 4c7c775c passes over the whole range, including the workflow checks.
  • check_registered_tests.py passes; on main it fails, naming the nightly=True file.
  • Suite selection was simulated with run_suite.filter_tests, and compute_partitions.py gives the same shard counts for pr-test, pr-test-extra, nightly and weekly.
  • update_est_time.py dry run: 1402 updates, up from 1372; the old 1372 keep their values.
  • The diffusion partition planner reports 0 missing estimates, and both newly scheduled 2-GPU files ran on 2×GB300.
  • An import walk over sglang.{kernels,multimodal_gen,srt,test} finds no new import failures.
  • After merging main 851642ad, the 18 test files touched by the newly merged upstream commits pass on main and on this branch.

Not validated:

  • true SM100 / B200 (no B200 machine was available; sm_103 stood in for every major == 10 gate);
  • ROCm, NPU, XPU and MUSA (those changes are mechanical and import-checked);
  • the VMoBA CUDA path (FlashAttention 2 is not in the image);
  • FLUX.3 e2e (gated weights);
  • HiDream-O1, HunyuanImage-3 and Nomas e2e (unit tests only).

Speed Tests and Profiling

No perf-path changes intended:

  • Default-tier fused paths are unchanged.
  • The removed kernels and branches were not selected in production.
  • The mla_qk_prep default variant is byte-identical.

Checklist

🤖 Generated with Claude Code


CI States

Latest PR Test (Base): ⏳ Run #42123536075
Latest PR Test (Extra): ❌ Run #42123536315
Latest PR Test (AMD ROCm 10): ⏳ Run #42123536434

Tip

Base CI only covers the core cases, and its workload may not exercise your change. Use /rerun-test to run the tests your change touches, or copy this prompt for an agent:

Find the registered tests this PR affects and run them with /rerun-test, following https://github.com/sgl-project/sglang/blob/main/.agents/skills/rerun-test/SKILL.md

BBuf and others added 30 commits November 4, 2026 13:28
…d disagg flags

The encoder connection timeout flag stored into a name no field has, so the
documented value never reached the encoder health check; it also parsed as int
while the default is 3.05. Seven disaggregation flags were parsed into fields
nothing reads (the transfer backend is always Mooncake); remove them with
their parse-only tests.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- gemm: drop the HIP unrolledx4 block-FP8 kernel; its selector never
  returned True, so only the tuner ran it.
- embeddings: drop the never-wired ngram decode fast paths (Python, CUDA
  kernels, two benches); the test now checks the general kernels against
  a reference.
- gemm: drop the unused MuseSpark router_gemm JIT kernel.
- sampling: drop the standalone TextSeal selector, create_argmax_workspace;
  memory: drop the device-descriptor small_field_copy entry (tests and
  bench now exercise the SmallFieldCopyPlan path production uses).
- lora: drop the unregistered fused_moe_lora shrink/expand custom ops and
  the unloadable default MoE LoRA plan table.
- lplb: keep only the H100 shmem budget check the solver calls.
- speculative: drop row_argmax (no caller); activation: drop the
  gelu_and_mul_triton quantize path that could only raise.
- moe: drop a dead if/elif-pass platform branch and a commented print.
- srt: delete cutlass_moe_params.py, only a test used it.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- environ: declare SGLANG_ROCM_SMALLM_FP8_BS, SGLANG_DISABLE_LFM_FUSED_CONV
  and SGLANG_LORA_CONFIG_DIR (same names and defaults) and read them via
  envs; the smallm FP8 block-scale build failure logs a warning instead
  of printing.
- hicache: page_unified admission now also probes the load-back JIT
  kernel, so a load-back build failure refuses the layout up front
  instead of surfacing on the first load.
- dsa: dsa_indexer_kpool picks act_quant with the same TileLang/Triton
  predicate as dsa_indexer and dsa_prefill_cuda_graph (gfx1250 now takes
  Triton there too).
- lora: bound the CuTeDSL grouped-GEMM specs to SM90..SM10x and pick the
  tcgen05 class only on SM10x (SM120 now raises instead of picking it).
- moe topk: import the in-repo moe_topk_sigmoid without a silent
  ImportError guard.
- Drop knobs nothing uses: verify_tree_greedy_func's topk, softmax_jit's
  config override, the never-set softmax branches of MuseSpark routing,
  and the SGLANG_HIP_MHC_BOUNDARY_WIDE_TILE kill switch of a bitwise
  identical tile.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
SGLANG_ENABLE_FP8_BLOCK_GEMV, SGLANG_ENABLE_FUSED_KDA_OUTPUT_QUANT
(Ling-3.1-flash), SGLANG_DSA_TRITON_PREFILL (GLM-5.1 MI355X),
SGLANG_K3_ATTN_RES_MODE, SGLANG_MOE_FUSED_GATE_RADIX (Kimi-K3 H100/H200)
and SGLANG_OPT_USE_JIT_NORM (DeepSeek-V4 B200) have no reader in python/,
rust/, sgl-model-gateway/, scripts/ or docker/; setting them changes
nothing.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The CuTe DSL decode wrappers accepted scale, use_qk_l2norm_in_kernel,
softplus_beta and softplus_threshold but always ran the kernel compiled with
K**-0.5 / True / 1.0 / 20.0. Production passes those values, so nothing is
wrong today, but any other value was silently dropped. Raise instead.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
gemm.dsv3_fused_a_gemm registered sgl_kernel:dsv3_fused_a_gemm, which the AOT
package does not export. The inventory test only checked that sglang.kernels.ops
target modules exist; it now resolves every target statically (module file,
top-level name, class method), with sgl_kernel targets checked against the
in-tree AOT package, so a missing attribute fails on the CPU lane.

Also drop the capabilities_satisfied branch for a bare CapabilityRequirement;
every caller passes a set.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…orwarding wrapper

- moe_route_warp, argmax/softmax JIT, mhc, the MoE LoRA PDL switch and
  the MoE LoRA fp8 scale arithmetic query the device through
  current_arch() instead of raw torch.cuda / srt SM helpers. argmax and
  softmax take the vector width from max_vec_bytes; the fp8 constant
  stays exactly SM 9.0. covered() / muse_spark_route_warp_covered() lose
  their per-device index, like every other current_arch() caller.
- moe: callers import situ_and_mul_masked_post_quant from its module;
  the package-level wrapper only forwarded its arguments.
- quantization.fp8_utils: drop the device / cuda_capability arguments no
  caller passes.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…arch

Migrate the hand-rolled capability / SM-count probes the arch refactor missed
to current_arch(), keeping each probe's semantics. Import-time probes (the
CuTe DSL MTP ring's NUM_SMS / packed-FMA gate, extend/prefill attention's
CUDA_CAPABILITY) now run at call time. Probes that took the gfx major on ROCm
(prefill attention block size, vision RoPE, lightning attention) read
props.major so ROCm keeps its current branch. _is_fa3_supported no longer
takes a device argument it ignored.

Behavior change: the Unlimited-OCR DASH decode and the fused Q/K
RMSNorm+RoPE+gate kernel took PDL whenever the device major was >= 9, which
includes ROCm gfx9; they now use current_arch().supports_pdl, false on HIP.
Not validated on ROCm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
SGLANG_ASM_VERIFY_ATTN, SGLANG_ROCM_KDA_HIP_PREFILL,
SGLANG_ROCM_K3_KDA_FUSED_BACKEND and SGLANG_K3_FLYDSL_SOURCE were read with
raw os.environ. Declare them in Envs with the same names, defaults and
string parsing, read them through envs, and document them.

unified_attention_3d_mtp wrapped the in-tree vattn_asm_gfx950 import in
except ImportError, which only hid a broken module behind a silent Triton
fallback; import it directly. Its CU count now comes from current_arch().

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- jit/utils/arch.py resolved CUDA_HOME itself; use toolchain.cuda_home().
- hipcc_path() read ROCM_PATH only while toolchain.rocm_home() prefers
  ROCM_HOME, so the two could pick different hipcc; use rocm_home().
- check_hip prefixed every HIP error with "smallm_moe:" although six gemm,
  moe and attention modules share it.
- vattn_asm_gfx950 carried a copy of the hipcc ctypes binding and error check;
  use hip_lib()/check_hip(). Its 3-D launcher stays. Not validated on ROCm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- jit/csrc/moe/expert_specialization/*.cuh (865 lines): a JIT copy of the AOT
  ES-MXFP8 grouped GEMM/quant that no load_jit, #include or build file reaches;
  production uses the sgl_kernel AOT ops.
- math.cuh fma/mul/add_f32x2: left behind by a reverted change, no caller.
- dsv4 fp4_indexer finish_flat_indexer_topk and its kernel: no caller outside
  its own two tests, which go with it.
- helion_config_bucket and its ContextVar: never used, not even by tests.
- unused use_tma parameters of three MiniMax sparse attention entry points.
- the GDN-only ReplaySSM aliases (callers now use the real name),
  h_v2b_cutedsl, a stale cache dict in kda_nvidia_prefill, six unused DSV4
  layout constants, setup_sail's all_common_sources, and five
  jit.utils.compile re-exports nobody imports.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
absorbed_bmm_concat_cast_q_fp8 kept every K=192 candidate from its A/B
(loop, three_dot, pad, single_k) plus block_k/num_stages knobs; production
only reaches the single-dot path (power-of-2 K) and two_dot ("auto" for
other K). Keep those two and the "cuda" variant, and drop the A/B from the
microbenchmark.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… one module

_add, _encode_groups, _quantize_input, _unpack_words, _word_bytes and
_weight_offset were pasted verbatim into up to seven fused-MoE family
files; each family now imports the one copy from gluon_moe_gfx950/common.py.
Variants whose body differs stay local to their family. Every family
compiles to the same gfx950 AMDGCN as before.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
smallm_fp8_blockscale, smallm_fp8, smallm_moe and smallm_router each
pasted the same hipcc --genco + hipModuleLoad sequence; it now lives in
jit/utils/hipcc.py as build_gfx950_module. Failure handling per caller is
unchanged (m31_router keeps its own cached build). Not validated on ROCm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…6 sibling

GroupedGemmKernelSm90Fp8 carried verbatim copies of the BF16 class's
host helpers (_setup_attributes, smem/TMA layout builders, the
can_implement checks); it now subclasses GroupedGemmKernelSm90Bf16 like
the SM100 FP8 class already does, overriding only __init__, the dtype
check, __call__ and the kernel body.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
It reads other ranks' shards through peer (NVLink) mappings, so under the
operator-group rules it belongs to communication; it keeps importing the
shared row dequantizer from embeddings.engram_gather.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
attention/msa_v1/compact_dequant.py gathers and dequantizes KV-cache
pages; ops that read or write KV-cache contents belong to kvcache. Moved
as-is to kvcache/msa_v1_compact_dequant.py.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The dispatch docstring called the torch implementation a reference for CPU
unit tests. DeepseekV4AttnBackend also serves XPU, whose tensors are not
is_cuda, so the torch path is production code there.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…own groups

diffusion.mxfp8_quantize_swizzled and diffusion.silu_mul_mxfp8 already
lived in quantization / activation; their op ids now say so
(quantization.mxfp8_quantize_swizzled, activation.silu_mul_mxfp8), and
the diffusion facade no longer forwards them or their can_use_*
predicates. Diffusion callers import the defining modules.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…aller

qwen_moe_topk_output is Qwen2MoeSparseMoeBlock dispatch (layer
attributes, MoE runner backend, EP and expert-distribution recording),
not a kernel; it now lives in srt/models/qwen2_moe.py, and
kernels/ops/moe/smallm_router_gfx950 keeps only the kernel and its
support check. Not validated on ROCm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… quantization

per_token_quant_fp8_gfx1250 is a pure format conversion; it moves from
gemm/fp8_gfx1250.py to quantization/per_token_quant_fp8_gfx1250.py. The
gfx1250 GEMM tile stays in gemm. Not validated on ROCm (no gfx1250 runner).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- core/test_gated_launch.py: drop nightly=True. No CUDA job passes
  --nightly, so the file was filtered out of nightly-test-1-gpu-large.
  check_registered_tests.py now rejects nightly=True on CUDA/CPU.
- layers/mamba/test_mamba_ssm.py: also register on base-b 4-gpu-b200. The
  stochastic-rounding cases need SM100 and skipped on the 5090 lane.
- attention/unittests/dense/test_extend_init_contract.py: also register on
  base-b 1-gpu-large. The FA3 cases skip on SM100+, i.e. on the 5090.
- hicache/test_hicache_storage_3fs_backend.py: the AMD registration named
  the CUDA suite base-b-test-2-gpu-large, which no AMD job runs; mirror
  the CUDA extra-a tier with extra-a-test-2-gpu-large-amd.
- amd/test_minimax_hd128_asm_prefill.py: gfx950-only, move from the MI300
  stage-b suite to stage-b-test-1-gpu-small-amd-mi35x.
- _pr-test-stage.yml: hand HF_MUSE_SPARK_TOKEN to every 8-gpu-b300 run so
  the nightly Muse Spark test can download its gated checkpoint.
- Inkling-NVFP4 nightly and Mistral-Large-3 weekly: drop the 8-gpu-h200
  registrations, every test in both files skips off Blackwell.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The regex only matched single-line `register_X_ci(est_time=N, ...)` calls
with an integer value, so 150 of 2459 CUDA/CPU registrations (black-wrapped
calls, positional `register_cpu_ci(5, "suite")`, float values, a comment
between stage and runner_config) never had est_time refreshed. Read the
call the same way ci_register.py does instead.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- nightly-test-intel.yml: drop the 8-GPU XPU job; no file registers
  nightly-xpu-kernel-main-8-gpu.
- pr-test-jit-kernel.yml: drop the unit-test-8-gpu-h200 row; no file
  registers base-b-kernel-unit-test-8-gpu-h200, so every kernel PR booked
  an 8xH200 runner to log "No tests found".
- run_suite.py: drop those two suites plus base-c-test-8-gpu-b200,
  nightly-amd-vlm and nightly-acc-16-npu-a3-cann910 (no registrations,
  no job).
- compute_partitions.py: drop the python/sglang/jit_kernel globs; the
  directory no longer exists.
- stage_models_overrides.json: drop four suite_labels keys that are not
  suite names.
- rerun_test_groups.json: drop the deleted test_unified_radix_lock_ref.py.
- ci_register.py: list register_musa_ci once.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Leftovers of the same pattern sgl-project#47603 removed:
- test_generation_models.py: drop ALL_MODELS and MAMBA_MODEL_PATHS (no
  reader since test_all_models went away) and the skipIf that made a
  local run do nothing; point the new-model docs at CI_MODELS.
- test_transformers_models.py: drop test_others, which returned early in
  CI and needed tp=2 on a 1-GPU lane.
- test_quark_mxfp4.py: the four local-only classes and the TP=2 class
  can't run on the 1-GPU MI355X nightly runner; move them to
  test/manual/quantization/test_quark_mxfp4_multi_gpu.py. Drop the
  duplicate `import os`.
- test_qwen3_next_models_mtp.py: move the two @unittest.skip manual-only
  classes into the existing test/manual archive for this file.
- test_inkling_unified.py: move the INKLING_UNIFIED_PARITY-gated parity
  class (nothing sets the env) to test/manual/models/.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- utils/test_bench_typebaseddispatcher.py: a timeit script on an MI300
  lane that benchmarked TypeBasedDispatcher against a copy of the old list
  dispatcher, and whose edge-case checks only printed. Replace it with a
  CPU test of exact/inherited dispatch, the unknown-type error, the
  fallback and +=.
- core/test_srt_endpoint.py: drop the two unconditionally skipped
  stateful custom-logit-processor tests, their helper and the Rust-server
  skip stubs for them.
- hicache/test_hicache_l3_object_layout.py: drop a test that pops
  msgspec_to_builtins from the scheduler module, which imports it inside
  the handler; its asserts repeat the envelope test above it.
- debug_utils/comparator/testing_helpers.py: a helper module with a
  disabled registration only to satisfy the lint; move it to
  sglang.test.comparator_test_utils and drop the sys.path setup.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…p gates

Several kernel tests were registered only on lanes where their capability
gate always skips, so the production path they cover had no CI:

- SM120 FlashMLA fallback, SANA-Video depthwise GLU and Joy MLP epilogue
  (H100 lane) -> RTX 5090 lane; the SM120 fast-path cases move out of
  test_model_fast_paths.py into test_model_fast_paths_sm120.py.
- Blackwell paths of fp4_indexer_quant and the FlashInfer KDA input pack /
  prefill graph -> add B200.
- SM103 Qwen-Image output-bias absorption and FLUX.2 token-cat NVFP4
  -> B300 (the latter split into test_model_fast_paths_sm103.py).
- gfx950-only hc_mix FlyDSL, smallm router, MiniMax-M3 MXFP8, sparse-MLA
  topk_length and the gfx95 GDN launch config -> MI35x suites.
- ROCm-only aten_head_rmsnorm cases had no AMD registration -> own AMD file;
  moe_topk_sigmoid, reached only on ROCm, gets an AMD registration.
- The 2-GPU device_guard case -> 2-GPU lane.

Also drop registrations that test nothing: softmax on SM120 compares
torch.softmax with itself, and the DSV3.2 indexer fusion tests all skip
on HIP.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- test_kpool_topk_transform_fused: unittest.main() never collected the two
  module-level pytest functions (4 cases); run the file under pytest.
- test_q8kv8_sparse_prefill_backend: the request-window cases need
  Blackwell but ran only on H100; register B200 and fix their stale
  expectation that c1/c2/c4 history uses a c_max stride per request (it is
  packed per request; only c128 keeps the stride).
- test_rmsnorm: register B200, where the JIT rmsnorm is used, and make the
  strided-rows padding check compare against the untouched buffer instead
  of itself.
- test_moe_wna16_marlin: the GEMM test only compared against an AOT op that
  no longer exists, so all cases skipped; compare against the dequantized
  reference instead.
- test_awq_marlin_(moe_)repack: drop the jit_vs_aot cases for the same
  missing AOT ops.
- test_fp4_moe: __main__ ran one case of the grid; run the grid and drop
  four unused helpers.
- Delete test_amd_nccl_allreduce_determinism: it exercises torch's RCCL
  all-reduce, asserts nothing and ignores worker exit codes.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
bench_fp8_blockwise_gemm needs SM120 and bench_softmax needs SM100, but
both run only on the H100/AMD benchmark lanes and exit early there;
bench_bf16xfp32_router_gemm and test_subblock_sage_fp8_sm90 need packages
CI does not install (HPC-Ops, SpargeAttention).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@mintlify

mintlify Bot commented Nov 4, 2026 •

Copy link
Copy Markdown
Contributor

Preview approval required

This pull request comes from a fork, so Mintlify won't build a preview until someone with write access approves commit 41ac5ac.

Approve preview

@mickqian

mickqian commented Nov 4, 2026

Copy link
Copy Markdown
Collaborator

/tag-and-rerun-ci

@github-actions github-actions Bot added the run-ci CI: run the baseline test suite on this PR label Nov 4, 2026
@BBuf

BBuf commented Nov 4, 2026

Copy link
Copy Markdown
Collaborator Author

/tag-and-rerun-ci

BBuf and others added 8 commits November 4, 2026 22:36
…tail case on sm_100

- test_moe_topk_sigmoid.py fails every case on the MI300 jit-kernel lane; its
  AMD registration was never run on AMD hardware, so it goes back to CUDA only.
- TestSubBlockNumerics::test_ragged_tail_reproduces_dense returns NaN on the
  B200 lane (sm_100) while it passes on sm_103 and sm_120; the kernel is
  unchanged here, so skip it on sm_100 and leave the fix to the owner.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@github-actions

This comment has been minimized.

@BBuf
BBuf merged commit 944549c into sgl-project:main Nov 6, 2026
159 of 193 checks passed
@BBuf
BBuf deleted the bbuf/cleanup-round2 branch November 6, 2026 01:34
DarkSharpness added a commit to DarkSharpness/sglang that referenced this pull request Nov 6, 2026
…y the LiteTopK indexers

main pruned fma_f32x2 and mul_f32x2 from math.cuh as unused (sgl-project#48366), while this branch's fp8 / fp4 indexers call them; the merge dropped them and the indexers no longer compile.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
DarkSharpness added a commit to DarkSharpness/sglang that referenced this pull request Nov 6, 2026
…nalize kernel

> (auto generated by Claude)

sgl-project#48366 removed the packed fp32x2 FMA / multiply helpers from sgl_kernel/math.cuh as unused, but moe_finalize_fuse_shared.cuh uses math::fma_f32x2. Restore both helpers verbatim, as they were before sgl-project#48366.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
DarkSharpness added a commit to DarkSharpness/sglang that referenced this pull request Nov 6, 2026
…nalize kernel

> (auto generated by Claude)

sgl-project#48366 removed the packed fp32x2 FMA / multiply helpers from sgl_kernel/math.cuh as unused, but moe_finalize_fuse_shared.cuh uses math::fma_f32x2. Restore both helpers verbatim, as they were before sgl-project#48366.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

amd blackwell SM100/SM120 deepseek dependencies Pull requests that update a dependency file diffusion SGLang Diffusion documentation Improvements or additions to documentation hicache Hierarchical Caching for SGLang jit-kernel lora mooncake Multi-modal multi-modal language model npu quant LLM Quantization run-ci CI: run the baseline test suite on this PR sgl-kernel

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants