Repository navigation
[Refactor] kernels/diffusion/tests: Prune dead code and fix tests that never ran - #44690
Merged
BBuf merged 80 commits intoOct 19, 2026
Merged
Conversation
…ve paths SenseNovaU1GenerationStage only calls NEOChatModel.t2i_generate. Remove the upstream chat/batch_chat/generate/it2i/interleave entry points and the helpers only they reached (reference-image feature cache, thw indexes, image loading utils), plus unused fm-module classes, HF task heads, and compat helpers. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…se estimators estimate_layerwise_layer_uses, estimate_default_workload_timing and estimate_workload_phase_peaks (plus their private helpers) were added for a residency planner that never landed; only unit tests imported them. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The torch-native FusedGroupNorm3D compatibility wrapper was only selected by an undeclared env var that defaults to false and nothing sets; the default GroupNorm/SpatialNorm3D + norm_silu path is unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…a stray print The diffusion registry no longer passes per-spec descriptions or format signatures (KernelSpec defaults apply). flux3_image.pack_cfg has no caller, and causal_wanvideo logs its unsupported-QK-norm error like wanvideo does instead of printing it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The SANAWM_FORK_DUMP_DIR/SANAWM_RT_DUMP_DIR/SANAWM_BLOCK_PROBE/SANAWM_INJECT_DIR debug hooks were undeclared and no-op by default, yet forward_long read the environment on every call and the offline stage could torch.load tensors over its inputs. Drop the sana_wm_parity module, its hooks, and the on_noise_pred callback that only fed it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…cheduler module The MiniMax-H3 pipeline declares no scheduler (model_index carries scheduler=null) and its denoise loop applies the update in place, so the adapter's EntryClass is never resolved. Keep the two reference formulas the denoise-loop unit test compares against inside the test. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… ops
Keep the KernelSpec inventory but remove the metadata nothing reads:
FormatSignature, KernelSpec.format_signature/description/group/name,
BaseFusedOp.format_signature/descriptions/dispatch_forward,
PlatformInfo.is_hip, selector.clear_cache and
fused_op.clear_platform_caches. Name the CUDA-only capability set once
(spec.CUDA_ONLY) instead of spelling frozenset({CapabilityRequirement.CUDA})
at every registration, and let fused_op share selector's cached platform
probe instead of keeping a second copy.
Fixes:
- BaseFusedOp had no forward_flydsl stub although BACKEND_METHODS maps
FLYDSL to it, so SGLANG_FORCE_FUSED_OP_BACKEND=flydsl raised
AttributeError instead of falling back.
- Fp8ScaledMMOp declared its SM90 performance heuristic as the TORCH
backend's capability, so the registry claimed forward_native only runs
on SM90. The heuristic is now checked directly; dispatch is unchanged.
Remove the BaseFusedOp layers production never calls (srt and
multimodal_gen call the kernels directly): the four RMSNorm ops in
ops.layernorm and the gelu/gelu_tanh/relu2/gelu_quick ops in
ops.activation, plus their tests. Keep silu_and_mul. Drop the unused
facade wrappers gemm.{dsv3_fused_a_gemm,tiny_gemm_bf16,n128k512_gemm_bf16,
n32k5120_gemm_bf16} and quantization.{per_token_group_quant,
sgl_per_token_group_quant_fp8,sgl_per_token_group_quant_int8};
gemm.dsv3_fused_a_gemm and quantization.per_token_group_quant were also
shadowed by same-named submodules once those were imported. Drop the dead
softcap_out Triton kernel.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…dler table Each request-gated fusion family now exports its QualityGatedFusion object (with its static reject_reason) instead of four one-line mark/active/mount/unmount wrappers. The denoising stage's handler table holds (tier, fusion) pairs and a shared toggle_quality_fusions, which the Helios chunked stage now reuses instead of its own mount toggle. The Helios and LingBot Video gated-residual sites, identical apart from names, share one module and try_gated_residual. Also drops the dead common/platform.has_triton. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…rs and platforms - EncodingStage (only re-exported), SlidingTileAttnConfig, six unused StageValidators methods - BitsAndBytes4BitLinear / swap_linears_to_* (callers removed in sgl-project#27393), register_quantization_config, _nvfp4_uses_tma_scale_layout - the subprocess inspect path of the model registry (its command pointed at a nonexistent module) - realtime delta-gzip build/restore helpers (the server only sends raw or preview-encoded frames since sgl-project#27297) and the docs describing that encoding - bundled model-overlay lookup (model_overlays/ was removed in sgl-project#22182) - Platform.is_async_output_supported, the musa __main__ print block - DiffGenerator.generate_with_lora (it always passed sampling_params= into SamplingParams and could not run), flux3_image_quantization_name, BCG _env_float, align_hunyuan_image3_resolution, load_safetensors_state_dict, safetensors_keys, ServerArgs.is_component_attention_backend_automatic and derive_pool_control_* plus the write-only state behind them Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…onvert_hf_to_fp8 tool Platform.verify_quantization/supported_quantization came with the initial multimodal_gen import and never had a call site; the diffusion plugin docs do not list it as an active hook (the plugin.mdx row describes SRT DeviceMixin). convert_hf_to_fp8.py lost its last docs reference in sgl-project#21356/sgl-project#22672, which point users at the ModelOpt FP8 builder instead. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Delete code with no caller anywhere in the repo: - seg_la_kernel and its commented-out launch block (linear/seg_la.py) - fused_ape_pool_norm_rope (dsv4/fused_compress_triton.py) - fp8_index / fp8_index_kernel (dsa/tilelang_kernel.py) - kda_nvidia_prefill timing scaffolding, unused beta-absorb helpers and side stream, standalone-script import fallbacks and sys.path hacks - nvlink_comm all_reduce/reduce_scatter push/pull wrappers and their JIT entry points (only all_gather is used) - inkling_two_shot_all_reduce and its unfenced v1 kernel - gather_dsa_kv_scales, gather_dequant_requant_fp8_paged_ref - sp_collective.get_tuning, shmem_budget.max_nc_for_nv - unused cute_utils.cvt conversions - gdn_replayssm_circular_commit_kernel, reset_gdn_replayssm_spec_cursors - unused top_k/top_p_renorm_probs facade wrappers in ops.sampling Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- Drop the stale top_p_renorm_triton copy (production uses renorm_triton.top_p_renorm_probs_triton, which has the int64 row offset). - Drop the T=1 GDN decode entry (gated_delta_rule, its frozen wide-vec kernel and wrapper) from cutedsl_gdn_mtp_ring; production only calls gated_delta_rule_mtp. - Drop the plain FlyDSL Kimi-K3 KDA decode; production uses the _with_f_b specialization, whose tests are kept. - Drop dsa index_buf_accessor GetK/GetS and their Triton kernels; the manual test now checks GetKAndS against local torch gathers. - Drop adler32_checksum/adler32_regions_checksum and the whole-tensor kernel; the test now covers the same cases through adler32_strided_checksum. - Drop the unread _use_aiter flag in dsa index_buf_accessor. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- Remove the SGLANG_INKLING_FA4_USE_PIP debug branch that routed FA4 through the pip flash-attn package, along with the out= signature probe and the None fallbacks it required in flash_attention_v4_sm120. - Report vattn_asm_gfx950 build/load status through the module logger. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Reference implementations that only tests and benchmarks use now live next
to the other test helpers:
- dsv4 quant_k_cache -> test/kernels/deepseek_v4/quant_k_cache.py
- dsv4 V41 fp8 KV quantize/dequantize reference -> dsv41_kv_quant_reference
- LPLB solve_ipm / dispatch_probability torch references -> test/kernels/lplb.py
- DSA transform_index_page_table_{decode,prefill}_ref -> dsa_transform_index.py
- DCP LSE combine CPU reference -> test/kernels/dcp.py
- extend attention redundant_attention -> test/kernels/extend_attention.py
- DSpark CompactRowIndex/CompactVerifyIds parity pairs -> test/kernels/dspark.py
- torchrun multigpu_launch -> test/kernels/multigpu.py (mp.py keeps the
production register_comm_cleanup)
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
triton_sparse_mla_decode carried copies of the fused and split-K kernels from triton_sparse_mla. The fused copy equals _sparse_mla_fused_kernel with USE_TOPK_LENGTH=False and three pipeline stages. The split-K copy differs only in the LSE written for splits with no valid index (-2**30 instead of -FLT_MAX); the reduce kernel maps both to a zero weight, so the decode output is unchanged. Launch the shared kernels instead. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- test_swapab_attention: move CUDA registration from 1-gpu-large (H100) to
4-gpu-b200; the class only runs on SM100/SM103.
- bench_{custom_all_reduce,symm_mem_all_gather,tp_qknorm}: drop AMD
registrations; the 1-GPU MI300 runner skips every N>1 world size.
- test_per_token_quant_fp8: drop the nightly entry that re-ran the
per-commit grid.
- HPC-Ops tests: disable registrations; CI never installs HPC-Ops so every
case was skipped.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- Remove twelve envs.py entries nothing reads (vLLM install-time vars, CONFIG_ROOT, LOGGING_LEVEL/PREFIX, ATTENTION_BACKEND, SERVER_DEV_MODE, LD_LIBRARY_PATH, CUDA_VISIBLE_DEVICES) and the docs presenting them as working knobs. - Read SGLANG_DIFFUSION_TORCH_PROFILER_DIR through envs. - Declare SGLANG_ENABLE_FUSED_QKNORM_ROPE and the MiniMax-H3 VAE FUSED_NORM / VIT_FP32_NORM / STREAM_TEMPORAL_CAT switches in envs.py with their existing parsing and defaults, and switch every read site to them. - Drop the MiniMax-H3 VAE dev-only knobs from the initial bundle that default off and nothing sets (ViT FeedForward/RoPE torch.compile toggles with their *_FATAL/_BACKEND/_MODE/_FULLGRAPH/_DYNAMIC suffixes, TEMPORAL_CAT_DTYPE) together with the code paths they gated. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…paths Promote MOSS-Image's _fused_or_native into BitExactFusionGate.run: skip when disabled, verify an unseen signature outside compile tracing and CUDA graph capture, and fall back (disabling the gate) on a launch failure or mismatch. 24 hand-written sites that follow exactly this sequence now call it (MOSS-Image x5, ERNIE x3, GLM x2, Ideogram x2, Joy x2, LongCat x2, FLUX 3 x2, FLUX.2 strided QK, Ming, Anima, SANA-Video, Wan temb slices, LTX-2.5 decoder RoPE). The FLUX 3, GLM and LongCat LayerNorm+modulate wrappers share layernorm_modulate_bitexact. Every hand-written mismatch_msg restated the default warning, so the parameter is gone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… env knobs
PHASE_B_D_SPLITS / PHASE_B_DTILE_{NW,NS,ACC} and FUSED_GDN_PRECISION were
unprefixed, undocumented tuning knobs with no tests. Drop the env reads and keep
the defaults: per-arch Phase-B D-tile selection and bf16 dots for the camera
scan. The launch-config helpers that only served the precision override go too.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…stry has_registered_diffusion_model_path re-implemented the exact / short-name / HF-cache-path steps of _get_config_info; both now use one helper. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…en-Image VAE Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The decoder attention had two fused QK RMSNorm+RoPE paths for the same op. Measured on H200 against the eager norm + native NeoX RoPE chain: h3_vae_qk_rmsnorm_rope (attention.py, on by default for quality=exact) rounds about a quarter of the rotated channels differently and takes 41.1 us per call at [1, 4096, 16, 64] inside a CUDA graph, while the quality-gated fused_inplace_qknorm_rope path in minimax_h3_vae_cuda_opt matches the rotation bit for bit (only ~1e-4 of the elements differ, in the RMSNorm reduction) and takes 21.7 us. Keep the latter and drop the h3 kernel, its registry/export entries and tests; quality=exact now runs the eager chain. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…irectly Joy, Krea 2, FLUX 3, MiniMax-H3 and Cosmos3's packed-KV variant cannot route through layernorm.apply_qk_norm_rope without changing their reference numerics or layout; say so at each fork. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…rnels Each symbol below has no runtime caller; tests and benchmarks that only used it as an oracle now use a torch or production reference instead. - quantization: mxfp8 v-cache seq-blocked update + from_mxfp8; fp8_kernel ue8m0/hopper-moe-mn-major/per_group_transpose helpers; the per_token_group_quant_8bit dispatchers (moved into sglang.test.kernels.per_token_group_quant for the AOT test/bench); awq_gemm_triton; minimax per_token_quant_fp8_ue8m0 (the scatter test now compares against the production UE8M0 group quant); nvfp4 cvt_sf_*. - gemm: the test-only w8a8_block_fp8_matmul dispatcher and the test that never ran on CUDA (its guard returned on every sm>=90 GPU), splitk autotune_tactics, TGV tactic getters. - moe: rocm tkw1 helpers, the CUDA JIT moe_fused_gate (production uses the Triton router; its test now checks the torch reference at model shapes), Inkling gate v1, post_reorder_triton_kernel. - layernorm/elementwise: fused_inplace_qknorm_across_heads, the unused rmsnorm autotune path and silu_and_mul_triton copy, hc_mix_shape_ok; qwen4_gate reuses numerics.round_bf16_to_fp32; the mHC torch sinkhorn oracle moves to sglang.test.kernels.mhc. - aot: the AOT merge_state test imports the production Triton kernel, drop _get_cache_buf and a stale post-reorder benchmark, trim the per-token-group quant test grid from 3632 to 320 cases. Also fix the per_token_group_quant_8bit_v2 docstring (production still routes fp32 UE8M0 scales to it) and route gfx950 small-M kernel status messages through the logger instead of print. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…K norm; point LongCat tests at the raw LN-modulate launch layernorm_modulate_bitexact launches the raw kernel outside compile tracing, so the LongCat gate tests patch that symbol. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The realtime server only sends raw RGB or webp/jpeg preview frames, so the web UI's delta-gzip restore path (and the previous-frame state it needed) is unreachable. The transport option that omits realtime_output_format now says it gets the server's raw RGB default. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… fields pool_control_endpoint / pool_control_advertised_endpoint had no CLI flag, no docs and no reader; the derive_pool_control_* helpers that matched them are already gone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The dsv4 attention backends selected by create_dsv4_backend (DeepseekV4AttnBackend and its Trtllm/Mega/MultiStep subclasses on CUDA, DeepseekV4HipRadixBackend on ROCm) all mix in compressor_v2's CompressorBackendMixin, and the Ascend backend has its own. The older mixin in dsv4/compressor.py and its helpers (FusedCompressMetadata, is_overlap_compress, make_compressor_plan, create_paged_compressor_data) were only referenced under TYPE_CHECKING, so the code they reached is dead: - compress_old.py's plans, compress_forward and compress_fused_norm_rope_inplace - the HIP Triton compressor in fused_compress_triton.py - CompressorHip in srt dsv4/compress_hip.py (never imported) and the fused_softmax_pool_triton kernel only it used fused_norm_rope_inplace, still used by DeepSeek-V4 and the memory pool, moves next to the other norm/rope kernels in dsv4/elementwise.py. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…shim - dsv4 SetKAndS.torch and dequantize_k_cache_paged_ref were only used by tests; they now live with the dsv4 FP8 quantizer reference in sglang.test.kernels.deepseek_v4.k_cache_reference. - Delete lplb/cublasdx_solver.py, a back-compat re-export nothing imports. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
It only fed the commented-out fallback launch removed earlier. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
5 tasks done
This was referenced Oct 19, 2026
5 tasks done
This was referenced Oct 21, 2026
This was referenced Oct 23, 2026
rohithrr9
added a commit
to rohithrr9/sglang
that referenced
this pull request
Oct 26, 2026
…view comments Registry: gemm.dsv3_fused_a_gemm_fp8 passed format_signature and description to KernelSpec. FormatSignature was pruned from the registry spec upstream (sgl-project#44690), so `import sglang.kernels` raised NameError on every platform. Drop both arguments and add a registry test: the op resolves and is eligible on SM103 only. Kernel: remove the unused kBatch (Config::batch, the template parameter and the stage[kBatch] loops); the device code is unchanged (SASS identical in 15 of 16 functions, the 16th differs in register naming only). Comments no longer quote measurements or claim that the PDL wait covers only writes before the trigger, and the FP16 operand range is stated exactly (2^-17 <= |x| <= 65504). Tests: the dispatch tests need CUDA only, only the tests that run the kernel need SM103; drop the 4-gpu-b200 registration and restore the benchmark's est_time. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
rohithrr9
added a commit
to rohithrr9/sglang
that referenced
this pull request
Oct 26, 2026
…view comments Registry: gemm.dsv3_fused_a_gemm_fp8 passed format_signature and description to KernelSpec. FormatSignature was pruned from the registry spec upstream (sgl-project#44690), so `import sglang.kernels` raised NameError on every platform. Drop both arguments and add a registry test: the op resolves and is eligible on SM103 only. Kernel: remove the unused kBatch (Config::batch, the template parameter and the stage[kBatch] loops); the device code is unchanged (SASS identical in 15 of 16 functions, the 16th differs in register naming only). Comments no longer quote measurements or claim that the PDL wait covers only writes before the trigger, and the FP16 operand range is stated exactly (2^-17 <= |x| <= 65504). Tests: the dispatch tests need CUDA only, only the tests that run the kernel need SM103; drop the 4-gpu-b200 registration and restore the benchmark's est_time. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
Venkat2811
added a commit
to datacrunch-research/sglang
that referenced
this pull request
Oct 27, 2026
…gl-project#45544) Layer 4 of the Kimi-K3 nightly line on upstream main 5884e4f. Under expert parallelism each rank's trtllm-gen routed MoE computes only its own experts (112 of 896 at EP8), but FlashInfer picks its routing kernel from num_experts. K3's 896 experts select the Tier<896, 16> specialisation, which is above the dyn-block kernel's 512-expert limit, so decode steps of 5 to 16 tokens run routingIndicesClusterKernel, about 11 to 12.6 us per MoE layer on 8x B300 (once per MoE layer, so about 1 ms per decode step). The fused route+quant launch (route_quant_fused, decode batches up to 64 tokens) already writes the packed (id << 16 | bf16 weight) ids the routed MoE consumes. When the model stages the rank's expert range, its pack epilogue now writes the rank's experts as 0..count-1 and any other expert in top-k slot k as count + k (unique per row), and the runner passes num_experts = round_up(count + top_k, 4) with local_expert_offset 0 (128 at EP8: Tier<256, 16>, dyn-block kernel up to 16 tokens, a smaller cluster kernel above). The global ids (out_i, which TopKOutput carries on) are unchanged, and the same (token, local expert, weight) entries reach the expert GEMMs in the same order, so the MoE output is bitwise unchanged for the same GEMM tactic. Plumbing: - route_radix.cuh / route_quant_fused.cuh / moe_route_quant_fused.py: optional (local_expert_offset, num_local_experts) for the packed output only; 0 keeps the global ids. MAX_TOKENS (the fused launch's token cap) is public. - route_quant_handoff.py: stage(x, local_experts) takes the rank's range; rank_local_num_experts() decides whether it applies (a proper sub-range that shrinks the routing problem); the range is recorded with the packed tensor and read back once with take_local_experts(routing), keyed by that tensor. - kimi_k3.py: stages the range Mxfp4MoEMethod hands the runner, read per forward because bind_full_expert_weights can change it after construction. - flashinfer_trtllm.py: passes the rank-local num_experts and offset 0 for such packed ids, after asserting the range equals quant_info's. Autotune coverage: FlashInfer's MoE tuner key includes num_experts. When the decode-shaped autotune forward runs above the fused route's 64-token cap (decode CUDA graphs past 64 rows), it takes the unfused routing and would only tune the global key, leaving every rank-local decode call on FlashInfer's untuned default tactic. During tuning, when the runner sees unfused routing and rank-local ids apply, it also runs FlashInfer on up to 64 of that forward's routes, relabelled as the fused route would, so the decode key is tuned from real routes. Serving paths are unchanged. Measured on 8x B300, Kimi-K3 TP8 DCP8 EP8, one node: routing kernel per MoE layer 11.1 / 12.2 / 12.6 us -> 4.3 / 4.8 / 7.1 us at 8 / 16 / 24 decode rows; decode output throughput (bench_one_batch_server, 8192 in / 512 out) +6.1 / +4.2 / +1.0 % at batch 8 / 16 / 32 with the expert-GEMM tactics held at main's fresh picks, and +6.3 / +10.7 / -5.0 % with a fresh autotune per arm (the tuner picks different GEMM tactics for the new key). Tests: the kernel's rank-local ids against a torch oracle (EP 2/8/16, first, middle and last rank, 1 to 64 tokens) with everything else it writes unchanged; the handoff bookkeeping (keyed, consumed once, global ids when the range does not apply); and FlashInfer's routed MXFP4 MoE output bitwise equal between global and rank-local ids (EP 2/4/8/16, 1 to 64 tokens, SM100). Source: e251e3d on the kimi-k3-nightly-3c560089 line (fork PR #20), a mirror of upstream sgl-project#45544 (head fc02c94). One conflict, in moe_route_quant_fused.covered(): upstream sgl-project#44690 dropped the x.shape[0] == scores.shape[0] check (a mismatch is a caller bug the launcher reports); kept that, and took the PR's public MAX_TOKENS rename. Checked against sgl-project#44480 (merged 20 Oct, StandardDispatcher maps expert ids for runners that need local ids under EP): flashinfer_trtllm and flashinfer_mxfp4 stay in its keep-global-ids list, so nothing else remaps for this runner and the packed ids are made rank-local once, in the fused route epilogue. The range the model stages (moe_ep_rank * num_local_experts, num_local_experts) is the one the MXFP4 quant methods hand the runner at 5884e4f (mxfp4.py, mxfp4_flashinfer_trtllm_moe.py), and the runner asserts the two are equal. The runner function this changes is untouched by upstream since 3c56008. Tests: GPU only (test_rank_local_routed_moe.py, test_route_quant_prerequisites.py launch the JIT kernels; on a CPU host the new file does not import on the parent and skips or fails for want of CUDA here, as the existing cases in test_route_quant_prerequisites.py do on the base). On one B300 GPU on the 3c56008 line: test_rank_local_routed_moe.py 1/1 at this layer (2/2 on 23 Oct counted layer 6's padded-row case) and test_route_quant_prerequisites.py 5/5. They rerun on GPU with this branch's image. Upstream: sgl-project#45544 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Venkat2811
added a commit
to datacrunch-research/sglang
that referenced
this pull request
Oct 27, 2026
…gl-project#45544) Layer 4 of the Kimi-K3 nightly line on upstream main 5884e4f. Under expert parallelism each rank's trtllm-gen routed MoE computes only its own experts (112 of 896 at EP8), but FlashInfer picks its routing kernel from num_experts. K3's 896 experts select the Tier<896, 16> specialisation, which is above the dyn-block kernel's 512-expert limit, so decode steps of 5 to 16 tokens run routingIndicesClusterKernel, about 11 to 12.6 us per MoE layer on 8x B300 (once per MoE layer, so about 1 ms per decode step). The fused route+quant launch (route_quant_fused, decode batches up to 64 tokens) already writes the packed (id << 16 | bf16 weight) ids the routed MoE consumes. When the model stages the rank's expert range, its pack epilogue now writes the rank's experts as 0..count-1 and any other expert in top-k slot k as count + k (unique per row), and the runner passes num_experts = round_up(count + top_k, 4) with local_expert_offset 0 (128 at EP8: Tier<256, 16>, dyn-block kernel up to 16 tokens, a smaller cluster kernel above). The global ids (out_i, which TopKOutput carries on) are unchanged, and the same (token, local expert, weight) entries reach the expert GEMMs in the same order, so the MoE output is bitwise unchanged for the same GEMM tactic. Plumbing: - route_radix.cuh / route_quant_fused.cuh / moe_route_quant_fused.py: optional (local_expert_offset, num_local_experts) for the packed output only; 0 keeps the global ids. MAX_TOKENS (the fused launch's token cap) is public. - route_quant_handoff.py: stage(x, local_experts) takes the rank's range; rank_local_num_experts() decides whether it applies (a proper sub-range that shrinks the routing problem); the range is recorded with the packed tensor and read back once with take_local_experts(routing), keyed by that tensor. - kimi_k3.py: stages the range Mxfp4MoEMethod hands the runner, read per forward because bind_full_expert_weights can change it after construction. - flashinfer_trtllm.py: passes the rank-local num_experts and offset 0 for such packed ids, after asserting the range equals quant_info's. Autotune coverage: FlashInfer's MoE tuner key includes num_experts. When the decode-shaped autotune forward runs above the fused route's 64-token cap (decode CUDA graphs past 64 rows), it takes the unfused routing and would only tune the global key, leaving every rank-local decode call on FlashInfer's untuned default tactic. During tuning, when the runner sees unfused routing and rank-local ids apply, it also runs FlashInfer on up to 64 of that forward's routes, relabelled as the fused route would, so the decode key is tuned from real routes. Serving paths are unchanged. Measured on 8x B300, Kimi-K3 TP8 DCP8 EP8, one node: routing kernel per MoE layer 11.1 / 12.2 / 12.6 us -> 4.3 / 4.8 / 7.1 us at 8 / 16 / 24 decode rows; decode output throughput (bench_one_batch_server, 8192 in / 512 out) +6.1 / +4.2 / +1.0 % at batch 8 / 16 / 32 with the expert-GEMM tactics held at main's fresh picks, and +6.3 / +10.7 / -5.0 % with a fresh autotune per arm (the tuner picks different GEMM tactics for the new key). Tests: the kernel's rank-local ids against a torch oracle (EP 2/8/16, first, middle and last rank, 1 to 64 tokens) with everything else it writes unchanged; the handoff bookkeeping (keyed, consumed once, global ids when the range does not apply); and FlashInfer's routed MXFP4 MoE output bitwise equal between global and rank-local ids (EP 2/4/8/16, 1 to 64 tokens, SM100). Source: e251e3d on the kimi-k3-nightly-3c560089 line (fork PR #20), a mirror of upstream sgl-project#45544 (head fc02c94). One conflict, in moe_route_quant_fused.covered(): upstream sgl-project#44690 dropped the x.shape[0] == scores.shape[0] check (a mismatch is a caller bug the launcher reports); kept that, and took the PR's public MAX_TOKENS rename. Checked against sgl-project#44480 (merged 20 Oct, StandardDispatcher maps expert ids for runners that need local ids under EP): flashinfer_trtllm and flashinfer_mxfp4 stay in its keep-global-ids list, so nothing else remaps for this runner and the packed ids are made rank-local once, in the fused route epilogue. The range the model stages (moe_ep_rank * num_local_experts, num_local_experts) is the one the MXFP4 quant methods hand the runner at 5884e4f (mxfp4.py, mxfp4_flashinfer_trtllm_moe.py), and the runner asserts the two are equal. The runner function this changes is untouched by upstream since 3c56008. Tests: GPU only (test_rank_local_routed_moe.py, test_route_quant_prerequisites.py launch the JIT kernels; on a CPU host the new file does not import on the parent and skips or fails for want of CUDA here, as the existing cases in test_route_quant_prerequisites.py do on the base). On one B300 GPU on the 3c56008 line: test_rank_local_routed_moe.py 1/1 at this layer (2/2 on 23 Oct counted layer 6's padded-row case) and test_route_quant_prerequisites.py 5/5. They rerun on GPU with this branch's image. Upstream: sgl-project#45544 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
rohithrr9
added a commit
to rohithrr9/sglang
that referenced
this pull request
Oct 27, 2026
…view comments Registry: gemm.dsv3_fused_a_gemm_fp8 passed format_signature and description to KernelSpec. FormatSignature was pruned from the registry spec upstream (sgl-project#44690), so `import sglang.kernels` raised NameError on every platform. Drop both arguments and add a registry test: the op resolves and is eligible on SM103 only. Kernel: remove the unused kBatch (Config::batch, the template parameter and the stage[kBatch] loops); the device code is unchanged (SASS identical in 15 of 16 functions, the 16th differs in register naming only). Comments no longer quote measurements or claim that the PDL wait covers only writes before the trigger, and the FP16 operand range is stated exactly (2^-17 <= |x| <= 65504). Tests: the dispatch tests need CUDA only, only the tests that run the kernel need SM103; drop the 4-gpu-b200 registration and restore the benchmark's est_time. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
detain
pushed a commit
to detain/sglang
that referenced
this pull request
Oct 29, 2026
…ts only (sgl-project#44904) cherry-pick -x 48ea486. Landed second of the renorm pair: mainline makes sgl-project#44904 build ON TOP of sgl-project#44901 (its diff preimage contains fused_shared_ expert_weight/is_routed_scaling_factor_fused_in_topk), so the family follows the mainline order (see sgl-project#44901's note here). Wrong-output class: before this, the fused per-rank shared-slot paths asked the gate for top_k while num_fused_shared_experts markers also went through, renormalizing over one routed expert too many (routed sum 0.85-0.97 on CUDA). Conflict resolution (1): test/registered/kernels/ops/moe/test_moe_fused_gate.py carries pre-sgl-project#44690 docstring drift here (this fork never took sgl-project#44518's nemotron rework of that file's header); took the upstream post-image's new one-line docstring, kept this fork's file otherwise. The _reference_gate renorm-order fix, the renorm x apply_scale parametrization, and the production-impl comparison changes all landed. Fidelity: topk.py, hash_topk.py, state_capturer/routed_experts.py, test_fused_append_remap_per_rank_shared_slots.py and test_fused_shared_expert_scaling.py are BYTE-IDENTICAL to the upstream post-image; _router_triton_kernel (kernels/ops/moe/moe_fused_gate.py) verified byte-identical function body against upstream (the file's OTHER regions carry this fork's jit-loader drift and were not touched); fused_moe_triton_kernels.py hunks are docstring-only. Battery (CVD=9) vs sgl-project#44901 commit: test_fused_shared_expert_scaling 12P -> 13P/25sub; waterfill_eplb 4P -> 5P; scattered_a2a 3P -> 5P/2sub (the new routed-width capture case); hash_topk 4P held; append 6S -> 9S, padded 10S -> 11S (CUDA-gated adds); moe_fused_gate 46F -> 52F = the file's pre-existing no-CUDA class +6 new parametrized members (collects clean; needs hardware); kernels_namespace 76P (inventory guard, nothing added/removed); every other pin exact. ruff + format clean on all 11 files. Remap-on-CUDA numerics (torch.cat width change + padded-row remask) remain GPU-window debt, pinned at contract level by the fold-predicate probe commit that follows. (cherry picked from commit 48ea486) Co-authored-by: port-moe-contracts subagent (cherry picked from commit 0432a8fce0b4919b4a85fc6d3e724c6de30189be)
detain
added a commit
to detain/sglang
that referenced
this pull request
Oct 30, 2026
…traverse — entries 102,103,119-supersede,120,194,195,261-264 3 markers consumed (eagle_utils RE-PLANTED; spec_utils re-plant; eagle_info VERIFIED-ABSORBED). B4 tree count now 12. PER-FILE: - eagle_utils.py (entry 103): RE-PLANTED the tensors-after-early-return graft (num_needed_tokens == 0 returns before any torch.tensor -- the 17F were the fail-fast + alloc-path rebuilding CPU tensors every round; topk>1 max_alloc_len back to the list-builtin max) MERGED with upstream's in-range mirrorless block: it rides AFTER the tensors now (it reads cur_kv_lens_cpu) with a getattr-guarded seq_lens_cpu read -- identical semantics for the real ScheduleBatch (attr always present), stale-stub-proof. - spec_utils.py (entries 102,120,195; 261 deduped): * entry 261 sgl-project#45491: DEDUPED with evidence -- git log 12dc623..2a6c204 proves 573b571 IN RANGE; its verify-window semantics SURVIVE upstream's later refactors (sgl-project#45504 kernel move + sgl-project#46990 typed views collapsed the triton assign_extend_cache_locs pair) as fill_front_out_cache_loc_func(num_correct_drafts+1, batch.out_cache_loc, ...) carrying sgl-project#45491's OWN comment text; the fork-hunks window-form is thus superseded-in-place (origin/main has neither old name -- the replacement landed within the same range). * entry 195 sgl-project#44317: RE-PLANTED the traverse_tree host walk (tolist + numpy bitmask view; int words) + the needs_vocab_mask() skip (the backend-side methods ALREADY landed upstream in constrained/ -- consumer half was the dropped one) + the leftover-grammar_mask clear on the grammars-present-but-unconstrained path (the merged tree had the clear on the SUCCESS path only -- the stale-mask class this entry exists for). traverse_tree 3F -> green. * entry 102 (PLE on the ReplaySSM branches) + entry 120 (circular-ring fold before a plain extend): RE-PLANTED, RETARGETED to sgl-project#45504/sgl-project#46980 reality: kernel import path attention.fla.* -> attention.linear.* (upstream's kernels regroup), and the fold reads mamba_pool.get_speculative_mamba2_params_all_layers() (upstream sgl-project#46980 deleted the req-pool delegate; same route the verify-commit uses now). test_replayssm_extend_fold collection 1err -> 2P (the documented 'pair': its sibling test_replayssm_ring_fold.py was PRUNED by upstream's own sgl-project#44690 (file gone; a stale .pyc remains) and the 'gdn cutedsl probe' err belongs to kernels/ops/attention (6 CUDA-import collection errs there -- B1/B5 zone, not this dir): the spec-dir population is now 0err. test_replayssm_verify_ple_commit 1F -> green (stub re-anchored: router part_view seam + PLE pools off-pin restored -- the entry-102 pin the upstream file rewrite dropped, present in the 309-tip test verbatim). - eagle_info.py (entry 194): VERIFIED-ABSORBED -- the call sites pass num_token_blocks=None + kv_lens_sum, and runner_ids.pack_read_stream auto-sizes with the fork's OWN helper (spec_kv_index_token_blocks, now hosted in kernels/ops/attention/kv_indices) -- the 'AUTO-SIZED upgrade' B3-4 recorded; identical tiling math. - Entry-119 supersession chain CLOSED as briefed: sgl-project#45435's upstream form (tails padded -1 + mamba_lazy_undo_prealloc_at_boundary + untracked_decode_batch_idx gate) is the tree's reality (B4-2 audit); entry 119's production half stays dead. - test_ngram_mamba_verify_update.py (B2-marked file, B2 marker INTACT): restored only the two PLE off-pins + the entry-102 comment that the 309-tip fork version carried and upstream's rewrite dropped -- my re-plant re-creates the condition they exist for. ACTUALS (damage -> after, CVD=9, prod venv, per-file): test_eagle_decode_allocation 17F -> ** 18P ** (16-matrix + capacity + 1) test_spec_utils_traverse_tree 3F -> green (file pin incl.) test_replayssm_extend_fold 1err -> ** 2P ** test_replayssm_verify_ple_commit 1F -> green DIR unit/spec 363P/20F/21S/238sub/2err -> ** 388P/0F/21S/238sub/0err ** (two consecutive full-dir re-runs EXACT; DoD(b): zero non-env F). Lint: ruff format + F821/E9/F401 clean on the six files. No GPU; no push. Co-Authored-By: OpenCode <noreply@opencode.ai>
detain
added a commit
to detain/sglang
that referenced
this pull request
Oct 30, 2026
…re-plant) SESSION 2B COMPLETE Entry 312 records: the 5 cluster commits (9bd9007..5cde7a5) with per- cluster batteries; THE sgl-project#46977 RECONCILIATION DECISION TABLE (per-hunk survive-as-is / re-adapt / supersede, incl. the git-log proof that the spec-mamba four arrived in-range while the admission six NEVER merged -- correcting the briefing's dedupe list, and sgl-project#46643 fix(1)'s reconcile as already-merged base-PrefillBudget form = the reason prefill_budget.py got ZERO B4 edits); the entry-200 CAPACITY-PAIR restoration (our configurator's allocator.size branch survived -> upstream's consumer-side multiplications double-counted -> model_runner allocator= kwarg + tp_worker/scheduler subtractive forms re-planted; invariant-checker hunk OBSOLETE by upstream refactor); the io_struct NEEDS-ATTENTION VERDICT (upstream's own pointer_head guard vs stale fork stub -- test-side one-field re-anchor, matching sibling stubs); the -523sub PRUNE-vs-BREAKAGE audit (all three vanished managers files died to upstream's OWN sgl-project#45672 prune, two byte- identical to origin/main, the third's fork delta subsumed by sgl-project#46643 -- recorded, not re-homed); the B3 cross-note closures (MambaCowAllocError CATCHER landed; 216/217 admission interplay VERIFIED absorbed); the 2err population correction (replayssm_ring_fold pruned by sgl-project#44690, gdn cutedsl probe lives in the kernels zone, the spec dir's real err was merge damage, now 0err); the runner-zone finds the damage map missed (_init_model_hooks BOOT-INTEGRITY loss -- real decode-runner builds AttributeError, and component_manager's dropped residency half that parallel_executor already CALLS: both re-planted, 7F+crash-class closed); the true marker census (23 tree-carried, NOT the briefing's 20 -- B3 lesson applied; ALL consumed, B4=0, B1/B2/B5 diff-audited intact). GATES: trailing_endpoint ** 17P = 17P ** re-verified post-B4 (DoD-d, B3's gate unmoved). DoD(e) diff audit: ZERO mem_cache/ files changed, ZERO in qsa/moe-quant zones (moe_hook diff is marker-removal-only). NEW PINS: unit/managers 1009P/81F/9S/412sub/2err -> 1056P/2F/9S/474sub (2F = the two documented order-dependent mm_process_config ENV F's; the dir's 2err never reproduced -- population is the spec row's); unit/spec 363P/20F/21S/238sub/2err -> 388P/0F/21S/238sub/0err EXACT over 7 consecutive full-dir runs. OWED RESTART-WATCH REAFFIRMED: sgl-project#45472's sync-free planning (verified present, upstream form) + sgl-project#46643's unconditional per-round admission re-runs (re-planted) ride the entry-201 restart; entry-198's TTFT/starvation measure STILL OWED. Header: STACK -> qwennext-stack-312+merge5-b4-2026-10-29; +merge5-b4 wave tag; tip -> 5cde7a5; entries 1-312; INTEGRATION OWED annotated B4 LANDED (boot ban STAYS). wanted_stack: UNTOUCHED (no S-item changed by B4 findings -- recorded here). ENV debts recorded NOT-fixed: pcie_vocab_gather GPU suite (27F identical at the merge commit), openai 2.6.1<2.8.0 serving_responses import (upstream's own beam.py), imageio (whole diffusion import chain). NOTHING PUSHED; tree clean; zero GPU execution; zero prod boots all session. Co-Authored-By: OpenCode <noreply@opencode.ai>
This was referenced Nov 3, 2026
Merged
5 tasks done
This was referenced Nov 5, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Six months of fast kernel and diffusion work left a lot behind in
python/sglang/kernels,python/sglang/multimodal_genand the kernel/diffusion test trees: kernels nothing calls, registry metadata nothing reads, per-model copies of the same fused-op glue, env knobs that do nothing, and tests that pass without testing anything. This PR removes that, and fixes the real bugs the audit turned up along the way (several were silently shipping).408 files, +3,629 / -26,346. About 1.3K of the additions are test oracles and helpers moved out of production modules into
python/sglang/test/kernels/.Bugs fixed (please review these first)
vae.state_dict()beforeload_state_dict;AutoencoderDCbuilds its inner diffusers model lazily, so that call built it with random init and the later load returned early. All 375 VAE tensors were random; every SANA T2I image on main is noise.load_state_dictloads into an already-built inner model. Regression test intest_vae_loader.py(fails on main).query[:, t:] = q_imgwrite-back after an in-place RoPE; torch 2.14 rejects the self-copy, so every GLM-Image request fails in the first DiT block.test_glm_image_attention.py;test_ar_modelsre-enabled (fails on main, passes here in 175 s on 2xH200).h3_vae_qk_rmsnorm_ropepath ran inquality=exactbut was not bit-exact (~25% of rotated channels differ). A second, faster, quality-gated fused path for the same op already existed.STANDALONE_FILESmoved toserver/gpu_cases.pyin #24630 but the CUDA partition planner kept reading the runner, so no CUDA standalone file has run since June (12 2-GPU files + the 1-GPU ones).SGL_*names while docs/environ.pydefineSGLANG_*, and the rewrite is one-directional, so the documented names did nothing:SGLANG_USE_DEEPGEMM_BMM(5 models),SGLANG_DG_USE_NVRTC,SGLANG_CPU_QUANTIZATION,SGLANG_DSA_TOPK_ROW_SPLIT.envs; legacy names still work.BaseFusedOpKernelBackend.FLYDSLmapped to a missingforward_flydsl, soSGLANG_FORCE_FUSED_OP_BACKEND=flydslraisedAttributeError.test_fp8_kernelreturned early on every CUDA GPU with sm>=90;test_swapab_attentionwas registered on H100 but gated to sm_10x (production SM100 kernel had no CUDA coverage); sm_100 / sm_120 / 2-GPU / 4-GPU diffusion unit tests sat on a 1-GPU H100 lane;test_qwen_image21cache tests exercised a test-local re-implementation of the fix.FlowMatchingHeadcalledSimpleMLPAdaLNwith a signature the class doesn't have (crash forfm_head_layers > 2).Modifications
Kernels (
python/sglang/kernels, -13.8K)fused_ape_pool_norm_rope, tilelangfp8_index, KDA prefill timing scaffolding, nvlink push/pull all-reduce wrappers, inkling two-shot all-reduce,cutedsl_gdn_mtp_ringT=1 entry, non-f_bFlyDSL KDA decode, dsaGetK/GetS, adler32 whole-tensor entry points, ReplaySSM circular commit, cvt helpers, mxfp8 v-cache update, fp8 ue8m0/hopper/transpose helpers,w8a8_block_fp8_matmul,awq_gemm_triton, rocmtkw1,moe_fused_gate_jit, inkling gate v1,post_reorder_triton_kernel,fused_inplace_qknorm_across_heads, elementwisesilu_and_mul_triton/autotune path,rms_normalize_triton, staletop_p_renorm_tritoncopy, and others.compress_old.py,fused_compress_triton.py, srtdsv4/compress_hip.py, the oldCompressorBackendMixin): every CUDA/HIP backend usescompressor_v2;fused_norm_rope_inplace(live) moved todsv4/elementwise.py. AOT dsv4 q/k norm+RoPE ops (duplicates of the JIT ones) removed.FormatSignature,KernelSpec.format_signature/description/group/name,BaseFusedOp.format_signature/descriptions/dispatch_forward, unused selector/platform helpers).spec.CUDA_ONLYreplaces ~44 inline frozensets. UnusedBaseFusedOps (4 RMSNorm, Gelu/GeluTanh/ReLU2/QuickGELU) removed; facade wrappers nobody imports removed (this also fixes facade functions shadowed by same-named submodules).Fp8ScaledMMOpno longer declares an SM90 heuristic as a capability.mxfp8_quant. Forks that are not bit-exact stay, each with a one-line comment (grok RMSNorm, gfx1250 qk-RMSNorm, int8 group quant on ROCm, marlin SiTU, minimax Gemma RMSNorm).sglang.test.kernels(LPLB, dsa transform index, DCP LSE combine, dsv4 K-cache quant refs, extend attention, dspark window, multigpu launcher, mhc, per-token-group quant dispatch).Diffusion (
python/sglang/multimodal_gen, -9.4K)SANAWM_*env hooks incl.torch.loadinput injection),MINIMAX_H3_USE_FUSED_NORMpath, unresolved MiniMax-H3 euler scheduler, registry subprocess-inspect path (pointed at a nonexistent module), delta-gzip realtime encoding (server + web UI), bundled-overlay lookup,EncodingStage, bitsandbytes 4-bit swap,generate_with_lora(could never run),convert_hf_to_fp8.py, unused platform hooks, unreadServerArgsfields, etc.envs.pyentries nothing read removed (+ the docs rows presenting 5 of them as working knobs); raw reads ofSGLANG_ENABLE_FUSED_QKNORM_ROPE/MINIMAX_H3_VAE_*declared inenvs.pywith identical parsing; unprefixed GDN dev knobs dropped.BitExactFusionGate.run()replaces 24 hand-written verify-once/fallback blocks and themismatch_msgboilerplate; one LayerNorm+modulate helper for FLUX.3/GLM/LongCat;sites/exports oneQualityGatedFusionper family instead of four wrappers each (~50 facade entries, 28 imports gone). Sites that are hand-tuned for launch overhead, re-raise, verify inside the try, or differ in reference were left as they are (listed in the commit messages).Tests and CI
skipif(True), hard-coded dev paths), tautologies (default/literal echoes,inspect.getsourceshape tests, log-format echoes, torch-behavior tests), duplicate cases and duplicate helpers; trimmed grids with no extra coverage (fused_fp8_qkv_kv_cache1,344 -> 288 cases, AOTmoe_align4,368 -> 144, AOT per-token-group quant 3,632 -> 320). CI-tooling meta-tests are intentionally kept.quant-videogeninstalled for its tests; comm benches no longer register no-op AMD runs; HPC-Ops tests (always skipped: the package isn't in the CI venv) disabled explicitly.STANDALONE_FILE_EST_TIMESbefore merge.Accuracy Tests
End-to-end, same GPU, fixed seed, base vs this branch (fused paths confirmed active by a gate tracer on both sides; no gate disabled):
BitExactFusionGate.run, LN+modulate helper, QK-norm+RoPE, handler-table sites, gated residualKernel dedups were checked bit-for-bit on H200 and GB300 (scripts in the commit messages' descriptions).
Speed Tests and Profiling
No perf-path changes intended.
BitExactFusionGate.runwas not applied to the launch-overhead-tuned sites (the closures add ~160-180 ns per call on the host).Validation
test/registered/kernels/ops,test/registered/unitandmultimodal_gen/test/unit(1,624 files). Failures were re-run on main: identical failing test ids (rust extensions not built in the source tree, NPU-only, multi-GPU tests on one GPU, pre-existing diffusion kernel failures); 2 port collisions from parallel runs pass alone.test/registered/kernels/ops(334 files): failing ids identical to main.test_ar_modelsandtest_diffusion_bcg_tp2_zimage_turboat the final head); Kandinsky 6 4-GPU cases on 4xH200.pytest --collect-onlyover the test trees: no new collection errors; import walk oversglang.kernels,sglang.multimodal_gen,sglang.srt.{layers,models,...}: no new import failures.Checklist
🤖 Generated with Claude Code
CI States
Latest PR Test (Base): 🚫 Run #39111150246
Latest PR Test (Extra): ✅ Run #39111150188
Latest PR Test (AMD ROCm 10): ❌ Run #39111150233