From fa65f7c30cef89157ead525ad8530885a3cb2ac2 Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 09:04:57 +0800 Subject: [PATCH 1/6] [Kernel] Reclassify kernel tests by ops group under test/registered/kernels/ops/ (RFC #29630) Follow-up to RFC #29630: organize all kernel tests to mirror the sglang.kernels.ops. structure, and get the last test files out of the sglang package. - Move test/registered/jit/ (102) + test/registered/kernels/ (21) test & benchmark files into test/registered/kernels/ops// (tests) and ...//benchmark/ (benches), classified by the operator each exercises (attention/moe/layernorm/quantization/kvcache/gemm/diffusion/kv_canary/ speculative/communication/mamba/activation/embeddings/model). - Relocate shared test helpers out of the package's odd python/sglang/kernels/jit/tests/ folder into a proper support package python/sglang/kernels/testing/ (utils, kv_canary/, deepseek_v4/); rewrite all sglang.kernels.jit.tests imports -> sglang.kernels.testing. - Register the 6 previously-orphaned in-package tests (minimax_m3 x2, sconv metadata x2, moe preprocess/topk_sigmoid) that were never run by CI: CUDA ones -> base-b-kernel-unit; ROCm-only minimax_m3 ones -> jit-kernel-unit (amd). - Rename colliding test_sigmoid_gate_mul.py (broadcast variant) -> test_sigmoid_gate_mul_broadcast.py. CI discovery is recursive over test/registered/** and keyed on register_*_ci stage, so these moves don't change which suite each test runs in. python/sglang/ now has zero CI-registered tests. Validated: registry-validation + no-registered-tests-in-package hooks pass, testing package imports, namespace stays metadata-only, pre-commit clean. Co-Authored-By: Claude Opus 4.8 --- python/sglang/kernels/ops/communication/mp.py | 2 +- .../tests/deepseek_v4 => testing}/__init__.py | 0 .../deepseek_v4}/__init__.py | 0 .../tests => testing}/deepseek_v4/common.py | 0 .../kernels/testing/kv_canary/__init__.py | 0 .../kv_canary/_canary_helpers.py | 12 +++--- .../tests => testing}/kv_canary/_constants.py | 0 .../kv_canary/_differential.py | 12 +++--- .../tests => testing}/kv_canary/_fixtures.py | 2 +- .../kv_canary/_fuzz_driver.py | 2 +- .../kv_canary/_hand_oracle.py | 0 .../kv_canary/_invariants.py | 2 +- .../kernels/{jit/tests => testing}/utils.py | 0 .../activation}/benchmark/bench_activation.py | 0 .../ops/activation}/test_activation.py | 0 .../benchmark/bench_add_constant.py | 0 .../benchmark/bench_clamp_position.py | 0 .../attention}/benchmark/bench_concat_mla.py | 0 .../benchmark/bench_dsv4_fp4_indexer.py | 0 .../benchmark/bench_fused_qknorm_rope.py | 0 .../attention}/benchmark/bench_hadamard.py | 0 .../benchmark}/bench_minimax_decode_topk.py | 0 .../benchmark}/bench_minimax_qknorm_rope.py | 0 .../bench_mla_kv_pack_quantize_fp8.py | 0 .../benchmark/bench_online_c128_mtp.py | 0 .../ops/attention}/benchmark/bench_rope.py | 0 .../bench_sparse_mla_q8kv8_prefill_sm90.py | 0 .../ops/attention}/benchmark/bench_topk.py | 0 .../ops/attention}/test_add_constant.py | 0 .../ops/attention}/test_c128_v2.py | 4 +- .../ops/attention}/test_c4_v2.py | 4 +- .../ops/attention}/test_clamp_position.py | 0 .../ops/attention}/test_concat_mla.py | 0 .../test_cp_prefix_len_fa3_parity.py | 0 .../test_cute_dsl_fp8_paged_mqa_logits.py | 0 .../ops/attention}/test_cutedsl_gdn.py | 0 .../test_deepgemm_paged_mqa_logits.py | 0 ...epseek_v4_compress_state_runtime_shapes.py | 8 ++-- .../{ => ops/attention}/test_dsa_indexer.py | 0 .../{ => ops/attention}/test_dsa_metadata.py | 0 .../attention}/test_dsa_transform_index.py | 0 .../attention}/test_dsv32_indexer_fusion.py | 0 .../test_flash_attention_3_only_qv.py | 0 .../ops/attention}/test_flash_attention_4.py | 0 .../attention}/test_flash_mla_backends.py | 0 .../ops/attention}/test_fp4_indexer.py | 0 .../ops/attention}/test_fp8_wo_a.py | 0 .../attention}/test_fused_metadata_copy.py | 0 .../test_fused_store_index_cache.py | 0 .../test_fused_verify_triton_gdn.py | 0 .../ops/attention}/test_hadamard_jit.py | 0 .../test_inkling_attn_prologue_tau.py | 0 .../ops/attention}/test_inkling_row_scale.py | 0 .../attention}/test_minimax_decode_topk.py | 0 .../test_minimax_decode_topk_page_table.py | 0 .../test_minimax_m3_qk_norm_rope.py | 0 .../attention}/test_minimax_qknorm_rope.py | 0 .../attention}/test_mla_cp_fa3_parity.py | 0 .../test_mla_kv_pack_quantize_fp8.py | 0 .../ops/attention}/test_pos_enc.py | 0 .../ops/attention}/test_rope.py | 0 .../attention}/test_sm120_paged_mqa_logits.py | 0 .../test_sparse_mla_q8kv8_prefill_sm90.py | 0 .../ops/attention}/test_topk_v2.py | 0 .../benchmark/bench_custom_all_reduce.py | 0 .../benchmark/bench_symm_mem_all_gather.py | 0 .../benchmark/bench_tp_qknorm.py | 0 ...test_amd_deterministic_custom_allreduce.py | 0 .../test_amd_nccl_allreduce_determinism.py | 0 .../communication}/test_custom_all_reduce.py | 2 +- .../test_symm_mem_all_gather.py | 2 +- .../ops/communication}/test_tp_qknorm.py | 2 +- .../benchmark}/bench_causal_conv3d_cat_pad.py | 0 .../bench_diffusion_nvfp4_scaled_mm.py | 0 .../bench_fused_norm_scale_shift.py | 0 .../benchmark}/bench_group_norm_silu.py | 0 .../bench_ltx2_qknorm_split_rope.py | 0 .../diffusion/benchmark}/bench_norm_impls.py | 0 .../diffusion/benchmark}/bench_qknorm_rope.py | 0 .../benchmark}/bench_qwen_image_modulation.py | 0 .../benchmark}/bench_residual_gate_add.py | 0 .../diffusion/test_causal_conv3d_cat_pad.py | 0 .../test_diffusion_modelopt_fp8_scaled_mm.py | 0 .../test_diffusion_nvfp4_scaled_mm.py | 0 .../ops}/diffusion/test_flydsl_fused_norm.py | 0 .../diffusion/test_fused_norm_scale_shift.py | 0 .../ops}/diffusion/test_group_norm_silu.py | 0 .../ops}/diffusion/test_ltx2_ada_values.py | 0 .../diffusion/test_ltx2_qknorm_split_rope.py | 0 .../ops}/diffusion/test_qknorm_rope.py | 0 .../diffusion/test_qwen_image_modulation.py | 0 .../ops}/diffusion/test_residual_gate_add.py | 0 .../ops/diffusion}/test_timestep_embedding.py | 0 .../ops}/diffusion/test_varlen_pack_pad.py | 0 .../test_varlen_uspattn_equivalence.py | 0 .../bench_vocab_parallel_embedding.py | 0 .../test_vocab_parallel_embedding.py | 0 .../benchmark/bench_dsv3_fused_a_gemm.py | 0 .../gemm}/benchmark/bench_dsv3_router_gemm.py | 0 .../benchmark/bench_fp8_blockwise_gemm.py | 0 .../ops/gemm}/test_cutedsl_bf16_gemm.py | 0 .../gemm}/test_cutedsl_dsv3_fused_a_gemm.py | 0 .../ops/gemm}/test_dsv3_fused_a_gemm.py | 0 .../ops/gemm}/test_dsv3_router_gemm.py | 0 .../ops/gemm}/test_fp8_blockwise_gemm.py | 0 .../test_minimax_fused_qkv_index_gemm.py | 0 .../ops/kv_canary/benchmark}/bench_plan.py | 0 .../benchmark}/bench_scatter_req_token_ids.py | 0 .../ops/kv_canary/benchmark}/bench_verify.py | 0 .../ops/kv_canary/benchmark}/bench_write.py | 0 .../ops}/kv_canary/test_const_sync.py | 0 .../ops}/kv_canary/test_kernel_config.py | 22 +++++----- .../ops}/kv_canary/test_pipeline_e2e.py | 28 ++++++------- .../ops}/kv_canary/test_plan_fuzz.py | 8 ++-- .../ops}/kv_canary/test_plan_hand.py | 18 ++++----- .../kv_canary/test_scatter_req_token_ids.py | 0 .../ops}/kv_canary/test_utils.py | 0 .../ops}/kv_canary/test_verify_fuzz.py | 22 +++++----- .../ops}/kv_canary/test_verify_hand.py | 40 +++++++++---------- .../ops}/kv_canary/test_write_fuzz.py | 22 +++++----- .../ops}/kv_canary/test_write_hand.py | 34 ++++++++-------- .../benchmark/bench_fused_fp8_qkv_kv_cache.py | 0 .../ops/kvcache}/benchmark/bench_hicache.py | 0 .../ops/kvcache}/benchmark/bench_hisparse.py | 0 .../bench_minimax_store_kv_index.py | 0 .../benchmark/bench_set_mla_kv_buffer.py | 0 .../kvcache}/benchmark/bench_store_cache.py | 0 .../kvcache}/test_fused_fp8_qkv_kv_cache.py | 0 .../ops/kvcache}/test_hicache.py | 0 .../test_hicache_page_first_write_back.py | 0 .../ops/kvcache}/test_hisparse.py | 0 .../ops/kvcache}/test_kvcacheio_asymmetric.py | 0 .../kvcache}/test_minimax_store_kv_index.py | 0 .../ops/kvcache}/test_set_mla_kv_buffer.py | 0 .../ops/kvcache}/test_store_cache.py | 0 .../benchmark/bench_fused_eh_norm.py | 0 .../ops/layernorm}/benchmark/bench_norm.py | 0 .../ops/layernorm}/benchmark/bench_qknorm.py | 0 .../benchmark/bench_qknorm_across_heads.py | 0 .../ops/layernorm}/test_fused_add_rmsnorm.py | 0 .../ops/layernorm}/test_fused_eh_norm.py | 0 .../{ => ops/layernorm}/test_fused_op.py | 0 .../layernorm}/test_fused_op_gpu_parity.py | 0 .../test_fused_qk_gemma_rmsnorm_gate.py | 0 .../layernorm}/test_gemma4_fused_routing.py | 0 .../layernorm}/test_kernels_namespace.py | 0 .../{ => ops/layernorm}/test_mhc_kernels.py | 0 .../ops/layernorm}/test_minimax_m3_rmsnorm.py | 3 ++ .../ops/layernorm}/test_qknorm.py | 0 .../layernorm}/test_qknorm_across_heads.py | 0 .../ops/layernorm}/test_rmsnorm.py | 0 .../ops/layernorm}/test_rmsnorm_hf.py | 0 .../ops/mamba}/test_sconv_decode_metadata.py | 3 ++ .../ops/mamba}/test_sconv_extend_metadata.py | 3 ++ .../ops/mamba}/test_transfer_mamba.py | 0 .../ops/model}/test_inkling_rel_proj.py | 0 .../moe}/benchmark/bench_moe_fused_gate.py | 0 .../benchmark/bench_post_reorder_deepgemm.py | 0 .../ops/moe}/benchmark/bench_renorm.py | 0 .../kernels/{ => ops/moe}/test_fp4_moe.py | 0 .../{ => ops/moe}/test_fused_topk_deepseek.py | 0 .../ops/moe}/test_kpool_topk_transform.py | 0 .../moe}/test_marlin_packed_topk_unpack.py | 0 .../kernels/ops/moe}/test_minimax_m3_mxfp8.py | 5 ++- .../ops/moe}/test_minimax_quant_scatter.py | 0 .../ops/moe}/test_moe_align_block_size.py | 0 .../ops/moe}/test_moe_fused_gate.py | 0 .../moe}/test_moe_lora_align_block_size.py | 0 .../kernels/ops/moe}/test_moe_preprocess.py | 3 ++ .../kernels/ops/moe}/test_moe_topk_sigmoid.py | 3 ++ .../ops/moe}/test_moe_wna16_marlin.py | 0 .../ops/moe}/test_post_reorder_deepgemm.py | 0 .../{jit => kernels/ops/moe}/test_renorm.py | 0 .../ops/moe}/test_sigmoid_gate_mul.py | 0 .../moe/test_sigmoid_gate_mul_broadcast.py} | 0 .../benchmark/bench_awq_dequantize.py | 0 .../benchmark/bench_mxfp8_moe.py | 0 .../benchmark/bench_per_tensor_quant_fp8.py | 0 .../benchmark/bench_per_token_group_quant.py | 0 .../bench_per_token_group_quant_8bit_v2.py | 0 .../bench_per_token_group_quant_masked.py | 0 .../ops/quantization}/test_awq_dequantize.py | 0 .../test_awq_marlin_moe_repack.py | 0 .../quantization}/test_awq_marlin_repack.py | 0 .../ops/quantization}/test_gptq_marlin.py | 0 .../quantization}/test_gptq_marlin_repack.py | 0 .../ops/quantization}/test_mxfp8_moe.py | 0 .../test_per_tensor_quant_fp8.py | 0 .../test_per_token_group_quant.py | 0 .../test_per_token_group_quant_8bit_v2.py | 0 .../benchmark/bench_ngram_compute_decode.py | 0 .../bench_ngram_update_token_table.py | 0 .../bench_resolve_future_token_ids.py | 0 .../benchmark/bench_spec_topk1.py | 0 .../test_boundary_kv_fix_kernels.py | 0 .../speculative}/test_gather_spec_extras.py | 0 .../ops/speculative}/test_ngram_embedding.py | 0 .../test_resolve_future_token_ids.py | 0 .../{ => ops/speculative}/test_spec_topk1.py | 0 .../unit/distributed/test_vmm_utils.py | 2 +- 200 files changed, 144 insertions(+), 126 deletions(-) rename python/sglang/kernels/{jit/tests/deepseek_v4 => testing}/__init__.py (100%) rename python/sglang/kernels/{jit/tests/kv_canary => testing/deepseek_v4}/__init__.py (100%) rename python/sglang/kernels/{jit/tests => testing}/deepseek_v4/common.py (100%) create mode 100644 python/sglang/kernels/testing/kv_canary/__init__.py rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_canary_helpers.py (98%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_constants.py (100%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_differential.py (99%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_fixtures.py (98%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_fuzz_driver.py (95%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_hand_oracle.py (100%) rename python/sglang/kernels/{jit/tests => testing}/kv_canary/_invariants.py (99%) rename python/sglang/kernels/{jit/tests => testing}/utils.py (100%) rename test/registered/{jit => kernels/ops/activation}/benchmark/bench_activation.py (100%) rename test/registered/{jit => kernels/ops/activation}/test_activation.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_add_constant.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_clamp_position.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_concat_mla.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_dsv4_fp4_indexer.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_fused_qknorm_rope.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_hadamard.py (100%) rename test/registered/{jit/benchmark/minimax => kernels/ops/attention/benchmark}/bench_minimax_decode_topk.py (100%) rename test/registered/{jit/benchmark/minimax => kernels/ops/attention/benchmark}/bench_minimax_qknorm_rope.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_mla_kv_pack_quantize_fp8.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_online_c128_mtp.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_rope.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py (100%) rename test/registered/{jit => kernels/ops/attention}/benchmark/bench_topk.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_add_constant.py (100%) rename test/registered/{jit/deepseek_v4 => kernels/ops/attention}/test_c128_v2.py (99%) rename test/registered/{jit/deepseek_v4 => kernels/ops/attention}/test_c4_v2.py (99%) rename test/registered/{jit => kernels/ops/attention}/test_clamp_position.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_concat_mla.py (100%) rename test/registered/kernels/{ => ops/attention}/test_cp_prefix_len_fa3_parity.py (100%) rename test/registered/kernels/{ => ops/attention}/test_cute_dsl_fp8_paged_mqa_logits.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_cutedsl_gdn.py (100%) rename test/registered/kernels/{ => ops/attention}/test_deepgemm_paged_mqa_logits.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_deepseek_v4_compress_state_runtime_shapes.py (99%) rename test/registered/kernels/{ => ops/attention}/test_dsa_indexer.py (100%) rename test/registered/kernels/{ => ops/attention}/test_dsa_metadata.py (100%) rename test/registered/kernels/{ => ops/attention}/test_dsa_transform_index.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_dsv32_indexer_fusion.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_flash_attention_3_only_qv.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_flash_attention_4.py (100%) rename test/registered/kernels/{ => ops/attention}/test_flash_mla_backends.py (100%) rename test/registered/{jit/deepseek_v4 => kernels/ops/attention}/test_fp4_indexer.py (100%) rename test/registered/{jit/deepseek_v4 => kernels/ops/attention}/test_fp8_wo_a.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_fused_metadata_copy.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_fused_store_index_cache.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_fused_verify_triton_gdn.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_hadamard_jit.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_inkling_attn_prologue_tau.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_inkling_row_scale.py (100%) rename test/registered/{jit/minimax => kernels/ops/attention}/test_minimax_decode_topk.py (100%) rename test/registered/{jit/minimax => kernels/ops/attention}/test_minimax_decode_topk_page_table.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_minimax_m3_qk_norm_rope.py (100%) rename test/registered/{jit/minimax => kernels/ops/attention}/test_minimax_qknorm_rope.py (100%) rename test/registered/kernels/{ => ops/attention}/test_mla_cp_fa3_parity.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_mla_kv_pack_quantize_fp8.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_pos_enc.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_rope.py (100%) rename test/registered/kernels/{ => ops/attention}/test_sm120_paged_mqa_logits.py (100%) rename test/registered/{jit => kernels/ops/attention}/test_sparse_mla_q8kv8_prefill_sm90.py (100%) rename test/registered/{jit/deepseek_v4 => kernels/ops/attention}/test_topk_v2.py (100%) rename test/registered/{jit => kernels/ops/communication}/benchmark/bench_custom_all_reduce.py (100%) rename test/registered/{jit => kernels/ops/communication}/benchmark/bench_symm_mem_all_gather.py (100%) rename test/registered/{jit => kernels/ops/communication}/benchmark/bench_tp_qknorm.py (100%) rename test/registered/{jit => kernels/ops/communication}/test_amd_deterministic_custom_allreduce.py (100%) rename test/registered/{jit => kernels/ops/communication}/test_amd_nccl_allreduce_determinism.py (100%) rename test/registered/{jit => kernels/ops/communication}/test_custom_all_reduce.py (99%) rename test/registered/{jit => kernels/ops/communication}/test_symm_mem_all_gather.py (98%) rename test/registered/{jit => kernels/ops/communication}/test_tp_qknorm.py (99%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_causal_conv3d_cat_pad.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_diffusion_nvfp4_scaled_mm.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_fused_norm_scale_shift.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_group_norm_silu.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_ltx2_qknorm_split_rope.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_norm_impls.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_qknorm_rope.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_qwen_image_modulation.py (100%) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/bench_residual_gate_add.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_causal_conv3d_cat_pad.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_diffusion_nvfp4_scaled_mm.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_flydsl_fused_norm.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_fused_norm_scale_shift.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_group_norm_silu.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_ltx2_ada_values.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_ltx2_qknorm_split_rope.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_qknorm_rope.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_qwen_image_modulation.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_residual_gate_add.py (100%) rename test/registered/{jit => kernels/ops/diffusion}/test_timestep_embedding.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_varlen_pack_pad.py (100%) rename test/registered/{jit => kernels/ops}/diffusion/test_varlen_uspattn_equivalence.py (100%) rename test/registered/{jit => kernels/ops/embeddings}/benchmark/bench_vocab_parallel_embedding.py (100%) rename test/registered/kernels/{ => ops/embeddings}/test_vocab_parallel_embedding.py (100%) rename test/registered/{jit => kernels/ops/gemm}/benchmark/bench_dsv3_fused_a_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/benchmark/bench_dsv3_router_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/benchmark/bench_fp8_blockwise_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/test_cutedsl_bf16_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/test_cutedsl_dsv3_fused_a_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/test_dsv3_fused_a_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/test_dsv3_router_gemm.py (100%) rename test/registered/{jit => kernels/ops/gemm}/test_fp8_blockwise_gemm.py (100%) rename test/registered/{jit/minimax => kernels/ops/gemm}/test_minimax_fused_qkv_index_gemm.py (100%) rename test/registered/{jit/benchmark/kv_canary => kernels/ops/kv_canary/benchmark}/bench_plan.py (100%) rename test/registered/{jit/benchmark/kv_canary => kernels/ops/kv_canary/benchmark}/bench_scatter_req_token_ids.py (100%) rename test/registered/{jit/benchmark/kv_canary => kernels/ops/kv_canary/benchmark}/bench_verify.py (100%) rename test/registered/{jit/benchmark/kv_canary => kernels/ops/kv_canary/benchmark}/bench_write.py (100%) rename test/registered/{jit => kernels/ops}/kv_canary/test_const_sync.py (100%) rename test/registered/{jit => kernels/ops}/kv_canary/test_kernel_config.py (98%) rename test/registered/{jit => kernels/ops}/kv_canary/test_pipeline_e2e.py (99%) rename test/registered/{jit => kernels/ops}/kv_canary/test_plan_fuzz.py (96%) rename test/registered/{jit => kernels/ops}/kv_canary/test_plan_hand.py (99%) rename test/registered/{jit => kernels/ops}/kv_canary/test_scatter_req_token_ids.py (100%) rename test/registered/{jit => kernels/ops}/kv_canary/test_utils.py (100%) rename test/registered/{jit => kernels/ops}/kv_canary/test_verify_fuzz.py (95%) rename test/registered/{jit => kernels/ops}/kv_canary/test_verify_hand.py (99%) rename test/registered/{jit => kernels/ops}/kv_canary/test_write_fuzz.py (96%) rename test/registered/{jit => kernels/ops}/kv_canary/test_write_hand.py (99%) rename test/registered/{jit => kernels/ops/kvcache}/benchmark/bench_fused_fp8_qkv_kv_cache.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/benchmark/bench_hicache.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/benchmark/bench_hisparse.py (100%) rename test/registered/{jit/benchmark/minimax => kernels/ops/kvcache/benchmark}/bench_minimax_store_kv_index.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/benchmark/bench_set_mla_kv_buffer.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/benchmark/bench_store_cache.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_fused_fp8_qkv_kv_cache.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_hicache.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_hicache_page_first_write_back.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_hisparse.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_kvcacheio_asymmetric.py (100%) rename test/registered/{jit/minimax => kernels/ops/kvcache}/test_minimax_store_kv_index.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_set_mla_kv_buffer.py (100%) rename test/registered/{jit => kernels/ops/kvcache}/test_store_cache.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/benchmark/bench_fused_eh_norm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/benchmark/bench_norm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/benchmark/bench_qknorm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/benchmark/bench_qknorm_across_heads.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_fused_add_rmsnorm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_fused_eh_norm.py (100%) rename test/registered/kernels/{ => ops/layernorm}/test_fused_op.py (100%) rename test/registered/kernels/{ => ops/layernorm}/test_fused_op_gpu_parity.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_fused_qk_gemma_rmsnorm_gate.py (100%) rename test/registered/kernels/{ => ops/layernorm}/test_gemma4_fused_routing.py (100%) rename test/registered/kernels/{ => ops/layernorm}/test_kernels_namespace.py (100%) rename test/registered/kernels/{ => ops/layernorm}/test_mhc_kernels.py (100%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/layernorm}/test_minimax_m3_rmsnorm.py (95%) rename test/registered/{jit => kernels/ops/layernorm}/test_qknorm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_qknorm_across_heads.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_rmsnorm.py (100%) rename test/registered/{jit => kernels/ops/layernorm}/test_rmsnorm_hf.py (100%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/mamba}/test_sconv_decode_metadata.py (94%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/mamba}/test_sconv_extend_metadata.py (97%) rename test/registered/{jit => kernels/ops/mamba}/test_transfer_mamba.py (100%) rename test/registered/{jit => kernels/ops/model}/test_inkling_rel_proj.py (100%) rename test/registered/{jit => kernels/ops/moe}/benchmark/bench_moe_fused_gate.py (100%) rename test/registered/{jit => kernels/ops/moe}/benchmark/bench_post_reorder_deepgemm.py (100%) rename test/registered/{jit => kernels/ops/moe}/benchmark/bench_renorm.py (100%) rename test/registered/kernels/{ => ops/moe}/test_fp4_moe.py (100%) rename test/registered/kernels/{ => ops/moe}/test_fused_topk_deepseek.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_kpool_topk_transform.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_marlin_packed_topk_unpack.py (100%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/moe}/test_minimax_m3_mxfp8.py (97%) rename test/registered/{jit/minimax => kernels/ops/moe}/test_minimax_quant_scatter.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_moe_align_block_size.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_moe_fused_gate.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_moe_lora_align_block_size.py (100%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/moe}/test_moe_preprocess.py (95%) rename {python/sglang/kernels/jit/tests => test/registered/kernels/ops/moe}/test_moe_topk_sigmoid.py (99%) rename test/registered/{jit => kernels/ops/moe}/test_moe_wna16_marlin.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_post_reorder_deepgemm.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_renorm.py (100%) rename test/registered/{jit => kernels/ops/moe}/test_sigmoid_gate_mul.py (100%) rename test/registered/kernels/{test_sigmoid_gate_mul.py => ops/moe/test_sigmoid_gate_mul_broadcast.py} (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_awq_dequantize.py (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_mxfp8_moe.py (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_per_tensor_quant_fp8.py (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_per_token_group_quant.py (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_per_token_group_quant_8bit_v2.py (100%) rename test/registered/{jit => kernels/ops/quantization}/benchmark/bench_per_token_group_quant_masked.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_awq_dequantize.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_awq_marlin_moe_repack.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_awq_marlin_repack.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_gptq_marlin.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_gptq_marlin_repack.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_mxfp8_moe.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_per_tensor_quant_fp8.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_per_token_group_quant.py (100%) rename test/registered/{jit => kernels/ops/quantization}/test_per_token_group_quant_8bit_v2.py (100%) rename test/registered/{jit => kernels/ops/speculative}/benchmark/bench_ngram_compute_decode.py (100%) rename test/registered/{jit => kernels/ops/speculative}/benchmark/bench_ngram_update_token_table.py (100%) rename test/registered/{jit => kernels/ops/speculative}/benchmark/bench_resolve_future_token_ids.py (100%) rename test/registered/{jit => kernels/ops/speculative}/benchmark/bench_spec_topk1.py (100%) rename test/registered/kernels/{ => ops/speculative}/test_boundary_kv_fix_kernels.py (100%) rename test/registered/kernels/{ => ops/speculative}/test_gather_spec_extras.py (100%) rename test/registered/{jit => kernels/ops/speculative}/test_ngram_embedding.py (100%) rename test/registered/{jit => kernels/ops/speculative}/test_resolve_future_token_ids.py (100%) rename test/registered/kernels/{ => ops/speculative}/test_spec_topk1.py (100%) diff --git a/python/sglang/kernels/ops/communication/mp.py b/python/sglang/kernels/ops/communication/mp.py index 5a961ddb54b0..7b6361e63cd4 100644 --- a/python/sglang/kernels/ops/communication/mp.py +++ b/python/sglang/kernels/ops/communication/mp.py @@ -1,6 +1,6 @@ """Multi-process / multi-GPU launching utilities (torchrun-based). -Shared `multigpu_launch` helper that both `sglang.kernels.jit.tests.utils` and +Shared `multigpu_launch` helper that both `sglang.kernels.testing.utils` and `sglang.kernels.jit.benchmark.utils` build their domain-specific entry points on top of (`multigpu_pytest_main`, `multigpu_bench_main`). diff --git a/python/sglang/kernels/jit/tests/deepseek_v4/__init__.py b/python/sglang/kernels/testing/__init__.py similarity index 100% rename from python/sglang/kernels/jit/tests/deepseek_v4/__init__.py rename to python/sglang/kernels/testing/__init__.py diff --git a/python/sglang/kernels/jit/tests/kv_canary/__init__.py b/python/sglang/kernels/testing/deepseek_v4/__init__.py similarity index 100% rename from python/sglang/kernels/jit/tests/kv_canary/__init__.py rename to python/sglang/kernels/testing/deepseek_v4/__init__.py diff --git a/python/sglang/kernels/jit/tests/deepseek_v4/common.py b/python/sglang/kernels/testing/deepseek_v4/common.py similarity index 100% rename from python/sglang/kernels/jit/tests/deepseek_v4/common.py rename to python/sglang/kernels/testing/deepseek_v4/common.py diff --git a/python/sglang/kernels/testing/kv_canary/__init__.py b/python/sglang/kernels/testing/kv_canary/__init__.py new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/python/sglang/kernels/jit/tests/kv_canary/_canary_helpers.py b/python/sglang/kernels/testing/kv_canary/_canary_helpers.py similarity index 98% rename from python/sglang/kernels/jit/tests/kv_canary/_canary_helpers.py rename to python/sglang/kernels/testing/kv_canary/_canary_helpers.py index 9aa1e595cf1c..fbbae1c58e8a 100644 --- a/python/sglang/kernels/jit/tests/kv_canary/_canary_helpers.py +++ b/python/sglang/kernels/testing/kv_canary/_canary_helpers.py @@ -5,21 +5,21 @@ import torch -from sglang.kernels.jit.tests.kv_canary._constants import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 +from sglang.kernels.ops.kv_canary.verify import VerifyPlan +from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._constants import ( _I64_SIGN_BIT, _U64_MASK, DEFAULT_NUM_SLOTS, DEFAULT_RING_CAPACITY, DEFAULT_SLOT_STRIDE_BYTES, ) -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._fixtures import ( make_real_kv_source, make_real_kv_sources, ) -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 -from sglang.kernels.ops.kv_canary.verify import VerifyPlan -from sglang.kernels.ops.kv_canary.write import WritePlan __all__ = [ "FakeViolationLog", diff --git a/python/sglang/kernels/jit/tests/kv_canary/_constants.py b/python/sglang/kernels/testing/kv_canary/_constants.py similarity index 100% rename from python/sglang/kernels/jit/tests/kv_canary/_constants.py rename to python/sglang/kernels/testing/kv_canary/_constants.py diff --git a/python/sglang/kernels/jit/tests/kv_canary/_differential.py b/python/sglang/kernels/testing/kv_canary/_differential.py similarity index 99% rename from python/sglang/kernels/jit/tests/kv_canary/_differential.py rename to python/sglang/kernels/testing/kv_canary/_differential.py index fd2c26148801..97d08abaa3b2 100644 --- a/python/sglang/kernels/jit/tests/kv_canary/_differential.py +++ b/python/sglang/kernels/testing/kv_canary/_differential.py @@ -5,12 +5,6 @@ import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( - FakeViolationLog, - assert_canary_buf_equal, - assert_canary_state_equal, - make_log_pair, -) from sglang.kernels.ops.kv_canary import consts from sglang.kernels.ops.kv_canary.plan import launch_canary_plan_kernels from sglang.kernels.ops.kv_canary.plan_ref import ( @@ -30,6 +24,12 @@ from sglang.kernels.ops.kv_canary.write_ref import ( launch_canary_write_kernel_torch_reference, ) +from sglang.kernels.testing.kv_canary._canary_helpers import ( + FakeViolationLog, + assert_canary_buf_equal, + assert_canary_state_equal, + make_log_pair, +) _DEVICE = torch.device("cuda") diff --git a/python/sglang/kernels/jit/tests/kv_canary/_fixtures.py b/python/sglang/kernels/testing/kv_canary/_fixtures.py similarity index 98% rename from python/sglang/kernels/jit/tests/kv_canary/_fixtures.py rename to python/sglang/kernels/testing/kv_canary/_fixtures.py index 70fbb7587454..167a0bd297f0 100644 --- a/python/sglang/kernels/jit/tests/kv_canary/_fixtures.py +++ b/python/sglang/kernels/testing/kv_canary/_fixtures.py @@ -5,12 +5,12 @@ import torch -from sglang.kernels.jit.tests.kv_canary._constants import DEFAULT_NUM_SLOTS from sglang.kernels.ops.kv_canary.verify import ( RealKvSource, VerifyPlan, ) from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._constants import DEFAULT_NUM_SLOTS _DEVICE = torch.device("cuda") diff --git a/python/sglang/kernels/jit/tests/kv_canary/_fuzz_driver.py b/python/sglang/kernels/testing/kv_canary/_fuzz_driver.py similarity index 95% rename from python/sglang/kernels/jit/tests/kv_canary/_fuzz_driver.py rename to python/sglang/kernels/testing/kv_canary/_fuzz_driver.py index 49919e0683d4..7ba38f0c3352 100644 --- a/python/sglang/kernels/jit/tests/kv_canary/_fuzz_driver.py +++ b/python/sglang/kernels/testing/kv_canary/_fuzz_driver.py @@ -3,7 +3,7 @@ import random from typing import Any, Callable -from sglang.kernels.jit.tests.kv_canary._differential import ( +from sglang.kernels.testing.kv_canary._differential import ( ShrinkResult, shrink_inputs, ) diff --git a/python/sglang/kernels/jit/tests/kv_canary/_hand_oracle.py b/python/sglang/kernels/testing/kv_canary/_hand_oracle.py similarity index 100% rename from python/sglang/kernels/jit/tests/kv_canary/_hand_oracle.py rename to python/sglang/kernels/testing/kv_canary/_hand_oracle.py diff --git a/python/sglang/kernels/jit/tests/kv_canary/_invariants.py b/python/sglang/kernels/testing/kv_canary/_invariants.py similarity index 99% rename from python/sglang/kernels/jit/tests/kv_canary/_invariants.py rename to python/sglang/kernels/testing/kv_canary/_invariants.py index 192c79152cd9..79cdcf27f2a0 100644 --- a/python/sglang/kernels/jit/tests/kv_canary/_invariants.py +++ b/python/sglang/kernels/testing/kv_canary/_invariants.py @@ -11,10 +11,10 @@ import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import FakeViolationLog from sglang.kernels.ops.kv_canary import consts from sglang.kernels.ops.kv_canary.verify import CanaryLaunchTag, VerifyPlan from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._canary_helpers import FakeViolationLog class PlanInvariants: diff --git a/python/sglang/kernels/jit/tests/utils.py b/python/sglang/kernels/testing/utils.py similarity index 100% rename from python/sglang/kernels/jit/tests/utils.py rename to python/sglang/kernels/testing/utils.py diff --git a/test/registered/jit/benchmark/bench_activation.py b/test/registered/kernels/ops/activation/benchmark/bench_activation.py similarity index 100% rename from test/registered/jit/benchmark/bench_activation.py rename to test/registered/kernels/ops/activation/benchmark/bench_activation.py diff --git a/test/registered/jit/test_activation.py b/test/registered/kernels/ops/activation/test_activation.py similarity index 100% rename from test/registered/jit/test_activation.py rename to test/registered/kernels/ops/activation/test_activation.py diff --git a/test/registered/jit/benchmark/bench_add_constant.py b/test/registered/kernels/ops/attention/benchmark/bench_add_constant.py similarity index 100% rename from test/registered/jit/benchmark/bench_add_constant.py rename to test/registered/kernels/ops/attention/benchmark/bench_add_constant.py diff --git a/test/registered/jit/benchmark/bench_clamp_position.py b/test/registered/kernels/ops/attention/benchmark/bench_clamp_position.py similarity index 100% rename from test/registered/jit/benchmark/bench_clamp_position.py rename to test/registered/kernels/ops/attention/benchmark/bench_clamp_position.py diff --git a/test/registered/jit/benchmark/bench_concat_mla.py b/test/registered/kernels/ops/attention/benchmark/bench_concat_mla.py similarity index 100% rename from test/registered/jit/benchmark/bench_concat_mla.py rename to test/registered/kernels/ops/attention/benchmark/bench_concat_mla.py diff --git a/test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py b/test/registered/kernels/ops/attention/benchmark/bench_dsv4_fp4_indexer.py similarity index 100% rename from test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py rename to test/registered/kernels/ops/attention/benchmark/bench_dsv4_fp4_indexer.py diff --git a/test/registered/jit/benchmark/bench_fused_qknorm_rope.py b/test/registered/kernels/ops/attention/benchmark/bench_fused_qknorm_rope.py similarity index 100% rename from test/registered/jit/benchmark/bench_fused_qknorm_rope.py rename to test/registered/kernels/ops/attention/benchmark/bench_fused_qknorm_rope.py diff --git a/test/registered/jit/benchmark/bench_hadamard.py b/test/registered/kernels/ops/attention/benchmark/bench_hadamard.py similarity index 100% rename from test/registered/jit/benchmark/bench_hadamard.py rename to test/registered/kernels/ops/attention/benchmark/bench_hadamard.py diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py b/test/registered/kernels/ops/attention/benchmark/bench_minimax_decode_topk.py similarity index 100% rename from test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py rename to test/registered/kernels/ops/attention/benchmark/bench_minimax_decode_topk.py diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py b/test/registered/kernels/ops/attention/benchmark/bench_minimax_qknorm_rope.py similarity index 100% rename from test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py rename to test/registered/kernels/ops/attention/benchmark/bench_minimax_qknorm_rope.py diff --git a/test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py b/test/registered/kernels/ops/attention/benchmark/bench_mla_kv_pack_quantize_fp8.py similarity index 100% rename from test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py rename to test/registered/kernels/ops/attention/benchmark/bench_mla_kv_pack_quantize_fp8.py diff --git a/test/registered/jit/benchmark/bench_online_c128_mtp.py b/test/registered/kernels/ops/attention/benchmark/bench_online_c128_mtp.py similarity index 100% rename from test/registered/jit/benchmark/bench_online_c128_mtp.py rename to test/registered/kernels/ops/attention/benchmark/bench_online_c128_mtp.py diff --git a/test/registered/jit/benchmark/bench_rope.py b/test/registered/kernels/ops/attention/benchmark/bench_rope.py similarity index 100% rename from test/registered/jit/benchmark/bench_rope.py rename to test/registered/kernels/ops/attention/benchmark/bench_rope.py diff --git a/test/registered/jit/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py b/test/registered/kernels/ops/attention/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py similarity index 100% rename from test/registered/jit/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py rename to test/registered/kernels/ops/attention/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py diff --git a/test/registered/jit/benchmark/bench_topk.py b/test/registered/kernels/ops/attention/benchmark/bench_topk.py similarity index 100% rename from test/registered/jit/benchmark/bench_topk.py rename to test/registered/kernels/ops/attention/benchmark/bench_topk.py diff --git a/test/registered/jit/test_add_constant.py b/test/registered/kernels/ops/attention/test_add_constant.py similarity index 100% rename from test/registered/jit/test_add_constant.py rename to test/registered/kernels/ops/attention/test_add_constant.py diff --git a/test/registered/jit/deepseek_v4/test_c128_v2.py b/test/registered/kernels/ops/attention/test_c128_v2.py similarity index 99% rename from test/registered/jit/deepseek_v4/test_c128_v2.py rename to test/registered/kernels/ops/attention/test_c128_v2.py index 373fd2a247b7..5690402ab08e 100644 --- a/test/registered/jit/deepseek_v4/test_c128_v2.py +++ b/test/registered/kernels/ops/attention/test_c128_v2.py @@ -7,7 +7,8 @@ import torch import triton -from sglang.kernels.jit.tests.deepseek_v4.common import ( +from sglang.kernels.ops.attention.dsv4 import compress_forward +from sglang.kernels.testing.deepseek_v4.common import ( LegacyContext, PagedContext, make_legacy_context, @@ -15,7 +16,6 @@ make_state_pool, to_seq_extend, ) -from sglang.kernels.ops.attention.dsv4 import compress_forward from sglang.srt.utils import get_device from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci diff --git a/test/registered/jit/deepseek_v4/test_c4_v2.py b/test/registered/kernels/ops/attention/test_c4_v2.py similarity index 99% rename from test/registered/jit/deepseek_v4/test_c4_v2.py rename to test/registered/kernels/ops/attention/test_c4_v2.py index a9a7eb9d43e7..ae28242925e7 100644 --- a/test/registered/jit/deepseek_v4/test_c4_v2.py +++ b/test/registered/kernels/ops/attention/test_c4_v2.py @@ -7,7 +7,8 @@ import torch import triton -from sglang.kernels.jit.tests.deepseek_v4.common import ( +from sglang.kernels.ops.attention.dsv4 import compress_forward +from sglang.kernels.testing.deepseek_v4.common import ( LegacyContext, PagedContext, make_legacy_context, @@ -15,7 +16,6 @@ make_state_pool, to_seq_extend, ) -from sglang.kernels.ops.attention.dsv4 import compress_forward from sglang.srt.utils import get_device from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci diff --git a/test/registered/jit/test_clamp_position.py b/test/registered/kernels/ops/attention/test_clamp_position.py similarity index 100% rename from test/registered/jit/test_clamp_position.py rename to test/registered/kernels/ops/attention/test_clamp_position.py diff --git a/test/registered/jit/test_concat_mla.py b/test/registered/kernels/ops/attention/test_concat_mla.py similarity index 100% rename from test/registered/jit/test_concat_mla.py rename to test/registered/kernels/ops/attention/test_concat_mla.py diff --git a/test/registered/kernels/test_cp_prefix_len_fa3_parity.py b/test/registered/kernels/ops/attention/test_cp_prefix_len_fa3_parity.py similarity index 100% rename from test/registered/kernels/test_cp_prefix_len_fa3_parity.py rename to test/registered/kernels/ops/attention/test_cp_prefix_len_fa3_parity.py diff --git a/test/registered/kernels/test_cute_dsl_fp8_paged_mqa_logits.py b/test/registered/kernels/ops/attention/test_cute_dsl_fp8_paged_mqa_logits.py similarity index 100% rename from test/registered/kernels/test_cute_dsl_fp8_paged_mqa_logits.py rename to test/registered/kernels/ops/attention/test_cute_dsl_fp8_paged_mqa_logits.py diff --git a/test/registered/jit/test_cutedsl_gdn.py b/test/registered/kernels/ops/attention/test_cutedsl_gdn.py similarity index 100% rename from test/registered/jit/test_cutedsl_gdn.py rename to test/registered/kernels/ops/attention/test_cutedsl_gdn.py diff --git a/test/registered/kernels/test_deepgemm_paged_mqa_logits.py b/test/registered/kernels/ops/attention/test_deepgemm_paged_mqa_logits.py similarity index 100% rename from test/registered/kernels/test_deepgemm_paged_mqa_logits.py rename to test/registered/kernels/ops/attention/test_deepgemm_paged_mqa_logits.py diff --git a/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py similarity index 99% rename from test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py rename to test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py index 2658de98ad58..f2a279f0cd89 100644 --- a/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py +++ b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py @@ -89,15 +89,15 @@ import pytest import torch -from sglang.kernels.jit.tests.deepseek_v4.common import ( - make_legacy_context, - to_seq_extend, -) from sglang.kernels.ops.attention.dsv4 import ( CompressorDecodePlan, CompressorPrefillPlan, compress_forward, ) +from sglang.kernels.testing.deepseek_v4.common import ( + make_legacy_context, + to_seq_extend, +) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.utils import is_in_ci diff --git a/test/registered/kernels/test_dsa_indexer.py b/test/registered/kernels/ops/attention/test_dsa_indexer.py similarity index 100% rename from test/registered/kernels/test_dsa_indexer.py rename to test/registered/kernels/ops/attention/test_dsa_indexer.py diff --git a/test/registered/kernels/test_dsa_metadata.py b/test/registered/kernels/ops/attention/test_dsa_metadata.py similarity index 100% rename from test/registered/kernels/test_dsa_metadata.py rename to test/registered/kernels/ops/attention/test_dsa_metadata.py diff --git a/test/registered/kernels/test_dsa_transform_index.py b/test/registered/kernels/ops/attention/test_dsa_transform_index.py similarity index 100% rename from test/registered/kernels/test_dsa_transform_index.py rename to test/registered/kernels/ops/attention/test_dsa_transform_index.py diff --git a/test/registered/jit/test_dsv32_indexer_fusion.py b/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py similarity index 100% rename from test/registered/jit/test_dsv32_indexer_fusion.py rename to test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py diff --git a/test/registered/jit/test_flash_attention_3_only_qv.py b/test/registered/kernels/ops/attention/test_flash_attention_3_only_qv.py similarity index 100% rename from test/registered/jit/test_flash_attention_3_only_qv.py rename to test/registered/kernels/ops/attention/test_flash_attention_3_only_qv.py diff --git a/test/registered/jit/test_flash_attention_4.py b/test/registered/kernels/ops/attention/test_flash_attention_4.py similarity index 100% rename from test/registered/jit/test_flash_attention_4.py rename to test/registered/kernels/ops/attention/test_flash_attention_4.py diff --git a/test/registered/kernels/test_flash_mla_backends.py b/test/registered/kernels/ops/attention/test_flash_mla_backends.py similarity index 100% rename from test/registered/kernels/test_flash_mla_backends.py rename to test/registered/kernels/ops/attention/test_flash_mla_backends.py diff --git a/test/registered/jit/deepseek_v4/test_fp4_indexer.py b/test/registered/kernels/ops/attention/test_fp4_indexer.py similarity index 100% rename from test/registered/jit/deepseek_v4/test_fp4_indexer.py rename to test/registered/kernels/ops/attention/test_fp4_indexer.py diff --git a/test/registered/jit/deepseek_v4/test_fp8_wo_a.py b/test/registered/kernels/ops/attention/test_fp8_wo_a.py similarity index 100% rename from test/registered/jit/deepseek_v4/test_fp8_wo_a.py rename to test/registered/kernels/ops/attention/test_fp8_wo_a.py diff --git a/test/registered/jit/test_fused_metadata_copy.py b/test/registered/kernels/ops/attention/test_fused_metadata_copy.py similarity index 100% rename from test/registered/jit/test_fused_metadata_copy.py rename to test/registered/kernels/ops/attention/test_fused_metadata_copy.py diff --git a/test/registered/jit/test_fused_store_index_cache.py b/test/registered/kernels/ops/attention/test_fused_store_index_cache.py similarity index 100% rename from test/registered/jit/test_fused_store_index_cache.py rename to test/registered/kernels/ops/attention/test_fused_store_index_cache.py diff --git a/test/registered/jit/test_fused_verify_triton_gdn.py b/test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py similarity index 100% rename from test/registered/jit/test_fused_verify_triton_gdn.py rename to test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py diff --git a/test/registered/jit/test_hadamard_jit.py b/test/registered/kernels/ops/attention/test_hadamard_jit.py similarity index 100% rename from test/registered/jit/test_hadamard_jit.py rename to test/registered/kernels/ops/attention/test_hadamard_jit.py diff --git a/test/registered/jit/test_inkling_attn_prologue_tau.py b/test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py similarity index 100% rename from test/registered/jit/test_inkling_attn_prologue_tau.py rename to test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py diff --git a/test/registered/jit/test_inkling_row_scale.py b/test/registered/kernels/ops/attention/test_inkling_row_scale.py similarity index 100% rename from test/registered/jit/test_inkling_row_scale.py rename to test/registered/kernels/ops/attention/test_inkling_row_scale.py diff --git a/test/registered/jit/minimax/test_minimax_decode_topk.py b/test/registered/kernels/ops/attention/test_minimax_decode_topk.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_decode_topk.py rename to test/registered/kernels/ops/attention/test_minimax_decode_topk.py diff --git a/test/registered/jit/minimax/test_minimax_decode_topk_page_table.py b/test/registered/kernels/ops/attention/test_minimax_decode_topk_page_table.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_decode_topk_page_table.py rename to test/registered/kernels/ops/attention/test_minimax_decode_topk_page_table.py diff --git a/test/registered/jit/test_minimax_m3_qk_norm_rope.py b/test/registered/kernels/ops/attention/test_minimax_m3_qk_norm_rope.py similarity index 100% rename from test/registered/jit/test_minimax_m3_qk_norm_rope.py rename to test/registered/kernels/ops/attention/test_minimax_m3_qk_norm_rope.py diff --git a/test/registered/jit/minimax/test_minimax_qknorm_rope.py b/test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_qknorm_rope.py rename to test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py diff --git a/test/registered/kernels/test_mla_cp_fa3_parity.py b/test/registered/kernels/ops/attention/test_mla_cp_fa3_parity.py similarity index 100% rename from test/registered/kernels/test_mla_cp_fa3_parity.py rename to test/registered/kernels/ops/attention/test_mla_cp_fa3_parity.py diff --git a/test/registered/jit/test_mla_kv_pack_quantize_fp8.py b/test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py similarity index 100% rename from test/registered/jit/test_mla_kv_pack_quantize_fp8.py rename to test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py diff --git a/test/registered/jit/test_pos_enc.py b/test/registered/kernels/ops/attention/test_pos_enc.py similarity index 100% rename from test/registered/jit/test_pos_enc.py rename to test/registered/kernels/ops/attention/test_pos_enc.py diff --git a/test/registered/jit/test_rope.py b/test/registered/kernels/ops/attention/test_rope.py similarity index 100% rename from test/registered/jit/test_rope.py rename to test/registered/kernels/ops/attention/test_rope.py diff --git a/test/registered/kernels/test_sm120_paged_mqa_logits.py b/test/registered/kernels/ops/attention/test_sm120_paged_mqa_logits.py similarity index 100% rename from test/registered/kernels/test_sm120_paged_mqa_logits.py rename to test/registered/kernels/ops/attention/test_sm120_paged_mqa_logits.py diff --git a/test/registered/jit/test_sparse_mla_q8kv8_prefill_sm90.py b/test/registered/kernels/ops/attention/test_sparse_mla_q8kv8_prefill_sm90.py similarity index 100% rename from test/registered/jit/test_sparse_mla_q8kv8_prefill_sm90.py rename to test/registered/kernels/ops/attention/test_sparse_mla_q8kv8_prefill_sm90.py diff --git a/test/registered/jit/deepseek_v4/test_topk_v2.py b/test/registered/kernels/ops/attention/test_topk_v2.py similarity index 100% rename from test/registered/jit/deepseek_v4/test_topk_v2.py rename to test/registered/kernels/ops/attention/test_topk_v2.py diff --git a/test/registered/jit/benchmark/bench_custom_all_reduce.py b/test/registered/kernels/ops/communication/benchmark/bench_custom_all_reduce.py similarity index 100% rename from test/registered/jit/benchmark/bench_custom_all_reduce.py rename to test/registered/kernels/ops/communication/benchmark/bench_custom_all_reduce.py diff --git a/test/registered/jit/benchmark/bench_symm_mem_all_gather.py b/test/registered/kernels/ops/communication/benchmark/bench_symm_mem_all_gather.py similarity index 100% rename from test/registered/jit/benchmark/bench_symm_mem_all_gather.py rename to test/registered/kernels/ops/communication/benchmark/bench_symm_mem_all_gather.py diff --git a/test/registered/jit/benchmark/bench_tp_qknorm.py b/test/registered/kernels/ops/communication/benchmark/bench_tp_qknorm.py similarity index 100% rename from test/registered/jit/benchmark/bench_tp_qknorm.py rename to test/registered/kernels/ops/communication/benchmark/bench_tp_qknorm.py diff --git a/test/registered/jit/test_amd_deterministic_custom_allreduce.py b/test/registered/kernels/ops/communication/test_amd_deterministic_custom_allreduce.py similarity index 100% rename from test/registered/jit/test_amd_deterministic_custom_allreduce.py rename to test/registered/kernels/ops/communication/test_amd_deterministic_custom_allreduce.py diff --git a/test/registered/jit/test_amd_nccl_allreduce_determinism.py b/test/registered/kernels/ops/communication/test_amd_nccl_allreduce_determinism.py similarity index 100% rename from test/registered/jit/test_amd_nccl_allreduce_determinism.py rename to test/registered/kernels/ops/communication/test_amd_nccl_allreduce_determinism.py diff --git a/test/registered/jit/test_custom_all_reduce.py b/test/registered/kernels/ops/communication/test_custom_all_reduce.py similarity index 99% rename from test/registered/jit/test_custom_all_reduce.py rename to test/registered/kernels/ops/communication/test_custom_all_reduce.py index 38b2b0568626..d609f0ecfbdf 100644 --- a/test/registered/jit/test_custom_all_reduce.py +++ b/test/registered/kernels/ops/communication/test_custom_all_reduce.py @@ -28,13 +28,13 @@ import torch.distributed as dist import sglang.srt.distributed.parallel_state as ps -from sglang.kernels.jit.tests.utils import multigpu_pytest_main from sglang.kernels.jit.utils import cache_once, get_ci_test_range from sglang.kernels.ops.communication.all_reduce import ( AllReduceAlgo, get_all_reduce_module, ) from sglang.kernels.ops.communication.mp import register_comm_cleanup +from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.custom_all_reduce_v2 import ( CustomAllReduceV2, ) diff --git a/test/registered/jit/test_symm_mem_all_gather.py b/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py similarity index 98% rename from test/registered/jit/test_symm_mem_all_gather.py rename to test/registered/kernels/ops/communication/test_symm_mem_all_gather.py index 816964d9b48a..f38235ef1ee0 100644 --- a/test/registered/jit/test_symm_mem_all_gather.py +++ b/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py @@ -26,8 +26,8 @@ import torch.distributed as dist import sglang.srt.distributed.parallel_state as ps -from sglang.kernels.jit.tests.utils import multigpu_pytest_main from sglang.kernels.jit.utils import cache_once, get_ci_test_range +from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.triton_symm_mem_ag import ( all_gather_inner, create_state, diff --git a/test/registered/jit/test_tp_qknorm.py b/test/registered/kernels/ops/communication/test_tp_qknorm.py similarity index 99% rename from test/registered/jit/test_tp_qknorm.py rename to test/registered/kernels/ops/communication/test_tp_qknorm.py index 31f715eb0cc7..c7895c7a2f03 100644 --- a/test/registered/jit/test_tp_qknorm.py +++ b/test/registered/kernels/ops/communication/test_tp_qknorm.py @@ -14,7 +14,6 @@ import triton import sglang.srt.distributed.parallel_state as ps -from sglang.kernels.jit.tests.utils import multigpu_pytest_main from sglang.kernels.jit.utils import cache_once from sglang.kernels.ops.communication.all_reduce import ( fused_parallel_qknorm, @@ -22,6 +21,7 @@ get_fused_parallel_qknorm_module, ) from sglang.kernels.ops.communication.mp import register_comm_cleanup +from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.custom_all_reduce_v2 import ( CustomAllReduceV2, ) diff --git a/test/registered/jit/benchmark/diffusion/bench_causal_conv3d_cat_pad.py b/test/registered/kernels/ops/diffusion/benchmark/bench_causal_conv3d_cat_pad.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_causal_conv3d_cat_pad.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_causal_conv3d_cat_pad.py diff --git a/test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py b/test/registered/kernels/ops/diffusion/benchmark/bench_diffusion_nvfp4_scaled_mm.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_diffusion_nvfp4_scaled_mm.py diff --git a/test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py b/test/registered/kernels/ops/diffusion/benchmark/bench_fused_norm_scale_shift.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_fused_norm_scale_shift.py diff --git a/test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py b/test/registered/kernels/ops/diffusion/benchmark/bench_group_norm_silu.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_group_norm_silu.py diff --git a/test/registered/jit/benchmark/diffusion/bench_ltx2_qknorm_split_rope.py b/test/registered/kernels/ops/diffusion/benchmark/bench_ltx2_qknorm_split_rope.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_ltx2_qknorm_split_rope.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_ltx2_qknorm_split_rope.py diff --git a/test/registered/jit/benchmark/diffusion/bench_norm_impls.py b/test/registered/kernels/ops/diffusion/benchmark/bench_norm_impls.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_norm_impls.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_norm_impls.py diff --git a/test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py b/test/registered/kernels/ops/diffusion/benchmark/bench_qknorm_rope.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_qknorm_rope.py diff --git a/test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py b/test/registered/kernels/ops/diffusion/benchmark/bench_qwen_image_modulation.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_qwen_image_modulation.py diff --git a/test/registered/jit/benchmark/diffusion/bench_residual_gate_add.py b/test/registered/kernels/ops/diffusion/benchmark/bench_residual_gate_add.py similarity index 100% rename from test/registered/jit/benchmark/diffusion/bench_residual_gate_add.py rename to test/registered/kernels/ops/diffusion/benchmark/bench_residual_gate_add.py diff --git a/test/registered/jit/diffusion/test_causal_conv3d_cat_pad.py b/test/registered/kernels/ops/diffusion/test_causal_conv3d_cat_pad.py similarity index 100% rename from test/registered/jit/diffusion/test_causal_conv3d_cat_pad.py rename to test/registered/kernels/ops/diffusion/test_causal_conv3d_cat_pad.py diff --git a/test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py b/test/registered/kernels/ops/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py similarity index 100% rename from test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py rename to test/registered/kernels/ops/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py diff --git a/test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py b/test/registered/kernels/ops/diffusion/test_diffusion_nvfp4_scaled_mm.py similarity index 100% rename from test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py rename to test/registered/kernels/ops/diffusion/test_diffusion_nvfp4_scaled_mm.py diff --git a/test/registered/jit/diffusion/test_flydsl_fused_norm.py b/test/registered/kernels/ops/diffusion/test_flydsl_fused_norm.py similarity index 100% rename from test/registered/jit/diffusion/test_flydsl_fused_norm.py rename to test/registered/kernels/ops/diffusion/test_flydsl_fused_norm.py diff --git a/test/registered/jit/diffusion/test_fused_norm_scale_shift.py b/test/registered/kernels/ops/diffusion/test_fused_norm_scale_shift.py similarity index 100% rename from test/registered/jit/diffusion/test_fused_norm_scale_shift.py rename to test/registered/kernels/ops/diffusion/test_fused_norm_scale_shift.py diff --git a/test/registered/jit/diffusion/test_group_norm_silu.py b/test/registered/kernels/ops/diffusion/test_group_norm_silu.py similarity index 100% rename from test/registered/jit/diffusion/test_group_norm_silu.py rename to test/registered/kernels/ops/diffusion/test_group_norm_silu.py diff --git a/test/registered/jit/diffusion/test_ltx2_ada_values.py b/test/registered/kernels/ops/diffusion/test_ltx2_ada_values.py similarity index 100% rename from test/registered/jit/diffusion/test_ltx2_ada_values.py rename to test/registered/kernels/ops/diffusion/test_ltx2_ada_values.py diff --git a/test/registered/jit/diffusion/test_ltx2_qknorm_split_rope.py b/test/registered/kernels/ops/diffusion/test_ltx2_qknorm_split_rope.py similarity index 100% rename from test/registered/jit/diffusion/test_ltx2_qknorm_split_rope.py rename to test/registered/kernels/ops/diffusion/test_ltx2_qknorm_split_rope.py diff --git a/test/registered/jit/diffusion/test_qknorm_rope.py b/test/registered/kernels/ops/diffusion/test_qknorm_rope.py similarity index 100% rename from test/registered/jit/diffusion/test_qknorm_rope.py rename to test/registered/kernels/ops/diffusion/test_qknorm_rope.py diff --git a/test/registered/jit/diffusion/test_qwen_image_modulation.py b/test/registered/kernels/ops/diffusion/test_qwen_image_modulation.py similarity index 100% rename from test/registered/jit/diffusion/test_qwen_image_modulation.py rename to test/registered/kernels/ops/diffusion/test_qwen_image_modulation.py diff --git a/test/registered/jit/diffusion/test_residual_gate_add.py b/test/registered/kernels/ops/diffusion/test_residual_gate_add.py similarity index 100% rename from test/registered/jit/diffusion/test_residual_gate_add.py rename to test/registered/kernels/ops/diffusion/test_residual_gate_add.py diff --git a/test/registered/jit/test_timestep_embedding.py b/test/registered/kernels/ops/diffusion/test_timestep_embedding.py similarity index 100% rename from test/registered/jit/test_timestep_embedding.py rename to test/registered/kernels/ops/diffusion/test_timestep_embedding.py diff --git a/test/registered/jit/diffusion/test_varlen_pack_pad.py b/test/registered/kernels/ops/diffusion/test_varlen_pack_pad.py similarity index 100% rename from test/registered/jit/diffusion/test_varlen_pack_pad.py rename to test/registered/kernels/ops/diffusion/test_varlen_pack_pad.py diff --git a/test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py b/test/registered/kernels/ops/diffusion/test_varlen_uspattn_equivalence.py similarity index 100% rename from test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py rename to test/registered/kernels/ops/diffusion/test_varlen_uspattn_equivalence.py diff --git a/test/registered/jit/benchmark/bench_vocab_parallel_embedding.py b/test/registered/kernels/ops/embeddings/benchmark/bench_vocab_parallel_embedding.py similarity index 100% rename from test/registered/jit/benchmark/bench_vocab_parallel_embedding.py rename to test/registered/kernels/ops/embeddings/benchmark/bench_vocab_parallel_embedding.py diff --git a/test/registered/kernels/test_vocab_parallel_embedding.py b/test/registered/kernels/ops/embeddings/test_vocab_parallel_embedding.py similarity index 100% rename from test/registered/kernels/test_vocab_parallel_embedding.py rename to test/registered/kernels/ops/embeddings/test_vocab_parallel_embedding.py diff --git a/test/registered/jit/benchmark/bench_dsv3_fused_a_gemm.py b/test/registered/kernels/ops/gemm/benchmark/bench_dsv3_fused_a_gemm.py similarity index 100% rename from test/registered/jit/benchmark/bench_dsv3_fused_a_gemm.py rename to test/registered/kernels/ops/gemm/benchmark/bench_dsv3_fused_a_gemm.py diff --git a/test/registered/jit/benchmark/bench_dsv3_router_gemm.py b/test/registered/kernels/ops/gemm/benchmark/bench_dsv3_router_gemm.py similarity index 100% rename from test/registered/jit/benchmark/bench_dsv3_router_gemm.py rename to test/registered/kernels/ops/gemm/benchmark/bench_dsv3_router_gemm.py diff --git a/test/registered/jit/benchmark/bench_fp8_blockwise_gemm.py b/test/registered/kernels/ops/gemm/benchmark/bench_fp8_blockwise_gemm.py similarity index 100% rename from test/registered/jit/benchmark/bench_fp8_blockwise_gemm.py rename to test/registered/kernels/ops/gemm/benchmark/bench_fp8_blockwise_gemm.py diff --git a/test/registered/jit/test_cutedsl_bf16_gemm.py b/test/registered/kernels/ops/gemm/test_cutedsl_bf16_gemm.py similarity index 100% rename from test/registered/jit/test_cutedsl_bf16_gemm.py rename to test/registered/kernels/ops/gemm/test_cutedsl_bf16_gemm.py diff --git a/test/registered/jit/test_cutedsl_dsv3_fused_a_gemm.py b/test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py similarity index 100% rename from test/registered/jit/test_cutedsl_dsv3_fused_a_gemm.py rename to test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py diff --git a/test/registered/jit/test_dsv3_fused_a_gemm.py b/test/registered/kernels/ops/gemm/test_dsv3_fused_a_gemm.py similarity index 100% rename from test/registered/jit/test_dsv3_fused_a_gemm.py rename to test/registered/kernels/ops/gemm/test_dsv3_fused_a_gemm.py diff --git a/test/registered/jit/test_dsv3_router_gemm.py b/test/registered/kernels/ops/gemm/test_dsv3_router_gemm.py similarity index 100% rename from test/registered/jit/test_dsv3_router_gemm.py rename to test/registered/kernels/ops/gemm/test_dsv3_router_gemm.py diff --git a/test/registered/jit/test_fp8_blockwise_gemm.py b/test/registered/kernels/ops/gemm/test_fp8_blockwise_gemm.py similarity index 100% rename from test/registered/jit/test_fp8_blockwise_gemm.py rename to test/registered/kernels/ops/gemm/test_fp8_blockwise_gemm.py diff --git a/test/registered/jit/minimax/test_minimax_fused_qkv_index_gemm.py b/test/registered/kernels/ops/gemm/test_minimax_fused_qkv_index_gemm.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_fused_qkv_index_gemm.py rename to test/registered/kernels/ops/gemm/test_minimax_fused_qkv_index_gemm.py diff --git a/test/registered/jit/benchmark/kv_canary/bench_plan.py b/test/registered/kernels/ops/kv_canary/benchmark/bench_plan.py similarity index 100% rename from test/registered/jit/benchmark/kv_canary/bench_plan.py rename to test/registered/kernels/ops/kv_canary/benchmark/bench_plan.py diff --git a/test/registered/jit/benchmark/kv_canary/bench_scatter_req_token_ids.py b/test/registered/kernels/ops/kv_canary/benchmark/bench_scatter_req_token_ids.py similarity index 100% rename from test/registered/jit/benchmark/kv_canary/bench_scatter_req_token_ids.py rename to test/registered/kernels/ops/kv_canary/benchmark/bench_scatter_req_token_ids.py diff --git a/test/registered/jit/benchmark/kv_canary/bench_verify.py b/test/registered/kernels/ops/kv_canary/benchmark/bench_verify.py similarity index 100% rename from test/registered/jit/benchmark/kv_canary/bench_verify.py rename to test/registered/kernels/ops/kv_canary/benchmark/bench_verify.py diff --git a/test/registered/jit/benchmark/kv_canary/bench_write.py b/test/registered/kernels/ops/kv_canary/benchmark/bench_write.py similarity index 100% rename from test/registered/jit/benchmark/kv_canary/bench_write.py rename to test/registered/kernels/ops/kv_canary/benchmark/bench_write.py diff --git a/test/registered/jit/kv_canary/test_const_sync.py b/test/registered/kernels/ops/kv_canary/test_const_sync.py similarity index 100% rename from test/registered/jit/kv_canary/test_const_sync.py rename to test/registered/kernels/ops/kv_canary/test_const_sync.py diff --git a/test/registered/jit/kv_canary/test_kernel_config.py b/test/registered/kernels/ops/kv_canary/test_kernel_config.py similarity index 98% rename from test/registered/jit/kv_canary/test_kernel_config.py rename to test/registered/kernels/ops/kv_canary/test_kernel_config.py index 2917da076a0e..32f7f61320f4 100644 --- a/test/registered/jit/kv_canary/test_kernel_config.py +++ b/test/registered/kernels/ops/kv_canary/test_kernel_config.py @@ -3,7 +3,15 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary.verify import ( + CanaryLaunchTag, + VerifyOrWriteContext, + VerifyPlan, + launch_canary_verify_kernel, +) +from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_buf_equal, assert_canary_state_equal, @@ -15,24 +23,16 @@ make_write_plan_pair, stamp_clean_chain, ) -from sglang.kernels.jit.tests.kv_canary._differential import ( +from sglang.kernels.testing.kv_canary._differential import ( _assert_plans_byte_equal, _run_both_plan, _run_both_verify, _run_both_write, ) -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._fixtures import ( dummy_pseudo_tensors, empty_extras, ) -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary.verify import ( - CanaryLaunchTag, - VerifyOrWriteContext, - VerifyPlan, - launch_canary_verify_kernel, -) -from sglang.kernels.ops.kv_canary.write import WritePlan from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_pipeline_e2e.py b/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py similarity index 99% rename from test/registered/jit/kv_canary/test_pipeline_e2e.py rename to test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py index 24fa49b2a838..932b389c7236 100644 --- a/test/registered/jit/kv_canary/test_pipeline_e2e.py +++ b/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py @@ -5,20 +5,6 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( - FakeViolationLog, - assert_canary_buf_equal, - assert_canary_state_equal, - make_canary_buf, - make_real_kv_sources, - stamp_clean_chain, - write_slot_fields, -) -from sglang.kernels.jit.tests.kv_canary._fixtures import ( - clone_real_kv_sources, - empty_extras, - make_req_to_token, -) from sglang.kernels.ops.kv_canary import consts from sglang.kernels.ops.kv_canary.plan import launch_canary_plan_kernels from sglang.kernels.ops.kv_canary.plan_ref import ( @@ -38,6 +24,20 @@ from sglang.kernels.ops.kv_canary.write_ref import ( launch_canary_write_kernel_torch_reference, ) +from sglang.kernels.testing.kv_canary._canary_helpers import ( + FakeViolationLog, + assert_canary_buf_equal, + assert_canary_state_equal, + make_canary_buf, + make_real_kv_sources, + stamp_clean_chain, + write_slot_fields, +) +from sglang.kernels.testing.kv_canary._fixtures import ( + clone_real_kv_sources, + empty_extras, + make_req_to_token, +) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_plan_fuzz.py b/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py similarity index 96% rename from test/registered/jit/kv_canary/test_plan_fuzz.py rename to test/registered/kernels/ops/kv_canary/test_plan_fuzz.py index 25e3df6fc82f..137506423d62 100644 --- a/test/registered/jit/kv_canary/test_plan_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py @@ -7,19 +7,19 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._differential import _run_both_plan -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._differential import _run_both_plan +from sglang.kernels.testing.kv_canary._fixtures import ( allocate_plan_pair, derive_plan_capacity, make_lut, make_padding_mask, make_req_to_token, ) -from sglang.kernels.jit.tests.kv_canary._fuzz_driver import ( +from sglang.kernels.testing.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.jit.tests.kv_canary._invariants import PlanInvariants +from sglang.kernels.testing.kv_canary._invariants import PlanInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_plan_hand.py b/test/registered/kernels/ops/kv_canary/test_plan_hand.py similarity index 99% rename from test/registered/jit/kv_canary/test_plan_hand.py rename to test/registered/kernels/ops/kv_canary/test_plan_hand.py index bbf543cec577..b175b2c3ce4d 100644 --- a/test/registered/jit/kv_canary/test_plan_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_hand.py @@ -5,21 +5,21 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._differential import run_plan_diff -from sglang.kernels.jit.tests.kv_canary._fixtures import ( - allocate_plan_pair, - derive_plan_capacity, - empty_extras, - make_lut, - make_req_to_token, -) -from sglang.kernels.jit.tests.kv_canary._invariants import PlanInvariants from sglang.kernels.ops.kv_canary.plan import launch_canary_plan_kernels from sglang.kernels.ops.kv_canary.plan_ref import ( launch_canary_plan_kernels_torch_reference, ) from sglang.kernels.ops.kv_canary.verify import VerifyPlan from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._differential import run_plan_diff +from sglang.kernels.testing.kv_canary._fixtures import ( + allocate_plan_pair, + derive_plan_capacity, + empty_extras, + make_lut, + make_req_to_token, +) +from sglang.kernels.testing.kv_canary._invariants import PlanInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_scatter_req_token_ids.py b/test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py similarity index 100% rename from test/registered/jit/kv_canary/test_scatter_req_token_ids.py rename to test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py diff --git a/test/registered/jit/kv_canary/test_utils.py b/test/registered/kernels/ops/kv_canary/test_utils.py similarity index 100% rename from test/registered/jit/kv_canary/test_utils.py rename to test/registered/kernels/ops/kv_canary/test_utils.py diff --git a/test/registered/jit/kv_canary/test_verify_fuzz.py b/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py similarity index 95% rename from test/registered/jit/kv_canary/test_verify_fuzz.py rename to test/registered/kernels/ops/kv_canary/test_verify_fuzz.py index 56fa62eba0e1..ad8bf4921727 100644 --- a/test/registered/jit/kv_canary/test_verify_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py @@ -6,29 +6,29 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary.verify import ( + CanaryLaunchTag, + RealKvSource, + VerifyPlan, +) +from sglang.kernels.testing.kv_canary._canary_helpers import ( FakeViolationLog, make_canary_buf, make_log_pair, make_verify_plan_pair, stamp_clean_chain, ) -from sglang.kernels.jit.tests.kv_canary._differential import _run_both_verify -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._differential import _run_both_verify +from sglang.kernels.testing.kv_canary._fixtures import ( clone_real_kv_sources, make_real_kv_sources, ) -from sglang.kernels.jit.tests.kv_canary._fuzz_driver import ( +from sglang.kernels.testing.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.jit.tests.kv_canary._invariants import VerifyInvariants -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary.verify import ( - CanaryLaunchTag, - RealKvSource, - VerifyPlan, -) +from sglang.kernels.testing.kv_canary._invariants import VerifyInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_verify_hand.py b/test/registered/kernels/ops/kv_canary/test_verify_hand.py similarity index 99% rename from test/registered/jit/kv_canary/test_verify_hand.py rename to test/registered/kernels/ops/kv_canary/test_verify_hand.py index ea68821e8e64..de9fbe740ff2 100644 --- a/test/registered/jit/kv_canary/test_verify_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_hand.py @@ -8,7 +8,23 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 +from sglang.kernels.ops.kv_canary.verify import ( + CanaryLaunchTag, + RealKvSource, + VerifyOrWriteContext, + VerifyPlan, + launch_canary_verify_kernel, +) +from sglang.kernels.ops.kv_canary.verify_ref import ( + _compute_real_kv_hash_scalar, + launch_canary_verify_kernel_torch_reference, +) +from sglang.kernels.ops.kv_canary.write_ref import ( + launch_canary_write_kernel_torch_reference, +) +from sglang.kernels.testing.kv_canary._canary_helpers import ( FakeViolationLog, assert_only_bits_set, chain_anchor_signed, @@ -26,31 +42,15 @@ to_signed_int64, write_slot_fields, ) -from sglang.kernels.jit.tests.kv_canary._differential import ( +from sglang.kernels.testing.kv_canary._differential import ( _run_both_verify, run_verify_diff, ) -from sglang.kernels.jit.tests.kv_canary._fixtures import clone_real_kv_sources -from sglang.kernels.jit.tests.kv_canary._hand_oracle import ( +from sglang.kernels.testing.kv_canary._fixtures import clone_real_kv_sources +from sglang.kernels.testing.kv_canary._hand_oracle import ( _hand_fold_all, _hand_fold_partial, ) -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 -from sglang.kernels.ops.kv_canary.verify import ( - CanaryLaunchTag, - RealKvSource, - VerifyOrWriteContext, - VerifyPlan, - launch_canary_verify_kernel, -) -from sglang.kernels.ops.kv_canary.verify_ref import ( - _compute_real_kv_hash_scalar, - launch_canary_verify_kernel_torch_reference, -) -from sglang.kernels.ops.kv_canary.write_ref import ( - launch_canary_write_kernel_torch_reference, -) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_write_fuzz.py b/test/registered/kernels/ops/kv_canary/test_write_fuzz.py similarity index 96% rename from test/registered/jit/kv_canary/test_write_fuzz.py rename to test/registered/kernels/ops/kv_canary/test_write_fuzz.py index 87d8427f7814..18b4f12ab318 100644 --- a/test/registered/jit/kv_canary/test_write_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_write_fuzz.py @@ -6,29 +6,29 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary.verify import ( + CanaryLaunchTag, + RealKvSource, +) +from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._canary_helpers import ( FakeViolationLog, make_canary_buf, make_log_pair, make_write_plan_pair, stamp_pair, ) -from sglang.kernels.jit.tests.kv_canary._differential import _run_both_write -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._differential import _run_both_write +from sglang.kernels.testing.kv_canary._fixtures import ( clone_real_kv_sources, make_real_kv_sources, ) -from sglang.kernels.jit.tests.kv_canary._fuzz_driver import ( +from sglang.kernels.testing.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.jit.tests.kv_canary._invariants import WriteInvariants -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary.verify import ( - CanaryLaunchTag, - RealKvSource, -) -from sglang.kernels.ops.kv_canary.write import WritePlan +from sglang.kernels.testing.kv_canary._invariants import WriteInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/kv_canary/test_write_hand.py b/test/registered/kernels/ops/kv_canary/test_write_hand.py similarity index 99% rename from test/registered/jit/kv_canary/test_write_hand.py rename to test/registered/kernels/ops/kv_canary/test_write_hand.py index 9cc763dafd91..933fee9d1550 100644 --- a/test/registered/jit/kv_canary/test_write_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_write_hand.py @@ -7,7 +7,20 @@ import pytest import torch -from sglang.kernels.jit.tests.kv_canary._canary_helpers import ( +from sglang.kernels.ops.kv_canary import consts +from sglang.kernels.ops.kv_canary import write as write_module +from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 +from sglang.kernels.ops.kv_canary.verify import ( + CANARY_SLOT_BYTES, + CanaryLaunchTag, + RealKvSource, + VerifyOrWriteContext, + launch_canary_verify_kernel, +) +from sglang.kernels.ops.kv_canary.write import ( + launch_canary_write_kernel, +) +from sglang.kernels.testing.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_state_equal, assert_only_bits_set, @@ -24,31 +37,18 @@ stamp_pair, to_signed_int64, ) -from sglang.kernels.jit.tests.kv_canary._differential import ( +from sglang.kernels.testing.kv_canary._differential import ( _run_both_write, run_write_diff, ) -from sglang.kernels.jit.tests.kv_canary._fixtures import ( +from sglang.kernels.testing.kv_canary._fixtures import ( clone_real_kv_sources, dummy_pseudo_tensors, ) -from sglang.kernels.jit.tests.kv_canary._hand_oracle import ( +from sglang.kernels.testing.kv_canary._hand_oracle import ( _hand_fold_all, _hand_fold_partial, ) -from sglang.kernels.ops.kv_canary import consts -from sglang.kernels.ops.kv_canary import write as write_module -from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 -from sglang.kernels.ops.kv_canary.verify import ( - CANARY_SLOT_BYTES, - CanaryLaunchTag, - RealKvSource, - VerifyOrWriteContext, - launch_canary_verify_kernel, -) -from sglang.kernels.ops.kv_canary.write import ( - launch_canary_write_kernel, -) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/jit/benchmark/bench_fused_fp8_qkv_kv_cache.py b/test/registered/kernels/ops/kvcache/benchmark/bench_fused_fp8_qkv_kv_cache.py similarity index 100% rename from test/registered/jit/benchmark/bench_fused_fp8_qkv_kv_cache.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_fused_fp8_qkv_kv_cache.py diff --git a/test/registered/jit/benchmark/bench_hicache.py b/test/registered/kernels/ops/kvcache/benchmark/bench_hicache.py similarity index 100% rename from test/registered/jit/benchmark/bench_hicache.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_hicache.py diff --git a/test/registered/jit/benchmark/bench_hisparse.py b/test/registered/kernels/ops/kvcache/benchmark/bench_hisparse.py similarity index 100% rename from test/registered/jit/benchmark/bench_hisparse.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_hisparse.py diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py b/test/registered/kernels/ops/kvcache/benchmark/bench_minimax_store_kv_index.py similarity index 100% rename from test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_minimax_store_kv_index.py diff --git a/test/registered/jit/benchmark/bench_set_mla_kv_buffer.py b/test/registered/kernels/ops/kvcache/benchmark/bench_set_mla_kv_buffer.py similarity index 100% rename from test/registered/jit/benchmark/bench_set_mla_kv_buffer.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_set_mla_kv_buffer.py diff --git a/test/registered/jit/benchmark/bench_store_cache.py b/test/registered/kernels/ops/kvcache/benchmark/bench_store_cache.py similarity index 100% rename from test/registered/jit/benchmark/bench_store_cache.py rename to test/registered/kernels/ops/kvcache/benchmark/bench_store_cache.py diff --git a/test/registered/jit/test_fused_fp8_qkv_kv_cache.py b/test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py similarity index 100% rename from test/registered/jit/test_fused_fp8_qkv_kv_cache.py rename to test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py diff --git a/test/registered/jit/test_hicache.py b/test/registered/kernels/ops/kvcache/test_hicache.py similarity index 100% rename from test/registered/jit/test_hicache.py rename to test/registered/kernels/ops/kvcache/test_hicache.py diff --git a/test/registered/jit/test_hicache_page_first_write_back.py b/test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py similarity index 100% rename from test/registered/jit/test_hicache_page_first_write_back.py rename to test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py diff --git a/test/registered/jit/test_hisparse.py b/test/registered/kernels/ops/kvcache/test_hisparse.py similarity index 100% rename from test/registered/jit/test_hisparse.py rename to test/registered/kernels/ops/kvcache/test_hisparse.py diff --git a/test/registered/jit/test_kvcacheio_asymmetric.py b/test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py similarity index 100% rename from test/registered/jit/test_kvcacheio_asymmetric.py rename to test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py diff --git a/test/registered/jit/minimax/test_minimax_store_kv_index.py b/test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_store_kv_index.py rename to test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py diff --git a/test/registered/jit/test_set_mla_kv_buffer.py b/test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py similarity index 100% rename from test/registered/jit/test_set_mla_kv_buffer.py rename to test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py diff --git a/test/registered/jit/test_store_cache.py b/test/registered/kernels/ops/kvcache/test_store_cache.py similarity index 100% rename from test/registered/jit/test_store_cache.py rename to test/registered/kernels/ops/kvcache/test_store_cache.py diff --git a/test/registered/jit/benchmark/bench_fused_eh_norm.py b/test/registered/kernels/ops/layernorm/benchmark/bench_fused_eh_norm.py similarity index 100% rename from test/registered/jit/benchmark/bench_fused_eh_norm.py rename to test/registered/kernels/ops/layernorm/benchmark/bench_fused_eh_norm.py diff --git a/test/registered/jit/benchmark/bench_norm.py b/test/registered/kernels/ops/layernorm/benchmark/bench_norm.py similarity index 100% rename from test/registered/jit/benchmark/bench_norm.py rename to test/registered/kernels/ops/layernorm/benchmark/bench_norm.py diff --git a/test/registered/jit/benchmark/bench_qknorm.py b/test/registered/kernels/ops/layernorm/benchmark/bench_qknorm.py similarity index 100% rename from test/registered/jit/benchmark/bench_qknorm.py rename to test/registered/kernels/ops/layernorm/benchmark/bench_qknorm.py diff --git a/test/registered/jit/benchmark/bench_qknorm_across_heads.py b/test/registered/kernels/ops/layernorm/benchmark/bench_qknorm_across_heads.py similarity index 100% rename from test/registered/jit/benchmark/bench_qknorm_across_heads.py rename to test/registered/kernels/ops/layernorm/benchmark/bench_qknorm_across_heads.py diff --git a/test/registered/jit/test_fused_add_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py similarity index 100% rename from test/registered/jit/test_fused_add_rmsnorm.py rename to test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py diff --git a/test/registered/jit/test_fused_eh_norm.py b/test/registered/kernels/ops/layernorm/test_fused_eh_norm.py similarity index 100% rename from test/registered/jit/test_fused_eh_norm.py rename to test/registered/kernels/ops/layernorm/test_fused_eh_norm.py diff --git a/test/registered/kernels/test_fused_op.py b/test/registered/kernels/ops/layernorm/test_fused_op.py similarity index 100% rename from test/registered/kernels/test_fused_op.py rename to test/registered/kernels/ops/layernorm/test_fused_op.py diff --git a/test/registered/kernels/test_fused_op_gpu_parity.py b/test/registered/kernels/ops/layernorm/test_fused_op_gpu_parity.py similarity index 100% rename from test/registered/kernels/test_fused_op_gpu_parity.py rename to test/registered/kernels/ops/layernorm/test_fused_op_gpu_parity.py diff --git a/test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py b/test/registered/kernels/ops/layernorm/test_fused_qk_gemma_rmsnorm_gate.py similarity index 100% rename from test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py rename to test/registered/kernels/ops/layernorm/test_fused_qk_gemma_rmsnorm_gate.py diff --git a/test/registered/kernels/test_gemma4_fused_routing.py b/test/registered/kernels/ops/layernorm/test_gemma4_fused_routing.py similarity index 100% rename from test/registered/kernels/test_gemma4_fused_routing.py rename to test/registered/kernels/ops/layernorm/test_gemma4_fused_routing.py diff --git a/test/registered/kernels/test_kernels_namespace.py b/test/registered/kernels/ops/layernorm/test_kernels_namespace.py similarity index 100% rename from test/registered/kernels/test_kernels_namespace.py rename to test/registered/kernels/ops/layernorm/test_kernels_namespace.py diff --git a/test/registered/kernels/test_mhc_kernels.py b/test/registered/kernels/ops/layernorm/test_mhc_kernels.py similarity index 100% rename from test/registered/kernels/test_mhc_kernels.py rename to test/registered/kernels/ops/layernorm/test_mhc_kernels.py diff --git a/python/sglang/kernels/jit/tests/test_minimax_m3_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py similarity index 95% rename from python/sglang/kernels/jit/tests/test_minimax_m3_rmsnorm.py rename to test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py index 7eb477fecf21..e8f5abb8c31e 100644 --- a/python/sglang/kernels/jit/tests/test_minimax_m3_rmsnorm.py +++ b/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py @@ -18,6 +18,9 @@ gemma_fused_add_rmsnorm, gemma_rmsnorm, ) +from sglang.test.ci.ci_register import register_amd_ci + +register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" EPS = 1e-6 diff --git a/test/registered/jit/test_qknorm.py b/test/registered/kernels/ops/layernorm/test_qknorm.py similarity index 100% rename from test/registered/jit/test_qknorm.py rename to test/registered/kernels/ops/layernorm/test_qknorm.py diff --git a/test/registered/jit/test_qknorm_across_heads.py b/test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py similarity index 100% rename from test/registered/jit/test_qknorm_across_heads.py rename to test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py diff --git a/test/registered/jit/test_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_rmsnorm.py similarity index 100% rename from test/registered/jit/test_rmsnorm.py rename to test/registered/kernels/ops/layernorm/test_rmsnorm.py diff --git a/test/registered/jit/test_rmsnorm_hf.py b/test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py similarity index 100% rename from test/registered/jit/test_rmsnorm_hf.py rename to test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py diff --git a/python/sglang/kernels/jit/tests/test_sconv_decode_metadata.py b/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py similarity index 94% rename from python/sglang/kernels/jit/tests/test_sconv_decode_metadata.py rename to test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py index 6ac9f4c14cf2..3a3cb43fced1 100644 --- a/python/sglang/kernels/jit/tests/test_sconv_decode_metadata.py +++ b/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py @@ -13,6 +13,9 @@ fused_decode_sconv_metadata, precompute_helion_decode_metadata, ) +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/python/sglang/kernels/jit/tests/test_sconv_extend_metadata.py b/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py similarity index 97% rename from python/sglang/kernels/jit/tests/test_sconv_extend_metadata.py rename to test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py index cecb81ff5104..de7e36143534 100644 --- a/python/sglang/kernels/jit/tests/test_sconv_extend_metadata.py +++ b/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py @@ -18,6 +18,9 @@ fused_extend_sconv_metadata, precompute_helion_extend_metadata, ) +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/jit/test_transfer_mamba.py b/test/registered/kernels/ops/mamba/test_transfer_mamba.py similarity index 100% rename from test/registered/jit/test_transfer_mamba.py rename to test/registered/kernels/ops/mamba/test_transfer_mamba.py diff --git a/test/registered/jit/test_inkling_rel_proj.py b/test/registered/kernels/ops/model/test_inkling_rel_proj.py similarity index 100% rename from test/registered/jit/test_inkling_rel_proj.py rename to test/registered/kernels/ops/model/test_inkling_rel_proj.py diff --git a/test/registered/jit/benchmark/bench_moe_fused_gate.py b/test/registered/kernels/ops/moe/benchmark/bench_moe_fused_gate.py similarity index 100% rename from test/registered/jit/benchmark/bench_moe_fused_gate.py rename to test/registered/kernels/ops/moe/benchmark/bench_moe_fused_gate.py diff --git a/test/registered/jit/benchmark/bench_post_reorder_deepgemm.py b/test/registered/kernels/ops/moe/benchmark/bench_post_reorder_deepgemm.py similarity index 100% rename from test/registered/jit/benchmark/bench_post_reorder_deepgemm.py rename to test/registered/kernels/ops/moe/benchmark/bench_post_reorder_deepgemm.py diff --git a/test/registered/jit/benchmark/bench_renorm.py b/test/registered/kernels/ops/moe/benchmark/bench_renorm.py similarity index 100% rename from test/registered/jit/benchmark/bench_renorm.py rename to test/registered/kernels/ops/moe/benchmark/bench_renorm.py diff --git a/test/registered/kernels/test_fp4_moe.py b/test/registered/kernels/ops/moe/test_fp4_moe.py similarity index 100% rename from test/registered/kernels/test_fp4_moe.py rename to test/registered/kernels/ops/moe/test_fp4_moe.py diff --git a/test/registered/kernels/test_fused_topk_deepseek.py b/test/registered/kernels/ops/moe/test_fused_topk_deepseek.py similarity index 100% rename from test/registered/kernels/test_fused_topk_deepseek.py rename to test/registered/kernels/ops/moe/test_fused_topk_deepseek.py diff --git a/test/registered/jit/test_kpool_topk_transform.py b/test/registered/kernels/ops/moe/test_kpool_topk_transform.py similarity index 100% rename from test/registered/jit/test_kpool_topk_transform.py rename to test/registered/kernels/ops/moe/test_kpool_topk_transform.py diff --git a/test/registered/jit/test_marlin_packed_topk_unpack.py b/test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py similarity index 100% rename from test/registered/jit/test_marlin_packed_topk_unpack.py rename to test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py diff --git a/python/sglang/kernels/jit/tests/test_minimax_m3_mxfp8.py b/test/registered/kernels/ops/moe/test_minimax_m3_mxfp8.py similarity index 97% rename from python/sglang/kernels/jit/tests/test_minimax_m3_mxfp8.py rename to test/registered/kernels/ops/moe/test_minimax_m3_mxfp8.py index 3038d840539f..bc9f18d94933 100644 --- a/python/sglang/kernels/jit/tests/test_minimax_m3_mxfp8.py +++ b/test/registered/kernels/ops/moe/test_minimax_m3_mxfp8.py @@ -12,7 +12,7 @@ ``dot_scaled`` linear/MoE tests are gated to CDNA4 gfx95x (the hardware microscaling matrix cores) -- gfx942 has no native ``dot_scaled`` MX path. -Run: pytest python/sglang/kernels/jit/tests/test_minimax_m3_mxfp8.py -v +Run: pytest python/sglang/kernels/testing/test_minimax_m3_mxfp8.py -v """ import pytest @@ -33,6 +33,9 @@ _mxfp8_e4m3_quantize_triton, dequant_mxfp8_to_bf16, ) +from sglang.test.ci.ci_register import register_amd_ci + +register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" diff --git a/test/registered/jit/minimax/test_minimax_quant_scatter.py b/test/registered/kernels/ops/moe/test_minimax_quant_scatter.py similarity index 100% rename from test/registered/jit/minimax/test_minimax_quant_scatter.py rename to test/registered/kernels/ops/moe/test_minimax_quant_scatter.py diff --git a/test/registered/jit/test_moe_align_block_size.py b/test/registered/kernels/ops/moe/test_moe_align_block_size.py similarity index 100% rename from test/registered/jit/test_moe_align_block_size.py rename to test/registered/kernels/ops/moe/test_moe_align_block_size.py diff --git a/test/registered/jit/test_moe_fused_gate.py b/test/registered/kernels/ops/moe/test_moe_fused_gate.py similarity index 100% rename from test/registered/jit/test_moe_fused_gate.py rename to test/registered/kernels/ops/moe/test_moe_fused_gate.py diff --git a/test/registered/jit/test_moe_lora_align_block_size.py b/test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py similarity index 100% rename from test/registered/jit/test_moe_lora_align_block_size.py rename to test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py diff --git a/python/sglang/kernels/jit/tests/test_moe_preprocess.py b/test/registered/kernels/ops/moe/test_moe_preprocess.py similarity index 95% rename from python/sglang/kernels/jit/tests/test_moe_preprocess.py rename to test/registered/kernels/ops/moe/test_moe_preprocess.py index 121811e82a18..43e12c59e811 100644 --- a/python/sglang/kernels/jit/tests/test_moe_preprocess.py +++ b/test/registered/kernels/ops/moe/test_moe_preprocess.py @@ -13,6 +13,9 @@ get_src2dst, grouped_gemm_triton, ) +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/python/sglang/kernels/jit/tests/test_moe_topk_sigmoid.py b/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py similarity index 99% rename from python/sglang/kernels/jit/tests/test_moe_topk_sigmoid.py rename to test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py index 769dacd9232e..785ef7ec3eb4 100644 --- a/python/sglang/kernels/jit/tests/test_moe_topk_sigmoid.py +++ b/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py @@ -14,6 +14,9 @@ import torch from sglang.kernels.ops.moe.moe_topk_sigmoid import topk_sigmoid +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") try: from sgl_kernel import topk_sigmoid as topk_sigmoid_aot diff --git a/test/registered/jit/test_moe_wna16_marlin.py b/test/registered/kernels/ops/moe/test_moe_wna16_marlin.py similarity index 100% rename from test/registered/jit/test_moe_wna16_marlin.py rename to test/registered/kernels/ops/moe/test_moe_wna16_marlin.py diff --git a/test/registered/jit/test_post_reorder_deepgemm.py b/test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py similarity index 100% rename from test/registered/jit/test_post_reorder_deepgemm.py rename to test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py diff --git a/test/registered/jit/test_renorm.py b/test/registered/kernels/ops/moe/test_renorm.py similarity index 100% rename from test/registered/jit/test_renorm.py rename to test/registered/kernels/ops/moe/test_renorm.py diff --git a/test/registered/jit/test_sigmoid_gate_mul.py b/test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py similarity index 100% rename from test/registered/jit/test_sigmoid_gate_mul.py rename to test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py diff --git a/test/registered/kernels/test_sigmoid_gate_mul.py b/test/registered/kernels/ops/moe/test_sigmoid_gate_mul_broadcast.py similarity index 100% rename from test/registered/kernels/test_sigmoid_gate_mul.py rename to test/registered/kernels/ops/moe/test_sigmoid_gate_mul_broadcast.py diff --git a/test/registered/jit/benchmark/bench_awq_dequantize.py b/test/registered/kernels/ops/quantization/benchmark/bench_awq_dequantize.py similarity index 100% rename from test/registered/jit/benchmark/bench_awq_dequantize.py rename to test/registered/kernels/ops/quantization/benchmark/bench_awq_dequantize.py diff --git a/test/registered/jit/benchmark/bench_mxfp8_moe.py b/test/registered/kernels/ops/quantization/benchmark/bench_mxfp8_moe.py similarity index 100% rename from test/registered/jit/benchmark/bench_mxfp8_moe.py rename to test/registered/kernels/ops/quantization/benchmark/bench_mxfp8_moe.py diff --git a/test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py b/test/registered/kernels/ops/quantization/benchmark/bench_per_tensor_quant_fp8.py similarity index 100% rename from test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py rename to test/registered/kernels/ops/quantization/benchmark/bench_per_tensor_quant_fp8.py diff --git a/test/registered/jit/benchmark/bench_per_token_group_quant.py b/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant.py similarity index 100% rename from test/registered/jit/benchmark/bench_per_token_group_quant.py rename to test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant.py diff --git a/test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py b/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_8bit_v2.py similarity index 100% rename from test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py rename to test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_8bit_v2.py diff --git a/test/registered/jit/benchmark/bench_per_token_group_quant_masked.py b/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_masked.py similarity index 100% rename from test/registered/jit/benchmark/bench_per_token_group_quant_masked.py rename to test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_masked.py diff --git a/test/registered/jit/test_awq_dequantize.py b/test/registered/kernels/ops/quantization/test_awq_dequantize.py similarity index 100% rename from test/registered/jit/test_awq_dequantize.py rename to test/registered/kernels/ops/quantization/test_awq_dequantize.py diff --git a/test/registered/jit/test_awq_marlin_moe_repack.py b/test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py similarity index 100% rename from test/registered/jit/test_awq_marlin_moe_repack.py rename to test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py diff --git a/test/registered/jit/test_awq_marlin_repack.py b/test/registered/kernels/ops/quantization/test_awq_marlin_repack.py similarity index 100% rename from test/registered/jit/test_awq_marlin_repack.py rename to test/registered/kernels/ops/quantization/test_awq_marlin_repack.py diff --git a/test/registered/jit/test_gptq_marlin.py b/test/registered/kernels/ops/quantization/test_gptq_marlin.py similarity index 100% rename from test/registered/jit/test_gptq_marlin.py rename to test/registered/kernels/ops/quantization/test_gptq_marlin.py diff --git a/test/registered/jit/test_gptq_marlin_repack.py b/test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py similarity index 100% rename from test/registered/jit/test_gptq_marlin_repack.py rename to test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py diff --git a/test/registered/jit/test_mxfp8_moe.py b/test/registered/kernels/ops/quantization/test_mxfp8_moe.py similarity index 100% rename from test/registered/jit/test_mxfp8_moe.py rename to test/registered/kernels/ops/quantization/test_mxfp8_moe.py diff --git a/test/registered/jit/test_per_tensor_quant_fp8.py b/test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py similarity index 100% rename from test/registered/jit/test_per_tensor_quant_fp8.py rename to test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py diff --git a/test/registered/jit/test_per_token_group_quant.py b/test/registered/kernels/ops/quantization/test_per_token_group_quant.py similarity index 100% rename from test/registered/jit/test_per_token_group_quant.py rename to test/registered/kernels/ops/quantization/test_per_token_group_quant.py diff --git a/test/registered/jit/test_per_token_group_quant_8bit_v2.py b/test/registered/kernels/ops/quantization/test_per_token_group_quant_8bit_v2.py similarity index 100% rename from test/registered/jit/test_per_token_group_quant_8bit_v2.py rename to test/registered/kernels/ops/quantization/test_per_token_group_quant_8bit_v2.py diff --git a/test/registered/jit/benchmark/bench_ngram_compute_decode.py b/test/registered/kernels/ops/speculative/benchmark/bench_ngram_compute_decode.py similarity index 100% rename from test/registered/jit/benchmark/bench_ngram_compute_decode.py rename to test/registered/kernels/ops/speculative/benchmark/bench_ngram_compute_decode.py diff --git a/test/registered/jit/benchmark/bench_ngram_update_token_table.py b/test/registered/kernels/ops/speculative/benchmark/bench_ngram_update_token_table.py similarity index 100% rename from test/registered/jit/benchmark/bench_ngram_update_token_table.py rename to test/registered/kernels/ops/speculative/benchmark/bench_ngram_update_token_table.py diff --git a/test/registered/jit/benchmark/bench_resolve_future_token_ids.py b/test/registered/kernels/ops/speculative/benchmark/bench_resolve_future_token_ids.py similarity index 100% rename from test/registered/jit/benchmark/bench_resolve_future_token_ids.py rename to test/registered/kernels/ops/speculative/benchmark/bench_resolve_future_token_ids.py diff --git a/test/registered/jit/benchmark/bench_spec_topk1.py b/test/registered/kernels/ops/speculative/benchmark/bench_spec_topk1.py similarity index 100% rename from test/registered/jit/benchmark/bench_spec_topk1.py rename to test/registered/kernels/ops/speculative/benchmark/bench_spec_topk1.py diff --git a/test/registered/kernels/test_boundary_kv_fix_kernels.py b/test/registered/kernels/ops/speculative/test_boundary_kv_fix_kernels.py similarity index 100% rename from test/registered/kernels/test_boundary_kv_fix_kernels.py rename to test/registered/kernels/ops/speculative/test_boundary_kv_fix_kernels.py diff --git a/test/registered/kernels/test_gather_spec_extras.py b/test/registered/kernels/ops/speculative/test_gather_spec_extras.py similarity index 100% rename from test/registered/kernels/test_gather_spec_extras.py rename to test/registered/kernels/ops/speculative/test_gather_spec_extras.py diff --git a/test/registered/jit/test_ngram_embedding.py b/test/registered/kernels/ops/speculative/test_ngram_embedding.py similarity index 100% rename from test/registered/jit/test_ngram_embedding.py rename to test/registered/kernels/ops/speculative/test_ngram_embedding.py diff --git a/test/registered/jit/test_resolve_future_token_ids.py b/test/registered/kernels/ops/speculative/test_resolve_future_token_ids.py similarity index 100% rename from test/registered/jit/test_resolve_future_token_ids.py rename to test/registered/kernels/ops/speculative/test_resolve_future_token_ids.py diff --git a/test/registered/kernels/test_spec_topk1.py b/test/registered/kernels/ops/speculative/test_spec_topk1.py similarity index 100% rename from test/registered/kernels/test_spec_topk1.py rename to test/registered/kernels/ops/speculative/test_spec_topk1.py diff --git a/test/registered/unit/distributed/test_vmm_utils.py b/test/registered/unit/distributed/test_vmm_utils.py index fd3e6f5c642d..55511e01d1a6 100644 --- a/test/registered/unit/distributed/test_vmm_utils.py +++ b/test/registered/unit/distributed/test_vmm_utils.py @@ -20,8 +20,8 @@ import torch.distributed as dist from cuda.bindings import driver as drv -from sglang.kernels.jit.tests.utils import multigpu_pytest_main from sglang.kernels.jit.utils import cache_once +from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.vmm_utils import ( check_drv, exchange_posix_fds, From 9c34507d596d2776db880030cf0c865d305c7a19 Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 09:17:46 +0800 Subject: [PATCH 2/6] [Kernel] Follow-up: move nvfp4 bench data fixture + repoint CI change-filters to kernels/ops - Move test/registered/jit/benchmark/diffusion/diffusion_nvfp4_shapes.json (a non-.py data fixture the earlier .py-only move missed) into the nvfp4 bench's new dir; bench loads it via Path(__file__).parent / "...json". - Repoint stale CI change-filters test/registered/jit/** -> test/registered/kernels/ops/** (and the diffusion-specific globs) in pr-test-amd, pr-test-amd-rocm720, _pr-test-check-changes, so editing a moved test/bench still triggers its suite. test/registered/jit/ now fully empty. Benchmarks themselves keep running unchanged: run_suite.py discovers recursively over test/registered/** keyed on register_*_ci (path-independent). Co-Authored-By: Claude Opus 4.8 --- .github/workflows/_pr-test-check-changes.yml | 6 +++--- .github/workflows/pr-test-amd-rocm720.yml | 6 +++--- .github/workflows/pr-test-amd.yml | 6 +++--- .../ops/diffusion/benchmark}/diffusion_nvfp4_shapes.json | 0 4 files changed, 9 insertions(+), 9 deletions(-) rename test/registered/{jit/benchmark/diffusion => kernels/ops/diffusion/benchmark}/diffusion_nvfp4_shapes.json (100%) diff --git a/.github/workflows/_pr-test-check-changes.yml b/.github/workflows/_pr-test-check-changes.yml index f65e1a721bb1..b99744a223df 100644 --- a/.github/workflows/_pr-test-check-changes.yml +++ b/.github/workflows/_pr-test-check-changes.yml @@ -94,14 +94,14 @@ jobs: - "python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/**" - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - - "test/registered/jit/diffusion/**" - - "test/registered/jit/benchmark/diffusion/**" + - "test/registered/kernels/ops/diffusion/**" + - "test/registered/kernels/ops/diffusion/benchmark/**" - "python/sglang/cli/**" jit_kernel: - ".github/workflows/pr-test.yml" - ".github/workflows/pr-test-jit-kernel.yml" - "python/pyproject.toml" - - "test/registered/jit/**" + - "test/registered/kernels/ops/**" # sglang.kernels is the migrated kernel namespace (RFC #29630 / #30044); the # base-b-kernel suites import it directly, so kernel edits must run them. - "python/sglang/kernels/**" diff --git a/.github/workflows/pr-test-amd-rocm720.yml b/.github/workflows/pr-test-amd-rocm720.yml index 4ba8b30b510f..8a4ddfdd5409 100644 --- a/.github/workflows/pr-test-amd-rocm720.yml +++ b/.github/workflows/pr-test-amd-rocm720.yml @@ -199,15 +199,15 @@ jobs: - ".github/workflows/pr-test-amd-rocm720.yml" jit_kernel: - "python/sglang/kernels/**" - - "test/registered/jit/**" + - "test/registered/kernels/ops/**" - ".github/workflows/pr-test-amd-rocm720.yml" multimodal_gen: - "python/sglang/multimodal_gen/**/!(*.md|*.ipynb)" - "python/sglang/cli/**" - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - - "test/registered/jit/diffusion/**" - - "test/registered/jit/benchmark/diffusion/**" + - "test/registered/kernels/ops/diffusion/**" + - "test/registered/kernels/ops/diffusion/benchmark/**" - "python/pyproject_rocm.toml" - "python/pyproject_other.toml" diff --git a/.github/workflows/pr-test-amd.yml b/.github/workflows/pr-test-amd.yml index 7b4b15a263ae..3300d76d910c 100644 --- a/.github/workflows/pr-test-amd.yml +++ b/.github/workflows/pr-test-amd.yml @@ -187,15 +187,15 @@ jobs: - ".github/workflows/pr-test-amd.yml" jit_kernel: - "python/sglang/kernels/**" - - "test/registered/jit/**" + - "test/registered/kernels/ops/**" - ".github/workflows/pr-test-amd.yml" multimodal_gen: - "python/sglang/multimodal_gen/**/!(*.md|*.ipynb)" - "python/sglang/cli/**" - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - - "test/registered/jit/diffusion/**" - - "test/registered/jit/benchmark/diffusion/**" + - "test/registered/kernels/ops/diffusion/**" + - "test/registered/kernels/ops/diffusion/benchmark/**" - "python/pyproject_rocm.toml" - "python/pyproject_other.toml" diff --git a/test/registered/jit/benchmark/diffusion/diffusion_nvfp4_shapes.json b/test/registered/kernels/ops/diffusion/benchmark/diffusion_nvfp4_shapes.json similarity index 100% rename from test/registered/jit/benchmark/diffusion/diffusion_nvfp4_shapes.json rename to test/registered/kernels/ops/diffusion/benchmark/diffusion_nvfp4_shapes.json From 8c7935fcf617066611984c2cd320a651590cd2bd Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 09:28:10 +0800 Subject: [PATCH 3/6] [Kernel] Separate benchmarks from tests: test/registered/kernels/benchmark// Per review: keep benchmarks out of the ops/ test dirs, but still classified by group. Move test/registered/kernels/ops//benchmark/bench_*.py -> test/registered/kernels/benchmark// (57 benches + the nvfp4 shapes JSON fixture). Tests stay in test/registered/kernels/ops//. Benchmarks stay under test/registered/ (NOT the package): they are CI-registered and run_suite.py only globs test/registered/**, plus the check-no-registered-tests-in-package hook forbids register_*_ci inside the wheel. So this is the closest layout to "benchmarks under a jit-style dir" that keeps them runnable in CI. CI change-filters updated: general test path -> test/registered/kernels/** (covers both ops/ tests and benchmark/ benches); diffusion bench filter -> test/registered/kernels/benchmark/diffusion/**. Co-Authored-By: Claude Opus 4.8 --- .github/workflows/_pr-test-check-changes.yml | 4 ++-- .github/workflows/pr-test-amd-rocm720.yml | 4 ++-- .github/workflows/pr-test-amd.yml | 4 ++-- .../benchmark => benchmark/activation}/bench_activation.py | 0 .../benchmark => benchmark/attention}/bench_add_constant.py | 0 .../benchmark => benchmark/attention}/bench_clamp_position.py | 0 .../benchmark => benchmark/attention}/bench_concat_mla.py | 0 .../attention}/bench_dsv4_fp4_indexer.py | 0 .../attention}/bench_fused_qknorm_rope.py | 0 .../benchmark => benchmark/attention}/bench_hadamard.py | 0 .../attention}/bench_minimax_decode_topk.py | 0 .../attention}/bench_minimax_qknorm_rope.py | 0 .../attention}/bench_mla_kv_pack_quantize_fp8.py | 0 .../attention}/bench_online_c128_mtp.py | 0 .../attention/benchmark => benchmark/attention}/bench_rope.py | 0 .../attention}/bench_sparse_mla_q8kv8_prefill_sm90.py | 0 .../attention/benchmark => benchmark/attention}/bench_topk.py | 0 .../communication}/bench_custom_all_reduce.py | 0 .../communication}/bench_symm_mem_all_gather.py | 0 .../benchmark => benchmark/communication}/bench_tp_qknorm.py | 0 .../diffusion}/bench_causal_conv3d_cat_pad.py | 0 .../diffusion}/bench_diffusion_nvfp4_scaled_mm.py | 0 .../diffusion}/bench_fused_norm_scale_shift.py | 0 .../diffusion}/bench_group_norm_silu.py | 0 .../diffusion}/bench_ltx2_qknorm_split_rope.py | 0 .../benchmark => benchmark/diffusion}/bench_norm_impls.py | 0 .../benchmark => benchmark/diffusion}/bench_qknorm_rope.py | 0 .../diffusion}/bench_qwen_image_modulation.py | 0 .../diffusion}/bench_residual_gate_add.py | 0 .../diffusion}/diffusion_nvfp4_shapes.json | 0 .../embeddings}/bench_vocab_parallel_embedding.py | 0 .../benchmark => benchmark/gemm}/bench_dsv3_fused_a_gemm.py | 0 .../benchmark => benchmark/gemm}/bench_dsv3_router_gemm.py | 0 .../benchmark => benchmark/gemm}/bench_fp8_blockwise_gemm.py | 0 .../kv_canary/benchmark => benchmark/kv_canary}/bench_plan.py | 0 .../kv_canary}/bench_scatter_req_token_ids.py | 0 .../benchmark => benchmark/kv_canary}/bench_verify.py | 0 .../benchmark => benchmark/kv_canary}/bench_write.py | 0 .../kvcache}/bench_fused_fp8_qkv_kv_cache.py | 0 .../kvcache/benchmark => benchmark/kvcache}/bench_hicache.py | 0 .../kvcache/benchmark => benchmark/kvcache}/bench_hisparse.py | 0 .../kvcache}/bench_minimax_store_kv_index.py | 0 .../kvcache}/bench_set_mla_kv_buffer.py | 0 .../benchmark => benchmark/kvcache}/bench_store_cache.py | 0 .../benchmark => benchmark/layernorm}/bench_fused_eh_norm.py | 0 .../layernorm/benchmark => benchmark/layernorm}/bench_norm.py | 0 .../benchmark => benchmark/layernorm}/bench_qknorm.py | 0 .../layernorm}/bench_qknorm_across_heads.py | 0 .../moe/benchmark => benchmark/moe}/bench_moe_fused_gate.py | 0 .../moe}/bench_post_reorder_deepgemm.py | 0 .../{ops/moe/benchmark => benchmark/moe}/bench_renorm.py | 0 .../quantization}/bench_awq_dequantize.py | 0 .../benchmark => benchmark/quantization}/bench_mxfp8_moe.py | 0 .../quantization}/bench_per_tensor_quant_fp8.py | 0 .../quantization}/bench_per_token_group_quant.py | 0 .../quantization}/bench_per_token_group_quant_8bit_v2.py | 0 .../quantization}/bench_per_token_group_quant_masked.py | 0 .../speculative}/bench_ngram_compute_decode.py | 0 .../speculative}/bench_ngram_update_token_table.py | 0 .../speculative}/bench_resolve_future_token_ids.py | 0 .../benchmark => benchmark/speculative}/bench_spec_topk1.py | 0 61 files changed, 6 insertions(+), 6 deletions(-) rename test/registered/kernels/{ops/activation/benchmark => benchmark/activation}/bench_activation.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_add_constant.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_clamp_position.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_concat_mla.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_dsv4_fp4_indexer.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_fused_qknorm_rope.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_hadamard.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_minimax_decode_topk.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_minimax_qknorm_rope.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_mla_kv_pack_quantize_fp8.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_online_c128_mtp.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_rope.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_sparse_mla_q8kv8_prefill_sm90.py (100%) rename test/registered/kernels/{ops/attention/benchmark => benchmark/attention}/bench_topk.py (100%) rename test/registered/kernels/{ops/communication/benchmark => benchmark/communication}/bench_custom_all_reduce.py (100%) rename test/registered/kernels/{ops/communication/benchmark => benchmark/communication}/bench_symm_mem_all_gather.py (100%) rename test/registered/kernels/{ops/communication/benchmark => benchmark/communication}/bench_tp_qknorm.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_causal_conv3d_cat_pad.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_diffusion_nvfp4_scaled_mm.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_fused_norm_scale_shift.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_group_norm_silu.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_ltx2_qknorm_split_rope.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_norm_impls.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_qknorm_rope.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_qwen_image_modulation.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/bench_residual_gate_add.py (100%) rename test/registered/kernels/{ops/diffusion/benchmark => benchmark/diffusion}/diffusion_nvfp4_shapes.json (100%) rename test/registered/kernels/{ops/embeddings/benchmark => benchmark/embeddings}/bench_vocab_parallel_embedding.py (100%) rename test/registered/kernels/{ops/gemm/benchmark => benchmark/gemm}/bench_dsv3_fused_a_gemm.py (100%) rename test/registered/kernels/{ops/gemm/benchmark => benchmark/gemm}/bench_dsv3_router_gemm.py (100%) rename test/registered/kernels/{ops/gemm/benchmark => benchmark/gemm}/bench_fp8_blockwise_gemm.py (100%) rename test/registered/kernels/{ops/kv_canary/benchmark => benchmark/kv_canary}/bench_plan.py (100%) rename test/registered/kernels/{ops/kv_canary/benchmark => benchmark/kv_canary}/bench_scatter_req_token_ids.py (100%) rename test/registered/kernels/{ops/kv_canary/benchmark => benchmark/kv_canary}/bench_verify.py (100%) rename test/registered/kernels/{ops/kv_canary/benchmark => benchmark/kv_canary}/bench_write.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_fused_fp8_qkv_kv_cache.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_hicache.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_hisparse.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_minimax_store_kv_index.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_set_mla_kv_buffer.py (100%) rename test/registered/kernels/{ops/kvcache/benchmark => benchmark/kvcache}/bench_store_cache.py (100%) rename test/registered/kernels/{ops/layernorm/benchmark => benchmark/layernorm}/bench_fused_eh_norm.py (100%) rename test/registered/kernels/{ops/layernorm/benchmark => benchmark/layernorm}/bench_norm.py (100%) rename test/registered/kernels/{ops/layernorm/benchmark => benchmark/layernorm}/bench_qknorm.py (100%) rename test/registered/kernels/{ops/layernorm/benchmark => benchmark/layernorm}/bench_qknorm_across_heads.py (100%) rename test/registered/kernels/{ops/moe/benchmark => benchmark/moe}/bench_moe_fused_gate.py (100%) rename test/registered/kernels/{ops/moe/benchmark => benchmark/moe}/bench_post_reorder_deepgemm.py (100%) rename test/registered/kernels/{ops/moe/benchmark => benchmark/moe}/bench_renorm.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_awq_dequantize.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_mxfp8_moe.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_per_tensor_quant_fp8.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_per_token_group_quant.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_per_token_group_quant_8bit_v2.py (100%) rename test/registered/kernels/{ops/quantization/benchmark => benchmark/quantization}/bench_per_token_group_quant_masked.py (100%) rename test/registered/kernels/{ops/speculative/benchmark => benchmark/speculative}/bench_ngram_compute_decode.py (100%) rename test/registered/kernels/{ops/speculative/benchmark => benchmark/speculative}/bench_ngram_update_token_table.py (100%) rename test/registered/kernels/{ops/speculative/benchmark => benchmark/speculative}/bench_resolve_future_token_ids.py (100%) rename test/registered/kernels/{ops/speculative/benchmark => benchmark/speculative}/bench_spec_topk1.py (100%) diff --git a/.github/workflows/_pr-test-check-changes.yml b/.github/workflows/_pr-test-check-changes.yml index b99744a223df..ae3b060dbd04 100644 --- a/.github/workflows/_pr-test-check-changes.yml +++ b/.github/workflows/_pr-test-check-changes.yml @@ -95,13 +95,13 @@ jobs: - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - "test/registered/kernels/ops/diffusion/**" - - "test/registered/kernels/ops/diffusion/benchmark/**" + - "test/registered/kernels/benchmark/diffusion/**" - "python/sglang/cli/**" jit_kernel: - ".github/workflows/pr-test.yml" - ".github/workflows/pr-test-jit-kernel.yml" - "python/pyproject.toml" - - "test/registered/kernels/ops/**" + - "test/registered/kernels/**" # sglang.kernels is the migrated kernel namespace (RFC #29630 / #30044); the # base-b-kernel suites import it directly, so kernel edits must run them. - "python/sglang/kernels/**" diff --git a/.github/workflows/pr-test-amd-rocm720.yml b/.github/workflows/pr-test-amd-rocm720.yml index 8a4ddfdd5409..afdf9662f1d9 100644 --- a/.github/workflows/pr-test-amd-rocm720.yml +++ b/.github/workflows/pr-test-amd-rocm720.yml @@ -199,7 +199,7 @@ jobs: - ".github/workflows/pr-test-amd-rocm720.yml" jit_kernel: - "python/sglang/kernels/**" - - "test/registered/kernels/ops/**" + - "test/registered/kernels/**" - ".github/workflows/pr-test-amd-rocm720.yml" multimodal_gen: - "python/sglang/multimodal_gen/**/!(*.md|*.ipynb)" @@ -207,7 +207,7 @@ jobs: - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - "test/registered/kernels/ops/diffusion/**" - - "test/registered/kernels/ops/diffusion/benchmark/**" + - "test/registered/kernels/benchmark/diffusion/**" - "python/pyproject_rocm.toml" - "python/pyproject_other.toml" diff --git a/.github/workflows/pr-test-amd.yml b/.github/workflows/pr-test-amd.yml index 3300d76d910c..566200d8ed56 100644 --- a/.github/workflows/pr-test-amd.yml +++ b/.github/workflows/pr-test-amd.yml @@ -187,7 +187,7 @@ jobs: - ".github/workflows/pr-test-amd.yml" jit_kernel: - "python/sglang/kernels/**" - - "test/registered/kernels/ops/**" + - "test/registered/kernels/**" - ".github/workflows/pr-test-amd.yml" multimodal_gen: - "python/sglang/multimodal_gen/**/!(*.md|*.ipynb)" @@ -195,7 +195,7 @@ jobs: - "python/sglang/srt/observability/**" - "python/sglang/kernels/ops/diffusion/**" - "test/registered/kernels/ops/diffusion/**" - - "test/registered/kernels/ops/diffusion/benchmark/**" + - "test/registered/kernels/benchmark/diffusion/**" - "python/pyproject_rocm.toml" - "python/pyproject_other.toml" diff --git a/test/registered/kernels/ops/activation/benchmark/bench_activation.py b/test/registered/kernels/benchmark/activation/bench_activation.py similarity index 100% rename from test/registered/kernels/ops/activation/benchmark/bench_activation.py rename to test/registered/kernels/benchmark/activation/bench_activation.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_add_constant.py b/test/registered/kernels/benchmark/attention/bench_add_constant.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_add_constant.py rename to test/registered/kernels/benchmark/attention/bench_add_constant.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_clamp_position.py b/test/registered/kernels/benchmark/attention/bench_clamp_position.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_clamp_position.py rename to test/registered/kernels/benchmark/attention/bench_clamp_position.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_concat_mla.py b/test/registered/kernels/benchmark/attention/bench_concat_mla.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_concat_mla.py rename to test/registered/kernels/benchmark/attention/bench_concat_mla.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_dsv4_fp4_indexer.py b/test/registered/kernels/benchmark/attention/bench_dsv4_fp4_indexer.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_dsv4_fp4_indexer.py rename to test/registered/kernels/benchmark/attention/bench_dsv4_fp4_indexer.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_fused_qknorm_rope.py b/test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_fused_qknorm_rope.py rename to test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_hadamard.py b/test/registered/kernels/benchmark/attention/bench_hadamard.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_hadamard.py rename to test/registered/kernels/benchmark/attention/bench_hadamard.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_minimax_decode_topk.py b/test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_minimax_decode_topk.py rename to test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_minimax_qknorm_rope.py b/test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_minimax_qknorm_rope.py rename to test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_mla_kv_pack_quantize_fp8.py b/test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_mla_kv_pack_quantize_fp8.py rename to test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_online_c128_mtp.py b/test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_online_c128_mtp.py rename to test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_rope.py b/test/registered/kernels/benchmark/attention/bench_rope.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_rope.py rename to test/registered/kernels/benchmark/attention/bench_rope.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py b/test/registered/kernels/benchmark/attention/bench_sparse_mla_q8kv8_prefill_sm90.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_sparse_mla_q8kv8_prefill_sm90.py rename to test/registered/kernels/benchmark/attention/bench_sparse_mla_q8kv8_prefill_sm90.py diff --git a/test/registered/kernels/ops/attention/benchmark/bench_topk.py b/test/registered/kernels/benchmark/attention/bench_topk.py similarity index 100% rename from test/registered/kernels/ops/attention/benchmark/bench_topk.py rename to test/registered/kernels/benchmark/attention/bench_topk.py diff --git a/test/registered/kernels/ops/communication/benchmark/bench_custom_all_reduce.py b/test/registered/kernels/benchmark/communication/bench_custom_all_reduce.py similarity index 100% rename from test/registered/kernels/ops/communication/benchmark/bench_custom_all_reduce.py rename to test/registered/kernels/benchmark/communication/bench_custom_all_reduce.py diff --git a/test/registered/kernels/ops/communication/benchmark/bench_symm_mem_all_gather.py b/test/registered/kernels/benchmark/communication/bench_symm_mem_all_gather.py similarity index 100% rename from test/registered/kernels/ops/communication/benchmark/bench_symm_mem_all_gather.py rename to test/registered/kernels/benchmark/communication/bench_symm_mem_all_gather.py diff --git a/test/registered/kernels/ops/communication/benchmark/bench_tp_qknorm.py b/test/registered/kernels/benchmark/communication/bench_tp_qknorm.py similarity index 100% rename from test/registered/kernels/ops/communication/benchmark/bench_tp_qknorm.py rename to test/registered/kernels/benchmark/communication/bench_tp_qknorm.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_causal_conv3d_cat_pad.py b/test/registered/kernels/benchmark/diffusion/bench_causal_conv3d_cat_pad.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_causal_conv3d_cat_pad.py rename to test/registered/kernels/benchmark/diffusion/bench_causal_conv3d_cat_pad.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_diffusion_nvfp4_scaled_mm.py b/test/registered/kernels/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_diffusion_nvfp4_scaled_mm.py rename to test/registered/kernels/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_fused_norm_scale_shift.py b/test/registered/kernels/benchmark/diffusion/bench_fused_norm_scale_shift.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_fused_norm_scale_shift.py rename to test/registered/kernels/benchmark/diffusion/bench_fused_norm_scale_shift.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_group_norm_silu.py b/test/registered/kernels/benchmark/diffusion/bench_group_norm_silu.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_group_norm_silu.py rename to test/registered/kernels/benchmark/diffusion/bench_group_norm_silu.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_ltx2_qknorm_split_rope.py b/test/registered/kernels/benchmark/diffusion/bench_ltx2_qknorm_split_rope.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_ltx2_qknorm_split_rope.py rename to test/registered/kernels/benchmark/diffusion/bench_ltx2_qknorm_split_rope.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_norm_impls.py b/test/registered/kernels/benchmark/diffusion/bench_norm_impls.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_norm_impls.py rename to test/registered/kernels/benchmark/diffusion/bench_norm_impls.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_qknorm_rope.py b/test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_qknorm_rope.py rename to test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_qwen_image_modulation.py b/test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_qwen_image_modulation.py rename to test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/bench_residual_gate_add.py b/test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/bench_residual_gate_add.py rename to test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py diff --git a/test/registered/kernels/ops/diffusion/benchmark/diffusion_nvfp4_shapes.json b/test/registered/kernels/benchmark/diffusion/diffusion_nvfp4_shapes.json similarity index 100% rename from test/registered/kernels/ops/diffusion/benchmark/diffusion_nvfp4_shapes.json rename to test/registered/kernels/benchmark/diffusion/diffusion_nvfp4_shapes.json diff --git a/test/registered/kernels/ops/embeddings/benchmark/bench_vocab_parallel_embedding.py b/test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py similarity index 100% rename from test/registered/kernels/ops/embeddings/benchmark/bench_vocab_parallel_embedding.py rename to test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py diff --git a/test/registered/kernels/ops/gemm/benchmark/bench_dsv3_fused_a_gemm.py b/test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py similarity index 100% rename from test/registered/kernels/ops/gemm/benchmark/bench_dsv3_fused_a_gemm.py rename to test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py diff --git a/test/registered/kernels/ops/gemm/benchmark/bench_dsv3_router_gemm.py b/test/registered/kernels/benchmark/gemm/bench_dsv3_router_gemm.py similarity index 100% rename from test/registered/kernels/ops/gemm/benchmark/bench_dsv3_router_gemm.py rename to test/registered/kernels/benchmark/gemm/bench_dsv3_router_gemm.py diff --git a/test/registered/kernels/ops/gemm/benchmark/bench_fp8_blockwise_gemm.py b/test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py similarity index 100% rename from test/registered/kernels/ops/gemm/benchmark/bench_fp8_blockwise_gemm.py rename to test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py diff --git a/test/registered/kernels/ops/kv_canary/benchmark/bench_plan.py b/test/registered/kernels/benchmark/kv_canary/bench_plan.py similarity index 100% rename from test/registered/kernels/ops/kv_canary/benchmark/bench_plan.py rename to test/registered/kernels/benchmark/kv_canary/bench_plan.py diff --git a/test/registered/kernels/ops/kv_canary/benchmark/bench_scatter_req_token_ids.py b/test/registered/kernels/benchmark/kv_canary/bench_scatter_req_token_ids.py similarity index 100% rename from test/registered/kernels/ops/kv_canary/benchmark/bench_scatter_req_token_ids.py rename to test/registered/kernels/benchmark/kv_canary/bench_scatter_req_token_ids.py diff --git a/test/registered/kernels/ops/kv_canary/benchmark/bench_verify.py b/test/registered/kernels/benchmark/kv_canary/bench_verify.py similarity index 100% rename from test/registered/kernels/ops/kv_canary/benchmark/bench_verify.py rename to test/registered/kernels/benchmark/kv_canary/bench_verify.py diff --git a/test/registered/kernels/ops/kv_canary/benchmark/bench_write.py b/test/registered/kernels/benchmark/kv_canary/bench_write.py similarity index 100% rename from test/registered/kernels/ops/kv_canary/benchmark/bench_write.py rename to test/registered/kernels/benchmark/kv_canary/bench_write.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_fused_fp8_qkv_kv_cache.py b/test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_fused_fp8_qkv_kv_cache.py rename to test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_hicache.py b/test/registered/kernels/benchmark/kvcache/bench_hicache.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_hicache.py rename to test/registered/kernels/benchmark/kvcache/bench_hicache.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_hisparse.py b/test/registered/kernels/benchmark/kvcache/bench_hisparse.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_hisparse.py rename to test/registered/kernels/benchmark/kvcache/bench_hisparse.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_minimax_store_kv_index.py b/test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_minimax_store_kv_index.py rename to test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_set_mla_kv_buffer.py b/test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_set_mla_kv_buffer.py rename to test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py diff --git a/test/registered/kernels/ops/kvcache/benchmark/bench_store_cache.py b/test/registered/kernels/benchmark/kvcache/bench_store_cache.py similarity index 100% rename from test/registered/kernels/ops/kvcache/benchmark/bench_store_cache.py rename to test/registered/kernels/benchmark/kvcache/bench_store_cache.py diff --git a/test/registered/kernels/ops/layernorm/benchmark/bench_fused_eh_norm.py b/test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py similarity index 100% rename from test/registered/kernels/ops/layernorm/benchmark/bench_fused_eh_norm.py rename to test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py diff --git a/test/registered/kernels/ops/layernorm/benchmark/bench_norm.py b/test/registered/kernels/benchmark/layernorm/bench_norm.py similarity index 100% rename from test/registered/kernels/ops/layernorm/benchmark/bench_norm.py rename to test/registered/kernels/benchmark/layernorm/bench_norm.py diff --git a/test/registered/kernels/ops/layernorm/benchmark/bench_qknorm.py b/test/registered/kernels/benchmark/layernorm/bench_qknorm.py similarity index 100% rename from test/registered/kernels/ops/layernorm/benchmark/bench_qknorm.py rename to test/registered/kernels/benchmark/layernorm/bench_qknorm.py diff --git a/test/registered/kernels/ops/layernorm/benchmark/bench_qknorm_across_heads.py b/test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py similarity index 100% rename from test/registered/kernels/ops/layernorm/benchmark/bench_qknorm_across_heads.py rename to test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py diff --git a/test/registered/kernels/ops/moe/benchmark/bench_moe_fused_gate.py b/test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py similarity index 100% rename from test/registered/kernels/ops/moe/benchmark/bench_moe_fused_gate.py rename to test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py diff --git a/test/registered/kernels/ops/moe/benchmark/bench_post_reorder_deepgemm.py b/test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py similarity index 100% rename from test/registered/kernels/ops/moe/benchmark/bench_post_reorder_deepgemm.py rename to test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py diff --git a/test/registered/kernels/ops/moe/benchmark/bench_renorm.py b/test/registered/kernels/benchmark/moe/bench_renorm.py similarity index 100% rename from test/registered/kernels/ops/moe/benchmark/bench_renorm.py rename to test/registered/kernels/benchmark/moe/bench_renorm.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_awq_dequantize.py b/test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_awq_dequantize.py rename to test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_mxfp8_moe.py b/test/registered/kernels/benchmark/quantization/bench_mxfp8_moe.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_mxfp8_moe.py rename to test/registered/kernels/benchmark/quantization/bench_mxfp8_moe.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_per_tensor_quant_fp8.py b/test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_per_tensor_quant_fp8.py rename to test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant.py rename to test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_8bit_v2.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_8bit_v2.py rename to test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py diff --git a/test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_masked.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py similarity index 100% rename from test/registered/kernels/ops/quantization/benchmark/bench_per_token_group_quant_masked.py rename to test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py diff --git a/test/registered/kernels/ops/speculative/benchmark/bench_ngram_compute_decode.py b/test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py similarity index 100% rename from test/registered/kernels/ops/speculative/benchmark/bench_ngram_compute_decode.py rename to test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py diff --git a/test/registered/kernels/ops/speculative/benchmark/bench_ngram_update_token_table.py b/test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py similarity index 100% rename from test/registered/kernels/ops/speculative/benchmark/bench_ngram_update_token_table.py rename to test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py diff --git a/test/registered/kernels/ops/speculative/benchmark/bench_resolve_future_token_ids.py b/test/registered/kernels/benchmark/speculative/bench_resolve_future_token_ids.py similarity index 100% rename from test/registered/kernels/ops/speculative/benchmark/bench_resolve_future_token_ids.py rename to test/registered/kernels/benchmark/speculative/bench_resolve_future_token_ids.py diff --git a/test/registered/kernels/ops/speculative/benchmark/bench_spec_topk1.py b/test/registered/kernels/benchmark/speculative/bench_spec_topk1.py similarity index 100% rename from test/registered/kernels/ops/speculative/benchmark/bench_spec_topk1.py rename to test/registered/kernels/benchmark/speculative/bench_spec_topk1.py From bf900a9fd1b059ba7052a70ff9d0c2484cf81e56 Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 09:51:37 +0800 Subject: [PATCH 4/6] [Kernel] Move kernel test helpers to sglang.test.kernels (out of the runtime kernels package) The shared test-support helpers do not belong under the runtime `kernels/` package. Move python/sglang/kernels/testing/ -> python/sglang/test/kernels/ (utils, kv_canary/, deepseek_v4/), the established in-package test-support namespace (alongside test_utils, ci_register, runners, mock_model, ...), and rewrite sglang.kernels.testing -> sglang.test.kernels imports. Kept separate from the existing srt-level sglang/test/kv_canary/ (e2e fixtures): these are the kernel-level fuzz/oracle helpers, so they live under sglang/test/kernels/kv_canary/ mirroring the kernel test tree test/registered/kernels/. Still an importable installed package (required: test/registered/ tests import it by absolute path). Co-Authored-By: Claude Opus 4.8 --- python/sglang/kernels/ops/communication/mp.py | 2 +- .../{kernels/testing => test/kernels}/__init__.py | 0 .../testing => test/kernels}/deepseek_v4/__init__.py | 0 .../testing => test/kernels}/deepseek_v4/common.py | 0 .../testing => test/kernels}/kv_canary/__init__.py | 0 .../kernels}/kv_canary/_canary_helpers.py | 4 ++-- .../testing => test/kernels}/kv_canary/_constants.py | 0 .../kernels}/kv_canary/_differential.py | 2 +- .../testing => test/kernels}/kv_canary/_fixtures.py | 2 +- .../kernels}/kv_canary/_fuzz_driver.py | 2 +- .../kernels}/kv_canary/_hand_oracle.py | 0 .../kernels}/kv_canary/_invariants.py | 2 +- .../{kernels/testing => test/kernels}/utils.py | 0 .../registered/kernels/ops/attention/test_c128_v2.py | 6 +++--- test/registered/kernels/ops/attention/test_c4_v2.py | 6 +++--- ...test_deepseek_v4_compress_state_runtime_shapes.py | 4 ++-- .../ops/communication/test_custom_all_reduce.py | 2 +- .../ops/communication/test_symm_mem_all_gather.py | 2 +- .../kernels/ops/communication/test_tp_qknorm.py | 2 +- .../kernels/ops/kv_canary/test_kernel_config.py | 8 ++++---- .../kernels/ops/kv_canary/test_pipeline_e2e.py | 6 +++--- .../kernels/ops/kv_canary/test_plan_fuzz.py | 10 +++++----- .../kernels/ops/kv_canary/test_plan_hand.py | 8 ++++---- .../kernels/ops/kv_canary/test_verify_fuzz.py | 12 ++++++------ .../kernels/ops/kv_canary/test_verify_hand.py | 10 +++++----- .../kernels/ops/kv_canary/test_write_fuzz.py | 12 ++++++------ .../kernels/ops/kv_canary/test_write_hand.py | 10 +++++----- test/registered/unit/distributed/test_vmm_utils.py | 2 +- 28 files changed, 57 insertions(+), 57 deletions(-) rename python/sglang/{kernels/testing => test/kernels}/__init__.py (100%) rename python/sglang/{kernels/testing => test/kernels}/deepseek_v4/__init__.py (100%) rename python/sglang/{kernels/testing => test/kernels}/deepseek_v4/common.py (100%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/__init__.py (100%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_canary_helpers.py (98%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_constants.py (100%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_differential.py (99%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_fixtures.py (98%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_fuzz_driver.py (95%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_hand_oracle.py (100%) rename python/sglang/{kernels/testing => test/kernels}/kv_canary/_invariants.py (99%) rename python/sglang/{kernels/testing => test/kernels}/utils.py (100%) diff --git a/python/sglang/kernels/ops/communication/mp.py b/python/sglang/kernels/ops/communication/mp.py index 7b6361e63cd4..957ac0647c91 100644 --- a/python/sglang/kernels/ops/communication/mp.py +++ b/python/sglang/kernels/ops/communication/mp.py @@ -1,6 +1,6 @@ """Multi-process / multi-GPU launching utilities (torchrun-based). -Shared `multigpu_launch` helper that both `sglang.kernels.testing.utils` and +Shared `multigpu_launch` helper that both `sglang.test.kernels.utils` and `sglang.kernels.jit.benchmark.utils` build their domain-specific entry points on top of (`multigpu_pytest_main`, `multigpu_bench_main`). diff --git a/python/sglang/kernels/testing/__init__.py b/python/sglang/test/kernels/__init__.py similarity index 100% rename from python/sglang/kernels/testing/__init__.py rename to python/sglang/test/kernels/__init__.py diff --git a/python/sglang/kernels/testing/deepseek_v4/__init__.py b/python/sglang/test/kernels/deepseek_v4/__init__.py similarity index 100% rename from python/sglang/kernels/testing/deepseek_v4/__init__.py rename to python/sglang/test/kernels/deepseek_v4/__init__.py diff --git a/python/sglang/kernels/testing/deepseek_v4/common.py b/python/sglang/test/kernels/deepseek_v4/common.py similarity index 100% rename from python/sglang/kernels/testing/deepseek_v4/common.py rename to python/sglang/test/kernels/deepseek_v4/common.py diff --git a/python/sglang/kernels/testing/kv_canary/__init__.py b/python/sglang/test/kernels/kv_canary/__init__.py similarity index 100% rename from python/sglang/kernels/testing/kv_canary/__init__.py rename to python/sglang/test/kernels/kv_canary/__init__.py diff --git a/python/sglang/kernels/testing/kv_canary/_canary_helpers.py b/python/sglang/test/kernels/kv_canary/_canary_helpers.py similarity index 98% rename from python/sglang/kernels/testing/kv_canary/_canary_helpers.py rename to python/sglang/test/kernels/kv_canary/_canary_helpers.py index fbbae1c58e8a..99bacfb51adb 100644 --- a/python/sglang/kernels/testing/kv_canary/_canary_helpers.py +++ b/python/sglang/test/kernels/kv_canary/_canary_helpers.py @@ -9,14 +9,14 @@ from sglang.kernels.ops.kv_canary.consts import splitmix64, splitmix64_mix3 from sglang.kernels.ops.kv_canary.verify import VerifyPlan from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._constants import ( +from sglang.test.kernels.kv_canary._constants import ( _I64_SIGN_BIT, _U64_MASK, DEFAULT_NUM_SLOTS, DEFAULT_RING_CAPACITY, DEFAULT_SLOT_STRIDE_BYTES, ) -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._fixtures import ( make_real_kv_source, make_real_kv_sources, ) diff --git a/python/sglang/kernels/testing/kv_canary/_constants.py b/python/sglang/test/kernels/kv_canary/_constants.py similarity index 100% rename from python/sglang/kernels/testing/kv_canary/_constants.py rename to python/sglang/test/kernels/kv_canary/_constants.py diff --git a/python/sglang/kernels/testing/kv_canary/_differential.py b/python/sglang/test/kernels/kv_canary/_differential.py similarity index 99% rename from python/sglang/kernels/testing/kv_canary/_differential.py rename to python/sglang/test/kernels/kv_canary/_differential.py index 97d08abaa3b2..b122a2e1213c 100644 --- a/python/sglang/kernels/testing/kv_canary/_differential.py +++ b/python/sglang/test/kernels/kv_canary/_differential.py @@ -24,7 +24,7 @@ from sglang.kernels.ops.kv_canary.write_ref import ( launch_canary_write_kernel_torch_reference, ) -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_buf_equal, assert_canary_state_equal, diff --git a/python/sglang/kernels/testing/kv_canary/_fixtures.py b/python/sglang/test/kernels/kv_canary/_fixtures.py similarity index 98% rename from python/sglang/kernels/testing/kv_canary/_fixtures.py rename to python/sglang/test/kernels/kv_canary/_fixtures.py index 167a0bd297f0..c3cd5bdb15e4 100644 --- a/python/sglang/kernels/testing/kv_canary/_fixtures.py +++ b/python/sglang/test/kernels/kv_canary/_fixtures.py @@ -10,7 +10,7 @@ VerifyPlan, ) from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._constants import DEFAULT_NUM_SLOTS +from sglang.test.kernels.kv_canary._constants import DEFAULT_NUM_SLOTS _DEVICE = torch.device("cuda") diff --git a/python/sglang/kernels/testing/kv_canary/_fuzz_driver.py b/python/sglang/test/kernels/kv_canary/_fuzz_driver.py similarity index 95% rename from python/sglang/kernels/testing/kv_canary/_fuzz_driver.py rename to python/sglang/test/kernels/kv_canary/_fuzz_driver.py index 7ba38f0c3352..566aae086536 100644 --- a/python/sglang/kernels/testing/kv_canary/_fuzz_driver.py +++ b/python/sglang/test/kernels/kv_canary/_fuzz_driver.py @@ -3,7 +3,7 @@ import random from typing import Any, Callable -from sglang.kernels.testing.kv_canary._differential import ( +from sglang.test.kernels.kv_canary._differential import ( ShrinkResult, shrink_inputs, ) diff --git a/python/sglang/kernels/testing/kv_canary/_hand_oracle.py b/python/sglang/test/kernels/kv_canary/_hand_oracle.py similarity index 100% rename from python/sglang/kernels/testing/kv_canary/_hand_oracle.py rename to python/sglang/test/kernels/kv_canary/_hand_oracle.py diff --git a/python/sglang/kernels/testing/kv_canary/_invariants.py b/python/sglang/test/kernels/kv_canary/_invariants.py similarity index 99% rename from python/sglang/kernels/testing/kv_canary/_invariants.py rename to python/sglang/test/kernels/kv_canary/_invariants.py index 79cdcf27f2a0..588a242baa75 100644 --- a/python/sglang/kernels/testing/kv_canary/_invariants.py +++ b/python/sglang/test/kernels/kv_canary/_invariants.py @@ -14,7 +14,7 @@ from sglang.kernels.ops.kv_canary import consts from sglang.kernels.ops.kv_canary.verify import CanaryLaunchTag, VerifyPlan from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._canary_helpers import FakeViolationLog +from sglang.test.kernels.kv_canary._canary_helpers import FakeViolationLog class PlanInvariants: diff --git a/python/sglang/kernels/testing/utils.py b/python/sglang/test/kernels/utils.py similarity index 100% rename from python/sglang/kernels/testing/utils.py rename to python/sglang/test/kernels/utils.py diff --git a/test/registered/kernels/ops/attention/test_c128_v2.py b/test/registered/kernels/ops/attention/test_c128_v2.py index 5690402ab08e..a5b0e74fab94 100644 --- a/test/registered/kernels/ops/attention/test_c128_v2.py +++ b/test/registered/kernels/ops/attention/test_c128_v2.py @@ -8,7 +8,9 @@ import triton from sglang.kernels.ops.attention.dsv4 import compress_forward -from sglang.kernels.testing.deepseek_v4.common import ( +from sglang.srt.utils import get_device +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.deepseek_v4.common import ( LegacyContext, PagedContext, make_legacy_context, @@ -16,8 +18,6 @@ make_state_pool, to_seq_extend, ) -from sglang.srt.utils import get_device -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/kernels/ops/attention/test_c4_v2.py b/test/registered/kernels/ops/attention/test_c4_v2.py index ae28242925e7..f05aaa9039ad 100644 --- a/test/registered/kernels/ops/attention/test_c4_v2.py +++ b/test/registered/kernels/ops/attention/test_c4_v2.py @@ -8,7 +8,9 @@ import triton from sglang.kernels.ops.attention.dsv4 import compress_forward -from sglang.kernels.testing.deepseek_v4.common import ( +from sglang.srt.utils import get_device +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.deepseek_v4.common import ( LegacyContext, PagedContext, make_legacy_context, @@ -16,8 +18,6 @@ make_state_pool, to_seq_extend, ) -from sglang.srt.utils import get_device -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py index f2a279f0cd89..9a2f392a00b6 100644 --- a/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py +++ b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_state_runtime_shapes.py @@ -94,11 +94,11 @@ CompressorPrefillPlan, compress_forward, ) -from sglang.kernels.testing.deepseek_v4.common import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.deepseek_v4.common import ( make_legacy_context, to_seq_extend, ) -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.utils import is_in_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/communication/test_custom_all_reduce.py b/test/registered/kernels/ops/communication/test_custom_all_reduce.py index d609f0ecfbdf..894b163344e3 100644 --- a/test/registered/kernels/ops/communication/test_custom_all_reduce.py +++ b/test/registered/kernels/ops/communication/test_custom_all_reduce.py @@ -34,11 +34,11 @@ get_all_reduce_module, ) from sglang.kernels.ops.communication.mp import register_comm_cleanup -from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.custom_all_reduce_v2 import ( CustomAllReduceV2, ) from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.kernels.utils import multigpu_pytest_main register_cuda_ci( est_time=300, diff --git a/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py b/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py index f38235ef1ee0..bc73dde3d1e7 100644 --- a/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py +++ b/test/registered/kernels/ops/communication/test_symm_mem_all_gather.py @@ -27,12 +27,12 @@ import sglang.srt.distributed.parallel_state as ps from sglang.kernels.jit.utils import cache_once, get_ci_test_range -from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.triton_symm_mem_ag import ( all_gather_inner, create_state, ) from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.kernels.utils import multigpu_pytest_main register_cuda_ci(est_time=240, stage="extra-b", runner_config="8-gpu-h200") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. diff --git a/test/registered/kernels/ops/communication/test_tp_qknorm.py b/test/registered/kernels/ops/communication/test_tp_qknorm.py index c7895c7a2f03..712c24ecd077 100644 --- a/test/registered/kernels/ops/communication/test_tp_qknorm.py +++ b/test/registered/kernels/ops/communication/test_tp_qknorm.py @@ -21,11 +21,11 @@ get_fused_parallel_qknorm_module, ) from sglang.kernels.ops.communication.mp import register_comm_cleanup -from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.custom_all_reduce_v2 import ( CustomAllReduceV2, ) from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.kernels.utils import multigpu_pytest_main register_cuda_ci( est_time=300, diff --git a/test/registered/kernels/ops/kv_canary/test_kernel_config.py b/test/registered/kernels/ops/kv_canary/test_kernel_config.py index 32f7f61320f4..e3d7b6d6cb89 100644 --- a/test/registered/kernels/ops/kv_canary/test_kernel_config.py +++ b/test/registered/kernels/ops/kv_canary/test_kernel_config.py @@ -11,7 +11,8 @@ launch_canary_verify_kernel, ) from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_buf_equal, assert_canary_state_equal, @@ -23,17 +24,16 @@ make_write_plan_pair, stamp_clean_chain, ) -from sglang.kernels.testing.kv_canary._differential import ( +from sglang.test.kernels.kv_canary._differential import ( _assert_plans_byte_equal, _run_both_plan, _run_both_verify, _run_both_write, ) -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._fixtures import ( dummy_pseudo_tensors, empty_extras, ) -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=60, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py b/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py index 932b389c7236..975495fcdec4 100644 --- a/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py +++ b/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py @@ -24,7 +24,8 @@ from sglang.kernels.ops.kv_canary.write_ref import ( launch_canary_write_kernel_torch_reference, ) -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_buf_equal, assert_canary_state_equal, @@ -33,12 +34,11 @@ stamp_clean_chain, write_slot_fields, ) -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._fixtures import ( clone_real_kv_sources, empty_extras, make_req_to_token, ) -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py b/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py index 137506423d62..091f406fd2f6 100644 --- a/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py @@ -7,20 +7,20 @@ import pytest import torch -from sglang.kernels.testing.kv_canary._differential import _run_both_plan -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._differential import _run_both_plan +from sglang.test.kernels.kv_canary._fixtures import ( allocate_plan_pair, derive_plan_capacity, make_lut, make_padding_mask, make_req_to_token, ) -from sglang.kernels.testing.kv_canary._fuzz_driver import ( +from sglang.test.kernels.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.testing.kv_canary._invariants import PlanInvariants -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._invariants import PlanInvariants register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_plan_hand.py b/test/registered/kernels/ops/kv_canary/test_plan_hand.py index b175b2c3ce4d..487f26d80e92 100644 --- a/test/registered/kernels/ops/kv_canary/test_plan_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_hand.py @@ -11,16 +11,16 @@ ) from sglang.kernels.ops.kv_canary.verify import VerifyPlan from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._differential import run_plan_diff -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._differential import run_plan_diff +from sglang.test.kernels.kv_canary._fixtures import ( allocate_plan_pair, derive_plan_capacity, empty_extras, make_lut, make_req_to_token, ) -from sglang.kernels.testing.kv_canary._invariants import PlanInvariants -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._invariants import PlanInvariants register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py b/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py index ad8bf4921727..b1405a9079b3 100644 --- a/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py @@ -12,24 +12,24 @@ RealKvSource, VerifyPlan, ) -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, make_canary_buf, make_log_pair, make_verify_plan_pair, stamp_clean_chain, ) -from sglang.kernels.testing.kv_canary._differential import _run_both_verify -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._differential import _run_both_verify +from sglang.test.kernels.kv_canary._fixtures import ( clone_real_kv_sources, make_real_kv_sources, ) -from sglang.kernels.testing.kv_canary._fuzz_driver import ( +from sglang.test.kernels.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.testing.kv_canary._invariants import VerifyInvariants -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._invariants import VerifyInvariants register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_verify_hand.py b/test/registered/kernels/ops/kv_canary/test_verify_hand.py index de9fbe740ff2..266b7ae610ee 100644 --- a/test/registered/kernels/ops/kv_canary/test_verify_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_hand.py @@ -24,7 +24,8 @@ from sglang.kernels.ops.kv_canary.write_ref import ( launch_canary_write_kernel_torch_reference, ) -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, assert_only_bits_set, chain_anchor_signed, @@ -42,16 +43,15 @@ to_signed_int64, write_slot_fields, ) -from sglang.kernels.testing.kv_canary._differential import ( +from sglang.test.kernels.kv_canary._differential import ( _run_both_verify, run_verify_diff, ) -from sglang.kernels.testing.kv_canary._fixtures import clone_real_kv_sources -from sglang.kernels.testing.kv_canary._hand_oracle import ( +from sglang.test.kernels.kv_canary._fixtures import clone_real_kv_sources +from sglang.test.kernels.kv_canary._hand_oracle import ( _hand_fold_all, _hand_fold_partial, ) -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_write_fuzz.py b/test/registered/kernels/ops/kv_canary/test_write_fuzz.py index 18b4f12ab318..cd78d6a3552c 100644 --- a/test/registered/kernels/ops/kv_canary/test_write_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_write_fuzz.py @@ -12,24 +12,24 @@ RealKvSource, ) from sglang.kernels.ops.kv_canary.write import WritePlan -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, make_canary_buf, make_log_pair, make_write_plan_pair, stamp_pair, ) -from sglang.kernels.testing.kv_canary._differential import _run_both_write -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._differential import _run_both_write +from sglang.test.kernels.kv_canary._fixtures import ( clone_real_kv_sources, make_real_kv_sources, ) -from sglang.kernels.testing.kv_canary._fuzz_driver import ( +from sglang.test.kernels.kv_canary._fuzz_driver import ( FUZZ_SEEDS_PR, run_fuzz_combo, ) -from sglang.kernels.testing.kv_canary._invariants import WriteInvariants -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._invariants import WriteInvariants register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_write_hand.py b/test/registered/kernels/ops/kv_canary/test_write_hand.py index 933fee9d1550..065633aa1c14 100644 --- a/test/registered/kernels/ops/kv_canary/test_write_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_write_hand.py @@ -20,7 +20,8 @@ from sglang.kernels.ops.kv_canary.write import ( launch_canary_write_kernel, ) -from sglang.kernels.testing.kv_canary._canary_helpers import ( +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kernels.kv_canary._canary_helpers import ( FakeViolationLog, assert_canary_state_equal, assert_only_bits_set, @@ -37,19 +38,18 @@ stamp_pair, to_signed_int64, ) -from sglang.kernels.testing.kv_canary._differential import ( +from sglang.test.kernels.kv_canary._differential import ( _run_both_write, run_write_diff, ) -from sglang.kernels.testing.kv_canary._fixtures import ( +from sglang.test.kernels.kv_canary._fixtures import ( clone_real_kv_sources, dummy_pseudo_tensors, ) -from sglang.kernels.testing.kv_canary._hand_oracle import ( +from sglang.test.kernels.kv_canary._hand_oracle import ( _hand_fold_all, _hand_fold_partial, ) -from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/unit/distributed/test_vmm_utils.py b/test/registered/unit/distributed/test_vmm_utils.py index 55511e01d1a6..23178cbe7ae8 100644 --- a/test/registered/unit/distributed/test_vmm_utils.py +++ b/test/registered/unit/distributed/test_vmm_utils.py @@ -21,7 +21,6 @@ from cuda.bindings import driver as drv from sglang.kernels.jit.utils import cache_once -from sglang.kernels.testing.utils import multigpu_pytest_main from sglang.srt.distributed.device_communicators.vmm_utils import ( check_drv, exchange_posix_fds, @@ -32,6 +31,7 @@ release_mappings, ) from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.kernels.utils import multigpu_pytest_main register_cuda_ci(est_time=60, stage="base-b", runner_config="2-gpu-large") From 7f22c9a7f9fba301408e837ca249af014dd40889 Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 09:56:36 +0800 Subject: [PATCH 5/6] [Kernel] Update kernels/README for the final RFC #29630 layout - Add the jit/ subpackage (shared JIT build/runtime infra) to the Layout. - All 19 operator groups are populated; drop the stale "reserved placeholders / implementations not moved yet" wording (migration complete via #31666/#32015/ #32045/#32072). - Note sglang.jit_kernel removed; ops live at kernels.ops.._jit_; tests/benchmarks under test/registered/kernels/{ops,benchmark}//; shared test helpers in sglang.test.kernels. - Fix the fused-op parity test path (now under ops/layernorm/). Co-Authored-By: Claude Opus 4.8 --- python/sglang/kernels/README.md | 31 +++++++++++++++++++------------ 1 file changed, 19 insertions(+), 12 deletions(-) diff --git a/python/sglang/kernels/README.md b/python/sglang/kernels/README.md index 817e34f45949..7b69d06e7df7 100644 --- a/python/sglang/kernels/README.md +++ b/python/sglang/kernels/README.md @@ -19,21 +19,28 @@ sglang/kernels/ selector.py # heuristic select_kernel() and cached get_kernel() fused_op.py # BaseFusedOp: per-operator multi-backend contract ops/ - / # one subpackage per operator group + / # one subpackage per operator group (see list below) + jit/ # shared JIT CUDA build/runtime infra: utils/, csrc/, + # include/, __main__ (KERNEL_PATH resolves here) ``` -Groups populated in this phase: `activation`, `gemm`, `kvcache`, `layernorm`, -`moe`, `quantization`. The remaining groups (`attention`, `communication`, -`diffusion`, `grammar`, `mamba`, `memory`, `sampling`, `spatial`, -`speculative`) are reserved package placeholders whose implementations still -live in `sglang.kernels.jit` / `sgl_kernel` / `triton_ops` and will migrate in -later phases. +Operator groups (all populated): `activation`, `attention`, `communication`, +`diffusion`, `embeddings`, `gemm`, `grammar`, `kv_canary`, `kvcache`, +`layernorm`, `lplb`, `mamba`, `memory`, `model`, `moe`, `quantization`, +`sampling`, `spatial`, `speculative`. + +As of the RFC #29630 finale (#32072) the legacy `sglang.jit_kernel` package has +been **removed**: its shared build/runtime infra moved to `sglang.kernels.jit` +and each JIT-backed operator into its group as +`sglang.kernels.ops.._jit_`. Tests and benchmarks live under +`test/registered/kernels/` (`ops//` for tests, `benchmark//` for +benchmarks); shared test helpers are in `sglang.test.kernels`. ## How it works -Implementations are not moved yet. Each `ops.` function is a thin -wrapper that forwards to a chosen backend, and every backend is described by a -`KernelSpec` in the registry so alternatives can be inventoried and compared: +Each `ops.` function is a thin wrapper that forwards to a chosen +backend, and every backend is described by a `KernelSpec` in the registry so +alternatives can be inventoried and compared: - `register_kernel(KernelSpec(...))` records metadata only — an operator id (`"."`), a backend, and an import path (`"module:attr"`). No @@ -85,8 +92,8 @@ What this buys (see the - **Unified correctness testing** — a generic harness enumerates `available_backends()` and asserts each one matches `forward_native` - (`test/registered/kernels/test_fused_op_gpu_parity.py`); new backends are - picked up automatically. + (`test/registered/kernels/ops/layernorm/test_fused_op_gpu_parity.py`); new + backends are picked up automatically. - **One-switch debugging** — `SGLANG_FORCE_FUSED_OP_BACKEND=torch` (or `set_fused_op_backend(KernelBackend.TORCH)`) flips *every* fused op to its reference implementation for numerical-bug bisection. From a6dabc26c1e1b7f12a2394395365aabfdbdeb6b9 Mon Sep 17 00:00:00 2001 From: BBuf Date: Thu, 23 Jul 2026 10:04:20 +0800 Subject: [PATCH 6/6] [Kernel] Add missing __main__ entry to test_minimax_m3_rmsnorm The newly-registered test_minimax_m3_rmsnorm.py (register_amd_ci) lacked the `if __name__ == "__main__":` entry that run_suite.py's collection requires for registered pytest-style tests (otherwise `python3 file.py -f` silently skips). Surfaced by stage-a-unit-test-mlx collection. Add the standard `sys.exit(pytest.main([__file__, "-v"]))` block. Co-Authored-By: Claude Opus 4.8 --- .../kernels/ops/layernorm/test_minimax_m3_rmsnorm.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py index e8f5abb8c31e..dc46ffd909e0 100644 --- a/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py +++ b/test/registered/kernels/ops/layernorm/test_minimax_m3_rmsnorm.py @@ -79,3 +79,9 @@ def test_gemma_fused_add_rmsnorm_matches_reference(shape, dtype): torch.testing.assert_close(residual_out, ref_residual, atol=2e-2, rtol=2e-2) torch.testing.assert_close(got, ref, atol=2e-2, rtol=2e-2) + + +if __name__ == "__main__": + import sys + + sys.exit(pytest.main([__file__, "-v"]))