From 9a5b21aa1a477ce1412c9acca11f7f0e16ca652d Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 25 Sep 2026 06:23:31 +0000 Subject: [PATCH 01/10] [CI] Move 1-gpu-large tests that pass on the 5090 to 1-gpu-small Run every base-b, base-b-kernel-unit and extra-a file registered on 1-gpu-large (H100) on an RTX 5090 (SM120, 32 GB) with the CI install and environment, and move the ones that pass without losing coverage. A file with tests skipped on the 5090 only moves when those tests also skip on H100 (ROCm-, SM100- or missing-op-gated). - base-b-kernel-unit: 154 of 176 files move. The kernel lane now runs 3 shards on 1-gpu-small and 1 job on 1-gpu-large (was 1 and 2). - base-b: 84 files move; extra-a: 15 files move. Both are sized from est_time by compute_partitions.py (base-b: 5090 8 -> 13, H100 4 -> 2). - test_nvfp4_marlin was already registered on 1-gpu-small; drop its 1-gpu-large duplicate. - Raise est_time to the measured 5090 duration where the H100 estimate was more than 1.5x low (24 files). - Point new kernel correctness tests at 1-gpu-small in test/README.md and the write-sglang-test / add-jit-kernel skills; 1-gpu-large stays for SM90-only kernels. About 189 min of per-commit H100 est_time moves to the 5090 pool. Kept on 1-gpu-large: FA3 / FlashMLA / DeepGEMM / HPC-Ops / SM90-only kernels; weights over ~26 GB or two resident models; H100-tuned perf thresholds; int8_scaled_mm (no SM120 kernel); lmms_eval (not installed on Blackwell runners); and files that fail on SM120 (Triton shared- memory limits, FlashInfer top-k, dsv3_fused_a_gemm / tiny_gemm mismatches, EAGLE3 deterministic parity). --- .claude/skills/add-jit-kernel/SKILL.md | 11 ++++++----- .claude/skills/write-sglang-test/SKILL.md | 4 ++-- .github/workflows/pr-test-jit-kernel.yml | 13 +++++++++---- test/README.md | 4 ++-- .../attention/test_chunk_gated_delta_rule.py | 2 +- .../attention/test_gdn_fused_split_head_ratios.py | 2 +- .../attention/test_gdn_noncontiguous_stride.py | 2 +- .../registered/attention/test_gdn_prefill_layout.py | 2 +- .../test_gemma4_unified_swa_virtual_ids.py | 2 +- test/registered/attention/test_kda_kernels.py | 2 +- .../attention/test_normal_decode_set_metadata.py | 2 +- .../attention/test_triton_sliding_window.py | 2 +- .../attention/test_unified_memory_deterministic.py | 2 +- .../attention/unittests/dense/test_fa4.py | 2 +- .../attention/unittests/dense/test_flashinfer.py | 2 +- .../unittests/dense/test_flex_attention.py | 2 +- .../attention/unittests/dense/test_hybrid_attn.py | 2 +- .../attention/unittests/dense/test_torch_native.py | 2 +- .../attention/unittests/dense/test_triton.py | 2 +- .../attention/unittests/dense/test_trtllm_mha.py | 2 +- .../unittests/gdn/test_gdn_cutedsl_ring_verify.py | 2 +- .../unittests/gdn/test_gdn_replayssm_spec_fold.py | 2 +- .../unittests/gdn/test_linear_replayssm_decode.py | 2 +- .../attention/unittests/gdn/test_torch_native.py | 2 +- .../attention/unittests/gdn/test_triton.py | 2 +- .../test_flashinfer_mla_chunk_metadata.py | 2 +- .../hybrid_linear/test_kda_fused_accept_indices.py | 2 +- .../hybrid_linear/test_verify_buffer_fixup_hook.py | 2 +- .../attention/unittests/kda/test_triton.py | 2 +- .../attention/unittests/lightning/test_triton.py | 2 +- .../attention/unittests/mamba/test_mamba2.py | 2 +- .../attention/unittests/mla/test_triton.py | 2 +- .../attention/unittests/swa/test_flashinfer.py | 2 +- .../attention/unittests/swa/test_torch_native.py | 2 +- .../attention/unittests/swa/test_triton.py | 2 +- .../chunked_prefill/test_scripted_swa_1gpu.py | 2 +- .../cuda_graph/breakable/test_bcg_with_lora.py | 2 +- .../breakable/test_breakable_cuda_graph.py | 2 +- .../e2e/entrypoints/test_engine_multi_tokenizer.py | 2 +- .../e2e/gdn/test_qwen35_gdn_multi_item_scoring.py | 2 +- .../e2e/hicache/test_hicache_storage_lora.py | 2 +- .../e2e/models/test_compressed_tensors_models.py | 2 +- .../kernels/ops/activation/test_activation.py | 2 +- .../kernels/ops/attention/test_c128_v2.py | 2 +- test/registered/kernels/ops/attention/test_c4_v2.py | 2 +- .../kernels/ops/attention/test_clamp_position.py | 2 +- .../ops/attention/test_combine_topk_swa_indices.py | 2 +- .../kernels/ops/attention/test_concat_mla.py | 2 +- .../kernels/ops/attention/test_dcp_lse_combine.py | 2 +- .../test_deepseek_v4_compress_plan_bounds.py | 2 +- ...est_deepseek_v4_compress_state_runtime_shapes.py | 2 +- .../ops/attention/test_dsa_kpool_metadata_fusion.py | 2 +- .../ops/attention/test_dsa_kpool_multi_pool.py | 2 +- .../kernels/ops/attention/test_dsa_metadata.py | 2 +- .../ops/attention/test_dsa_metadata_replay.py | 2 +- .../ops/attention/test_dsa_multi_ctas_counter.py | 2 +- .../ops/attention/test_dsa_sink_pad_cache.py | 2 +- .../ops/attention/test_dsa_transform_index.py | 2 +- .../ops/attention/test_dsv32_indexer_fusion.py | 2 +- .../ops/attention/test_dsv4_indexer_quant.py | 2 +- .../kernels/ops/attention/test_dsv4_q_rope_store.py | 2 +- .../kernels/ops/attention/test_flash_attention_4.py | 2 +- .../ops/attention/test_flash_mla_backends.py | 2 +- .../kernels/ops/attention/test_fp4_indexer.py | 2 +- .../test_fused_kda_conv_recurrent_verify.py | 2 +- .../ops/attention/test_fused_metadata_copy.py | 2 +- .../attention/test_fused_qk_rmsnorm_rope_gate.py | 2 +- .../ops/attention/test_fused_store_index_cache.py | 2 +- .../ops/attention/test_fused_verify_triton_gdn.py | 2 +- .../attention/test_gdn_decode_fused_proj_conv.py | 2 +- .../ops/attention/test_inkling_attn_prologue_tau.py | 2 +- .../kernels/ops/attention/test_inkling_rel_proj.py | 2 +- .../kernels/ops/attention/test_inkling_row_scale.py | 2 +- .../kernels/ops/attention/test_kda_fused_decode.py | 2 +- .../ops/attention/test_kda_fused_verify_backend.py | 2 +- .../ops/attention/test_kda_gate_beta_cumsum.py | 2 +- .../kernels/ops/attention/test_kda_helion.py | 2 +- .../ops/attention/test_kda_replayssm_fold.py | 2 +- .../attention/test_kda_replayssm_fold_batched.py | 2 +- .../ops/attention/test_kda_replayssm_ring_fused.py | 2 +- .../ops/attention/test_kda_replayssm_ring_ragged.py | 2 +- .../ops/attention/test_minimax_decode_topk.py | 2 +- .../ops/attention/test_minimax_qknorm_rope.py | 2 +- .../ops/attention/test_mla_kv_pack_quantize_fp8.py | 2 +- .../kernels/ops/attention/test_pos_enc.py | 2 +- test/registered/kernels/ops/attention/test_rope.py | 2 +- .../ops/attention/test_suffix_attention_merge.py | 2 +- .../ops/attention/test_tilelang_dsa_zero_rope.py | 2 +- .../kernels/ops/attention/test_topk_v2.py | 2 +- .../kernels/ops/cuda_graph/test_cuda_graph_dedup.py | 2 +- .../kernels/ops/diffusion/test_activation.py | 2 +- .../kernels/ops/diffusion/test_complex_rope.py | 2 +- .../kernels/ops/diffusion/test_gelu_tanh_cat.py | 2 +- .../kernels/ops/diffusion/test_helios_qk_rope.py | 2 +- .../ops/diffusion/test_interleaved_rope_fp64.py | 2 +- .../kernels/ops/diffusion/test_joint_qkv_cat.py | 2 +- .../ops/diffusion/test_layernorm_modulate.py | 2 +- .../ops/diffusion/test_lingbot_modulation.py | 2 +- .../diffusion/test_longcat_image_norm_modulate.py | 2 +- .../ops/diffusion/test_ltx25_decoder_rope.py | 2 +- .../kernels/ops/diffusion/test_modulate.py | 2 +- test/registered/kernels/ops/diffusion/test_norm.py | 2 +- .../kernels/ops/diffusion/test_qwen_qkv_epilogue.py | 2 +- .../diffusion/test_rmsnorm_preserve_reduction.py | 2 +- test/registered/kernels/ops/diffusion/test_rope.py | 2 +- .../kernels/ops/diffusion/test_routing.py | 2 +- .../kernels/ops/diffusion/test_sana_wm_conv_post.py | 2 +- .../ops/diffusion/test_sana_wm_reverse_scan.py | 2 +- test/registered/kernels/ops/diffusion/test_sites.py | 2 +- .../ops/diffusion/test_wan_norm_silu_post.py | 2 +- .../kernels/ops/disaggregation/test_kv_checksum.py | 2 +- .../kernels/ops/dllm/test_gemma4_sampling_cuda.py | 2 +- .../kernels/ops/elementwise/test_add_constant.py | 2 +- .../kernels/ops/elementwise/test_bias_gelu.py | 2 +- .../kernels/ops/elementwise/test_fast_topk.py | 2 +- .../kernels/ops/elementwise/test_hc_combine.py | 2 +- .../kernels/ops/elementwise/test_launch_carveout.py | 2 +- .../ops/embeddings/test_qwen4_ple_offload.py | 2 +- .../ops/gemm/test_chunked_sgmv_cuda_graph.py | 2 +- .../ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py | 2 +- .../ops/kimi_k3/test_kimi_k3_prerequisite_ops.py | 2 +- .../kernels/ops/kv_canary/test_const_sync.py | 2 +- .../kernels/ops/kv_canary/test_kernel_config.py | 2 +- .../kernels/ops/kv_canary/test_pipeline_e2e.py | 2 +- .../kernels/ops/kv_canary/test_plan_fuzz.py | 2 +- .../kernels/ops/kv_canary/test_plan_hand.py | 2 +- .../ops/kv_canary/test_scatter_req_token_ids.py | 2 +- .../kernels/ops/kv_canary/test_verify_fuzz.py | 2 +- .../kernels/ops/kv_canary/test_verify_hand.py | 2 +- .../kernels/ops/kv_canary/test_write_fuzz.py | 2 +- .../kernels/ops/kv_canary/test_write_hand.py | 2 +- .../ops/kvcache/test_fused_fp8_qkv_kv_cache.py | 2 +- test/registered/kernels/ops/kvcache/test_hicache.py | 2 +- .../kvcache/test_hicache_page_first_write_back.py | 2 +- .../registered/kernels/ops/kvcache/test_hisparse.py | 2 +- .../ops/kvcache/test_hisparse_slot_mapping.py | 2 +- .../kernels/ops/kvcache/test_hisparse_spec.py | 2 +- .../ops/kvcache/test_kvcacheio_asymmetric.py | 2 +- .../ops/kvcache/test_minimax_store_kv_index.py | 2 +- .../kernels/ops/kvcache/test_pd_dcp_gather.py | 2 +- .../kernels/ops/kvcache/test_set_mla_kv_buffer.py | 2 +- .../kernels/ops/kvcache/test_store_cache.py | 2 +- .../ops/kvcache/test_unified_swa_tail_allocation.py | 2 +- .../kernels/ops/layernorm/test_fused_add_rmsnorm.py | 2 +- .../kernels/ops/layernorm/test_fused_eh_norm.py | 2 +- .../kernels/ops/layernorm/test_hc_combine.py | 2 +- .../ops/layernorm/test_hy4_decode_kernels.py | 2 +- .../kernels/ops/layernorm/test_mhc_kernels.py | 2 +- .../registered/kernels/ops/layernorm/test_qknorm.py | 2 +- .../ops/layernorm/test_qknorm_across_heads.py | 2 +- .../kernels/ops/layernorm/test_rmsnorm.py | 2 +- .../kernels/ops/layernorm/test_rmsnorm_hf.py | 2 +- .../kernels/ops/mamba/test_causal_conv1d.py | 2 +- .../ops/mamba/test_fused_replay_state_indices.py | 2 +- .../kernels/ops/mamba/test_sconv_cache.py | 2 +- .../kernels/ops/mamba/test_sconv_decode_metadata.py | 2 +- .../kernels/ops/mamba/test_sconv_extend_metadata.py | 2 +- .../kernels/ops/mamba/test_transfer_mamba.py | 2 +- .../kernels/ops/memory/test_adler32_checksum.py | 2 +- .../kernels/ops/memory/test_swa_page_free.py | 2 +- .../ops/minicpm_sala/test_get_block_table.py | 2 +- .../ops/minicpm_sala/test_minicpm_compress_k.py | 2 +- .../ops/moe/test_deepep_v2_contig_scatter.py | 2 +- .../moe/test_fp8_per_token_to_per_tensor_quant.py | 2 +- .../kernels/ops/moe/test_fused_swiglu_epilogue.py | 2 +- .../kernels/ops/moe/test_inkling_silu_and_mul.py | 2 +- .../kernels/ops/moe/test_jit_grouped_topk.py | 2 +- .../ops/moe/test_marlin_packed_topk_unpack.py | 2 +- .../kernels/ops/moe/test_moe_align_block_size.py | 2 +- .../kernels/ops/moe/test_moe_align_small_numel.py | 2 +- .../kernels/ops/moe/test_moe_fused_gate.py | 2 +- .../ops/moe/test_moe_lora_align_block_size.py | 2 +- .../kernels/ops/moe/test_moe_topk_sigmoid.py | 2 +- .../kernels/ops/moe/test_mova_routed_linear.py | 2 +- .../kernels/ops/moe/test_pad_expert_counts.py | 2 +- .../kernels/ops/moe/test_post_reorder_deepgemm.py | 2 +- test/registered/kernels/ops/moe/test_renorm.py | 2 +- .../ops/moe/test_shuffle_rows_with_scales.py | 2 +- .../kernels/ops/moe/test_sigmoid_gate_mul.py | 2 +- test/registered/kernels/ops/qsa/test_qsa_indexer.py | 2 +- .../kernels/ops/qsa/test_qsa_strided_zero_fill.py | 2 +- .../kernels/ops/quantization/test_awq_dequantize.py | 2 +- .../ops/quantization/test_awq_marlin_moe_repack.py | 2 +- .../ops/quantization/test_awq_marlin_repack.py | 2 +- .../test_diffusion_modelopt_fp8_scaled_mm.py | 2 +- .../kernels/ops/quantization/test_gptq_marlin.py | 2 +- .../ops/quantization/test_gptq_marlin_repack.py | 2 +- .../kernels/ops/quantization/test_hadamard_jit.py | 2 +- .../kernels/ops/quantization/test_nvfp4_marlin.py | 3 +-- .../ops/quantization/test_per_tensor_quant_fp8.py | 2 +- .../ops/quantization/test_per_token_group_quant.py | 2 +- .../ops/quantization/test_per_token_quant_fp8.py | 2 +- .../sampling/test_apply_custom_logit_processor.py | 2 +- .../kernels/ops/speculative/test_dflash_domino.py | 2 +- .../ops/speculative/test_draft_decode_window.py | 2 +- .../kernels/ops/speculative/test_eagle_utils.py | 2 +- .../kernels/ops/speculative/test_ngram_embedding.py | 2 +- .../kernels/ops/speculative/test_ngram_utils.py | 2 +- .../ops/speculative/test_spec_kv_indices_grid.py | 2 +- .../ops/speculative/test_speculative_sampling.py | 2 +- .../ops/speculative/test_verify_commit_triton.py | 2 +- test/registered/layers/test_layernorm_fusion.py | 2 +- test/registered/lora/test_fused_moe_lora_kernel.py | 2 +- test/registered/lora/test_lora_overlap_loading.py | 2 +- .../lora/test_lora_qwen3_5_4b_logprob_diff.py | 2 +- .../lora/test_lora_qwen3_8b_logprob_diff.py | 2 +- test/registered/lora/test_lora_spec_decoding.py | 2 +- test/registered/lora/test_lora_update.py | 2 +- .../mem_cache/test_post_capture_kv_sizing.py | 2 +- .../model_loading/test_utils_update_weights.py | 2 +- ...test_fused_append_remap_per_rank_shared_slots.py | 2 +- .../moe/test_fused_append_shared_experts.py | 2 +- .../moe/test_fused_append_shared_experts_top6.py | 2 +- test/registered/moe/test_fused_moe.py | 2 +- test/registered/moe/test_topk_padded_region.py | 2 +- .../moe/test_topk_renormalize_degenerate.py | 2 +- test/registered/moe/test_triton_fused_moe.py | 2 +- .../moe/test_triton_moe_channel_fp8_kernel.py | 2 +- .../function_call/test_anthropic_tool_use.py | 2 +- .../function_call/test_openai_function_calling.py | 2 +- .../validation/test_large_max_new_tokens.py | 2 +- .../validation/test_request_length_validation.py | 2 +- test/registered/prefill_only/test_serving_rerank.py | 2 +- .../profiling/test_diffusion_stage_profiler_sync.py | 2 +- test/registered/quant/test_awq.py | 2 +- test/registered/quant/test_fp8_kernel.py | 2 +- test/registered/quant/test_fp8_utils.py | 2 +- test/registered/quant/test_fp8kv_triton.py | 2 +- test/registered/quant/test_gptqmodel_dynamic.py | 2 +- test/registered/quant/test_modelopt_fp8.py | 2 +- test/registered/rl/test_lora_load_from_tensor.py | 2 +- test/registered/sessions/test_session_control.py | 2 +- test/registered/sessions/test_session_latency.py | 2 +- test/registered/sessions/test_streaming_session.py | 2 +- .../sessions/test_streaming_session_extra.py | 2 +- .../sessions/test_streaming_session_swa.py | 2 +- .../sessions/test_streaming_session_swa_extra.py | 2 +- .../spec/eagle/test_adaptive_speculative.py | 2 +- .../spec/eagle/test_eagle_constrained_decoding.py | 2 +- .../registered/spec/eagle/test_spec_eagle_stress.py | 2 +- test/registered/spec/test_frozen_kv_mtp.py | 2 +- test/registered/spec/test_spec_ngram_extra.py | 2 +- .../unit/layers/moe/test_deepep_v2_masked_slab.py | 2 +- .../unit/layers/moe/test_flashinfer_a2a_wide_ep.py | 2 +- test/registered/unit/managers/test_io_struct.py | 2 +- test/registered/unit/mem_cache/test_swa_unittest.py | 2 +- .../runner/test_flashinfer_autotune_sync.py | 2 +- .../unit/model_loader/test_modelopt_export.py | 2 +- .../unit/models/test_kimi_k3_bfa_overlap.py | 2 +- .../unit/multimodal/test_cuda_ipc_transport.py | 2 +- .../multimodal/test_vit_cuda_graph_metadata_cuda.py | 2 +- .../unit/sampling/test_sampling_metadata_staging.py | 2 +- .../unit/spec/test_resolve_swa_kv_pool.py | 2 +- test/registered/vlm/test_paddleocr_vl_server.py | 2 +- test/registered/vlm/test_rust_native_mm_e2e.py | 2 +- test/registered/vlm/test_rust_native_mm_mmmu.py | 2 +- test/registered/vlm/test_token_id_retokenize_e2e.py | 2 +- 257 files changed, 272 insertions(+), 267 deletions(-) diff --git a/.claude/skills/add-jit-kernel/SKILL.md b/.claude/skills/add-jit-kernel/SKILL.md index 11fa1d02e6ac..007eb53390cb 100644 --- a/.claude/skills/add-jit-kernel/SKILL.md +++ b/.claude/skills/add-jit-kernel/SKILL.md @@ -524,7 +524,7 @@ if torch.cuda.get_device_capability()[0] < 9: JIT kernel correctness tests and benchmarks live under `test/registered/kernels/ops//` and `test/registered/kernels/benchmark//`, mirroring the wrapper's group under `python/sglang/kernels/ops/` (NOT inside the `sglang` package -- a `register_*_ci(...)` call anywhere under `python/sglang/` is rejected by the `check-no-registered-tests-in-package` pre-commit hook). Only their test-only helpers (e.g. `benchmark/marker.py`) stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path. **CI does not run `pytest` in those directories directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` and `bench_*.py` under `test/registered/`, collects `register_*_ci(...)` calls by **statically parsing each file's AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check. -- **PR / per-commit CUDA suites** (see `test/run_suite.py` → `PER_COMMIT_SUITES`): JIT unit tests use `base-b-kernel-unit-test-1-gpu-large` on H100 and `base-b-kernel-unit-test-4-gpu-b200` on B200/SM100 paths (see `.github/workflows/pr-test-jit-kernel.yml`). Multi-GPU JIT tests use `base-b-kernel-unit-test-8-gpu-h200`. +- **PR / per-commit CUDA suites** (see `test/run_suite.py` → `PER_COMMIT_SUITES`): JIT unit tests use `base-b-kernel-unit-test-1-gpu-small` on the RTX 5090 (SM120) by default, `base-b-kernel-unit-test-1-gpu-large` on H100 only for SM90-only kernels (or shapes that exceed SM120's ~99 KB shared memory), and `base-b-kernel-unit-test-4-gpu-b200` on B200/SM100 paths (see `.github/workflows/pr-test-jit-kernel.yml`). Multi-GPU JIT tests use `base-b-kernel-unit-test-8-gpu-h200`. - **Nightly kernel suite**: register with `stage="nightly"` plus the `runner_config` of the machine it needs (e.g. `1-gpu-large`), giving the `nightly-test-1-gpu-large` suite. `.github/workflows/nightly-test-nvidia.yml` sets `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` for the whole nightly run, so the expanded parameter grids apply automatically (see `python/sglang/kernels/jit/utils/common.py` → `should_run_full_tests` / `get_ci_test_range`). There is no separate kernel-only nightly job: every nightly test on one machine type shares that machine's suite. Registration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing): @@ -532,7 +532,8 @@ Registration pattern (module level, **literal** `est_time`, `stage`, and `runner ```python from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") +# SM90-only kernels register on runner_config="1-gpu-large" instead. # Optional B200/SM100 registration for tests that cover Blackwell-specific code paths # register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") # Optional second registration: same file also runs nightly, same form, @@ -540,7 +541,7 @@ register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-l # register_cuda_ci(est_time=120, stage="nightly", runner_config="1-gpu-large") ``` -CI generates the suite name as `{stage}-test-{runner_config}`, so `stage="base-b-kernel-unit", runner_config="1-gpu-large"` becomes the `base-b-kernel-unit-test-1-gpu-large` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage="nightly"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools. +CI generates the suite name as `{stage}-test-{runner_config}`, so `stage="base-b-kernel-unit", runner_config="1-gpu-small"` becomes the `base-b-kernel-unit-test-1-gpu-small` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage="nightly"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools. Keep `est_time`, `stage`, `runner_config`, and `suite` as literal values. `run_suite.py` collects them from the file AST, so computed values and helper wrappers can break CI discovery. @@ -549,7 +550,7 @@ Use `register_cuda_ci(..., disabled="reason")` if the file must stay in-tree but **Run like CI** (from repo root): ```bash -(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large) +(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-small) # For B200/SM100-specific coverage: (cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200) ``` @@ -564,7 +565,7 @@ import torch from sglang.kernels.ops.elementwise.scale import scale from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32]) diff --git a/.claude/skills/write-sglang-test/SKILL.md b/.claude/skills/write-sglang-test/SKILL.md index c932523dd098..5f156c77457c 100644 --- a/.claude/skills/write-sglang-test/SKILL.md +++ b/.claude/skills/write-sglang-test/SKILL.md @@ -101,7 +101,7 @@ Use the lightest suite that meets your test's needs: - **No GPU required** → `base-a-test-cpu` - **Most small GPU tests** → `base-b-test-1-gpu-small` (default choice) - **Need H100 memory or Hopper features** → `base-b-test-1-gpu-large` -- **JIT kernel correctness** → `base-b-kernel-unit-test-1-gpu-large` +- **JIT kernel correctness** → `base-b-kernel-unit-test-1-gpu-small` (default; SM90-only kernels → `base-b-kernel-unit-test-1-gpu-large`) - **JIT kernel correctness for B200 / SM100 paths** → `base-b-kernel-unit-test-4-gpu-b200` - **JIT kernel benchmarks** → `base-b-kernel-benchmark-test-1-gpu-large` - **Multi-GPU** → only when the test actually needs multiple GPUs @@ -335,7 +335,7 @@ They are ordinary registered tests; only their stage differs: from sglang.test.ci.ci_register import register_cuda_ci # Correctness tests in test/registered/kernels/ops// -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="8-gpu-h200") diff --git a/.github/workflows/pr-test-jit-kernel.yml b/.github/workflows/pr-test-jit-kernel.yml index e8685122fbd8..32f01d016e64 100644 --- a/.github/workflows/pr-test-jit-kernel.yml +++ b/.github/workflows/pr-test-jit-kernel.yml @@ -54,14 +54,19 @@ jobs: include: - suite: unit-test-1-gpu-large runner_config: 1-gpu-large + suite_args: --fork-worker-batch-size 20 + - suite: unit-test-1-gpu-small + runner_config: 1-gpu-small shard: " (0)" - suite_args: --auto-partition-id 0 --auto-partition-size 2 --fork-worker-batch-size 20 - - suite: unit-test-1-gpu-large - runner_config: 1-gpu-large + suite_args: --auto-partition-id 0 --auto-partition-size 3 --fork-worker-batch-size 20 + - suite: unit-test-1-gpu-small + runner_config: 1-gpu-small shard: " (1)" - suite_args: --auto-partition-id 1 --auto-partition-size 2 --fork-worker-batch-size 20 + suite_args: --auto-partition-id 1 --auto-partition-size 3 --fork-worker-batch-size 20 - suite: unit-test-1-gpu-small runner_config: 1-gpu-small + shard: " (2)" + suite_args: --auto-partition-id 2 --auto-partition-size 3 --fork-worker-batch-size 20 - suite: unit-test-4-gpu-b200 runner_config: 4-gpu-b200 - suite: unit-test-8-gpu-h200 diff --git a/test/README.md b/test/README.md index a99158328507..01bb6ba2ede0 100644 --- a/test/README.md +++ b/test/README.md @@ -90,7 +90,7 @@ root and mirror the operator group under `python/sglang/kernels/ops/`. Helpers stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path: -- Correctness tests: `test/registered/kernels/ops//test_*.py` → `base-b-kernel-unit-test-1-gpu-large` +- Correctness tests: `test/registered/kernels/ops//test_*.py` → `base-b-kernel-unit-test-1-gpu-small` (use `base-b-kernel-unit-test-1-gpu-large` only for SM90-only kernels or shapes that exceed the 5090's ~99 KB shared memory) - Benchmarks: `test/registered/kernels/benchmark//bench_*.py` → `base-b-kernel-benchmark-test-1-gpu-large` ## Choosing a Suite @@ -102,7 +102,7 @@ Use the lightest suite that meets your test's needs. Full suite tables are in th | No GPU required | `base-a-test-cpu` | | Small GPU (fits 5090, 32GB) | `base-b-test-1-gpu-small` (most tests go here) | | Large GPU memory or Hopper features | `base-b-test-1-gpu-large` | -| JIT kernel correctness | `base-b-kernel-unit-test-1-gpu-large` | +| JIT kernel correctness | `base-b-kernel-unit-test-1-gpu-small` (SM90-only kernels: `base-b-kernel-unit-test-1-gpu-large`) | | JIT kernel benchmarks | `base-b-kernel-benchmark-test-1-gpu-large` | | Multi-GPU (2/4/8) | `base-b-test-2-gpu-large`, `base-c-test-*` | | Long-running or experimental | `nightly-*` suites | diff --git a/test/registered/attention/test_chunk_gated_delta_rule.py b/test/registered/attention/test_chunk_gated_delta_rule.py index 83cff71ace47..6b6a2f8382f6 100644 --- a/test/registered/attention/test_chunk_gated_delta_rule.py +++ b/test/registered/attention/test_chunk_gated_delta_rule.py @@ -13,7 +13,7 @@ register_xpu_ci, ) -register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=79, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=11, stage="stage-b", runner_config="1-gpu-large-amd") register_xpu_ci(est_time=900, suite="stage-b-test-1-gpu-xpu") diff --git a/test/registered/attention/test_gdn_fused_split_head_ratios.py b/test/registered/attention/test_gdn_fused_split_head_ratios.py index 9c0d7bf8cd0d..da370c8be8ef 100644 --- a/test/registered/attention/test_gdn_fused_split_head_ratios.py +++ b/test/registered/attention/test_gdn_fused_split_head_ratios.py @@ -7,7 +7,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-small") def _reference_split(mixed_qkvz, mixed_ba, num_heads_qk, num_heads_v, head_qk, head_v): diff --git a/test/registered/attention/test_gdn_noncontiguous_stride.py b/test/registered/attention/test_gdn_noncontiguous_stride.py index 65b801d54999..18ddeb16df34 100644 --- a/test/registered/attention/test_gdn_noncontiguous_stride.py +++ b/test/registered/attention/test_gdn_noncontiguous_stride.py @@ -14,7 +14,7 @@ ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=7, stage="stage-b", runner_config="1-gpu-large-amd") diff --git a/test/registered/attention/test_gdn_prefill_layout.py b/test/registered/attention/test_gdn_prefill_layout.py index 497004b4ef54..fc83204b90eb 100644 --- a/test/registered/attention/test_gdn_prefill_layout.py +++ b/test/registered/attention/test_gdn_prefill_layout.py @@ -13,7 +13,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "Test requires CUDA") diff --git a/test/registered/attention/test_gemma4_unified_swa_virtual_ids.py b/test/registered/attention/test_gemma4_unified_swa_virtual_ids.py index 4b43774406ed..dea5d8fdc1bb 100644 --- a/test/registered/attention/test_gemma4_unified_swa_virtual_ids.py +++ b/test/registered/attention/test_gemma4_unified_swa_virtual_ids.py @@ -48,7 +48,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=72, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=72, stage="base-b", runner_config="1-gpu-small") # Healthy 0.865-0.880 (static pool 0.870); a died-mid-run server scores # 0.05-0.21. 1 sigma over 200 examples is ~0.024. diff --git a/test/registered/attention/test_kda_kernels.py b/test/registered/attention/test_kda_kernels.py index 83b8d6de4ca8..bacd410d92ed 100644 --- a/test/registered/attention/test_kda_kernels.py +++ b/test/registered/attention/test_kda_kernels.py @@ -19,7 +19,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=25, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=145, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=12, stage="stage-b", runner_config="1-gpu-large-amd") diff --git a/test/registered/attention/test_normal_decode_set_metadata.py b/test/registered/attention/test_normal_decode_set_metadata.py index d9d4f927ed77..a2ef36460f7c 100644 --- a/test/registered/attention/test_normal_decode_set_metadata.py +++ b/test/registered/attention/test_normal_decode_set_metadata.py @@ -21,7 +21,7 @@ from sglang.test.test_utils import CustomTestCase # Register this test for CUDA CI in base-b (fast attention/kernel tests) -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=17, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/test_triton_sliding_window.py b/test/registered/attention/test_triton_sliding_window.py index 24f88fcd1e2f..d42543190c89 100644 --- a/test/registered/attention/test_triton_sliding_window.py +++ b/test/registered/attention/test_triton_sliding_window.py @@ -15,7 +15,7 @@ ) # Sliding window attention with Triton backend (Gemma-3 model) -register_cuda_ci(est_time=80, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=80, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=200, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/attention/test_unified_memory_deterministic.py b/test/registered/attention/test_unified_memory_deterministic.py index 4cd894fa03b9..c1b2960b7b9c 100644 --- a/test/registered/attention/test_unified_memory_deterministic.py +++ b/test/registered/attention/test_unified_memory_deterministic.py @@ -27,7 +27,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=71, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=133, stage="base-b", runner_config="1-gpu-small") KIMI_LINEAR_MODEL = "yujiepan/kimi-linear-tiny-random" SERVER_ENV = {"SGLANG_BATCH_INVARIANT_OPS_ENABLE_MM_DEEPGEMM": "0"} diff --git a/test/registered/attention/unittests/dense/test_fa4.py b/test/registered/attention/unittests/dense/test_fa4.py index 2adbce5d6ee1..a52288dd74c1 100644 --- a/test/registered/attention/unittests/dense/test_fa4.py +++ b/test/registered/attention/unittests/dense/test_fa4.py @@ -30,7 +30,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=27, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=22, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=22, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/dense/test_flashinfer.py b/test/registered/attention/unittests/dense/test_flashinfer.py index 214bcecb74cb..1b761240e399 100644 --- a/test/registered/attention/unittests/dense/test_flashinfer.py +++ b/test/registered/attention/unittests/dense/test_flashinfer.py @@ -27,7 +27,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf( diff --git a/test/registered/attention/unittests/dense/test_flex_attention.py b/test/registered/attention/unittests/dense/test_flex_attention.py index 123243368561..f4e1e06ef603 100644 --- a/test/registered/attention/unittests/dense/test_flex_attention.py +++ b/test/registered/attention/unittests/dense/test_flex_attention.py @@ -15,7 +15,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/dense/test_hybrid_attn.py b/test/registered/attention/unittests/dense/test_hybrid_attn.py index 444dd5a412d3..1121ec2399f3 100644 --- a/test/registered/attention/unittests/dense/test_hybrid_attn.py +++ b/test/registered/attention/unittests/dense/test_hybrid_attn.py @@ -19,7 +19,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf( diff --git a/test/registered/attention/unittests/dense/test_torch_native.py b/test/registered/attention/unittests/dense/test_torch_native.py index 9eebe7c4c64f..2f92ec5b5aed 100644 --- a/test/registered/attention/unittests/dense/test_torch_native.py +++ b/test/registered/attention/unittests/dense/test_torch_native.py @@ -12,7 +12,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/dense/test_triton.py b/test/registered/attention/unittests/dense/test_triton.py index 64559579caa9..016ae11be9ff 100644 --- a/test/registered/attention/unittests/dense/test_triton.py +++ b/test/registered/attention/unittests/dense/test_triton.py @@ -31,7 +31,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=25, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/dense/test_trtllm_mha.py b/test/registered/attention/unittests/dense/test_trtllm_mha.py index 26678d9b2687..6d799682ea9d 100644 --- a/test/registered/attention/unittests/dense/test_trtllm_mha.py +++ b/test/registered/attention/unittests/dense/test_trtllm_mha.py @@ -29,7 +29,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=17, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf( diff --git a/test/registered/attention/unittests/gdn/test_gdn_cutedsl_ring_verify.py b/test/registered/attention/unittests/gdn/test_gdn_cutedsl_ring_verify.py index d0039b3166b9..958b1e24fb72 100644 --- a/test/registered/attention/unittests/gdn/test_gdn_cutedsl_ring_verify.py +++ b/test/registered/attention/unittests/gdn/test_gdn_cutedsl_ring_verify.py @@ -28,7 +28,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=23, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=23, stage="base-b", runner_config="1-gpu-small") T, H, HV, K, V, SLOTS = 4, 4, 16, 128, 128, 16 DEVICE = "cuda" diff --git a/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py b/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py index f5e531e2284e..5af352d8369c 100644 --- a/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py +++ b/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py @@ -20,7 +20,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") B, T = 3, 4 H, HV = 4, 8 diff --git a/test/registered/attention/unittests/gdn/test_linear_replayssm_decode.py b/test/registered/attention/unittests/gdn/test_linear_replayssm_decode.py index 0c3ef19e7335..b4b4d3d70d6e 100644 --- a/test/registered/attention/unittests/gdn/test_linear_replayssm_decode.py +++ b/test/registered/attention/unittests/gdn/test_linear_replayssm_decode.py @@ -45,7 +45,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=57, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/gdn/test_torch_native.py b/test/registered/attention/unittests/gdn/test_torch_native.py index ef516ab5b7ee..3959bfda6e37 100644 --- a/test/registered/attention/unittests/gdn/test_torch_native.py +++ b/test/registered/attention/unittests/gdn/test_torch_native.py @@ -16,7 +16,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/gdn/test_triton.py b/test/registered/attention/unittests/gdn/test_triton.py index 129e76a5c806..03b3fec90435 100644 --- a/test/registered/attention/unittests/gdn/test_triton.py +++ b/test/registered/attention/unittests/gdn/test_triton.py @@ -30,7 +30,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/hybrid_linear/test_flashinfer_mla_chunk_metadata.py b/test/registered/attention/unittests/hybrid_linear/test_flashinfer_mla_chunk_metadata.py index 9a078c277227..dd768b330e13 100644 --- a/test/registered/attention/unittests/hybrid_linear/test_flashinfer_mla_chunk_metadata.py +++ b/test/registered/attention/unittests/hybrid_linear/test_flashinfer_mla_chunk_metadata.py @@ -32,7 +32,7 @@ ) from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") _KV_LORA_RANK = DEFAULT_KV_LORA_RANK diff --git a/test/registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py b/test/registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py index aba244cfd524..38b29fd644cc 100644 --- a/test/registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py +++ b/test/registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py @@ -24,7 +24,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") class TestBuildFusedAcceptIndices(CustomTestCase): diff --git a/test/registered/attention/unittests/hybrid_linear/test_verify_buffer_fixup_hook.py b/test/registered/attention/unittests/hybrid_linear/test_verify_buffer_fixup_hook.py index 8aa979bc5138..0dc552a335c8 100644 --- a/test/registered/attention/unittests/hybrid_linear/test_verify_buffer_fixup_hook.py +++ b/test/registered/attention/unittests/hybrid_linear/test_verify_buffer_fixup_hook.py @@ -22,7 +22,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") _STALE = -555 _MAX_BS = 4 diff --git a/test/registered/attention/unittests/kda/test_triton.py b/test/registered/attention/unittests/kda/test_triton.py index d9f24bbabd60..2a25cad9bba1 100644 --- a/test/registered/attention/unittests/kda/test_triton.py +++ b/test/registered/attention/unittests/kda/test_triton.py @@ -23,7 +23,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/lightning/test_triton.py b/test/registered/attention/unittests/lightning/test_triton.py index 961014bb13d5..7c7b6fa37cfd 100644 --- a/test/registered/attention/unittests/lightning/test_triton.py +++ b/test/registered/attention/unittests/lightning/test_triton.py @@ -20,7 +20,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/mamba/test_mamba2.py b/test/registered/attention/unittests/mamba/test_mamba2.py index e635874c2bd3..0623ead82ff3 100644 --- a/test/registered/attention/unittests/mamba/test_mamba2.py +++ b/test/registered/attention/unittests/mamba/test_mamba2.py @@ -33,7 +33,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/mla/test_triton.py b/test/registered/attention/unittests/mla/test_triton.py index 5c8956506cdf..938eb2e523f3 100644 --- a/test/registered/attention/unittests/mla/test_triton.py +++ b/test/registered/attention/unittests/mla/test_triton.py @@ -29,7 +29,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/swa/test_flashinfer.py b/test/registered/attention/unittests/swa/test_flashinfer.py index 8b3c4b87e2db..83f12046488e 100644 --- a/test/registered/attention/unittests/swa/test_flashinfer.py +++ b/test/registered/attention/unittests/swa/test_flashinfer.py @@ -28,7 +28,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf( diff --git a/test/registered/attention/unittests/swa/test_torch_native.py b/test/registered/attention/unittests/swa/test_torch_native.py index f98a62100126..d179795228de 100644 --- a/test/registered/attention/unittests/swa/test_torch_native.py +++ b/test/registered/attention/unittests/swa/test_torch_native.py @@ -13,7 +13,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/swa/test_triton.py b/test/registered/attention/unittests/swa/test_triton.py index aeeb84c08527..03104761057e 100644 --- a/test/registered/attention/unittests/swa/test_triton.py +++ b/test/registered/attention/unittests/swa/test_triton.py @@ -24,7 +24,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/chunked_prefill/test_scripted_swa_1gpu.py b/test/registered/chunked_prefill/test_scripted_swa_1gpu.py index 7937f4a1573b..f267b93987cf 100644 --- a/test/registered/chunked_prefill/test_scripted_swa_1gpu.py +++ b/test/registered/chunked_prefill/test_scripted_swa_1gpu.py @@ -5,7 +5,7 @@ from sglang.test.scripted_runtime.test_case import ScriptedTestCase from sglang.test.scripted_runtime_chunked_helpers import base_engine_kwargs -register_cuda_ci(est_time=130, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=130, stage="extra-a", runner_config="1-gpu-small") _SWA_MODEL = "openai/gpt-oss-20b" diff --git a/test/registered/cuda_graph/breakable/test_bcg_with_lora.py b/test/registered/cuda_graph/breakable/test_bcg_with_lora.py index bbf70e6ed726..12379c47e6d5 100644 --- a/test/registered/cuda_graph/breakable/test_bcg_with_lora.py +++ b/test/registered/cuda_graph/breakable/test_bcg_with_lora.py @@ -26,7 +26,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=178, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=178, stage="base-b", runner_config="1-gpu-small") BASE_MODEL = "Qwen/Qwen3-4B" LORA_ADAPTER = "nissenj/Qwen3-4B-lora-v2" diff --git a/test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py b/test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py index 9c9faa1eb00d..40574c106446 100644 --- a/test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py +++ b/test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py @@ -24,7 +24,7 @@ ) # CI Registration — large suite to fit the integration test's server startup. -register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=141, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=200, suite="stage-c-test-large-8-gpu-amd-mi35x") diff --git a/test/registered/e2e/entrypoints/test_engine_multi_tokenizer.py b/test/registered/e2e/entrypoints/test_engine_multi_tokenizer.py index 3700663360ec..66f698c7f2c0 100644 --- a/test/registered/e2e/entrypoints/test_engine_multi_tokenizer.py +++ b/test/registered/e2e/entrypoints/test_engine_multi_tokenizer.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import DEFAULT_SMALL_MODEL_NAME_FOR_TEST, CustomTestCase -register_cuda_ci(est_time=180, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=180, stage="base-b", runner_config="1-gpu-small") PROMPTS = ["The capital of France is", "1 + 1 ="] SAMPLING = {"temperature": 0, "max_new_tokens": 8} diff --git a/test/registered/e2e/gdn/test_qwen35_gdn_multi_item_scoring.py b/test/registered/e2e/gdn/test_qwen35_gdn_multi_item_scoring.py index fe9737147e0d..a533dc87ace9 100644 --- a/test/registered/e2e/gdn/test_qwen35_gdn_multi_item_scoring.py +++ b/test/registered/e2e/gdn/test_qwen35_gdn_multi_item_scoring.py @@ -14,7 +14,7 @@ CustomTestCase, ) -register_cuda_ci(est_time=240, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=240, stage="extra-a", runner_config="1-gpu-small") class TestQwen35GDNMultiItemScoring(CustomTestCase): diff --git a/test/registered/e2e/hicache/test_hicache_storage_lora.py b/test/registered/e2e/hicache/test_hicache_storage_lora.py index fd8051b0905f..f6774698e275 100644 --- a/test/registered/e2e/hicache/test_hicache_storage_lora.py +++ b/test/registered/e2e/hicache/test_hicache_storage_lora.py @@ -21,7 +21,7 @@ terminate_and_kill_process_tree, ) -register_cuda_ci(est_time=300, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=300, stage="base-b", runner_config="1-gpu-small") LORA_NAME = "sql" LORA_PATH = "philschmid/code-llama-3-1-8b-text-to-sql-lora" diff --git a/test/registered/e2e/models/test_compressed_tensors_models.py b/test/registered/e2e/models/test_compressed_tensors_models.py index b9135d57c6b1..c0f9cd2f3853 100644 --- a/test/registered/e2e/models/test_compressed_tensors_models.py +++ b/test/registered/e2e/models/test_compressed_tensors_models.py @@ -13,7 +13,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=63, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=63, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=42, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/kernels/ops/activation/test_activation.py b/test/registered/kernels/ops/activation/test_activation.py index fb2562e0fca5..e82f734f64cc 100644 --- a/test/registered/kernels/ops/activation/test_activation.py +++ b/test/registered/kernels/ops/activation/test_activation.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=20, stage="nightly", runner_config="1-gpu-large") register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/attention/test_c128_v2.py b/test/registered/kernels/ops/attention/test_c128_v2.py index a5b0e74fab94..dc920f684209 100644 --- a/test/registered/kernels/ops/attention/test_c128_v2.py +++ b/test/registered/kernels/ops/attention/test_c128_v2.py @@ -19,7 +19,7 @@ to_seq_extend, ) -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) Context = Union[LegacyContext, PagedContext] diff --git a/test/registered/kernels/ops/attention/test_c4_v2.py b/test/registered/kernels/ops/attention/test_c4_v2.py index cd419ef76db6..a45f96882bf6 100644 --- a/test/registered/kernels/ops/attention/test_c4_v2.py +++ b/test/registered/kernels/ops/attention/test_c4_v2.py @@ -23,7 +23,7 @@ to_seq_extend, ) -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) Context = Union[LegacyContext, PagedContext] diff --git a/test/registered/kernels/ops/attention/test_clamp_position.py b/test/registered/kernels/ops/attention/test_clamp_position.py index 579e5b4d7b8d..b18d2c8a7f5e 100644 --- a/test/registered/kernels/ops/attention/test_clamp_position.py +++ b/test/registered/kernels/ops/attention/test_clamp_position.py @@ -10,7 +10,7 @@ ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=12, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/attention/test_combine_topk_swa_indices.py b/test/registered/kernels/ops/attention/test_combine_topk_swa_indices.py index ff7fbdee15df..b41c110d8b9c 100644 --- a/test/registered/kernels/ops/attention/test_combine_topk_swa_indices.py +++ b/test/registered/kernels/ops/attention/test_combine_topk_swa_indices.py @@ -20,7 +20,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available(), reason="combine_topk_swa_indices requires CUDA" diff --git a/test/registered/kernels/ops/attention/test_concat_mla.py b/test/registered/kernels/ops/attention/test_concat_mla.py index 1375061f24bd..fe65d1105f54 100644 --- a/test/registered/kernels/ops/attention/test_concat_mla.py +++ b/test/registered/kernels/ops/attention/test_concat_mla.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=17, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b-kernel-unit", runner_config="1-gpu-small") def torch_concat_mla_k( diff --git a/test/registered/kernels/ops/attention/test_dcp_lse_combine.py b/test/registered/kernels/ops/attention/test_dcp_lse_combine.py index f177178c1f6e..d4b855271ee4 100644 --- a/test/registered/kernels/ops/attention/test_dcp_lse_combine.py +++ b/test/registered/kernels/ops/attention/test_dcp_lse_combine.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") class TestLSECombineTritonVsCPU(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_deepseek_v4_compress_plan_bounds.py b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_plan_bounds.py index 37bd285bffc6..44618bc684f5 100644 --- a/test/registered/kernels/ops/attention/test_deepseek_v4_compress_plan_bounds.py +++ b/test/registered/kernels/ops/attention/test_deepseek_v4_compress_plan_bounds.py @@ -14,7 +14,7 @@ ) from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestCompressWritePlanBounds(CustomTestCase): 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 1d1328e4abdb..b3a9279f75b8 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 @@ -101,7 +101,7 @@ ) from sglang.utils import is_in_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=25, suite="nightly-amd-kernel-1-gpu", nightly=True) Mode = Literal["decode", "prefill"] diff --git a/test/registered/kernels/ops/attention/test_dsa_kpool_metadata_fusion.py b/test/registered/kernels/ops/attention/test_dsa_kpool_metadata_fusion.py index 021067027982..f85bb6bb15fb 100644 --- a/test/registered/kernels/ops/attention/test_dsa_kpool_metadata_fusion.py +++ b/test/registered/kernels/ops/attention/test_dsa_kpool_metadata_fusion.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestKPoolMetadataFusion(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_dsa_kpool_multi_pool.py b/test/registered/kernels/ops/attention/test_dsa_kpool_multi_pool.py index 26ce445333f0..31e13fc03d81 100644 --- a/test/registered/kernels/ops/attention/test_dsa_kpool_multi_pool.py +++ b/test/registered/kernels/ops/attention/test_dsa_kpool_multi_pool.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "Test requires CUDA") diff --git a/test/registered/kernels/ops/attention/test_dsa_metadata.py b/test/registered/kernels/ops/attention/test_dsa_metadata.py index 149e7c4e3c7b..b4c42c5f1e52 100644 --- a/test/registered/kernels/ops/attention/test_dsa_metadata.py +++ b/test/registered/kernels/ops/attention/test_dsa_metadata.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=15, stage="stage-b", runner_config="1-gpu-large-amd") diff --git a/test/registered/kernels/ops/attention/test_dsa_metadata_replay.py b/test/registered/kernels/ops/attention/test_dsa_metadata_replay.py index e47b7aefa737..31fab9254c1f 100644 --- a/test/registered/kernels/ops/attention/test_dsa_metadata_replay.py +++ b/test/registered/kernels/ops/attention/test_dsa_metadata_replay.py @@ -20,7 +20,7 @@ ) from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestDSAMetadataReplay(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_dsa_multi_ctas_counter.py b/test/registered/kernels/ops/attention/test_dsa_multi_ctas_counter.py index 4902e296fa46..5c035d0bf283 100644 --- a/test/registered/kernels/ops/attention/test_dsa_multi_ctas_counter.py +++ b/test/registered/kernels/ops/attention/test_dsa_multi_ctas_counter.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") _NUM_Q_HEADS = 128 _MAX_CTX_LEN = 64 diff --git a/test/registered/kernels/ops/attention/test_dsa_sink_pad_cache.py b/test/registered/kernels/ops/attention/test_dsa_sink_pad_cache.py index 2af86eb97d71..ff5e45169542 100644 --- a/test/registered/kernels/ops/attention/test_dsa_sink_pad_cache.py +++ b/test/registered/kernels/ops/attention/test_dsa_sink_pad_cache.py @@ -7,7 +7,7 @@ from sglang.srt.layers.attention.dsa_backend import DeepseekSparseAttnBackend from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required") diff --git a/test/registered/kernels/ops/attention/test_dsa_transform_index.py b/test/registered/kernels/ops/attention/test_dsa_transform_index.py index 231540540556..52ce01eef773 100644 --- a/test/registered/kernels/ops/attention/test_dsa_transform_index.py +++ b/test/registered/kernels/ops/attention/test_dsa_transform_index.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=9, suite="stage-b-test-1-gpu-small-amd-mi35x") TOPK = 2048 diff --git a/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py b/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py index 268a1c4dff7d..5204308dc7a3 100644 --- a/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py +++ b/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py @@ -31,7 +31,7 @@ _is_hip = is_hip() -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=45, suite="jit-kernel-unit-test-amd") HEAD_DIM = 128 diff --git a/test/registered/kernels/ops/attention/test_dsv4_indexer_quant.py b/test/registered/kernels/ops/attention/test_dsv4_indexer_quant.py index fe2d562b9a4e..7ddec586b557 100644 --- a/test/registered/kernels/ops/attention/test_dsv4_indexer_quant.py +++ b/test/registered/kernels/ops/attention/test_dsv4_indexer_quant.py @@ -29,7 +29,7 @@ _is_hip = is_hip() -register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-small") # the mi35x suite rather than the default AMD one: the ROCm case below is gfx95-only, and # everything else in here skips on HIP, so the mi300 registration only ever produced skips register_amd_ci(est_time=45, suite="stage-b-test-1-gpu-small-amd-mi35x") diff --git a/test/registered/kernels/ops/attention/test_dsv4_q_rope_store.py b/test/registered/kernels/ops/attention/test_dsv4_q_rope_store.py index c4943112a83b..7ed02c0fd501 100644 --- a/test/registered/kernels/ops/attention/test_dsv4_q_rope_store.py +++ b/test/registered/kernels/ops/attention/test_dsv4_q_rope_store.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestQRopeStore(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_flash_attention_4.py b/test/registered/kernels/ops/attention/test_flash_attention_4.py index 69132bba3fcb..fa2625edabac 100644 --- a/test/registered/kernels/ops/attention/test_flash_attention_4.py +++ b/test/registered/kernels/ops/attention/test_flash_attention_4.py @@ -18,7 +18,7 @@ from sglang.srt.utils import is_sm100_or_sm110_supported from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="4-gpu-b200") # Skip this test on Hopper machine diff --git a/test/registered/kernels/ops/attention/test_flash_mla_backends.py b/test/registered/kernels/ops/attention/test_flash_mla_backends.py index d105bdfc77a5..6781213eeb16 100644 --- a/test/registered/kernels/ops/attention/test_flash_mla_backends.py +++ b/test/registered/kernels/ops/attention/test_flash_mla_backends.py @@ -54,7 +54,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-small") # Per-token byte layout diff --git a/test/registered/kernels/ops/attention/test_fp4_indexer.py b/test/registered/kernels/ops/attention/test_fp4_indexer.py index 10509632c624..e67a946a5520 100644 --- a/test/registered/kernels/ops/attention/test_fp4_indexer.py +++ b/test/registered/kernels/ops/attention/test_fp4_indexer.py @@ -22,7 +22,7 @@ from sglang.srt.utils import get_device, is_xpu from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") _is_xpu = is_xpu() if _is_xpu: diff --git a/test/registered/kernels/ops/attention/test_fused_kda_conv_recurrent_verify.py b/test/registered/kernels/ops/attention/test_fused_kda_conv_recurrent_verify.py index 4d5957404deb..d9e9f1e900a0 100644 --- a/test/registered/kernels/ops/attention/test_fused_kda_conv_recurrent_verify.py +++ b/test/registered/kernels/ops/attention/test_fused_kda_conv_recurrent_verify.py @@ -14,7 +14,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-small") _DEVICE = "cuda" diff --git a/test/registered/kernels/ops/attention/test_fused_metadata_copy.py b/test/registered/kernels/ops/attention/test_fused_metadata_copy.py index a4f093c24703..685783e3c2ad 100644 --- a/test/registered/kernels/ops/attention/test_fused_metadata_copy.py +++ b/test/registered/kernels/ops/attention/test_fused_metadata_copy.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=100, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=100, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=100, suite="nightly-amd-kernel-1-gpu", nightly=True) # ============================================================================= diff --git a/test/registered/kernels/ops/attention/test_fused_qk_rmsnorm_rope_gate.py b/test/registered/kernels/ops/attention/test_fused_qk_rmsnorm_rope_gate.py index b9f3338f6ea9..6df0de1e80c7 100644 --- a/test/registered/kernels/ops/attention/test_fused_qk_rmsnorm_rope_gate.py +++ b/test/registered/kernels/ops/attention/test_fused_qk_rmsnorm_rope_gate.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-small") def gemma_rmsnorm(x: torch.Tensor, weight: torch.Tensor, eps: float) -> torch.Tensor: diff --git a/test/registered/kernels/ops/attention/test_fused_store_index_cache.py b/test/registered/kernels/ops/attention/test_fused_store_index_cache.py index a7ecb21cb99e..c1ab340cb29d 100644 --- a/test/registered/kernels/ops/attention/test_fused_store_index_cache.py +++ b/test/registered/kernels/ops/attention/test_fused_store_index_cache.py @@ -48,7 +48,7 @@ except ImportError: _is_fp8_fnuz = False -register_cuda_ci(est_time=24, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=24, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=24, suite="nightly-amd-kernel-1-gpu", nightly=True) PAGE_SIZE = 64 diff --git a/test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py b/test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py index f66a8932038a..bc6b25342900 100644 --- a/test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py +++ b/test/registered/kernels/ops/attention/test_fused_verify_triton_gdn.py @@ -28,7 +28,7 @@ except ImportError: KERNELS_AVAILABLE = False -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/kernels/ops/attention/test_gdn_decode_fused_proj_conv.py b/test/registered/kernels/ops/attention/test_gdn_decode_fused_proj_conv.py index 7eda705639aa..6b0e427dd7da 100644 --- a/test/registered/kernels/ops/attention/test_gdn_decode_fused_proj_conv.py +++ b/test/registered/kernels/ops/attention/test_gdn_decode_fused_proj_conv.py @@ -14,7 +14,7 @@ from sglang.kernels.ops.mamba.causal_conv1d_triton import causal_conv1d_update from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=7, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=7, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _reference( diff --git a/test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py b/test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py index a5e4a7804770..171422a5cc18 100644 --- a/test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py +++ b/test/registered/kernels/ops/attention/test_inkling_attn_prologue_tau.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="4-gpu-b200") HEAD = 128 diff --git a/test/registered/kernels/ops/attention/test_inkling_rel_proj.py b/test/registered/kernels/ops/attention/test_inkling_rel_proj.py index 3f90b250da01..b8862d710fab 100644 --- a/test/registered/kernels/ops/attention/test_inkling_rel_proj.py +++ b/test/registered/kernels/ops/attention/test_inkling_rel_proj.py @@ -9,7 +9,7 @@ from sglang.kernels.ops.attention.inkling_rel_proj import rel_proj_small_t from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") H, K, E, ROW = 16, 16, 1024, 2816 diff --git a/test/registered/kernels/ops/attention/test_inkling_row_scale.py b/test/registered/kernels/ops/attention/test_inkling_row_scale.py index 68b5febadc2e..ed5eaefbbd92 100644 --- a/test/registered/kernels/ops/attention/test_inkling_row_scale.py +++ b/test/registered/kernels/ops/attention/test_inkling_row_scale.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") diff --git a/test/registered/kernels/ops/attention/test_kda_fused_decode.py b/test/registered/kernels/ops/attention/test_kda_fused_decode.py index 3c4e6d665198..b830d616927c 100644 --- a/test/registered/kernels/ops/attention/test_kda_fused_decode.py +++ b/test/registered/kernels/ops/attention/test_kda_fused_decode.py @@ -20,7 +20,7 @@ from sglang.kernels.ops.mamba.causal_conv1d_triton import causal_conv1d_update from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=8, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=8, stage="base-b-kernel-unit", runner_config="1-gpu-small") _HEAD_DIM = 128 _CONV_STATE_W = 3 diff --git a/test/registered/kernels/ops/attention/test_kda_fused_verify_backend.py b/test/registered/kernels/ops/attention/test_kda_fused_verify_backend.py index e042c12fb595..ba21fce5067b 100644 --- a/test/registered/kernels/ops/attention/test_kda_fused_verify_backend.py +++ b/test/registered/kernels/ops/attention/test_kda_fused_verify_backend.py @@ -23,7 +23,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-small") # One bf16 ulp: the fused and unfused kernels reduce K in different orders. _OUTPUT_TOL = dict(rtol=2**-7, atol=1e-7) diff --git a/test/registered/kernels/ops/attention/test_kda_gate_beta_cumsum.py b/test/registered/kernels/ops/attention/test_kda_gate_beta_cumsum.py index 60dc2b38791f..cdc575a56222 100644 --- a/test/registered/kernels/ops/attention/test_kda_gate_beta_cumsum.py +++ b/test/registered/kernels/ops/attention/test_kda_gate_beta_cumsum.py @@ -5,7 +5,7 @@ from sglang.kernels.ops.attention.fla.kda import chunk_kda, kda_gate_chunk_cumsum from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "requires CUDA") diff --git a/test/registered/kernels/ops/attention/test_kda_helion.py b/test/registered/kernels/ops/attention/test_kda_helion.py index f4f0c558a09b..c1b0f2be105e 100644 --- a/test/registered/kernels/ops/attention/test_kda_helion.py +++ b/test/registered/kernels/ops/attention/test_kda_helion.py @@ -37,7 +37,7 @@ helion_fused_recurrent_kda_replayssm_decode, ) -register_cuda_ci(est_time=180, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=180, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not HELION_AVAILABLE, diff --git a/test/registered/kernels/ops/attention/test_kda_replayssm_fold.py b/test/registered/kernels/ops/attention/test_kda_replayssm_fold.py index 249b4007664f..fd6998cd4a02 100644 --- a/test/registered/kernels/ops/attention/test_kda_replayssm_fold.py +++ b/test/registered/kernels/ops/attention/test_kda_replayssm_fold.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") if not torch.cuda.is_available(): import pytest diff --git a/test/registered/kernels/ops/attention/test_kda_replayssm_fold_batched.py b/test/registered/kernels/ops/attention/test_kda_replayssm_fold_batched.py index 938c7b1d4bd5..760dc75754cb 100644 --- a/test/registered/kernels/ops/attention/test_kda_replayssm_fold_batched.py +++ b/test/registered/kernels/ops/attention/test_kda_replayssm_fold_batched.py @@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") if not torch.cuda.is_available(): pytest.skip( diff --git a/test/registered/kernels/ops/attention/test_kda_replayssm_ring_fused.py b/test/registered/kernels/ops/attention/test_kda_replayssm_ring_fused.py index afbba8dbb43b..8631f3bddb3c 100644 --- a/test/registered/kernels/ops/attention/test_kda_replayssm_ring_fused.py +++ b/test/registered/kernels/ops/attention/test_kda_replayssm_ring_fused.py @@ -20,7 +20,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") if not torch.cuda.is_available(): pytest.skip( diff --git a/test/registered/kernels/ops/attention/test_kda_replayssm_ring_ragged.py b/test/registered/kernels/ops/attention/test_kda_replayssm_ring_ragged.py index daef77a7041e..0b84ff281819 100644 --- a/test/registered/kernels/ops/attention/test_kda_replayssm_ring_ragged.py +++ b/test/registered/kernels/ops/attention/test_kda_replayssm_ring_ragged.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") if not torch.cuda.is_available(): pytest.skip( diff --git a/test/registered/kernels/ops/attention/test_minimax_decode_topk.py b/test/registered/kernels/ops/attention/test_minimax_decode_topk.py index aeaf103f0896..ccb776dc69cb 100644 --- a/test/registered/kernels/ops/attention/test_minimax_decode_topk.py +++ b/test/registered/kernels/ops/attention/test_minimax_decode_topk.py @@ -12,7 +12,7 @@ from sglang.kernels.ops.attention.minimax_decode_topk import minimax_decode_topk from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_amd_ci(est_time=15, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py b/test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py index 74c94d1dc2fa..4538b6637483 100644 --- a/test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py +++ b/test/registered/kernels/ops/attention/test_minimax_qknorm_rope.py @@ -14,7 +14,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") dev = "cuda" diff --git a/test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py b/test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py index 0bf8faed919a..d83f719d2515 100644 --- a/test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py +++ b/test/registered/kernels/ops/attention/test_mla_kv_pack_quantize_fp8.py @@ -9,7 +9,7 @@ ) 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_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=60, suite="nightly-amd-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/kernels/ops/attention/test_pos_enc.py b/test/registered/kernels/ops/attention/test_pos_enc.py index 58e3647c651c..d32f4e4f8d72 100644 --- a/test/registered/kernels/ops/attention/test_pos_enc.py +++ b/test/registered/kernels/ops/attention/test_pos_enc.py @@ -10,7 +10,7 @@ from sglang.kernels.ops.attention.rope import rotary_embedding from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=18, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=18, stage="base-b-kernel-unit", runner_config="1-gpu-small") @triton.jit diff --git a/test/registered/kernels/ops/attention/test_rope.py b/test/registered/kernels/ops/attention/test_rope.py index a60d94349e0b..b4cceceea2ff 100644 --- a/test/registered/kernels/ops/attention/test_rope.py +++ b/test/registered/kernels/ops/attention/test_rope.py @@ -9,7 +9,7 @@ from sglang.srt.utils import is_hip from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=24, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=103, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=390, stage="nightly", runner_config="1-gpu-large") register_amd_ci(est_time=64, suite="jit-kernel-unit-test-amd") diff --git a/test/registered/kernels/ops/attention/test_suffix_attention_merge.py b/test/registered/kernels/ops/attention/test_suffix_attention_merge.py index b5b426327aff..fce66f62ba9c 100644 --- a/test/registered/kernels/ops/attention/test_suffix_attention_merge.py +++ b/test/registered/kernels/ops/attention/test_suffix_attention_merge.py @@ -14,7 +14,7 @@ register_cuda_ci( est_time=15, stage="base-b-kernel-unit", - runner_config="1-gpu-large", + runner_config="1-gpu-small", ) diff --git a/test/registered/kernels/ops/attention/test_tilelang_dsa_zero_rope.py b/test/registered/kernels/ops/attention/test_tilelang_dsa_zero_rope.py index 1458bd73bd86..e71c6443314b 100644 --- a/test/registered/kernels/ops/attention/test_tilelang_dsa_zero_rope.py +++ b/test/registered/kernels/ops/attention/test_tilelang_dsa_zero_rope.py @@ -16,7 +16,7 @@ from sglang.test.test_utils import CustomTestCase # backend-specific: the zero-tail specialization only exists in the HIP TileLang kernels -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=180, suite="stage-b-test-1-gpu-small-amd-mi35x") diff --git a/test/registered/kernels/ops/attention/test_topk_v2.py b/test/registered/kernels/ops/attention/test_topk_v2.py index 3869c4d28ee4..d5291938a98b 100644 --- a/test/registered/kernels/ops/attention/test_topk_v2.py +++ b/test/registered/kernels/ops/attention/test_topk_v2.py @@ -43,7 +43,7 @@ from sglang.srt.utils import is_hip from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") PAGE_SIZE = 64 # c4 page size = 256 // 4 diff --git a/test/registered/kernels/ops/cuda_graph/test_cuda_graph_dedup.py b/test/registered/kernels/ops/cuda_graph/test_cuda_graph_dedup.py index 2976937a1454..606db27148ad 100644 --- a/test/registered/kernels/ops/cuda_graph/test_cuda_graph_dedup.py +++ b/test/registered/kernels/ops/cuda_graph/test_cuda_graph_dedup.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "requires CUDA") diff --git a/test/registered/kernels/ops/diffusion/test_activation.py b/test/registered/kernels/ops/diffusion/test_activation.py index 4a2802ad2719..50735be2c8b1 100644 --- a/test/registered/kernels/ops/diffusion/test_activation.py +++ b/test/registered/kernels/ops/diffusion/test_activation.py @@ -26,7 +26,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_complex_rope.py b/test/registered/kernels/ops/diffusion/test_complex_rope.py index 5ecf2eeb6d2a..800a03393b13 100644 --- a/test/registered/kernels/ops/diffusion/test_complex_rope.py +++ b/test/registered/kernels/ops/diffusion/test_complex_rope.py @@ -12,7 +12,7 @@ from sglang.multimodal_gen.runtime.models.dits import qwen_image21 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_gelu_tanh_cat.py b/test/registered/kernels/ops/diffusion/test_gelu_tanh_cat.py index 85dbacc658f3..7bcd9834675c 100644 --- a/test/registered/kernels/ops/diffusion/test_gelu_tanh_cat.py +++ b/test/registered/kernels/ops/diffusion/test_gelu_tanh_cat.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_helios_qk_rope.py b/test/registered/kernels/ops/diffusion/test_helios_qk_rope.py index 271719d33aaa..8188862ef021 100644 --- a/test/registered/kernels/ops/diffusion/test_helios_qk_rope.py +++ b/test/registered/kernels/ops/diffusion/test_helios_qk_rope.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_interleaved_rope_fp64.py b/test/registered/kernels/ops/diffusion/test_interleaved_rope_fp64.py index 5c77c3485161..1ebb6a497518 100644 --- a/test/registered/kernels/ops/diffusion/test_interleaved_rope_fp64.py +++ b/test/registered/kernels/ops/diffusion/test_interleaved_rope_fp64.py @@ -10,7 +10,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_joint_qkv_cat.py b/test/registered/kernels/ops/diffusion/test_joint_qkv_cat.py index ff254d0cc793..db8c8b387be8 100644 --- a/test/registered/kernels/ops/diffusion/test_joint_qkv_cat.py +++ b/test/registered/kernels/ops/diffusion/test_joint_qkv_cat.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") def make_inputs(batch, image_tokens, text_tokens, heads, dim, dtype): diff --git a/test/registered/kernels/ops/diffusion/test_layernorm_modulate.py b/test/registered/kernels/ops/diffusion/test_layernorm_modulate.py index 03c614a4d75b..fbe91297d3c0 100644 --- a/test/registered/kernels/ops/diffusion/test_layernorm_modulate.py +++ b/test/registered/kernels/ops/diffusion/test_layernorm_modulate.py @@ -11,7 +11,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_lingbot_modulation.py b/test/registered/kernels/ops/diffusion/test_lingbot_modulation.py index e4617c8e1c48..ec394d8553bc 100644 --- a/test/registered/kernels/ops/diffusion/test_lingbot_modulation.py +++ b/test/registered/kernels/ops/diffusion/test_lingbot_modulation.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-small") def make_sites(channels): diff --git a/test/registered/kernels/ops/diffusion/test_longcat_image_norm_modulate.py b/test/registered/kernels/ops/diffusion/test_longcat_image_norm_modulate.py index fd22bbd90ab2..32639393f9fc 100644 --- a/test/registered/kernels/ops/diffusion/test_longcat_image_norm_modulate.py +++ b/test/registered/kernels/ops/diffusion/test_longcat_image_norm_modulate.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestLongCatNormModulation(CustomTestCase): diff --git a/test/registered/kernels/ops/diffusion/test_ltx25_decoder_rope.py b/test/registered/kernels/ops/diffusion/test_ltx25_decoder_rope.py index f4ba23a52a20..3ed517f770ce 100644 --- a/test/registered/kernels/ops/diffusion/test_ltx25_decoder_rope.py +++ b/test/registered/kernels/ops/diffusion/test_ltx25_decoder_rope.py @@ -10,7 +10,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_modulate.py b/test/registered/kernels/ops/diffusion/test_modulate.py index c4bee5102960..8d77de0ba4fd 100644 --- a/test/registered/kernels/ops/diffusion/test_modulate.py +++ b/test/registered/kernels/ops/diffusion/test_modulate.py @@ -32,7 +32,7 @@ ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=75, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=75, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="4-gpu-b200") # Nightly is not redundant: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1, which # expands the get_ci_test_range sweeps below. diff --git a/test/registered/kernels/ops/diffusion/test_norm.py b/test/registered/kernels/ops/diffusion/test_norm.py index 261db1ae0702..0c11fe8c7a24 100644 --- a/test/registered/kernels/ops/diffusion/test_norm.py +++ b/test/registered/kernels/ops/diffusion/test_norm.py @@ -43,7 +43,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=70, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=70, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_qwen_qkv_epilogue.py b/test/registered/kernels/ops/diffusion/test_qwen_qkv_epilogue.py index 83ccc46aee63..f222f22256ca 100644 --- a/test/registered/kernels/ops/diffusion/test_qwen_qkv_epilogue.py +++ b/test/registered/kernels/ops/diffusion/test_qwen_qkv_epilogue.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="4-gpu-b200") -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.cuda.get_device_capability()[0] < 9, diff --git a/test/registered/kernels/ops/diffusion/test_rmsnorm_preserve_reduction.py b/test/registered/kernels/ops/diffusion/test_rmsnorm_preserve_reduction.py index aac5659801c2..a4c723d85cf0 100644 --- a/test/registered/kernels/ops/diffusion/test_rmsnorm_preserve_reduction.py +++ b/test/registered/kernels/ops/diffusion/test_rmsnorm_preserve_reduction.py @@ -10,7 +10,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_rope.py b/test/registered/kernels/ops/diffusion/test_rope.py index ee1fa089fc25..868bbf5563e6 100644 --- a/test/registered/kernels/ops/diffusion/test_rope.py +++ b/test/registered/kernels/ops/diffusion/test_rope.py @@ -31,7 +31,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=18, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=131, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1, which # expands the get_ci_test_range sweeps below. register_cuda_ci(est_time=220, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/diffusion/test_routing.py b/test/registered/kernels/ops/diffusion/test_routing.py index d82e4da6d936..129bcc293588 100644 --- a/test/registered/kernels/ops/diffusion/test_routing.py +++ b/test/registered/kernels/ops/diffusion/test_routing.py @@ -11,7 +11,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_sana_wm_conv_post.py b/test/registered/kernels/ops/diffusion/test_sana_wm_conv_post.py index 1f33b86258ff..0527c296d832 100644 --- a/test/registered/kernels/ops/diffusion/test_sana_wm_conv_post.py +++ b/test/registered/kernels/ops/diffusion/test_sana_wm_conv_post.py @@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestSanaWMConvPost(CustomTestCase): diff --git a/test/registered/kernels/ops/diffusion/test_sana_wm_reverse_scan.py b/test/registered/kernels/ops/diffusion/test_sana_wm_reverse_scan.py index 476e2949bacb..28d244c614a4 100644 --- a/test/registered/kernels/ops/diffusion/test_sana_wm_reverse_scan.py +++ b/test/registered/kernels/ops/diffusion/test_sana_wm_reverse_scan.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestSanaWMReverseScan(CustomTestCase): diff --git a/test/registered/kernels/ops/diffusion/test_sites.py b/test/registered/kernels/ops/diffusion/test_sites.py index a62595529447..ef4add8922a7 100644 --- a/test/registered/kernels/ops/diffusion/test_sites.py +++ b/test/registered/kernels/ops/diffusion/test_sites.py @@ -50,7 +50,7 @@ from sglang.test.test_utils import CustomTestCase register_cpu_ci(est_time=12, suite="base-a-test-cpu") -register_cuda_ci(est_time=38, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=38, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif( not torch.cuda.is_available(), reason="CUDA required" diff --git a/test/registered/kernels/ops/diffusion/test_wan_norm_silu_post.py b/test/registered/kernels/ops/diffusion/test_wan_norm_silu_post.py index 060765c9e247..d887e1294a7b 100644 --- a/test/registered/kernels/ops/diffusion/test_wan_norm_silu_post.py +++ b/test/registered/kernels/ops/diffusion/test_wan_norm_silu_post.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") def reference(x, gamma, bias, scale): diff --git a/test/registered/kernels/ops/disaggregation/test_kv_checksum.py b/test/registered/kernels/ops/disaggregation/test_kv_checksum.py index 70321df97573..3126722e06f6 100644 --- a/test/registered/kernels/ops/disaggregation/test_kv_checksum.py +++ b/test/registered/kernels/ops/disaggregation/test_kv_checksum.py @@ -17,7 +17,7 @@ from sglang.srt.environ import envs from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _ref_strided_adler32(tensors, indices, strides) -> int: diff --git a/test/registered/kernels/ops/dllm/test_gemma4_sampling_cuda.py b/test/registered/kernels/ops/dllm/test_gemma4_sampling_cuda.py index 30e2d2c1d43c..94c9ef3e7746 100644 --- a/test/registered/kernels/ops/dllm/test_gemma4_sampling_cuda.py +++ b/test/registered/kernels/ops/dllm/test_gemma4_sampling_cuda.py @@ -13,7 +13,7 @@ from sglang.srt.runtime_context import get_context, get_parallel from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/kernels/ops/elementwise/test_add_constant.py b/test/registered/kernels/ops/elementwise/test_add_constant.py index c5fc47fd84f2..58ad134fcf79 100644 --- a/test/registered/kernels/ops/elementwise/test_add_constant.py +++ b/test/registered/kernels/ops/elementwise/test_add_constant.py @@ -6,7 +6,7 @@ from sglang.kernels.ops.elementwise.add_constant import add_constant from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=8, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/elementwise/test_bias_gelu.py b/test/registered/kernels/ops/elementwise/test_bias_gelu.py index 6ff81595862e..6c5b0f85a019 100644 --- a/test/registered/kernels/ops/elementwise/test_bias_gelu.py +++ b/test/registered/kernels/ops/elementwise/test_bias_gelu.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.elementwise.bias_gelu import bias_gelu_tanh 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="4-gpu-b200") diff --git a/test/registered/kernels/ops/elementwise/test_fast_topk.py b/test/registered/kernels/ops/elementwise/test_fast_topk.py index c203acb49a46..412e8ef2a143 100644 --- a/test/registered/kernels/ops/elementwise/test_fast_topk.py +++ b/test/registered/kernels/ops/elementwise/test_fast_topk.py @@ -4,7 +4,7 @@ from sglang.kernels.ops.elementwise.fast_topk import fast_topk from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _check_topk_values(score, lengths, indices, topk, row_starts): diff --git a/test/registered/kernels/ops/elementwise/test_hc_combine.py b/test/registered/kernels/ops/elementwise/test_hc_combine.py index 10bc98113834..c26762530913 100644 --- a/test/registered/kernels/ops/elementwise/test_hc_combine.py +++ b/test/registered/kernels/ops/elementwise/test_hc_combine.py @@ -5,7 +5,7 @@ from sglang.kernels.ops.elementwise.hc_combine import hc_combine from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") HC_COUNT = 4 HIDDEN_SIZE = 2560 diff --git a/test/registered/kernels/ops/elementwise/test_launch_carveout.py b/test/registered/kernels/ops/elementwise/test_launch_carveout.py index d49461efb0f9..a38cce4963f5 100644 --- a/test/registered/kernels/ops/elementwise/test_launch_carveout.py +++ b/test/registered/kernels/ops/elementwise/test_launch_carveout.py @@ -9,7 +9,7 @@ from sglang.kernels.jit.utils import load_jit from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_launch_carveout(): diff --git a/test/registered/kernels/ops/embeddings/test_qwen4_ple_offload.py b/test/registered/kernels/ops/embeddings/test_qwen4_ple_offload.py index b95c4d8ae70b..0e7b22cc6bbe 100644 --- a/test/registered/kernels/ops/embeddings/test_qwen4_ple_offload.py +++ b/test/registered/kernels/ops/embeddings/test_qwen4_ple_offload.py @@ -20,7 +20,7 @@ from sglang.srt.utils import set_weight_attrs from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available(), reason="CUDA is required for this test." diff --git a/test/registered/kernels/ops/gemm/test_chunked_sgmv_cuda_graph.py b/test/registered/kernels/ops/gemm/test_chunked_sgmv_cuda_graph.py index b47f44481c7f..32454abcc120 100644 --- a/test/registered/kernels/ops/gemm/test_chunked_sgmv_cuda_graph.py +++ b/test/registered/kernels/ops/gemm/test_chunked_sgmv_cuda_graph.py @@ -37,7 +37,7 @@ from sglang.srt.runtime_context import get_context from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py b/test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py index 96a6004e41b9..fd9e6782371b 100644 --- a/test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py +++ b/test/registered/kernels/ops/gemm/test_cutedsl_dsv3_fused_a_gemm.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") if not torch.cuda.is_available(): pytest.skip("CUDA required", allow_module_level=True) diff --git a/test/registered/kernels/ops/kimi_k3/test_kimi_k3_prerequisite_ops.py b/test/registered/kernels/ops/kimi_k3/test_kimi_k3_prerequisite_ops.py index 6f47dc3d6a06..50698353b082 100644 --- a/test/registered/kernels/ops/kimi_k3/test_kimi_k3_prerequisite_ops.py +++ b/test/registered/kernels/ops/kimi_k3/test_kimi_k3_prerequisite_ops.py @@ -42,7 +42,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") NUM_EXPERTS = 896 TOPK = 16 diff --git a/test/registered/kernels/ops/kv_canary/test_const_sync.py b/test/registered/kernels/ops/kv_canary/test_const_sync.py index de39c8f743ba..42611c49b224 100644 --- a/test/registered/kernels/ops/kv_canary/test_const_sync.py +++ b/test/registered/kernels/ops/kv_canary/test_const_sync.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.kv_canary import consts from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=5, stage="jit-kernel-unit", runner_config="amd") 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 bbaad5a9ce15..84edc1c8252a 100644 --- a/test/registered/kernels/ops/kv_canary/test_kernel_config.py +++ b/test/registered/kernels/ops/kv_canary/test_kernel_config.py @@ -35,7 +35,7 @@ empty_extras, ) -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=60, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") 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 0178d2afd457..90c7cf64c17b 100644 --- a/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py +++ b/test/registered/kernels/ops/kv_canary/test_pipeline_e2e.py @@ -40,7 +40,7 @@ make_req_to_token, ) -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") 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 091f406fd2f6..73cee9670a0a 100644 --- a/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_fuzz.py @@ -22,7 +22,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") 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 89ebfafc83a9..c9fff8357f6e 100644 --- a/test/registered/kernels/ops/kv_canary/test_plan_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_plan_hand.py @@ -22,7 +22,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py b/test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py index 5c6fd2368394..69f832c27d38 100644 --- a/test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py +++ b/test/registered/kernels/ops/kv_canary/test_scatter_req_token_ids.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, 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 b1405a9079b3..1808d2caabe0 100644 --- a/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_fuzz.py @@ -31,7 +31,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") 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 a71c53ac6d45..8259a89f35e6 100644 --- a/test/registered/kernels/ops/kv_canary/test_verify_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_verify_hand.py @@ -54,7 +54,7 @@ _hand_fold_partial, ) -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") 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 33004395ac6f..3c72dbe32a67 100644 --- a/test/registered/kernels/ops/kv_canary/test_write_fuzz.py +++ b/test/registered/kernels/ops/kv_canary/test_write_fuzz.py @@ -31,7 +31,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") 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 79ba9f260265..243cd0578cb1 100644 --- a/test/registered/kernels/ops/kv_canary/test_write_hand.py +++ b/test/registered/kernels/ops/kv_canary/test_write_hand.py @@ -51,7 +51,7 @@ _hand_fold_partial, ) -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py b/test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py index c4e134a791ec..7fe6bd34360e 100644 --- a/test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py +++ b/test/registered/kernels/ops/kvcache/test_fused_fp8_qkv_kv_cache.py @@ -4,7 +4,7 @@ from sglang.kernels.ops.kvcache.fused_fp8_qkv_kv_cache import fused_fp8_qkv_kv_cache from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="4-gpu-b200") FP8 = torch.float8_e4m3fn diff --git a/test/registered/kernels/ops/kvcache/test_hicache.py b/test/registered/kernels/ops/kvcache/test_hicache.py index 4f2fe71c4cb7..81f5eec95aa9 100644 --- a/test/registered/kernels/ops/kvcache/test_hicache.py +++ b/test/registered/kernels/ops/kvcache/test_hicache.py @@ -19,7 +19,7 @@ from sglang.srt.utils import is_cuda, is_hip, is_npu, is_xpu from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=43, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=12, stage="jit-kernel-unit", runner_config="amd") pytestmark = pytest.mark.skipif( diff --git a/test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py b/test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py index 73499e35c530..adf13f5a1ae1 100644 --- a/test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py +++ b/test/registered/kernels/ops/kvcache/test_hicache_page_first_write_back.py @@ -30,7 +30,7 @@ from sglang.srt.utils import is_cuda, is_hip, is_npu, is_xpu from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=55, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") pytestmark = pytest.mark.skipif( diff --git a/test/registered/kernels/ops/kvcache/test_hisparse.py b/test/registered/kernels/ops/kvcache/test_hisparse.py index a41618222618..1f6b8ed1d06c 100644 --- a/test/registered/kernels/ops/kvcache/test_hisparse.py +++ b/test/registered/kernels/ops/kvcache/test_hisparse.py @@ -17,7 +17,7 @@ ) register_amd_ci(est_time=30, stage="stage-b", runner_config="1-gpu-small-amd") -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=50, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_xpu_ci(est_time=60, suite="stage-b-test-1-gpu-xpu") if is_xpu(): diff --git a/test/registered/kernels/ops/kvcache/test_hisparse_slot_mapping.py b/test/registered/kernels/ops/kvcache/test_hisparse_slot_mapping.py index 93ca37657ce7..2ed294015b04 100644 --- a/test/registered/kernels/ops/kvcache/test_hisparse_slot_mapping.py +++ b/test/registered/kernels/ops/kvcache/test_hisparse_slot_mapping.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kvcache/test_hisparse_spec.py b/test/registered/kernels/ops/kvcache/test_hisparse_spec.py index 7f46073a2b2b..eb228f02db53 100644 --- a/test/registered/kernels/ops/kvcache/test_hisparse_spec.py +++ b/test/registered/kernels/ops/kvcache/test_hisparse_spec.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or is_npu() or is_xpu(), diff --git a/test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py b/test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py index fa137f026408..b19be6a0cfe0 100644 --- a/test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py +++ b/test/registered/kernels/ops/kvcache/test_kvcacheio_asymmetric.py @@ -7,7 +7,7 @@ from sglang.srt.mem_cache.pool_host.mha import AsymmetricMHATokenToKVPoolHost from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) # These tests use AsymmetricMHATokenToKVPoolHost methods and let that class call diff --git a/test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py b/test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py index ae94a6b17db5..0dd4940d1151 100644 --- a/test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py +++ b/test/registered/kernels/ops/kvcache/test_minimax_store_kv_index.py @@ -11,7 +11,7 @@ from sglang.kernels.ops.kvcache.minimax_store_kv_index import store_kv_index 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/kernels/ops/kvcache/test_pd_dcp_gather.py b/test/registered/kernels/ops/kvcache/test_pd_dcp_gather.py index 723916e35c48..433ae539a936 100644 --- a/test/registered/kernels/ops/kvcache/test_pd_dcp_gather.py +++ b/test/registered/kernels/ops/kvcache/test_pd_dcp_gather.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestPdDcpGather(CustomTestCase): diff --git a/test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py b/test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py index ab7598db6976..d04ff2fa71b2 100644 --- a/test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py +++ b/test/registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py @@ -15,7 +15,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=15, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" diff --git a/test/registered/kernels/ops/kvcache/test_store_cache.py b/test/registered/kernels/ops/kvcache/test_store_cache.py index 17b898a2c6cd..cb12b1187705 100644 --- a/test/registered/kernels/ops/kvcache/test_store_cache.py +++ b/test/registered/kernels/ops/kvcache/test_store_cache.py @@ -8,7 +8,7 @@ from sglang.kernels.ops.kvcache.kvcache import can_use_store_cache, store_cache from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=78, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=40, stage="nightly", runner_config="1-gpu-large") register_amd_ci(est_time=55, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/kvcache/test_unified_swa_tail_allocation.py b/test/registered/kernels/ops/kvcache/test_unified_swa_tail_allocation.py index 866de4d95ecd..ba936079f4b0 100644 --- a/test/registered/kernels/ops/kvcache/test_unified_swa_tail_allocation.py +++ b/test/registered/kernels/ops/kvcache/test_unified_swa_tail_allocation.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestUnifiedSWATailAllocation(CustomTestCase): diff --git a/test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py index b1afd5e1f607..c87f2f2c8f60 100644 --- a/test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py +++ b/test/registered/kernels/ops/layernorm/test_fused_add_rmsnorm.py @@ -7,7 +7,7 @@ from sglang.kernels.jit.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=20, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/layernorm/test_fused_eh_norm.py b/test/registered/kernels/ops/layernorm/test_fused_eh_norm.py index b7f0ded07ee6..9858d189f758 100644 --- a/test/registered/kernels/ops/layernorm/test_fused_eh_norm.py +++ b/test/registered/kernels/ops/layernorm/test_fused_eh_norm.py @@ -6,7 +6,7 @@ from sglang.kernels.ops.layernorm.fused_eh_norm import fused_eh_norm from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.cuda is None, diff --git a/test/registered/kernels/ops/layernorm/test_hc_combine.py b/test/registered/kernels/ops/layernorm/test_hc_combine.py index 08173765a9c4..20382dc17c40 100644 --- a/test/registered/kernels/ops/layernorm/test_hc_combine.py +++ b/test/registered/kernels/ops/layernorm/test_hc_combine.py @@ -6,7 +6,7 @@ from sglang.kernels.ops.layernorm.mhc import hc_combine from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.cuda is None, diff --git a/test/registered/kernels/ops/layernorm/test_hy4_decode_kernels.py b/test/registered/kernels/ops/layernorm/test_hy4_decode_kernels.py index dfc0f72d864c..55bf6aac20e8 100644 --- a/test/registered/kernels/ops/layernorm/test_hy4_decode_kernels.py +++ b/test/registered/kernels/ops/layernorm/test_hy4_decode_kernels.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=35, stage="base-b-kernel-unit", runner_config="1-gpu-small") @triton.jit diff --git a/test/registered/kernels/ops/layernorm/test_mhc_kernels.py b/test/registered/kernels/ops/layernorm/test_mhc_kernels.py index b2010cd6ce56..1cd7d022e0d2 100644 --- a/test/registered/kernels/ops/layernorm/test_mhc_kernels.py +++ b/test/registered/kernels/ops/layernorm/test_mhc_kernels.py @@ -8,7 +8,7 @@ from sglang.kernels.ops.layernorm.mhc import mhc_fused_post_pre, mhc_post, mhc_pre from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=127, stage="base-b", runner_config="1-gpu-small") @pytest.fixture diff --git a/test/registered/kernels/ops/layernorm/test_qknorm.py b/test/registered/kernels/ops/layernorm/test_qknorm.py index 7b85ca67574a..2c28f50a27fc 100644 --- a/test/registered/kernels/ops/layernorm/test_qknorm.py +++ b/test/registered/kernels/ops/layernorm/test_qknorm.py @@ -8,7 +8,7 @@ from sglang.kernels.jit.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=37, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=37, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=130, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py b/test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py index 414d18abc0d7..ded311773f69 100644 --- a/test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py +++ b/test/registered/kernels/ops/layernorm/test_qknorm_across_heads.py @@ -8,7 +8,7 @@ from sglang.kernels.jit.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=20, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/layernorm/test_rmsnorm.py b/test/registered/kernels/ops/layernorm/test_rmsnorm.py index 53a5c64cfe37..e82a5a55bb44 100644 --- a/test/registered/kernels/ops/layernorm/test_rmsnorm.py +++ b/test/registered/kernels/ops/layernorm/test_rmsnorm.py @@ -8,7 +8,7 @@ from sglang.srt.utils import is_hip from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=160, stage="nightly", runner_config="1-gpu-large") register_amd_ci(est_time=45, suite="jit-kernel-unit-test-amd") diff --git a/test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py b/test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py index c676de355061..28035bdadc5f 100644 --- a/test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py +++ b/test/registered/kernels/ops/layernorm/test_rmsnorm_hf.py @@ -13,7 +13,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=100, stage="nightly", 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/mamba/test_causal_conv1d.py b/test/registered/kernels/ops/mamba/test_causal_conv1d.py index 38a916aa80fc..df6a1b76c596 100644 --- a/test/registered/kernels/ops/mamba/test_causal_conv1d.py +++ b/test/registered/kernels/ops/mamba/test_causal_conv1d.py @@ -33,7 +33,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_cuda_ci(est_time=180, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/mamba/test_fused_replay_state_indices.py b/test/registered/kernels/ops/mamba/test_fused_replay_state_indices.py index 9410aeea6824..c8414de31fdd 100644 --- a/test/registered/kernels/ops/mamba/test_fused_replay_state_indices.py +++ b/test/registered/kernels/ops/mamba/test_fused_replay_state_indices.py @@ -31,7 +31,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Guard tail appended to every buffer; must stay untouched by both paths. _GUARD = 8 diff --git a/test/registered/kernels/ops/mamba/test_sconv_cache.py b/test/registered/kernels/ops/mamba/test_sconv_cache.py index d4ccbbd675e4..058d6670af97 100644 --- a/test/registered/kernels/ops/mamba/test_sconv_cache.py +++ b/test/registered/kernels/ops/mamba/test_sconv_cache.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py b/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py index 228238a5b89f..ad2dee7ead38 100644 --- a/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py +++ b/test/registered/kernels/ops/mamba/test_sconv_decode_metadata.py @@ -15,7 +15,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py b/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py index 8a1fccfb65a3..fac703236c93 100644 --- a/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py +++ b/test/registered/kernels/ops/mamba/test_sconv_extend_metadata.py @@ -22,7 +22,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly expands the representative PR cases below to the complete matrix. register_cuda_ci(est_time=40, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/mamba/test_transfer_mamba.py b/test/registered/kernels/ops/mamba/test_transfer_mamba.py index ff552f30fb01..a9221935e1ff 100644 --- a/test/registered/kernels/ops/mamba/test_transfer_mamba.py +++ b/test/registered/kernels/ops/mamba/test_transfer_mamba.py @@ -15,7 +15,7 @@ from sglang.srt.mem_cache.pool_host.mamba import MambaPoolHost from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) pytestmark = pytest.mark.skipif( diff --git a/test/registered/kernels/ops/memory/test_adler32_checksum.py b/test/registered/kernels/ops/memory/test_adler32_checksum.py index 7f39dd7fbc4b..f1362da1e699 100644 --- a/test/registered/kernels/ops/memory/test_adler32_checksum.py +++ b/test/registered/kernels/ops/memory/test_adler32_checksum.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _ref_adler32(tensor: torch.Tensor) -> int: diff --git a/test/registered/kernels/ops/memory/test_swa_page_free.py b/test/registered/kernels/ops/memory/test_swa_page_free.py index adfa96196b19..862db5283314 100644 --- a/test/registered/kernels/ops/memory/test_swa_page_free.py +++ b/test/registered/kernels/ops/memory/test_swa_page_free.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/minicpm_sala/test_get_block_table.py b/test/registered/kernels/ops/minicpm_sala/test_get_block_table.py index ea7f8994e45e..1fdd239c5c23 100644 --- a/test/registered/kernels/ops/minicpm_sala/test_get_block_table.py +++ b/test/registered/kernels/ops/minicpm_sala/test_get_block_table.py @@ -4,7 +4,7 @@ from sglang.kernels.ops.minicpm_sala.get_block_table import get_block_table from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") _HEAD_GROUP = 2 _SPARSE_BLOCK_SIZE = 64 diff --git a/test/registered/kernels/ops/minicpm_sala/test_minicpm_compress_k.py b/test/registered/kernels/ops/minicpm_sala/test_minicpm_compress_k.py index a03ea9b7dd31..18fd07488493 100644 --- a/test/registered/kernels/ops/minicpm_sala/test_minicpm_compress_k.py +++ b/test/registered/kernels/ops/minicpm_sala/test_minicpm_compress_k.py @@ -7,7 +7,7 @@ from sglang.srt.layers.attention.minicpm.sparse_utils import compress_k_core_new from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_compress_k_writes_each_head_once(): diff --git a/test/registered/kernels/ops/moe/test_deepep_v2_contig_scatter.py b/test/registered/kernels/ops/moe/test_deepep_v2_contig_scatter.py index ce201ee6d0d4..c421b63a94ab 100644 --- a/test/registered/kernels/ops/moe/test_deepep_v2_contig_scatter.py +++ b/test/registered/kernels/ops/moe/test_deepep_v2_contig_scatter.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" HIDDEN = 256 diff --git a/test/registered/kernels/ops/moe/test_fp8_per_token_to_per_tensor_quant.py b/test/registered/kernels/ops/moe/test_fp8_per_token_to_per_tensor_quant.py index dcbe1eeb60f8..0d3ddceaa568 100644 --- a/test/registered/kernels/ops/moe/test_fp8_per_token_to_per_tensor_quant.py +++ b/test/registered/kernels/ops/moe/test_fp8_per_token_to_per_tensor_quant.py @@ -15,7 +15,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") dev = "cuda" FP8 = torch.float8_e4m3fn diff --git a/test/registered/kernels/ops/moe/test_fused_swiglu_epilogue.py b/test/registered/kernels/ops/moe/test_fused_swiglu_epilogue.py index 9e483f1ffa6b..5a1f6bab160c 100644 --- a/test/registered/kernels/ops/moe/test_fused_swiglu_epilogue.py +++ b/test/registered/kernels/ops/moe/test_fused_swiglu_epilogue.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import publish_build_topology -register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/moe/test_inkling_silu_and_mul.py b/test/registered/kernels/ops/moe/test_inkling_silu_and_mul.py index 58f6c7406faa..d7584f83060c 100644 --- a/test/registered/kernels/ops/moe/test_inkling_silu_and_mul.py +++ b/test/registered/kernels/ops/moe/test_inkling_silu_and_mul.py @@ -20,7 +20,7 @@ from sglang.kernels.ops.moe.inkling_moe import silu_and_mul from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/kernels/ops/moe/test_jit_grouped_topk.py b/test/registered/kernels/ops/moe/test_jit_grouped_topk.py index 08bb8866c03f..3f7a360dd802 100644 --- a/test/registered/kernels/ops/moe/test_jit_grouped_topk.py +++ b/test/registered/kernels/ops/moe/test_jit_grouped_topk.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="jit-kernel-unit-test-amd") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="needs a GPU") diff --git a/test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py b/test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py index 09062c8969ca..65193355318e 100644 --- a/test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py +++ b/test/registered/kernels/ops/moe/test_marlin_packed_topk_unpack.py @@ -17,7 +17,7 @@ from sglang.srt.layers.moe.moe_runner.marlin import _fused_unpack_packed_topk from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _torch_unpack(packed: torch.Tensor): diff --git a/test/registered/kernels/ops/moe/test_moe_align_block_size.py b/test/registered/kernels/ops/moe/test_moe_align_block_size.py index 326905ec61db..47f3189855e1 100644 --- a/test/registered/kernels/ops/moe/test_moe_align_block_size.py +++ b/test/registered/kernels/ops/moe/test_moe_align_block_size.py @@ -10,7 +10,7 @@ from sglang.kernels.ops.moe.moe_align import moe_align_block_size from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=50, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/moe/test_moe_align_small_numel.py b/test/registered/kernels/ops/moe/test_moe_align_small_numel.py index fd67fecab228..52afb875afb1 100644 --- a/test/registered/kernels/ops/moe/test_moe_align_small_numel.py +++ b/test/registered/kernels/ops/moe/test_moe_align_small_numel.py @@ -23,7 +23,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Bucket counts above this are not uniformly supported by the AOT sgl_kernel # path across wheel versions, so the cross-check against it stops here; the diff --git a/test/registered/kernels/ops/moe/test_moe_fused_gate.py b/test/registered/kernels/ops/moe/test_moe_fused_gate.py index 3e909fa958d1..c620a90eb923 100644 --- a/test/registered/kernels/ops/moe/test_moe_fused_gate.py +++ b/test/registered/kernels/ops/moe/test_moe_fused_gate.py @@ -25,7 +25,7 @@ from sglang.kernels.ops.moe.moe_fused_gate import moe_fused_gate, moe_fused_gate_jit from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=8, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=48, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" diff --git a/test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py b/test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py index 35cc3142707e..aaa9884795fd 100644 --- a/test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py +++ b/test/registered/kernels/ops/moe/test_moe_lora_align_block_size.py @@ -11,7 +11,7 @@ from sglang.kernels.ops.moe.moe_lora_align import moe_lora_align_block_size from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py b/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py index 247e7d8f008c..e242ba78236f 100644 --- a/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py +++ b/test/registered/kernels/ops/moe/test_moe_topk_sigmoid.py @@ -16,7 +16,7 @@ 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") try: from sgl_kernel import topk_sigmoid as topk_sigmoid_aot diff --git a/test/registered/kernels/ops/moe/test_mova_routed_linear.py b/test/registered/kernels/ops/moe/test_mova_routed_linear.py index e3e149358482..df3ea837fc9f 100644 --- a/test/registered/kernels/ops/moe/test_mova_routed_linear.py +++ b/test/registered/kernels/ops/moe/test_mova_routed_linear.py @@ -7,7 +7,7 @@ from sglang.srt.runtime_context import get_context from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/kernels/ops/moe/test_pad_expert_counts.py b/test/registered/kernels/ops/moe/test_pad_expert_counts.py index 3f901f36dc50..27c2dd5f4fd9 100644 --- a/test/registered/kernels/ops/moe/test_pad_expert_counts.py +++ b/test/registered/kernels/ops/moe/test_pad_expert_counts.py @@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "pad_expert_counts needs CUDA") diff --git a/test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py b/test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py index ef01918ce5dd..3b064c3a1ac7 100644 --- a/test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py +++ b/test/registered/kernels/ops/moe/test_post_reorder_deepgemm.py @@ -7,7 +7,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") dev = "cuda" diff --git a/test/registered/kernels/ops/moe/test_renorm.py b/test/registered/kernels/ops/moe/test_renorm.py index ae66e0ff9881..4c010a03aa54 100644 --- a/test/registered/kernels/ops/moe/test_renorm.py +++ b/test/registered/kernels/ops/moe/test_renorm.py @@ -9,7 +9,7 @@ from sglang.srt.utils import is_hip from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="stage-b-test-1-gpu-small-amd-mi35x") if is_hip(): diff --git a/test/registered/kernels/ops/moe/test_shuffle_rows_with_scales.py b/test/registered/kernels/ops/moe/test_shuffle_rows_with_scales.py index 767a32c00441..f6b259be3b08 100644 --- a/test/registered/kernels/ops/moe/test_shuffle_rows_with_scales.py +++ b/test/registered/kernels/ops/moe/test_shuffle_rows_with_scales.py @@ -20,7 +20,7 @@ from sglang.kernels.ops.moe.shuffle_rows_with_scales import shuffle_rows_with_scales 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") GROUP_SIZE = 128 diff --git a/test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py b/test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py index 4a0cad31f5d6..c0fee36ccd08 100644 --- a/test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py +++ b/test/registered/kernels/ops/moe/test_sigmoid_gate_mul.py @@ -5,7 +5,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=4, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=4, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=4, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" diff --git a/test/registered/kernels/ops/qsa/test_qsa_indexer.py b/test/registered/kernels/ops/qsa/test_qsa_indexer.py index 9664b83f9c5a..d7b93de2dfa0 100644 --- a/test/registered/kernels/ops/qsa/test_qsa_indexer.py +++ b/test/registered/kernels/ops/qsa/test_qsa_indexer.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") from sglang.srt.layers.attention.qsa.kernel import ( average_pool_qsa_keys, diff --git a/test/registered/kernels/ops/qsa/test_qsa_strided_zero_fill.py b/test/registered/kernels/ops/qsa/test_qsa_strided_zero_fill.py index 89973cd8a3dc..d08746b8bffe 100644 --- a/test/registered/kernels/ops/qsa/test_qsa_strided_zero_fill.py +++ b/test/registered/kernels/ops/qsa/test_qsa_strided_zero_fill.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") from sglang.srt.layers.attention.qsa.sparse_attn import ( qwen_sparse_fa2_cu_seqlens_triton, diff --git a/test/registered/kernels/ops/quantization/test_awq_dequantize.py b/test/registered/kernels/ops/quantization/test_awq_dequantize.py index 3425e7ba14a4..6bc8f3779d27 100644 --- a/test/registered/kernels/ops/quantization/test_awq_dequantize.py +++ b/test/registered/kernels/ops/quantization/test_awq_dequantize.py @@ -9,7 +9,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-small") def reverse_awq_order(t: torch.Tensor): diff --git a/test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py b/test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py index d407afee9331..7bd1668f72af 100644 --- a/test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py +++ b/test/registered/kernels/ops/quantization/test_awq_marlin_moe_repack.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_marlin_utils import awq_pack -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _has_aot_awq_marlin_moe_repack() -> bool: diff --git a/test/registered/kernels/ops/quantization/test_awq_marlin_repack.py b/test/registered/kernels/ops/quantization/test_awq_marlin_repack.py index 10b8efa83dfd..e642a9dc51fd 100644 --- a/test/registered/kernels/ops/quantization/test_awq_marlin_repack.py +++ b/test/registered/kernels/ops/quantization/test_awq_marlin_repack.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_marlin_utils import awq_pack, get_weight_perm, marlin_weights -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _has_aot_awq_marlin_repack() -> bool: diff --git a/test/registered/kernels/ops/quantization/test_diffusion_modelopt_fp8_scaled_mm.py b/test/registered/kernels/ops/quantization/test_diffusion_modelopt_fp8_scaled_mm.py index 0b060b8547cf..88d95761ea88 100644 --- a/test/registered/kernels/ops/quantization/test_diffusion_modelopt_fp8_scaled_mm.py +++ b/test/registered/kernels/ops/quantization/test_diffusion_modelopt_fp8_scaled_mm.py @@ -14,7 +14,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" DTYPE = torch.bfloat16 diff --git a/test/registered/kernels/ops/quantization/test_gptq_marlin.py b/test/registered/kernels/ops/quantization/test_gptq_marlin.py index eac5bbc8452e..d644e4dfccc2 100644 --- a/test/registered/kernels/ops/quantization/test_gptq_marlin.py +++ b/test/registered/kernels/ops/quantization/test_gptq_marlin.py @@ -14,7 +14,7 @@ marlin_quantize, ) -register_cuda_ci(est_time=13, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=141, stage="base-b-kernel-unit", runner_config="1-gpu-small") MNK_FACTORS = [ (1, 1, 1), diff --git a/test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py b/test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py index 6f633e0da5b7..01fe565b99e4 100644 --- a/test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py +++ b/test/registered/kernels/ops/quantization/test_gptq_marlin_repack.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_marlin_utils import get_weight_perm, marlin_weights -register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-small") MARLIN_K_CHUNKS = [128] MARLIN_N_CHUNKS = [64, 256] diff --git a/test/registered/kernels/ops/quantization/test_hadamard_jit.py b/test/registered/kernels/ops/quantization/test_hadamard_jit.py index 7860685da9ee..41fbe1179676 100644 --- a/test/registered/kernels/ops/quantization/test_hadamard_jit.py +++ b/test/registered/kernels/ops/quantization/test_hadamard_jit.py @@ -18,7 +18,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=32, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=81, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=128, stage="nightly", runner_config="1-gpu-large") # Exact M×N Hadamard matrices (±1 entries) copied from diff --git a/test/registered/kernels/ops/quantization/test_nvfp4_marlin.py b/test/registered/kernels/ops/quantization/test_nvfp4_marlin.py index 9465ddd0622e..9e7436acfc5e 100644 --- a/test/registered/kernels/ops/quantization/test_nvfp4_marlin.py +++ b/test/registered/kernels/ops/quantization/test_nvfp4_marlin.py @@ -19,8 +19,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_marlin_utils import make_nvfp4_weight_and_ref -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") -register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=142, stage="base-b", runner_config="1-gpu-small") @pytest.mark.skipif( diff --git a/test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py b/test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py index 1f013dc14b90..5dd7b367a41d 100644 --- a/test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py +++ b/test/registered/kernels/ops/quantization/test_per_tensor_quant_fp8.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-small") # Nightly is not redundant here: it sets SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1 to expand get_ci_test_range sweeps. register_cuda_ci(est_time=20, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/quantization/test_per_token_group_quant.py b/test/registered/kernels/ops/quantization/test_per_token_group_quant.py index d7d51271f55d..bae4b34f34aa 100644 --- a/test/registered/kernels/ops/quantization/test_per_token_group_quant.py +++ b/test/registered/kernels/ops/quantization/test_per_token_group_quant.py @@ -32,7 +32,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=65, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=177, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_cuda_ci(est_time=120, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/quantization/test_per_token_quant_fp8.py b/test/registered/kernels/ops/quantization/test_per_token_quant_fp8.py index 386c1d388df5..74302f4bba61 100644 --- a/test/registered/kernels/ops/quantization/test_per_token_quant_fp8.py +++ b/test/registered/kernels/ops/quantization/test_per_token_quant_fp8.py @@ -8,7 +8,7 @@ from sglang.kernels.ops.quantization.per_token_quant_fp8 import per_token_quant_fp8 from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_cuda_ci(est_time=30, stage="nightly", runner_config="1-gpu-large") diff --git a/test/registered/kernels/ops/sampling/test_apply_custom_logit_processor.py b/test/registered/kernels/ops/sampling/test_apply_custom_logit_processor.py index 4eab75fa5860..592f90b34708 100644 --- a/test/registered/kernels/ops/sampling/test_apply_custom_logit_processor.py +++ b/test/registered/kernels/ops/sampling/test_apply_custom_logit_processor.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestApplyCustomLogitProcessorCUDA(CustomTestCase): diff --git a/test/registered/kernels/ops/speculative/test_dflash_domino.py b/test/registered/kernels/ops/speculative/test_dflash_domino.py index 1d145ed48a57..730f7acd8fe4 100644 --- a/test/registered/kernels/ops/speculative/test_dflash_domino.py +++ b/test/registered/kernels/ops/speculative/test_dflash_domino.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/kernels/ops/speculative/test_draft_decode_window.py b/test/registered/kernels/ops/speculative/test_draft_decode_window.py index 0617d295f7d6..212beee272ef 100644 --- a/test/registered/kernels/ops/speculative/test_draft_decode_window.py +++ b/test/registered/kernels/ops/speculative/test_draft_decode_window.py @@ -59,7 +59,7 @@ from sglang.srt.utils import next_power_of_2 from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") _HAS_CUDA = torch.cuda.is_available() diff --git a/test/registered/kernels/ops/speculative/test_eagle_utils.py b/test/registered/kernels/ops/speculative/test_eagle_utils.py index d91f31eec00e..601adecaa06f 100644 --- a/test/registered/kernels/ops/speculative/test_eagle_utils.py +++ b/test/registered/kernels/ops/speculative/test_eagle_utils.py @@ -12,7 +12,7 @@ else: from sgl_kernel import verify_tree_greedy -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_verify_tree_greedy(): diff --git a/test/registered/kernels/ops/speculative/test_ngram_embedding.py b/test/registered/kernels/ops/speculative/test_ngram_embedding.py index 152c1fa1c1ea..2d686c304af5 100644 --- a/test/registered/kernels/ops/speculative/test_ngram_embedding.py +++ b/test/registered/kernels/ops/speculative/test_ngram_embedding.py @@ -11,7 +11,7 @@ ) 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_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=8, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/speculative/test_ngram_utils.py b/test/registered/kernels/ops/speculative/test_ngram_utils.py index a6fc76bdcc11..eab385ab4a7a 100644 --- a/test/registered/kernels/ops/speculative/test_ngram_utils.py +++ b/test/registered/kernels/ops/speculative/test_ngram_utils.py @@ -10,7 +10,7 @@ else: from sgl_kernel import reconstruct_indices_from_tree_mask -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_reconstruct_indices_from_tree_mask(): diff --git a/test/registered/kernels/ops/speculative/test_spec_kv_indices_grid.py b/test/registered/kernels/ops/speculative/test_spec_kv_indices_grid.py index d93ad83dee47..f5427e675d02 100644 --- a/test/registered/kernels/ops/speculative/test_spec_kv_indices_grid.py +++ b/test/registered/kernels/ops/speculative/test_spec_kv_indices_grid.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=25, stage="jit-kernel-unit", runner_config="amd") SENTINEL = 0x7EADBEEF diff --git a/test/registered/kernels/ops/speculative/test_speculative_sampling.py b/test/registered/kernels/ops/speculative/test_speculative_sampling.py index 77e16ade16e8..1782faf10d5e 100644 --- a/test/registered/kernels/ops/speculative/test_speculative_sampling.py +++ b/test/registered/kernels/ops/speculative/test_speculative_sampling.py @@ -13,7 +13,7 @@ else: from sgl_kernel import tree_speculative_sampling_target_only -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=46, stage="base-b-kernel-unit", runner_config="1-gpu-small") test_cases = [ ( diff --git a/test/registered/kernels/ops/speculative/test_verify_commit_triton.py b/test/registered/kernels/ops/speculative/test_verify_commit_triton.py index fed30ec8f21e..fedb13e843ec 100644 --- a/test/registered/kernels/ops/speculative/test_verify_commit_triton.py +++ b/test/registered/kernels/ops/speculative/test_verify_commit_triton.py @@ -8,7 +8,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") def _reference(accept_index, accept_lens, seq_lens, draft_token_num, track_interval): diff --git a/test/registered/layers/test_layernorm_fusion.py b/test/registered/layers/test_layernorm_fusion.py index 5e6bb7391153..7145626f727f 100644 --- a/test/registered/layers/test_layernorm_fusion.py +++ b/test/registered/layers/test_layernorm_fusion.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=2, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/lora/test_fused_moe_lora_kernel.py b/test/registered/lora/test_fused_moe_lora_kernel.py index 8b63796ce73a..131a49d80574 100644 --- a/test/registered/lora/test_fused_moe_lora_kernel.py +++ b/test/registered/lora/test_fused_moe_lora_kernel.py @@ -17,7 +17,7 @@ # ============================================================================== -register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-small") def round_up(x, base): diff --git a/test/registered/lora/test_lora_overlap_loading.py b/test/registered/lora/test_lora_overlap_loading.py index e88e5c6b7c2e..eb9369f7b8ca 100644 --- a/test/registered/lora/test_lora_overlap_loading.py +++ b/test/registered/lora/test_lora_overlap_loading.py @@ -31,7 +31,7 @@ ) from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=237, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=237, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=75, suite="stage-b-test-1-gpu-small-amd") DEVICE = get_device(0) diff --git a/test/registered/lora/test_lora_qwen3_5_4b_logprob_diff.py b/test/registered/lora/test_lora_qwen3_5_4b_logprob_diff.py index 21bb051a7d25..5791868331f0 100644 --- a/test/registered/lora/test_lora_qwen3_5_4b_logprob_diff.py +++ b/test/registered/lora/test_lora_qwen3_5_4b_logprob_diff.py @@ -34,7 +34,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=59, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=59, stage="extra-a", runner_config="1-gpu-small") BASE_MODEL = "Qwen/Qwen3.5-4B" LORA_HF_REPO = "opherlie/lora-test-case-Qwen3.5-4B" diff --git a/test/registered/lora/test_lora_qwen3_8b_logprob_diff.py b/test/registered/lora/test_lora_qwen3_8b_logprob_diff.py index c08ef5ffcbed..a1a83206fed6 100644 --- a/test/registered/lora/test_lora_qwen3_8b_logprob_diff.py +++ b/test/registered/lora/test_lora_qwen3_8b_logprob_diff.py @@ -37,7 +37,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=56, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=56, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=120, suite="stage-b-test-1-gpu-small-amd") BASE_MODEL = "Qwen/Qwen3-8B" diff --git a/test/registered/lora/test_lora_spec_decoding.py b/test/registered/lora/test_lora_spec_decoding.py index ab81cf988f76..b7087815e36d 100644 --- a/test/registered/lora/test_lora_spec_decoding.py +++ b/test/registered/lora/test_lora_spec_decoding.py @@ -35,7 +35,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=95, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=95, stage="base-b", runner_config="1-gpu-small") PROMPTS = [ "What is the capital of France? Answer in one sentence.", diff --git a/test/registered/lora/test_lora_update.py b/test/registered/lora/test_lora_update.py index ff6047823641..13fba2ca6f71 100644 --- a/test/registered/lora/test_lora_update.py +++ b/test/registered/lora/test_lora_update.py @@ -37,7 +37,7 @@ register_cuda_ci( est_time=487, stage="base-b", - runner_config="1-gpu-large", + runner_config="1-gpu-small", ) register_amd_ci( est_time=730, diff --git a/test/registered/mem_cache/test_post_capture_kv_sizing.py b/test/registered/mem_cache/test_post_capture_kv_sizing.py index bbe3445095cb..dd4a04bda1d0 100644 --- a/test/registered/mem_cache/test_post_capture_kv_sizing.py +++ b/test/registered/mem_cache/test_post_capture_kv_sizing.py @@ -28,7 +28,7 @@ ) # CI Registration -register_cuda_ci(est_time=75, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=75, stage="base-b", runner_config="1-gpu-small") STDOUT_FILENAME = "post_capture_kv_sizing_stdout.log" STDERR_FILENAME = "post_capture_kv_sizing_stderr.log" diff --git a/test/registered/model_loading/test_utils_update_weights.py b/test/registered/model_loading/test_utils_update_weights.py index a995e8b1f22d..61da69f0d0c0 100644 --- a/test/registered/model_loading/test_utils_update_weights.py +++ b/test/registered/model_loading/test_utils_update_weights.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import DEFAULT_SMALL_MODEL_NAME_FOR_TEST -register_cuda_ci(est_time=42, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=42, stage="base-b", runner_config="1-gpu-small") class AsyncEngine(Engine): diff --git a/test/registered/moe/test_fused_append_remap_per_rank_shared_slots.py b/test/registered/moe/test_fused_append_remap_per_rank_shared_slots.py index b5090f5a1420..7e49039bd8b4 100644 --- a/test/registered/moe/test_fused_append_remap_per_rank_shared_slots.py +++ b/test/registered/moe/test_fused_append_remap_per_rank_shared_slots.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/moe/test_fused_append_shared_experts.py b/test/registered/moe/test_fused_append_shared_experts.py index fbe98cbc15c9..e3d2eb0f7d2a 100644 --- a/test/registered/moe/test_fused_append_shared_experts.py +++ b/test/registered/moe/test_fused_append_shared_experts.py @@ -22,7 +22,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase, empty_gpu_cache -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/moe/test_fused_append_shared_experts_top6.py b/test/registered/moe/test_fused_append_shared_experts_top6.py index 47393b6915a4..26a3cbe72511 100644 --- a/test/registered/moe/test_fused_append_shared_experts_top6.py +++ b/test/registered/moe/test_fused_append_shared_experts_top6.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/moe/test_fused_moe.py b/test/registered/moe/test_fused_moe.py index ace128004cc9..dcbf64b4524c 100644 --- a/test/registered/moe/test_fused_moe.py +++ b/test/registered/moe/test_fused_moe.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase, empty_gpu_cache -register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=109, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=30, suite="stage-b-test-1-gpu-small-amd") _is_hip = is_hip() diff --git a/test/registered/moe/test_topk_padded_region.py b/test/registered/moe/test_topk_padded_region.py index 929c96d7a688..6aae4eaa94b9 100644 --- a/test/registered/moe/test_topk_padded_region.py +++ b/test/registered/moe/test_topk_padded_region.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=60, stage="stage-b", runner_config="1-gpu-small-amd") _IS_HIP = is_hip() diff --git a/test/registered/moe/test_topk_renormalize_degenerate.py b/test/registered/moe/test_topk_renormalize_degenerate.py index c48f9209a6d4..9491f5fc75bb 100644 --- a/test/registered/moe/test_topk_renormalize_degenerate.py +++ b/test/registered/moe/test_topk_renormalize_degenerate.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=60, stage="stage-b", runner_config="1-gpu-small-amd") torch.manual_seed(1234) diff --git a/test/registered/moe/test_triton_fused_moe.py b/test/registered/moe/test_triton_fused_moe.py index 3591507b091d..45aacfe6e6fd 100644 --- a/test/registered/moe/test_triton_fused_moe.py +++ b/test/registered/moe/test_triton_fused_moe.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=78, stage="base-b", runner_config="1-gpu-small") class TestFusedMOE(CustomTestCase): diff --git a/test/registered/moe/test_triton_moe_channel_fp8_kernel.py b/test/registered/moe/test_triton_moe_channel_fp8_kernel.py index 21b372a6af69..0057660069a3 100644 --- a/test/registered/moe/test_triton_moe_channel_fp8_kernel.py +++ b/test/registered/moe/test_triton_moe_channel_fp8_kernel.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") def native_w8a8_per_token_matmul(A, B, As, Bs, output_dtype=torch.float16): diff --git a/test/registered/openai_server/function_call/test_anthropic_tool_use.py b/test/registered/openai_server/function_call/test_anthropic_tool_use.py index 2ecfa26d3140..539ceaf3772e 100644 --- a/test/registered/openai_server/function_call/test_anthropic_tool_use.py +++ b/test/registered/openai_server/function_call/test_anthropic_tool_use.py @@ -29,7 +29,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=50, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=50, stage="base-b", runner_config="1-gpu-small") # Backend-specific: streaming tool-argument assembly is rebuilt from # incremental decode output, so ROCm chunk-boundary divergence surfaces here. register_amd_ci(est_time=140, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/openai_server/function_call/test_openai_function_calling.py b/test/registered/openai_server/function_call/test_openai_function_calling.py index ca1afd43415f..70f4d52e91f0 100644 --- a/test/registered/openai_server/function_call/test_openai_function_calling.py +++ b/test/registered/openai_server/function_call/test_openai_function_calling.py @@ -19,7 +19,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=100, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=195, stage="base-b", runner_config="1-gpu-small") # Backend-specific: the llama3 and pythonic parsers run on real decoded text, # so ROCm decode divergence surfaces as tool calls that no longer parse. register_amd_ci(est_time=73, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/openai_server/validation/test_large_max_new_tokens.py b/test/registered/openai_server/validation/test_large_max_new_tokens.py index 2ce6f44c4d54..600b81213fd5 100644 --- a/test/registered/openai_server/validation/test_large_max_new_tokens.py +++ b/test/registered/openai_server/validation/test_large_max_new_tokens.py @@ -25,7 +25,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=58, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=58, stage="base-b", runner_config="1-gpu-small") register_cpu_ci(est_time=101, suite="stage-b-test-cpu-intel") diff --git a/test/registered/openai_server/validation/test_request_length_validation.py b/test/registered/openai_server/validation/test_request_length_validation.py index f8810771c796..98d08800358a 100644 --- a/test/registered/openai_server/validation/test_request_length_validation.py +++ b/test/registered/openai_server/validation/test_request_length_validation.py @@ -13,7 +13,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=49, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=49, stage="base-b", runner_config="1-gpu-small") class TestRequestLengthValidation(CustomTestCase): diff --git a/test/registered/prefill_only/test_serving_rerank.py b/test/registered/prefill_only/test_serving_rerank.py index ba7f071082cd..acb9339a9035 100644 --- a/test/registered/prefill_only/test_serving_rerank.py +++ b/test/registered/prefill_only/test_serving_rerank.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci # Keep consistent with other openai_server/basic unit tests. -register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="stage-b-test-1-gpu-small-amd") try: diff --git a/test/registered/profiling/test_diffusion_stage_profiler_sync.py b/test/registered/profiling/test_diffusion_stage_profiler_sync.py index 82ba6fe290c7..53babef96e3f 100644 --- a/test/registered/profiling/test_diffusion_stage_profiler_sync.py +++ b/test/registered/profiling/test_diffusion_stage_profiler_sync.py @@ -16,7 +16,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA") diff --git a/test/registered/quant/test_awq.py b/test/registered/quant/test_awq.py index 4ef1bfa47195..ab38b3cecbb7 100644 --- a/test/registered/quant/test_awq.py +++ b/test/registered/quant/test_awq.py @@ -13,7 +13,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=220, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=431, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=200, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/quant/test_fp8_kernel.py b/test/registered/quant/test_fp8_kernel.py index 6ee0b8a14854..8e3a3cb1bebc 100644 --- a/test/registered/quant/test_fp8_kernel.py +++ b/test/registered/quant/test_fp8_kernel.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") from sglang.srt.utils import get_device, is_cuda, is_xpu diff --git a/test/registered/quant/test_fp8_utils.py b/test/registered/quant/test_fp8_utils.py index d0ed39ca3a1c..3fa1216a1d07 100644 --- a/test/registered/quant/test_fp8_utils.py +++ b/test/registered/quant/test_fp8_utils.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") class TestMxfp8MoeScaleLayout(CustomTestCase): diff --git a/test/registered/quant/test_fp8kv_triton.py b/test/registered/quant/test_fp8kv_triton.py index 895e33494c81..aed453536bf8 100644 --- a/test/registered/quant/test_fp8kv_triton.py +++ b/test/registered/quant/test_fp8kv_triton.py @@ -12,7 +12,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=64, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=64, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=94, suite="extra-a-test-1-gpu-large-amd") diff --git a/test/registered/quant/test_gptqmodel_dynamic.py b/test/registered/quant/test_gptqmodel_dynamic.py index 418e8c9264c9..a41f25683251 100644 --- a/test/registered/quant/test_gptqmodel_dynamic.py +++ b/test/registered/quant/test_gptqmodel_dynamic.py @@ -14,7 +14,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=50, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=50, stage="extra-a", runner_config="1-gpu-small") def check_quant_method(model_path: str): diff --git a/test/registered/quant/test_modelopt_fp8.py b/test/registered/quant/test_modelopt_fp8.py index 3958fbdde51a..cd9fe495ddd9 100644 --- a/test/registered/quant/test_modelopt_fp8.py +++ b/test/registered/quant/test_modelopt_fp8.py @@ -15,7 +15,7 @@ # Two classes run from this file: the default server plus the Rust-frontend # variant (when the embedded extension is built), each launches a server + eval. -register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-small") class TestModeloptFP8(CustomTestCase): diff --git a/test/registered/rl/test_lora_load_from_tensor.py b/test/registered/rl/test_lora_load_from_tensor.py index bb3a308ab2d2..4057996f566f 100644 --- a/test/registered/rl/test_lora_load_from_tensor.py +++ b/test/registered/rl/test_lora_load_from_tensor.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=104, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=104, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=90, suite="stage-b-test-1-gpu-small-amd") MODEL_PATH = "Qwen/Qwen3-0.6B" diff --git a/test/registered/sessions/test_session_control.py b/test/registered/sessions/test_session_control.py index fbe2498d40fa..590af9ea22cc 100644 --- a/test/registered/sessions/test_session_control.py +++ b/test/registered/sessions/test_session_control.py @@ -28,7 +28,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=137, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=137, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=87, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/sessions/test_session_latency.py b/test/registered/sessions/test_session_latency.py index 558ec07fceee..9e4a882e5592 100644 --- a/test/registered/sessions/test_session_latency.py +++ b/test/registered/sessions/test_session_latency.py @@ -27,7 +27,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=113, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=113, stage="extra-a", runner_config="1-gpu-small") NUM_TURNS = 150 INPUT_LEN = 16 diff --git a/test/registered/sessions/test_streaming_session.py b/test/registered/sessions/test_streaming_session.py index 06f15f136612..7c596d2f64e0 100644 --- a/test/registered/sessions/test_streaming_session.py +++ b/test/registered/sessions/test_streaming_session.py @@ -23,7 +23,7 @@ DEFAULT_TARGET_MODEL_EAGLE3, ) -register_cuda_ci(est_time=294, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=294, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=691, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/sessions/test_streaming_session_extra.py b/test/registered/sessions/test_streaming_session_extra.py index 96be186e5b9b..e50fb4e1d88b 100644 --- a/test/registered/sessions/test_streaming_session_extra.py +++ b/test/registered/sessions/test_streaming_session_extra.py @@ -10,7 +10,7 @@ DEFAULT_TARGET_MODEL_EAGLE3, ) -register_cuda_ci(est_time=466, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=466, stage="extra-a", runner_config="1-gpu-small") register_amd_ci(est_time=562, suite="extra-a-test-1-gpu-large-amd") diff --git a/test/registered/sessions/test_streaming_session_swa.py b/test/registered/sessions/test_streaming_session_swa.py index be61e8182f17..74bc0e549022 100644 --- a/test/registered/sessions/test_streaming_session_swa.py +++ b/test/registered/sessions/test_streaming_session_swa.py @@ -20,7 +20,7 @@ StreamingSessionServerBase, ) -register_cuda_ci(est_time=326, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=326, stage="base-b", runner_config="1-gpu-small") class TestStreamingSessionSWA(StreamingSessionServerBase, StreamingSessionKitMixin): diff --git a/test/registered/sessions/test_streaming_session_swa_extra.py b/test/registered/sessions/test_streaming_session_swa_extra.py index 52b06fbc6a5a..ecdaaacc0112 100644 --- a/test/registered/sessions/test_streaming_session_swa_extra.py +++ b/test/registered/sessions/test_streaming_session_swa_extra.py @@ -13,7 +13,7 @@ StreamingSessionServerBase, ) -register_cuda_ci(est_time=145, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=145, stage="extra-a", runner_config="1-gpu-small") register_cpu_ci(est_time=449, suite="stage-b-test-cpu-intel") diff --git a/test/registered/spec/eagle/test_adaptive_speculative.py b/test/registered/spec/eagle/test_adaptive_speculative.py index 1250369a01eb..d70063e407b0 100644 --- a/test/registered/spec/eagle/test_adaptive_speculative.py +++ b/test/registered/spec/eagle/test_adaptive_speculative.py @@ -18,7 +18,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=120, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=240, suite="stage-b-test-1-gpu-large-amd") HIGH_ACCEPT_PROMPT = ( diff --git a/test/registered/spec/eagle/test_eagle_constrained_decoding.py b/test/registered/spec/eagle/test_eagle_constrained_decoding.py index 1541354c63be..83b832f546d5 100644 --- a/test/registered/spec/eagle/test_eagle_constrained_decoding.py +++ b/test/registered/spec/eagle/test_eagle_constrained_decoding.py @@ -13,7 +13,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=107, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=161, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=165, stage="stage-b", runner_config="1-gpu-large-amd") diff --git a/test/registered/spec/eagle/test_spec_eagle_stress.py b/test/registered/spec/eagle/test_spec_eagle_stress.py index a6b229cae2c4..0e217ee31704 100644 --- a/test/registered/spec/eagle/test_spec_eagle_stress.py +++ b/test/registered/spec/eagle/test_spec_eagle_stress.py @@ -17,7 +17,7 @@ ) from sglang.test.server_fixtures.spec_eagle_fixture import Eagle3Base, EagleLlama2Base -register_cuda_ci(est_time=440, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=930, stage="extra-a", runner_config="1-gpu-small") class TestEagle3Perf(Eagle3Base, SpecPerfKit): diff --git a/test/registered/spec/test_frozen_kv_mtp.py b/test/registered/spec/test_frozen_kv_mtp.py index 96d23882d4ec..a6b220b6d4df 100644 --- a/test/registered/spec/test_frozen_kv_mtp.py +++ b/test/registered/spec/test_frozen_kv_mtp.py @@ -16,7 +16,7 @@ write_github_step_summary, ) -register_cuda_ci(est_time=120, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=450, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/spec/test_spec_ngram_extra.py b/test/registered/spec/test_spec_ngram_extra.py index 8a5f7096e32d..8edd88d80161 100644 --- a/test/registered/spec/test_spec_ngram_extra.py +++ b/test/registered/spec/test_spec_ngram_extra.py @@ -8,7 +8,7 @@ # Extra: Triton + Flashinfer NGRAM backends + non-overlap (sync V2) variant. # Sibling per-commit file (test_spec_ngram.py) keeps the Paged variant. -register_cuda_ci(est_time=197, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=197, stage="extra-a", runner_config="1-gpu-small") class TestNgramSpeculativeDecodingTriton(NgramServerBase, GSM8KMixin): diff --git a/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py b/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py index 61ed8ac97af5..67b4149c4b8e 100644 --- a/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py +++ b/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") DEVICE = "cuda" diff --git a/test/registered/unit/layers/moe/test_flashinfer_a2a_wide_ep.py b/test/registered/unit/layers/moe/test_flashinfer_a2a_wide_ep.py index 16197d843271..2bc60d496cd7 100644 --- a/test/registered/unit/layers/moe/test_flashinfer_a2a_wide_ep.py +++ b/test/registered/unit/layers/moe/test_flashinfer_a2a_wide_ep.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestFlashinferA2AWideEPPlumbing(CustomTestCase): diff --git a/test/registered/unit/managers/test_io_struct.py b/test/registered/unit/managers/test_io_struct.py index 5444ff5bcc60..59bb3c9283b9 100644 --- a/test/registered/unit/managers/test_io_struct.py +++ b/test/registered/unit/managers/test_io_struct.py @@ -42,7 +42,7 @@ CustomTestCase, ) -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=8, suite="stage-b-test-1-gpu-small-amd") register_cpu_ci(est_time=6, suite="stage-b-test-cpu-intel") diff --git a/test/registered/unit/mem_cache/test_swa_unittest.py b/test/registered/unit/mem_cache/test_swa_unittest.py index 7e6b9dd45e08..22999384ccaf 100644 --- a/test/registered/unit/mem_cache/test_swa_unittest.py +++ b/test/registered/unit/mem_cache/test_swa_unittest.py @@ -24,7 +24,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py b/test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py index 930450a83105..3aa56988b5fc 100644 --- a/test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py +++ b/test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_cpu_ci, register_cuda_ci register_cpu_ci(est_time=57, suite="base-a-test-cpu") -register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") import json import multiprocessing diff --git a/test/registered/unit/model_loader/test_modelopt_export.py b/test/registered/unit/model_loader/test_modelopt_export.py index 25313f28241c..695d5b3c496a 100644 --- a/test/registered/unit/model_loader/test_modelopt_export.py +++ b/test/registered/unit/model_loader/test_modelopt_export.py @@ -19,7 +19,7 @@ from sglang.srt.model_loader.loader import ModelOptModelLoader from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=9, suite="stage-b-test-1-gpu-small-amd") # Note: PYTHONPATH=python should be set when running tests diff --git a/test/registered/unit/models/test_kimi_k3_bfa_overlap.py b/test/registered/unit/models/test_kimi_k3_bfa_overlap.py index ade7a3546785..9cf0493bc22a 100644 --- a/test/registered/unit/models/test_kimi_k3_bfa_overlap.py +++ b/test/registered/unit/models/test_kimi_k3_bfa_overlap.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-small") _H = 7168 _QKVG = 6144 # q,k,v,g slices per rank at TP8 diff --git a/test/registered/unit/multimodal/test_cuda_ipc_transport.py b/test/registered/unit/multimodal/test_cuda_ipc_transport.py index e9a9ccebf214..227a584c3dbe 100644 --- a/test/registered/unit/multimodal/test_cuda_ipc_transport.py +++ b/test/registered/unit/multimodal/test_cuda_ipc_transport.py @@ -28,7 +28,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=37, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=37, stage="base-b", runner_config="1-gpu-small") def _produce_pooled_tensor(proxy_queue, consumer_done, result_queue): diff --git a/test/registered/unit/multimodal/test_vit_cuda_graph_metadata_cuda.py b/test/registered/unit/multimodal/test_vit_cuda_graph_metadata_cuda.py index 4b8280a09077..a2388c3117ca 100644 --- a/test/registered/unit/multimodal/test_vit_cuda_graph_metadata_cuda.py +++ b/test/registered/unit/multimodal/test_vit_cuda_graph_metadata_cuda.py @@ -7,7 +7,7 @@ from sglang.srt.multimodal.vit_cuda_graph_runner import ViTCudaGraphRunner from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") class _BoundaryBlock(nn.Module): diff --git a/test/registered/unit/sampling/test_sampling_metadata_staging.py b/test/registered/unit/sampling/test_sampling_metadata_staging.py index f469ae0c36b1..e33fab18b0ca 100644 --- a/test/registered/unit/sampling/test_sampling_metadata_staging.py +++ b/test/registered/unit/sampling/test_sampling_metadata_staging.py @@ -17,7 +17,7 @@ from sglang.test.test_utils import CustomTestCase register_cpu_ci(est_time=15, suite="base-a-test-cpu") -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") VOCAB_SIZE = 32 diff --git a/test/registered/unit/spec/test_resolve_swa_kv_pool.py b/test/registered/unit/spec/test_resolve_swa_kv_pool.py index 57e40d633134..27eecead67bc 100644 --- a/test/registered/unit/spec/test_resolve_swa_kv_pool.py +++ b/test/registered/unit/spec/test_resolve_swa_kv_pool.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=8, suite="stage-b-test-1-gpu-large-amd") _RESOLVERS = ( diff --git a/test/registered/vlm/test_paddleocr_vl_server.py b/test/registered/vlm/test_paddleocr_vl_server.py index 8dcdb59c35db..0ee45a6009eb 100644 --- a/test/registered/vlm/test_paddleocr_vl_server.py +++ b/test/registered/vlm/test_paddleocr_vl_server.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.vlm_utils import TestOpenAIMLLMServerBase -register_cuda_ci(est_time=63, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=63, stage="base-b", runner_config="1-gpu-small") class TestPaddleOCRVLServer(TestOpenAIMLLMServerBase): diff --git a/test/registered/vlm/test_rust_native_mm_e2e.py b/test/registered/vlm/test_rust_native_mm_e2e.py index 32dea3420500..1f98aeae04bb 100644 --- a/test/registered/vlm/test_rust_native_mm_e2e.py +++ b/test/registered/vlm/test_rust_native_mm_e2e.py @@ -25,7 +25,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=43, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=43, stage="base-b", runner_config="1-gpu-small") IMAGE_URL = "https://raw.githubusercontent.com/sgl-project/sgl-test-files/refs/heads/main/images/man_ironing_on_back_of_suv.png" VISION_BLOCK = "<|vision_start|><|image_pad|><|vision_end|>" diff --git a/test/registered/vlm/test_rust_native_mm_mmmu.py b/test/registered/vlm/test_rust_native_mm_mmmu.py index 9d8ba7f0a54a..f32a5e623088 100644 --- a/test/registered/vlm/test_rust_native_mm_mmmu.py +++ b/test/registered/vlm/test_rust_native_mm_mmmu.py @@ -30,7 +30,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=126, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=300, stage="base-b", runner_config="1-gpu-small") MODEL = "Qwen/Qwen3.5-0.8B" VISION_BLOCK = "<|vision_start|><|image_pad|><|vision_end|>" diff --git a/test/registered/vlm/test_token_id_retokenize_e2e.py b/test/registered/vlm/test_token_id_retokenize_e2e.py index 419d09a7a1a5..26e4b2c03ab1 100644 --- a/test/registered/vlm/test_token_id_retokenize_e2e.py +++ b/test/registered/vlm/test_token_id_retokenize_e2e.py @@ -38,7 +38,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=92, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=92, stage="base-b", runner_config="1-gpu-small") register_cpu_ci(est_time=123, suite="stage-b-test-cpu-intel") From b09fa62e1ab98bcca034be9af64221e9b4ca518c Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 16:42:17 +0000 Subject: [PATCH 02/10] [CI] Move newly added 1-gpu-large tests that pass on the 5090 to 1-gpu-small Re-run the 48 base-b, base-b-kernel-unit and extra-a files registered on 1-gpu-large since the last sweep on an RTX 5090 with the CI environment. - base-b-kernel-unit: 41 files move; test_convrot_int8 was already registered on 1-gpu-small, so drop its 1-gpu-large duplicate. - base-b: 4 files move (both test_fp8_kernel splits and the two rust_server test_mm_transport_1node cells); extra-a: test_lilicorr_cuda. - Raise est_time to the measured 5090 duration where the estimate was more than 1.5x low (9 files). - Kernel lane: 4 shards on 1-gpu-small (was 3) to keep each shard near 30 min of est_time. Kept on 1-gpu-large: test_deep_select (SM90/SM100/SM103 only, skips on SM120) and test_flashattention_return_lse (sgl-kernel FA3 is SM90-only). --- .github/workflows/pr-test-jit-kernel.yml | 10 +++++++--- .../e2e/rust_server/test_mm_transport_1node_cpu.py | 2 +- .../rust_server/test_mm_transport_1node_cuda_ipc.py | 2 +- .../ops/attention/dsv4/test_swapab_attention.py | 2 +- .../ops/attention/qsa/test_qsa_fused_kv_prepare.py | 2 +- .../attention/test_dsa_triton_sparse_mla_prefill.py | 2 +- .../ops/attention/test_dsv4_trtllm_q_metadata.py | 2 +- .../ops/attention/test_fused_norm_gate_strided.py | 2 +- .../ops/attention/test_inkling_checkpoint_indices.py | 2 +- .../ops/attention/test_iquest_q1_learned_sink.py | 2 +- .../kernels/ops/attention/test_kda_prefill_graph.py | 2 +- .../ops/attention/test_kpool_topk_transform_fused.py | 2 +- .../attention/test_minimax_sparse_decode_paged_tile.py | 2 +- .../kernels/ops/attention/test_mla_scatter_concat.py | 2 +- .../kernels/ops/attention/test_replayssm_ring_fold.py | 2 +- .../ops/attention/test_trtllm_kv_norm_rope_store.py | 2 +- .../kernels/ops/attention/test_vision_rope.py | 2 +- .../ops/diffusion/test_flux2_strided_qknorm_rope.py | 2 +- .../kernels/ops/diffusion/test_h3_vae_fused_norm.py | 2 +- .../kernels/ops/diffusion/test_matrix_rope.py | 2 +- .../kernels/ops/diffusion/test_residual_gate_fp32.py | 2 +- .../kernels/ops/diffusion/test_shape_specialization.py | 2 +- test/registered/kernels/ops/elementwise/test_add3.py | 2 +- test/registered/kernels/ops/gemm/test_convrot_int8.py | 1 - test/registered/kernels/ops/gemm/test_fp8_kernel.py | 2 +- .../kernels/ops/gemm/test_tiny_gemm_variants.py | 2 +- .../ops/kvcache/test_hicache_strided_device_rows.py | 2 +- .../ops/kvcache/test_masked_set_kv_buffer_strided.py | 2 +- .../kernels/ops/kvcache/test_prefix_valid_fp8.py | 2 +- .../kernels/ops/layernorm/test_gemma4_qkv_norm.py | 2 +- .../kernels/ops/mamba/test_sconv_tracking.py | 2 +- test/registered/kernels/ops/memory/test_row_compact.py | 2 +- test/registered/kernels/ops/memory/test_small_copy.py | 2 +- .../ops/memory/test_translate_token_ids_strided.py | 2 +- test/registered/kernels/ops/mm/test_image.py | 2 +- .../kernels/ops/moe/test_deepep_v2_expand_kernels.py | 2 +- .../kernels/ops/moe/test_gemma4_expert_activation.py | 2 +- .../kernels/ops/moe/test_round_robin_ep_balance.py | 2 +- .../kernels/ops/moe/test_route_quant_prerequisites.py | 2 +- .../kernels/ops/quantization/test_fp8_kernel.py | 2 +- .../kernels/ops/sampling/test_top_p_renorm.py | 2 +- .../kernels/ops/speculative/test_block_verification.py | 2 +- .../test_draft_decode_kv_indices_read_table.py | 2 +- .../kernels/ops/speculative/test_lilicorr_cuda.py | 2 +- .../kernels/ops/speculative/test_next_verify_prep.py | 2 +- .../unit/dllm/algorithm/test_gemma4_sampling_cuda.py | 2 +- .../unit/mem_cache/test_prefix_valid_fp8_pool.py | 2 +- 47 files changed, 52 insertions(+), 49 deletions(-) diff --git a/.github/workflows/pr-test-jit-kernel.yml b/.github/workflows/pr-test-jit-kernel.yml index 32f01d016e64..6479bd57f025 100644 --- a/.github/workflows/pr-test-jit-kernel.yml +++ b/.github/workflows/pr-test-jit-kernel.yml @@ -58,15 +58,19 @@ jobs: - suite: unit-test-1-gpu-small runner_config: 1-gpu-small shard: " (0)" - suite_args: --auto-partition-id 0 --auto-partition-size 3 --fork-worker-batch-size 20 + suite_args: --auto-partition-id 0 --auto-partition-size 4 --fork-worker-batch-size 20 - suite: unit-test-1-gpu-small runner_config: 1-gpu-small shard: " (1)" - suite_args: --auto-partition-id 1 --auto-partition-size 3 --fork-worker-batch-size 20 + suite_args: --auto-partition-id 1 --auto-partition-size 4 --fork-worker-batch-size 20 - suite: unit-test-1-gpu-small runner_config: 1-gpu-small shard: " (2)" - suite_args: --auto-partition-id 2 --auto-partition-size 3 --fork-worker-batch-size 20 + suite_args: --auto-partition-id 2 --auto-partition-size 4 --fork-worker-batch-size 20 + - suite: unit-test-1-gpu-small + runner_config: 1-gpu-small + shard: " (3)" + suite_args: --auto-partition-id 3 --auto-partition-size 4 --fork-worker-batch-size 20 - suite: unit-test-4-gpu-b200 runner_config: 4-gpu-b200 - suite: unit-test-8-gpu-h200 diff --git a/test/registered/e2e/rust_server/test_mm_transport_1node_cpu.py b/test/registered/e2e/rust_server/test_mm_transport_1node_cpu.py index 21cf811a569a..f7f2322460a3 100644 --- a/test/registered/e2e/rust_server/test_mm_transport_1node_cpu.py +++ b/test/registered/e2e/rust_server/test_mm_transport_1node_cpu.py @@ -12,7 +12,7 @@ RustMmTransportServerBase, ) -register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-small") class TestRustMmTransport1NodeCpu(RustMmTransportServerBase): diff --git a/test/registered/e2e/rust_server/test_mm_transport_1node_cuda_ipc.py b/test/registered/e2e/rust_server/test_mm_transport_1node_cuda_ipc.py index 7954f9925f29..04f4018d1e6e 100644 --- a/test/registered/e2e/rust_server/test_mm_transport_1node_cuda_ipc.py +++ b/test/registered/e2e/rust_server/test_mm_transport_1node_cuda_ipc.py @@ -13,7 +13,7 @@ RustMmTransportServerBase, ) -register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-small") class TestRustMmTransport1NodeCudaIpc(RustMmTransportServerBase): diff --git a/test/registered/kernels/ops/attention/dsv4/test_swapab_attention.py b/test/registered/kernels/ops/attention/dsv4/test_swapab_attention.py index f9ab0a9bb4fe..cc2c1fb81d43 100644 --- a/test/registered/kernels/ops/attention/dsv4/test_swapab_attention.py +++ b/test/registered/kernels/ops/attention/dsv4/test_swapab_attention.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") # backend-specific: the HIP swap-AB Gluon kernel runs on gfx950. register_amd_ci(est_time=30, stage="stage-b", runner_config="1-gpu-small-amd-mi35x") diff --git a/test/registered/kernels/ops/attention/qsa/test_qsa_fused_kv_prepare.py b/test/registered/kernels/ops/attention/qsa/test_qsa_fused_kv_prepare.py index 6ae2cf39477c..d825fff04c2a 100644 --- a/test/registered/kernels/ops/attention/qsa/test_qsa_fused_kv_prepare.py +++ b/test/registered/kernels/ops/attention/qsa/test_qsa_fused_kv_prepare.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.parametrize("requests", [1, 3, 8, 128]) diff --git a/test/registered/kernels/ops/attention/test_dsa_triton_sparse_mla_prefill.py b/test/registered/kernels/ops/attention/test_dsa_triton_sparse_mla_prefill.py index f37f0cd6a19c..78e1eccb94be 100644 --- a/test/registered/kernels/ops/attention/test_dsa_triton_sparse_mla_prefill.py +++ b/test/registered/kernels/ops/attention/test_dsa_triton_sparse_mla_prefill.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=42, stage="base-b-kernel-unit", runner_config="1-gpu-small") D_QK, D_V, SM_SCALE = 576, 512, 0.0625 diff --git a/test/registered/kernels/ops/attention/test_dsv4_trtllm_q_metadata.py b/test/registered/kernels/ops/attention/test_dsv4_trtllm_q_metadata.py index 1cf9bcd7152a..27aa9f42dd32 100644 --- a/test/registered/kernels/ops/attention/test_dsv4_trtllm_q_metadata.py +++ b/test/registered/kernels/ops/attention/test_dsv4_trtllm_q_metadata.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.attention.dsv4.trtllm_metadata import pack_sparse_tail from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.mark.parametrize("rows,heads", [(6, 16), (9, 128), (4097, 16)]) diff --git a/test/registered/kernels/ops/attention/test_fused_norm_gate_strided.py b/test/registered/kernels/ops/attention/test_fused_norm_gate_strided.py index 9c163b86672a..f993f1aff364 100644 --- a/test/registered/kernels/ops/attention/test_fused_norm_gate_strided.py +++ b/test/registered/kernels/ops/attention/test_fused_norm_gate_strided.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.attention.fla.fused_norm_gate import FusedRMSNormGated 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") _HEADS, _HEAD_DIM, _PROJ_WIDTH = 12, 128, 6144 diff --git a/test/registered/kernels/ops/attention/test_inkling_checkpoint_indices.py b/test/registered/kernels/ops/attention/test_inkling_checkpoint_indices.py index 554ef6bd4db5..16900c887039 100644 --- a/test/registered/kernels/ops/attention/test_inkling_checkpoint_indices.py +++ b/test/registered/kernels/ops/attention/test_inkling_checkpoint_indices.py @@ -27,7 +27,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") _BACKEND = "sglang.srt.layers.attention.linear.inkling_sconv_backend" diff --git a/test/registered/kernels/ops/attention/test_iquest_q1_learned_sink.py b/test/registered/kernels/ops/attention/test_iquest_q1_learned_sink.py index 5caafb0a9809..2ce0ce5d9e25 100644 --- a/test/registered/kernels/ops/attention/test_iquest_q1_learned_sink.py +++ b/test/registered/kernels/ops/attention/test_iquest_q1_learned_sink.py @@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") def torch_reference(query, sink_key, attn_output, lse, scale): diff --git a/test/registered/kernels/ops/attention/test_kda_prefill_graph.py b/test/registered/kernels/ops/attention/test_kda_prefill_graph.py index 098ee3c2083e..02f6cd888830 100644 --- a/test/registered/kernels/ops/attention/test_kda_prefill_graph.py +++ b/test/registered/kernels/ops/attention/test_kda_prefill_graph.py @@ -24,7 +24,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") HEADS, HEAD_DIM, CONV = 2, 128, 4 DIM = 3 * HEADS * HEAD_DIM diff --git a/test/registered/kernels/ops/attention/test_kpool_topk_transform_fused.py b/test/registered/kernels/ops/attention/test_kpool_topk_transform_fused.py index d36f29d617b8..ce1c5789276b 100644 --- a/test/registered/kernels/ops/attention/test_kpool_topk_transform_fused.py +++ b/test/registered/kernels/ops/attention/test_kpool_topk_transform_fused.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-small") # backend-specific: compile the HIP JIT path and exercise wave64 graph replay. register_amd_ci(est_time=90, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/attention/test_minimax_sparse_decode_paged_tile.py b/test/registered/kernels/ops/attention/test_minimax_sparse_decode_paged_tile.py index 317e0c2c8507..4d3b8708b672 100644 --- a/test/registered/kernels/ops/attention/test_minimax_sparse_decode_paged_tile.py +++ b/test/registered/kernels/ops/attention/test_minimax_sparse_decode_paged_tile.py @@ -14,7 +14,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="4-gpu-b200") BS, NQH, NKH, HD, BLK, TOPK = 4, 64, 1, 128, 128, 16 diff --git a/test/registered/kernels/ops/attention/test_mla_scatter_concat.py b/test/registered/kernels/ops/attention/test_mla_scatter_concat.py index fb2b380448c4..bc01f8b66daf 100644 --- a/test/registered/kernels/ops/attention/test_mla_scatter_concat.py +++ b/test/registered/kernels/ops/attention/test_mla_scatter_concat.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=36, stage="base-b-kernel-unit", runner_config="1-gpu-small") NOPE_DIM = 512 diff --git a/test/registered/kernels/ops/attention/test_replayssm_ring_fold.py b/test/registered/kernels/ops/attention/test_replayssm_ring_fold.py index 716b3eb1c7db..63c89eee614b 100644 --- a/test/registered/kernels/ops/attention/test_replayssm_ring_fold.py +++ b/test/registered/kernels/ops/attention/test_replayssm_ring_fold.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestReplaySSMRingFold(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_trtllm_kv_norm_rope_store.py b/test/registered/kernels/ops/attention/test_trtllm_kv_norm_rope_store.py index 6204405fdafc..88bdcd598d7f 100644 --- a/test/registered/kernels/ops/attention/test_trtllm_kv_norm_rope_store.py +++ b/test/registered/kernels/ops/attention/test_trtllm_kv_norm_rope_store.py @@ -11,7 +11,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") def inputs(n, dtype=torch.bfloat16, strided=False, index_dtype=torch.int64): diff --git a/test/registered/kernels/ops/attention/test_vision_rope.py b/test/registered/kernels/ops/attention/test_vision_rope.py index e504fae887aa..3dd457d600c1 100644 --- a/test/registered/kernels/ops/attention/test_vision_rope.py +++ b/test/registered/kernels/ops/attention/test_vision_rope.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestVisionRoPE(CustomTestCase): diff --git a/test/registered/kernels/ops/diffusion/test_flux2_strided_qknorm_rope.py b/test/registered/kernels/ops/diffusion/test_flux2_strided_qknorm_rope.py index 40ccd14fe476..4d4b9952b130 100644 --- a/test/registered/kernels/ops/diffusion/test_flux2_strided_qknorm_rope.py +++ b/test/registered/kernels/ops/diffusion/test_flux2_strided_qknorm_rope.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") def inputs(batch, tokens): diff --git a/test/registered/kernels/ops/diffusion/test_h3_vae_fused_norm.py b/test/registered/kernels/ops/diffusion/test_h3_vae_fused_norm.py index b7315f8a13a3..8ec05c116977 100644 --- a/test/registered/kernels/ops/diffusion/test_h3_vae_fused_norm.py +++ b/test/registered/kernels/ops/diffusion/test_h3_vae_fused_norm.py @@ -16,7 +16,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") @torch.inference_mode() diff --git a/test/registered/kernels/ops/diffusion/test_matrix_rope.py b/test/registered/kernels/ops/diffusion/test_matrix_rope.py index 0ba3c2331887..7cadcb1960b0 100644 --- a/test/registered/kernels/ops/diffusion/test_matrix_rope.py +++ b/test/registered/kernels/ops/diffusion/test_matrix_rope.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.diffusion import apply_matrix_rope from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_residual_gate_fp32.py b/test/registered/kernels/ops/diffusion/test_residual_gate_fp32.py index 1087928bfdcf..9b2fc49b5f10 100644 --- a/test/registered/kernels/ops/diffusion/test_residual_gate_fp32.py +++ b/test/registered/kernels/ops/diffusion/test_residual_gate_fp32.py @@ -7,7 +7,7 @@ from sglang.kernels.ops.diffusion import residual_gate_fp32 from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif( not torch.cuda.is_available() or torch.version.hip is not None, diff --git a/test/registered/kernels/ops/diffusion/test_shape_specialization.py b/test/registered/kernels/ops/diffusion/test_shape_specialization.py index 7811820f55be..bdf9add710e2 100644 --- a/test/registered/kernels/ops/diffusion/test_shape_specialization.py +++ b/test/registered/kernels/ops/diffusion/test_shape_specialization.py @@ -41,7 +41,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" HIDDEN = 512 diff --git a/test/registered/kernels/ops/elementwise/test_add3.py b/test/registered/kernels/ops/elementwise/test_add3.py index edb78dd38d5b..cbd7266426ff 100644 --- a/test/registered/kernels/ops/elementwise/test_add3.py +++ b/test/registered/kernels/ops/elementwise/test_add3.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestAdd3(CustomTestCase): diff --git a/test/registered/kernels/ops/gemm/test_convrot_int8.py b/test/registered/kernels/ops/gemm/test_convrot_int8.py index 6959cc59f251..105e6e94454c 100644 --- a/test/registered/kernels/ops/gemm/test_convrot_int8.py +++ b/test/registered/kernels/ops/gemm/test_convrot_int8.py @@ -16,7 +16,6 @@ from sglang.test.ci.ci_register import register_cuda_ci # One registration per GEMM path: H100 (sm_90a), B200 (sm_100a), RTX 5090 (CC 12.0). -register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="4-gpu-b200") register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-small") diff --git a/test/registered/kernels/ops/gemm/test_fp8_kernel.py b/test/registered/kernels/ops/gemm/test_fp8_kernel.py index 3042fc79259b..a8780c0f8eae 100644 --- a/test/registered/kernels/ops/gemm/test_fp8_kernel.py +++ b/test/registered/kernels/ops/gemm/test_fp8_kernel.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.kernels.fp8 import TestFP8Base -register_cuda_ci(est_time=5, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") _is_cuda = is_cuda() diff --git a/test/registered/kernels/ops/gemm/test_tiny_gemm_variants.py b/test/registered/kernels/ops/gemm/test_tiny_gemm_variants.py index 70bc7e7712a7..a687b1a655c3 100644 --- a/test/registered/kernels/ops/gemm/test_tiny_gemm_variants.py +++ b/test/registered/kernels/ops/gemm/test_tiny_gemm_variants.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestTinyGemmVariants(CustomTestCase): diff --git a/test/registered/kernels/ops/kvcache/test_hicache_strided_device_rows.py b/test/registered/kernels/ops/kvcache/test_hicache_strided_device_rows.py index c9f3b7476b94..5cdd3b23e2a8 100644 --- a/test/registered/kernels/ops/kvcache/test_hicache_strided_device_rows.py +++ b/test/registered/kernels/ops/kvcache/test_hicache_strided_device_rows.py @@ -18,7 +18,7 @@ from sglang.srt.mem_cache.pool_host.mha import MHATokenToKVPoolHost from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") _L, _H, _D, _ROWS = 2, 2, 64, 64 # One slot holds every layer's K and V plus a gap, as a unified-pool entry. diff --git a/test/registered/kernels/ops/kvcache/test_masked_set_kv_buffer_strided.py b/test/registered/kernels/ops/kvcache/test_masked_set_kv_buffer_strided.py index 9a4f17393d46..08cb36bd72ce 100644 --- a/test/registered/kernels/ops/kvcache/test_masked_set_kv_buffer_strided.py +++ b/test/registered/kernels/ops/kvcache/test_masked_set_kv_buffer_strided.py @@ -26,7 +26,7 @@ from sglang.srt.mem_cache.memory_pool import masked_set_kv_buffer_kernel from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestMaskedSetKvBufferStrided(unittest.TestCase): diff --git a/test/registered/kernels/ops/kvcache/test_prefix_valid_fp8.py b/test/registered/kernels/ops/kvcache/test_prefix_valid_fp8.py index 2beb9ab587b4..b6e9b064ae17 100644 --- a/test/registered/kernels/ops/kvcache/test_prefix_valid_fp8.py +++ b/test/registered/kernels/ops/kvcache/test_prefix_valid_fp8.py @@ -10,7 +10,7 @@ from sglang.test.kernels.prefix_valid import assert_prefix_commit, make_kv_cache from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") # backend-specific: ROCm's E4M3FNUZ pointer inference and conversion differ from CUDA. register_amd_ci(est_time=25, stage="jit-kernel-unit", runner_config="amd") diff --git a/test/registered/kernels/ops/layernorm/test_gemma4_qkv_norm.py b/test/registered/kernels/ops/layernorm/test_gemma4_qkv_norm.py index 023f87b2d9f2..cfab4b417e88 100644 --- a/test/registered/kernels/ops/layernorm/test_gemma4_qkv_norm.py +++ b/test/registered/kernels/ops/layernorm/test_gemma4_qkv_norm.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/kernels/ops/mamba/test_sconv_tracking.py b/test/registered/kernels/ops/mamba/test_sconv_tracking.py index be659c8f73a4..2cdbf9734e45 100644 --- a/test/registered/kernels/ops/mamba/test_sconv_tracking.py +++ b/test/registered/kernels/ops/mamba/test_sconv_tracking.py @@ -6,7 +6,7 @@ from sglang.kernels.ops.mamba.sconv_tracking import fill_track_conv_indices from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA only") diff --git a/test/registered/kernels/ops/memory/test_row_compact.py b/test/registered/kernels/ops/memory/test_row_compact.py index 4f9aa9bcc925..b220f8ce5d52 100644 --- a/test/registered/kernels/ops/memory/test_row_compact.py +++ b/test/registered/kernels/ops/memory/test_row_compact.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="4-gpu-b200") diff --git a/test/registered/kernels/ops/memory/test_small_copy.py b/test/registered/kernels/ops/memory/test_small_copy.py index 760dd6c42f79..880af39c864c 100644 --- a/test/registered/kernels/ops/memory/test_small_copy.py +++ b/test/registered/kernels/ops/memory/test_small_copy.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") _DTYPES = ( torch.bool, diff --git a/test/registered/kernels/ops/memory/test_translate_token_ids_strided.py b/test/registered/kernels/ops/memory/test_translate_token_ids_strided.py index 96ebca157243..30430d400eb7 100644 --- a/test/registered/kernels/ops/memory/test_translate_token_ids_strided.py +++ b/test/registered/kernels/ops/memory/test_translate_token_ids_strided.py @@ -28,7 +28,7 @@ from sglang.kernels.ops.memory.virtual_slot import translate_token_ids from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "needs CUDA") diff --git a/test/registered/kernels/ops/mm/test_image.py b/test/registered/kernels/ops/mm/test_image.py index f1e66754c4dd..da5868fa2f86 100644 --- a/test/registered/kernels/ops/mm/test_image.py +++ b/test/registered/kernels/ops/mm/test_image.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestImagePatchify(CustomTestCase): diff --git a/test/registered/kernels/ops/moe/test_deepep_v2_expand_kernels.py b/test/registered/kernels/ops/moe/test_deepep_v2_expand_kernels.py index 5192db0ef5c1..681194bdb227 100644 --- a/test/registered/kernels/ops/moe/test_deepep_v2_expand_kernels.py +++ b/test/registered/kernels/ops/moe/test_deepep_v2_expand_kernels.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" diff --git a/test/registered/kernels/ops/moe/test_gemma4_expert_activation.py b/test/registered/kernels/ops/moe/test_gemma4_expert_activation.py index 92044d79887c..c1ec8bbd78cc 100644 --- a/test/registered/kernels/ops/moe/test_gemma4_expert_activation.py +++ b/test/registered/kernels/ops/moe/test_gemma4_expert_activation.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=27, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/kernels/ops/moe/test_round_robin_ep_balance.py b/test/registered/kernels/ops/moe/test_round_robin_ep_balance.py index 08060747a36a..5878a586cf5e 100644 --- a/test/registered/kernels/ops/moe/test_round_robin_ep_balance.py +++ b/test/registered/kernels/ops/moe/test_round_robin_ep_balance.py @@ -15,7 +15,7 @@ from sglang.srt.runtime_context import get_forward, get_parallel 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") NUM_EXPERTS = 256 TOPK = 8 diff --git a/test/registered/kernels/ops/moe/test_route_quant_prerequisites.py b/test/registered/kernels/ops/moe/test_route_quant_prerequisites.py index 57ae938e73a7..525f470d49fc 100644 --- a/test/registered/kernels/ops/moe/test_route_quant_prerequisites.py +++ b/test/registered/kernels/ops/moe/test_route_quant_prerequisites.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="1-gpu-small") NUM_EXPERTS = 896 diff --git a/test/registered/kernels/ops/quantization/test_fp8_kernel.py b/test/registered/kernels/ops/quantization/test_fp8_kernel.py index 38afd872fa8e..3bed97449382 100644 --- a/test/registered/kernels/ops/quantization/test_fp8_kernel.py +++ b/test/registered/kernels/ops/quantization/test_fp8_kernel.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.kernels.fp8 import TestFP8Base -register_cuda_ci(est_time=4, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-small") _is_cuda = is_cuda() diff --git a/test/registered/kernels/ops/sampling/test_top_p_renorm.py b/test/registered/kernels/ops/sampling/test_top_p_renorm.py index 49a9538a5fd2..e586a6b0acb1 100644 --- a/test/registered/kernels/ops/sampling/test_top_p_renorm.py +++ b/test/registered/kernels/ops/sampling/test_top_p_renorm.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-small") class TestTopPRenorm(CustomTestCase): diff --git a/test/registered/kernels/ops/speculative/test_block_verification.py b/test/registered/kernels/ops/speculative/test_block_verification.py index 2e69343d7b5d..6c34a6fb9b40 100644 --- a/test/registered/kernels/ops/speculative/test_block_verification.py +++ b/test/registered/kernels/ops/speculative/test_block_verification.py @@ -15,7 +15,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cpu" if os.environ.get("TRITON_INTERPRET") == "1" else "cuda" diff --git a/test/registered/kernels/ops/speculative/test_draft_decode_kv_indices_read_table.py b/test/registered/kernels/ops/speculative/test_draft_decode_kv_indices_read_table.py index 97175e9c24f5..5f504b3317c3 100644 --- a/test/registered/kernels/ops/speculative/test_draft_decode_kv_indices_read_table.py +++ b/test/registered/kernels/ops/speculative/test_draft_decode_kv_indices_read_table.py @@ -29,7 +29,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") _SENTINEL = -7 diff --git a/test/registered/kernels/ops/speculative/test_lilicorr_cuda.py b/test/registered/kernels/ops/speculative/test_lilicorr_cuda.py index 9e305bf096c0..f2627e6ef2e2 100644 --- a/test/registered/kernels/ops/speculative/test_lilicorr_cuda.py +++ b/test/registered/kernels/ops/speculative/test_lilicorr_cuda.py @@ -12,7 +12,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="extra-a", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available(), reason="LiLiCorr Triton kernels require CUDA" diff --git a/test/registered/kernels/ops/speculative/test_next_verify_prep.py b/test/registered/kernels/ops/speculative/test_next_verify_prep.py index d518d39f2c2e..98277f394e23 100644 --- a/test/registered/kernels/ops/speculative/test_next_verify_prep.py +++ b/test/registered/kernels/ops/speculative/test_next_verify_prep.py @@ -3,7 +3,7 @@ from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") from sglang.kernels.ops.speculative.eagle import ( prepare_draft_extend_inputs, diff --git a/test/registered/unit/dllm/algorithm/test_gemma4_sampling_cuda.py b/test/registered/unit/dllm/algorithm/test_gemma4_sampling_cuda.py index 8b663051d07a..10d3e9263d67 100644 --- a/test/registered/unit/dllm/algorithm/test_gemma4_sampling_cuda.py +++ b/test/registered/unit/dllm/algorithm/test_gemma4_sampling_cuda.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=19, stage="base-b-kernel-unit", runner_config="1-gpu-small") @unittest.skipUnless(torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/unit/mem_cache/test_prefix_valid_fp8_pool.py b/test/registered/unit/mem_cache/test_prefix_valid_fp8_pool.py index 8b5b4f5bc600..da8d8402ed2b 100644 --- a/test/registered/unit/mem_cache/test_prefix_valid_fp8_pool.py +++ b/test/registered/unit/mem_cache/test_prefix_valid_fp8_pool.py @@ -13,7 +13,7 @@ from sglang.test.test_utils import CustomTestCase # Preserve the original GPU lanes; the kernel file retains the other 25 seconds. -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") # backend-specific: loaded FP8 scales are doubled for ROCm's E4M3FNUZ format. register_amd_ci(est_time=15, stage="jit-kernel-unit", runner_config="amd") From 5c0dfbe4e8dad85f59ec74396db9110029fb3b1c Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 17:10:57 +0000 Subject: [PATCH 03/10] [CI] Move JIT kernel benchmarks that run on the 5090 to 1-gpu-small Run all 67 base-b-kernel-benchmark files on an RTX 5090 with the pr-test-jit-kernel.yml env, and check each log for empty result tables as well as exit codes (marker.skip leaves a table with no rows). - 62 files move to a new base-b-kernel-benchmark-test-1-gpu-small job (25.8 min est_time); 1-gpu-large keeps 5 (4.8 min). - Raise est_time to the measured 5090 duration where the estimate was more than 1.5x low (24 files). - Point new benchmarks at 1-gpu-small in test/README.md and the write-sglang-test / add-jit-kernel / ci-workflow-guide skills. Kept on 1-gpu-large: bench_sparse_mla_q8kv8_prefill_sm90 (SM90a/SM100f only), bench_topk (TopKPageTableTransform not supported on SM120), bench_hicache (TMA path: CUDA invalid argument on SM120), bench_fp4_index_logits and bench_bf16xfp32_router_gemm (Hopper only, produce no rows elsewhere). --- .github/workflows/pr-test-jit-kernel.yml | 2 ++ test/README.md | 4 ++-- .../kernels/benchmark/activation/bench_activation.py | 2 +- .../kernels/benchmark/attention/bench_clamp_position.py | 2 +- .../kernels/benchmark/attention/bench_concat_mla.py | 2 +- .../attention/bench_dsa_triton_sparse_mla_prefill.py | 2 +- .../kernels/benchmark/attention/bench_dsv4_c2_decode.py | 2 +- .../kernels/benchmark/attention/bench_fused_qknorm_rope.py | 2 +- .../kernels/benchmark/attention/bench_minimax_decode_topk.py | 2 +- .../kernels/benchmark/attention/bench_minimax_qknorm_rope.py | 2 +- .../benchmark/attention/bench_mla_kv_pack_quantize_fp8.py | 2 +- .../kernels/benchmark/attention/bench_online_c128_mtp.py | 2 +- test/registered/kernels/benchmark/attention/bench_rope.py | 2 +- .../benchmark/attention/minicpm_sala/bench_get_block_table.py | 2 +- .../kernels/benchmark/diffusion/bench_cosmos3_edge_fusions.py | 2 +- .../kernels/benchmark/diffusion/bench_flux2_fp8_norm_quant.py | 2 +- .../kernels/benchmark/diffusion/bench_flux2_token_cat_fp8.py | 2 +- .../kernels/benchmark/diffusion/bench_fp8_rowwise.py | 2 +- .../kernels/benchmark/diffusion/bench_helios_qk_rope.py | 2 +- .../benchmark/diffusion/bench_interleaved_rope_fp64.py | 2 +- .../kernels/benchmark/diffusion/bench_ltx25_decoder_rope.py | 2 +- .../kernels/benchmark/diffusion/bench_qknorm_rope.py | 2 +- .../benchmark/diffusion/bench_qwen_image_modulation.py | 2 +- .../benchmark/diffusion/bench_qwen_image_norm_fp8_quant.py | 2 +- .../kernels/benchmark/diffusion/bench_residual_gate_add.py | 2 +- .../kernels/benchmark/diffusion/bench_scale_shift_4d.py | 2 +- .../kernels/benchmark/diffusion/bench_timestep_embedding.py | 2 +- .../benchmark/diffusion/bench_varlen_segmented_pack.py | 2 +- .../kernels/benchmark/diffusion/bench_vdn_delta_factors.py | 2 +- .../kernels/benchmark/elementwise/bench_add_constant.py | 2 +- .../kernels/benchmark/elementwise/bench_bias_gelu.py | 2 +- .../benchmark/embeddings/bench_vocab_parallel_embedding.py | 2 +- test/registered/kernels/benchmark/gemm/bench_convrot_int8.py | 2 +- .../kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py | 2 +- .../kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py | 4 ++-- test/registered/kernels/benchmark/gemm/bench_tiny_gemm.py | 2 +- test/registered/kernels/benchmark/kv_canary/bench_write.py | 2 +- .../kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py | 2 +- test/registered/kernels/benchmark/kvcache/bench_hisparse.py | 2 +- .../kernels/benchmark/kvcache/bench_hisparse_spec.py | 2 +- .../kernels/benchmark/kvcache/bench_minimax_store_kv_index.py | 2 +- .../kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py | 2 +- .../registered/kernels/benchmark/kvcache/bench_store_cache.py | 2 +- .../kernels/benchmark/layernorm/bench_fused_eh_norm.py | 2 +- test/registered/kernels/benchmark/layernorm/bench_norm.py | 2 +- test/registered/kernels/benchmark/layernorm/bench_qknorm.py | 2 +- .../kernels/benchmark/layernorm/bench_qknorm_across_heads.py | 2 +- .../registered/kernels/benchmark/mamba/bench_causal_conv1d.py | 2 +- .../benchmark/moe/bench_fp8_per_token_to_per_tensor_quant.py | 2 +- test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py | 2 +- .../kernels/benchmark/moe/bench_moe_topk_softmax.py | 2 +- .../kernels/benchmark/moe/bench_post_reorder_deepgemm.py | 2 +- .../kernels/benchmark/moe/bench_shuffle_rows_with_scales.py | 2 +- .../kernels/benchmark/quantization/bench_awq_dequantize.py | 2 +- .../kernels/benchmark/quantization/bench_hadamard.py | 2 +- .../benchmark/quantization/bench_per_tensor_quant_fp8.py | 2 +- .../benchmark/quantization/bench_per_token_group_quant.py | 2 +- .../quantization/bench_per_token_group_quant_8bit_v2.py | 2 +- .../quantization/bench_per_token_group_quant_masked.py | 2 +- .../benchmark/quantization/bench_per_token_quant_fp8.py | 2 +- test/registered/kernels/benchmark/sampling/bench_renorm.py | 2 +- .../benchmark/speculative/bench_ngram_compute_decode.py | 2 +- .../benchmark/speculative/bench_ngram_update_token_table.py | 2 +- .../kernels/benchmark/speculative/bench_spec_topk1.py | 2 +- test/run_suite.py | 1 + 65 files changed, 68 insertions(+), 65 deletions(-) diff --git a/.github/workflows/pr-test-jit-kernel.yml b/.github/workflows/pr-test-jit-kernel.yml index 6479bd57f025..356e3808baf5 100644 --- a/.github/workflows/pr-test-jit-kernel.yml +++ b/.github/workflows/pr-test-jit-kernel.yml @@ -79,6 +79,8 @@ jobs: skip_health_check: true - suite: benchmark-test-1-gpu-large runner_config: 1-gpu-large + - suite: benchmark-test-1-gpu-small + runner_config: 1-gpu-small runs-on: ${{ fromJson(inputs.runs_on_map)[matrix.runner_config] }} timeout-minutes: 240 steps: diff --git a/test/README.md b/test/README.md index d6b0944f217a..c1a16ffa3086 100644 --- a/test/README.md +++ b/test/README.md @@ -91,7 +91,7 @@ stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path: - Correctness tests: `test/registered/kernels/ops//test_*.py` → `base-b-kernel-unit-test-1-gpu-small` (use `base-b-kernel-unit-test-1-gpu-large` only for SM90-only kernels or shapes that exceed the 5090's ~99 KB shared memory) -- Benchmarks: `test/registered/kernels/benchmark//bench_*.py` → `base-b-kernel-benchmark-test-1-gpu-large` +- Benchmarks: `test/registered/kernels/benchmark//bench_*.py` → `base-b-kernel-benchmark-test-1-gpu-small` (use `base-b-kernel-benchmark-test-1-gpu-large` only for SM90-only kernels) ## Choosing a Suite @@ -103,7 +103,7 @@ Use the lightest suite that meets your test's needs. Full suite tables are in th | Small GPU (fits 5090, 32GB) | `base-b-test-1-gpu-small` (most tests go here) | | Large GPU memory or Hopper features | `base-b-test-1-gpu-large` | | JIT kernel correctness | `base-b-kernel-unit-test-1-gpu-small` (SM90-only kernels: `base-b-kernel-unit-test-1-gpu-large`) | -| JIT kernel benchmarks | `base-b-kernel-benchmark-test-1-gpu-large` | +| JIT kernel benchmarks | `base-b-kernel-benchmark-test-1-gpu-small` (SM90-only kernels: `base-b-kernel-benchmark-test-1-gpu-large`) | | Multi-GPU (2/4/8) | `base-b-test-2-gpu-large`, `base-c-test-*` | | Long-running or experimental | `nightly-*` suites | diff --git a/test/registered/kernels/benchmark/activation/bench_activation.py b/test/registered/kernels/benchmark/activation/bench_activation.py index 983ef269f75b..e84be76f69ed 100644 --- a/test/registered/kernels/benchmark/activation/bench_activation.py +++ b/test/registered/kernels/benchmark/activation/bench_activation.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=30, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/attention/bench_clamp_position.py b/test/registered/kernels/benchmark/attention/bench_clamp_position.py index 393479683be8..084f3dc237c1 100644 --- a/test/registered/kernels/benchmark/attention/bench_clamp_position.py +++ b/test/registered/kernels/benchmark/attention/bench_clamp_position.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=16, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/attention/bench_concat_mla.py b/test/registered/kernels/benchmark/attention/bench_concat_mla.py index fb6183f65fcd..216e2e46b9b5 100644 --- a/test/registered/kernels/benchmark/attention/bench_concat_mla.py +++ b/test/registered/kernels/benchmark/attention/bench_concat_mla.py @@ -13,7 +13,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) IS_CI = is_in_ci() diff --git a/test/registered/kernels/benchmark/attention/bench_dsa_triton_sparse_mla_prefill.py b/test/registered/kernels/benchmark/attention/bench_dsa_triton_sparse_mla_prefill.py index 00c5cdda5102..08b5e55c0f5e 100644 --- a/test/registered/kernels/benchmark/attention/bench_dsa_triton_sparse_mla_prefill.py +++ b/test/registered/kernels/benchmark/attention/bench_dsa_triton_sparse_mla_prefill.py @@ -46,7 +46,7 @@ def _flashmla_runnable() -> bool: HAS_FLASHMLA = _flashmla_runnable() register_cuda_ci( - est_time=90, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=90, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) IS_CI = is_in_ci() diff --git a/test/registered/kernels/benchmark/attention/bench_dsv4_c2_decode.py b/test/registered/kernels/benchmark/attention/bench_dsv4_c2_decode.py index 5bb050b59e3b..2edabca20773 100644 --- a/test/registered/kernels/benchmark/attention/bench_dsv4_c2_decode.py +++ b/test/registered/kernels/benchmark/attention/bench_dsv4_c2_decode.py @@ -30,7 +30,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) HEAD_DIM = 512 diff --git a/test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py b/test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py index 2d80a4dfc6c0..7b483beadbb6 100644 --- a/test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py +++ b/test/registered/kernels/benchmark/attention/bench_fused_qknorm_rope.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=32, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) try: diff --git a/test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py b/test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py index 215279cdd31a..4349b3d297b7 100644 --- a/test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py +++ b/test/registered/kernels/benchmark/attention/bench_minimax_decode_topk.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=27, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=8, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py b/test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py index fdb98a36583a..2bd950837922 100644 --- a/test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py +++ b/test/registered/kernels/benchmark/attention/bench_minimax_qknorm_rope.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) HEAD_DIM, ROTARY_DIM, BASE, EPS, MAXPOS = 128, 64, 5_000_000, 1e-6, 131072 diff --git a/test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py b/test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py index cc0da748d741..10d61c41f6b6 100644 --- a/test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py +++ b/test/registered/kernels/benchmark/attention/bench_mla_kv_pack_quantize_fp8.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=15, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py b/test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py index 80376b8dd5dd..a5bea27cd664 100644 --- a/test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py +++ b/test/registered/kernels/benchmark/attention/bench_online_c128_mtp.py @@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=10, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/attention/bench_rope.py b/test/registered/kernels/benchmark/attention/bench_rope.py index 8091e7e65476..dc292e11d129 100644 --- a/test/registered/kernels/benchmark/attention/bench_rope.py +++ b/test/registered/kernels/benchmark/attention/bench_rope.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=22, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) MAX_SEQ_LEN = 131072 diff --git a/test/registered/kernels/benchmark/attention/minicpm_sala/bench_get_block_table.py b/test/registered/kernels/benchmark/attention/minicpm_sala/bench_get_block_table.py index a83f1c7f0c03..4de273d131e2 100644 --- a/test/registered/kernels/benchmark/attention/minicpm_sala/bench_get_block_table.py +++ b/test/registered/kernels/benchmark/attention/minicpm_sala/bench_get_block_table.py @@ -5,7 +5,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) _HEAD_GROUP = 2 diff --git a/test/registered/kernels/benchmark/diffusion/bench_cosmos3_edge_fusions.py b/test/registered/kernels/benchmark/diffusion/bench_cosmos3_edge_fusions.py index 0a134c845b4f..537c10548d73 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_cosmos3_edge_fusions.py +++ b/test/registered/kernels/benchmark/diffusion/bench_cosmos3_edge_fusions.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_flux2_fp8_norm_quant.py b/test/registered/kernels/benchmark/diffusion/bench_flux2_fp8_norm_quant.py index 3d229725ac59..ee65d8bc3693 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_flux2_fp8_norm_quant.py +++ b/test/registered/kernels/benchmark/diffusion/bench_flux2_fp8_norm_quant.py @@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=68, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) DEVICE = "cuda" diff --git a/test/registered/kernels/benchmark/diffusion/bench_flux2_token_cat_fp8.py b/test/registered/kernels/benchmark/diffusion/bench_flux2_token_cat_fp8.py index 6b36d943495f..d55026fefbbd 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_flux2_token_cat_fp8.py +++ b/test/registered/kernels/benchmark/diffusion/bench_flux2_token_cat_fp8.py @@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_fp8_rowwise.py b/test/registered/kernels/benchmark/diffusion/bench_fp8_rowwise.py index 4fb565331677..f88e8dc6f7d7 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_fp8_rowwise.py +++ b/test/registered/kernels/benchmark/diffusion/bench_fp8_rowwise.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_helios_qk_rope.py b/test/registered/kernels/benchmark/diffusion/bench_helios_qk_rope.py index 9108e499ea81..1ae02d556cc9 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_helios_qk_rope.py +++ b/test/registered/kernels/benchmark/diffusion/bench_helios_qk_rope.py @@ -7,7 +7,7 @@ register_cuda_ci( est_time=12, stage="base-b-kernel-benchmark", - runner_config="1-gpu-large", + runner_config="1-gpu-small", ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_interleaved_rope_fp64.py b/test/registered/kernels/benchmark/diffusion/bench_interleaved_rope_fp64.py index 807a8176f0ec..b22af96f7eea 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_interleaved_rope_fp64.py +++ b/test/registered/kernels/benchmark/diffusion/bench_interleaved_rope_fp64.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_ltx25_decoder_rope.py b/test/registered/kernels/benchmark/diffusion/bench_ltx25_decoder_rope.py index 308e15278ce2..bbb0dc6419e8 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_ltx25_decoder_rope.py +++ b/test/registered/kernels/benchmark/diffusion/bench_ltx25_decoder_rope.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=37, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py b/test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py index 3a73d39f36a2..b3ee97f003d4 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py +++ b/test/registered/kernels/benchmark/diffusion/bench_qknorm_rope.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=27, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) MAX_SEQ_LEN = 131072 diff --git a/test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py b/test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py index 83d00aebdc01..61dcbce2b4c4 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py +++ b/test/registered/kernels/benchmark/diffusion/bench_qwen_image_modulation.py @@ -14,7 +14,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=13, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/diffusion/bench_qwen_image_norm_fp8_quant.py b/test/registered/kernels/benchmark/diffusion/bench_qwen_image_norm_fp8_quant.py index 2ff550f815ea..17f5eddb2043 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_qwen_image_norm_fp8_quant.py +++ b/test/registered/kernels/benchmark/diffusion/bench_qwen_image_norm_fp8_quant.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) DEVICE = "cuda" diff --git a/test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py b/test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py index 4e66961ce62b..6fb05fbcaa90 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py +++ b/test/registered/kernels/benchmark/diffusion/bench_residual_gate_add.py @@ -9,7 +9,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_scale_shift_4d.py b/test/registered/kernels/benchmark/diffusion/bench_scale_shift_4d.py index dfc02a892fea..8538a80eaac3 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_scale_shift_4d.py +++ b/test/registered/kernels/benchmark/diffusion/bench_scale_shift_4d.py @@ -10,7 +10,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/diffusion/bench_timestep_embedding.py b/test/registered/kernels/benchmark/diffusion/bench_timestep_embedding.py index 140c43622cd6..2c75d71e9673 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_timestep_embedding.py +++ b/test/registered/kernels/benchmark/diffusion/bench_timestep_embedding.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) BATCHES = [1, 2, 8, 63, 256, 512, 613, 1024, 1536] diff --git a/test/registered/kernels/benchmark/diffusion/bench_varlen_segmented_pack.py b/test/registered/kernels/benchmark/diffusion/bench_varlen_segmented_pack.py index 3c55ba019e88..3ec1b9da29de 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_varlen_segmented_pack.py +++ b/test/registered/kernels/benchmark/diffusion/bench_varlen_segmented_pack.py @@ -5,7 +5,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) DEVICE = "cuda" diff --git a/test/registered/kernels/benchmark/diffusion/bench_vdn_delta_factors.py b/test/registered/kernels/benchmark/diffusion/bench_vdn_delta_factors.py index ee84cebe34f0..457ced221d01 100644 --- a/test/registered/kernels/benchmark/diffusion/bench_vdn_delta_factors.py +++ b/test/registered/kernels/benchmark/diffusion/bench_vdn_delta_factors.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) HEAD_DIM = 128 diff --git a/test/registered/kernels/benchmark/elementwise/bench_add_constant.py b/test/registered/kernels/benchmark/elementwise/bench_add_constant.py index a9996b6733df..f9d2902b0c31 100644 --- a/test/registered/kernels/benchmark/elementwise/bench_add_constant.py +++ b/test/registered/kernels/benchmark/elementwise/bench_add_constant.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=15, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/elementwise/bench_bias_gelu.py b/test/registered/kernels/benchmark/elementwise/bench_bias_gelu.py index 5c622b11d655..a033add63b6e 100644 --- a/test/registered/kernels/benchmark/elementwise/bench_bias_gelu.py +++ b/test/registered/kernels/benchmark/elementwise/bench_bias_gelu.py @@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py b/test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py index 6b332f875c98..c0e1a424692e 100644 --- a/test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py +++ b/test/registered/kernels/benchmark/embeddings/bench_vocab_parallel_embedding.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=10, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/gemm/bench_convrot_int8.py b/test/registered/kernels/benchmark/gemm/bench_convrot_int8.py index b4afa1e60375..cc0d13208fd0 100644 --- a/test/registered/kernels/benchmark/gemm/bench_convrot_int8.py +++ b/test/registered/kernels/benchmark/gemm/bench_convrot_int8.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=120, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=120, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) GROUP_SIZE = 256 diff --git a/test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py b/test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py index 36b976ceabe4..f97c30617591 100644 --- a/test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py +++ b/test/registered/kernels/benchmark/gemm/bench_dsv3_fused_a_gemm.py @@ -18,7 +18,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) IS_CI = is_in_ci() diff --git a/test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py b/test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py index 1df2da96d2c2..1a98ed540ae2 100644 --- a/test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py +++ b/test/registered/kernels/benchmark/gemm/bench_fp8_blockwise_gemm.py @@ -11,9 +11,9 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=5, + est_time=97, stage="base-b-kernel-benchmark", - runner_config="1-gpu-large", + runner_config="1-gpu-small", ) register_amd_ci(est_time=5, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/gemm/bench_tiny_gemm.py b/test/registered/kernels/benchmark/gemm/bench_tiny_gemm.py index 9375344a729a..5d1631c8c2f0 100644 --- a/test/registered/kernels/benchmark/gemm/bench_tiny_gemm.py +++ b/test/registered/kernels/benchmark/gemm/bench_tiny_gemm.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/kv_canary/bench_write.py b/test/registered/kernels/benchmark/kv_canary/bench_write.py index ffe05617e900..b21d9623435f 100644 --- a/test/registered/kernels/benchmark/kv_canary/bench_write.py +++ b/test/registered/kernels/benchmark/kv_canary/bench_write.py @@ -31,7 +31,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) # AMD mirrors the CUDA nightly registration (nightly-only, no per-PR suite). # Note: amd_ci_exec.sh sets SGLANG_IS_IN_CI, so this runs the CI-reduced range diff --git a/test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py b/test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py index 8011e16ff374..86957f11b31a 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py +++ b/test/registered/kernels/benchmark/kvcache/bench_fused_fp8_qkv_kv_cache.py @@ -5,7 +5,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) FP8 = torch.float8_e4m3fn diff --git a/test/registered/kernels/benchmark/kvcache/bench_hisparse.py b/test/registered/kernels/benchmark/kvcache/bench_hisparse.py index 17a768dccae8..a3de9a65cbef 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_hisparse.py +++ b/test/registered/kernels/benchmark/kvcache/bench_hisparse.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=18, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=49, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=18, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/kvcache/bench_hisparse_spec.py b/test/registered/kernels/benchmark/kvcache/bench_hisparse_spec.py index 941a222cf3fa..9ad3cdf69a98 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_hisparse_spec.py +++ b/test/registered/kernels/benchmark/kvcache/bench_hisparse_spec.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=45, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=45, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) DEVICE = "cuda" diff --git a/test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py b/test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py index 7bdd4c1e0527..4346503f00bf 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py +++ b/test/registered/kernels/benchmark/kvcache/bench_minimax_store_kv_index.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=6, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py b/test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py index e1d31186cb67..1f8358c88300 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py +++ b/test/registered/kernels/benchmark/kvcache/bench_set_mla_kv_buffer.py @@ -22,7 +22,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=9, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=9, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/kvcache/bench_store_cache.py b/test/registered/kernels/benchmark/kvcache/bench_store_cache.py index ec9f7872508d..430a7be680c4 100644 --- a/test/registered/kernels/benchmark/kvcache/bench_store_cache.py +++ b/test/registered/kernels/benchmark/kvcache/bench_store_cache.py @@ -10,7 +10,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=9, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=44, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=9, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py b/test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py index 8848a5c78b91..a70e2a83c27e 100644 --- a/test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py +++ b/test/registered/kernels/benchmark/layernorm/bench_fused_eh_norm.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=6, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/layernorm/bench_norm.py b/test/registered/kernels/benchmark/layernorm/bench_norm.py index be1605992c19..3ebba2a83ee6 100644 --- a/test/registered/kernels/benchmark/layernorm/bench_norm.py +++ b/test/registered/kernels/benchmark/layernorm/bench_norm.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/layernorm/bench_qknorm.py b/test/registered/kernels/benchmark/layernorm/bench_qknorm.py index 78f149a13526..448bc6f2c89c 100644 --- a/test/registered/kernels/benchmark/layernorm/bench_qknorm.py +++ b/test/registered/kernels/benchmark/layernorm/bench_qknorm.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) alt_stream = torch.cuda.Stream() diff --git a/test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py b/test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py index 298e3961de17..46248f5b9cfe 100644 --- a/test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py +++ b/test/registered/kernels/benchmark/layernorm/bench_qknorm_across_heads.py @@ -13,7 +13,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=23, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) IS_CI = is_in_ci() diff --git a/test/registered/kernels/benchmark/mamba/bench_causal_conv1d.py b/test/registered/kernels/benchmark/mamba/bench_causal_conv1d.py index 9316b93de79e..0924bc4e60d8 100644 --- a/test/registered/kernels/benchmark/mamba/bench_causal_conv1d.py +++ b/test/registered/kernels/benchmark/mamba/bench_causal_conv1d.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) PAD_SLOT_ID = -1 diff --git a/test/registered/kernels/benchmark/moe/bench_fp8_per_token_to_per_tensor_quant.py b/test/registered/kernels/benchmark/moe/bench_fp8_per_token_to_per_tensor_quant.py index 75000671204d..cffa753fe3f5 100644 --- a/test/registered/kernels/benchmark/moe/bench_fp8_per_token_to_per_tensor_quant.py +++ b/test/registered/kernels/benchmark/moe/bench_fp8_per_token_to_per_tensor_quant.py @@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=45, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=45, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) FP8 = torch.float8_e4m3fn diff --git a/test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py b/test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py index e220cd490bc7..6b156e19ddb3 100644 --- a/test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py +++ b/test/registered/kernels/benchmark/moe/bench_moe_fused_gate.py @@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/moe/bench_moe_topk_softmax.py b/test/registered/kernels/benchmark/moe/bench_moe_topk_softmax.py index f43b38ba3ed3..904f1a34c809 100644 --- a/test/registered/kernels/benchmark/moe/bench_moe_topk_softmax.py +++ b/test/registered/kernels/benchmark/moe/bench_moe_topk_softmax.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py b/test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py index 1c8a4d5a4f90..9c4eeb698b35 100644 --- a/test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py +++ b/test/registered/kernels/benchmark/moe/bench_post_reorder_deepgemm.py @@ -8,7 +8,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=8, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/moe/bench_shuffle_rows_with_scales.py b/test/registered/kernels/benchmark/moe/bench_shuffle_rows_with_scales.py index daff5d63de0f..2ae06f31645a 100644 --- a/test/registered/kernels/benchmark/moe/bench_shuffle_rows_with_scales.py +++ b/test/registered/kernels/benchmark/moe/bench_shuffle_rows_with_scales.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) GROUP_SIZE = 128 diff --git a/test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py b/test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py index 35ca91b65855..e3bdf73c4f93 100644 --- a/test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py +++ b/test/registered/kernels/benchmark/quantization/bench_awq_dequantize.py @@ -12,7 +12,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) IS_CI = is_in_ci() diff --git a/test/registered/kernels/benchmark/quantization/bench_hadamard.py b/test/registered/kernels/benchmark/quantization/bench_hadamard.py index 7332dfa36cc5..bccc3825b387 100644 --- a/test/registered/kernels/benchmark/quantization/bench_hadamard.py +++ b/test/registered/kernels/benchmark/quantization/bench_hadamard.py @@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=33, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) # AOT kernel: might not be available in all environments. diff --git a/test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py b/test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py index 5f8c046cd796..155e8f7e817f 100644 --- a/test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py +++ b/test/registered/kernels/benchmark/quantization/bench_per_tensor_quant_fp8.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) try: diff --git a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py index 2d190f3a9dbb..5808d7f7e067 100644 --- a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py +++ b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant.py @@ -18,7 +18,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) HIDDEN = 2048 diff --git a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py index 47b9b00b4a0b..edfea7843c42 100644 --- a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py +++ b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_8bit_v2.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) G = 128 diff --git a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py index c9ca5d82411e..26466d710e48 100644 --- a/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py +++ b/test/registered/kernels/benchmark/quantization/bench_per_token_group_quant_masked.py @@ -22,7 +22,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=25, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) # name -> (moe_intermediate_size, topk, num_experts, group_size) diff --git a/test/registered/kernels/benchmark/quantization/bench_per_token_quant_fp8.py b/test/registered/kernels/benchmark/quantization/bench_per_token_quant_fp8.py index b0d8ec079508..d1141f9a0970 100644 --- a/test/registered/kernels/benchmark/quantization/bench_per_token_quant_fp8.py +++ b/test/registered/kernels/benchmark/quantization/bench_per_token_quant_fp8.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( - est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=20, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/sampling/bench_renorm.py b/test/registered/kernels/benchmark/sampling/bench_renorm.py index 9195c032e9d9..06c434d0d9d0 100644 --- a/test/registered/kernels/benchmark/sampling/bench_renorm.py +++ b/test/registered/kernels/benchmark/sampling/bench_renorm.py @@ -10,7 +10,7 @@ from sglang.utils import is_in_ci register_cuda_ci( - est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) diff --git a/test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py b/test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py index 7f60bb4c734c..d47e08299839 100644 --- a/test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py +++ b/test/registered/kernels/benchmark/speculative/bench_ngram_compute_decode.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=15, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py b/test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py index bb179fbac0eb..816ea49b5a4c 100644 --- a/test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py +++ b/test/registered/kernels/benchmark/speculative/bench_ngram_update_token_table.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=15, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/registered/kernels/benchmark/speculative/bench_spec_topk1.py b/test/registered/kernels/benchmark/speculative/bench_spec_topk1.py index d44cdaadd7a2..9ea62c707f76 100644 --- a/test/registered/kernels/benchmark/speculative/bench_spec_topk1.py +++ b/test/registered/kernels/benchmark/speculative/bench_spec_topk1.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci register_cuda_ci( - est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-small" ) register_amd_ci(est_time=30, stage="jit-kernel-benchmark", runner_config="amd") diff --git a/test/run_suite.py b/test/run_suite.py index fda94d65cfdd..2da4f915bd58 100644 --- a/test/run_suite.py +++ b/test/run_suite.py @@ -80,6 +80,7 @@ "base-b-kernel-unit-test-4-gpu-b200", "base-b-kernel-unit-test-8-gpu-h200", "base-b-kernel-benchmark-test-1-gpu-large", + "base-b-kernel-benchmark-test-1-gpu-small", # Diffusion keeps case-level pytest partitioning behind registered # bridge files while sharing this discovery and dispatch entry point. "base-b-test-diffusion-1-gpu-h100", From 77067ac4dbc68006062fe16d7d6f6d10371576d8 Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 17:11:34 +0000 Subject: [PATCH 04/10] [CI] Point new JIT kernel benchmarks at 1-gpu-small in the skills Follow-up to the benchmark move: write-sglang-test, add-jit-kernel and ci-workflow-guide live under .agents/skills (.claude/skills is a symlink). --- .agents/skills/add-jit-kernel/SKILL.md | 6 +++--- .agents/skills/ci-workflow-guide/SKILL.md | 2 ++ .agents/skills/write-sglang-test/SKILL.md | 4 ++-- 3 files changed, 7 insertions(+), 5 deletions(-) diff --git a/.agents/skills/add-jit-kernel/SKILL.md b/.agents/skills/add-jit-kernel/SKILL.md index 2dc4292347cf..41882600a229 100644 --- a/.agents/skills/add-jit-kernel/SKILL.md +++ b/.agents/skills/add-jit-kernel/SKILL.md @@ -640,7 +640,7 @@ if __name__ == "__main__": ## Step 5: Add a benchmark (required) -Benchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large`). +Benchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-small`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-small`); use `base-b-kernel-benchmark-test-1-gpu-large` only for SM90-only kernels. Every benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md). @@ -669,7 +669,7 @@ from sglang.kernels.jit.benchmark.utils import create_random from sglang.kernels.ops.elementwise.scale import scale as jit_scale from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-small") @torch.compile() @@ -723,7 +723,7 @@ python test/registered/kernels/benchmark/elementwise/bench_scale.py Run the benchmark suite the way CI does: ```bash -cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large +cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-small ``` --- diff --git a/.agents/skills/ci-workflow-guide/SKILL.md b/.agents/skills/ci-workflow-guide/SKILL.md index 1aec68cca2d4..0dca5adcb0c1 100644 --- a/.agents/skills/ci-workflow-guide/SKILL.md +++ b/.agents/skills/ci-workflow-guide/SKILL.md @@ -283,9 +283,11 @@ Large suites are split across matrix jobs using the **LPT (Longest Processing Ti | `base-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) | | `base-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — | | `base-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — | +| `base-b-kernel-unit-test-1-gpu-small` | 4 | `1-gpu-5090` | — | | `base-b-kernel-unit-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — | | `base-b-kernel-unit-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — | | `base-b-kernel-unit-test-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — | +| `base-b-kernel-benchmark-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — | | `base-b-kernel-benchmark-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — | | `base-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — | | `base-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — | diff --git a/.agents/skills/write-sglang-test/SKILL.md b/.agents/skills/write-sglang-test/SKILL.md index aad46fc03284..e2a7337570c0 100644 --- a/.agents/skills/write-sglang-test/SKILL.md +++ b/.agents/skills/write-sglang-test/SKILL.md @@ -107,7 +107,7 @@ Use the lightest suite that meets your test's needs: - **Need H100 memory or Hopper features** → `base-b-test-1-gpu-large` - **JIT kernel correctness** → `base-b-kernel-unit-test-1-gpu-small` (default; SM90-only kernels → `base-b-kernel-unit-test-1-gpu-large`) - **JIT kernel correctness for B200 / SM100 paths** → `base-b-kernel-unit-test-4-gpu-b200` -- **JIT kernel benchmarks** → `base-b-kernel-benchmark-test-1-gpu-large` +- **JIT kernel benchmarks** → `base-b-kernel-benchmark-test-1-gpu-small` (default; SM90-only kernels → `base-b-kernel-benchmark-test-1-gpu-large`) - **Multi-GPU** → only when the test actually needs multiple GPUs --- @@ -344,7 +344,7 @@ register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="8-gpu-h200") # Benchmarks in test/registered/kernels/benchmark// -register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-small") # Optional nightly registration — same form, stage is just "nightly" register_cuda_ci(est_time=120, stage="nightly", runner_config="1-gpu-large") From ef3599a6dd9608723a6b28aa22c5896d5f675b2b Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 17:56:54 +0000 Subject: [PATCH 05/10] [CI] Move diffusion BCG and SANA-Video 2 smoke tests to the 5090 Both pass on an RTX 5090 with the pr-test-multimodal-gen.yml env and compare against no GPU-specific ground truth or perf baseline. - bcg-diffusion job: 1-gpu-h100 -> 1-gpu-5090, suite renamed base-b-test-diffusion-bcg-1-gpu-5090 (2 passed, 2m21s). - test_sana_video2: diffusion-1-gpu-h100 -> diffusion-1-gpu-5090, so it runs in the existing 5090 canary job (2 passed, 2m44s; est_time 60 -> 165). --- .github/workflows/pr-test-multimodal-gen.yml | 4 ++-- test/registered/e2e/diffusion/test_diffusion_bcg.py | 2 +- test/registered/e2e/diffusion/test_sana_video2.py | 2 +- test/run_suite.py | 2 +- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/.github/workflows/pr-test-multimodal-gen.yml b/.github/workflows/pr-test-multimodal-gen.yml index ada973468af1..fae67bbdb66a 100644 --- a/.github/workflows/pr-test-multimodal-gen.yml +++ b/.github/workflows/pr-test-multimodal-gen.yml @@ -215,7 +215,7 @@ jobs: if: | ((github.event_name == 'schedule' || inputs.test_parallel_dispatch == 'true') || (inputs.caller_needs_failure != 'true' && !cancelled())) && inputs.multimodal_gen == 'true' - runs-on: 1-gpu-h100 + runs-on: 1-gpu-5090 timeout-minutes: 90 steps: - name: Checkout code @@ -249,7 +249,7 @@ jobs: run: | python3 test/run_suite.py \ --hw cuda \ - --suite base-b-test-diffusion-bcg-1-gpu-h100 \ + --suite base-b-test-diffusion-bcg-1-gpu-5090 \ --timeout-per-file 3600 - name: Upload BCG diffusion artifacts diff --git a/test/registered/e2e/diffusion/test_diffusion_bcg.py b/test/registered/e2e/diffusion/test_diffusion_bcg.py index dff5d62f67cf..e3ece588df1c 100644 --- a/test/registered/e2e/diffusion/test_diffusion_bcg.py +++ b/test/registered/e2e/diffusion/test_diffusion_bcg.py @@ -5,7 +5,7 @@ register_cuda_ci( est_time=3600, stage="base-b", - runner_config="diffusion-bcg-1-gpu-h100", + runner_config="diffusion-bcg-1-gpu-5090", ) diff --git a/test/registered/e2e/diffusion/test_sana_video2.py b/test/registered/e2e/diffusion/test_sana_video2.py index bf7b1d1c6b70..ad3d18ed4d12 100644 --- a/test/registered/e2e/diffusion/test_sana_video2.py +++ b/test/registered/e2e/diffusion/test_sana_video2.py @@ -12,7 +12,7 @@ from sglang.multimodal_gen import DiffGenerator from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b", runner_config="diffusion-1-gpu-h100") +register_cuda_ci(est_time=165, stage="base-b", runner_config="diffusion-1-gpu-5090") pytestmark = [ pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA"), diff --git a/test/run_suite.py b/test/run_suite.py index 2da4f915bd58..07ab9371a3fb 100644 --- a/test/run_suite.py +++ b/test/run_suite.py @@ -87,7 +87,7 @@ "base-b-test-diffusion-1-gpu-5090", "base-b-test-diffusion-1-gpu-b200", "base-b-test-diffusion-2-gpu-h100", - "base-b-test-diffusion-bcg-1-gpu-h100", + "base-b-test-diffusion-bcg-1-gpu-5090", "base-b-test-diffusion-component-2-gpu-h100", "base-b-test-diffusion-unit-1-gpu-h100", "base-c-test-4-gpu-h100", From 95e9bf3b28e2ce53142b3ae2ab72c3e9b69b7a4b Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 19:03:38 +0000 Subject: [PATCH 06/10] [CI] Run the diffusion unit suite on the 5090, keep H100-only cases on H100 The full multimodal_gen/test/unit suite passes on an RTX 5090 (4807 passed, 146 skipped, 0 failed, 30m). Eight cases only run on H100, so a new unit-h100 file suite keeps the three files that hold them on a short H100 job: - test_kandinsky6_bf16_dtype.py: FA backend cases (SM12.x resolves FA to SDPA) - test_mxfp8_online_gemm.py: pre-Blackwell channelwise fallback - test_subblock_sparse_attention.py: SM90 kernel cases, >=48 GiB free case multimodal-gen-unit-test now runs on 1-gpu-5090 (suite base-b-test-diffusion-unit-1-gpu-5090); multimodal-gen-unit-test-h100 runs base-b-test-diffusion-unit-1-gpu-h100 (unit-h100, ~25s of tests). The 5090 run also covers the SM120 SubBlock kernel and Blackwell MXFP8 cases. --- .github/workflows/pr-test-multimodal-gen.yml | 42 ++++++++++++++++++- .../multimodal_gen/test/server/gpu_cases.py | 9 +++- .../e2e/diffusion/test_diffusion_unit.py | 2 +- .../e2e/diffusion/test_diffusion_unit_h100.py | 11 +++++ test/run_suite.py | 1 + 5 files changed, 61 insertions(+), 4 deletions(-) create mode 100644 test/registered/e2e/diffusion/test_diffusion_unit_h100.py diff --git a/.github/workflows/pr-test-multimodal-gen.yml b/.github/workflows/pr-test-multimodal-gen.yml index fae67bbdb66a..269f6353bd06 100644 --- a/.github/workflows/pr-test-multimodal-gen.yml +++ b/.github/workflows/pr-test-multimodal-gen.yml @@ -465,7 +465,7 @@ jobs: if: | ((github.event_name == 'schedule' || inputs.test_parallel_dispatch == 'true') || (inputs.caller_needs_failure != 'true' && !cancelled())) && inputs.multimodal_gen == 'true' - runs-on: 1-gpu-h100 + runs-on: 1-gpu-5090 timeout-minutes: 120 steps: - name: Checkout code @@ -495,9 +495,47 @@ jobs: run: | python3 test/run_suite.py \ --hw cuda \ - --suite base-b-test-diffusion-unit-1-gpu-h100 \ + --suite base-b-test-diffusion-unit-1-gpu-5090 \ --timeout-per-file 3600 + multimodal-gen-unit-test-h100: + # Only the unit files with SM90 / H100-only cases; the full suite runs on the 5090. + if: | + ((github.event_name == 'schedule' || inputs.test_parallel_dispatch == 'true') || (inputs.caller_needs_failure != 'true' && !cancelled())) && + inputs.multimodal_gen == 'true' + runs-on: 1-gpu-h100 + timeout-minutes: 60 + steps: + - name: Checkout code + uses: actions/checkout@v4 + with: + ref: ${{ inputs.git_ref || github.sha }} + + - uses: ./.github/actions/check-pr-test-health + + - uses: ./.github/actions/check-maintenance + + - name: Download artifacts + if: inputs.sgl_kernel == 'true' + uses: actions/download-artifact@v4 + with: + path: python/sglang/kernels/aot/dist/ + merge-multiple: true + pattern: wheel-python3.10-cuda* + + - name: Install dependencies + timeout-minutes: 20 + run: | + CUSTOM_BUILD_SGL_KERNEL=${{inputs.sgl_kernel}} bash scripts/ci/cuda/ci_install_dependency.sh diffusion + + - name: Run H100-only diffusion unit tests + timeout-minutes: 20 + run: | + python3 test/run_suite.py \ + --hw cuda \ + --suite base-b-test-diffusion-unit-1-gpu-h100 \ + --timeout-per-file 1200 + diffusion-coverage-check: needs: [multimodal-gen-test-1-gpu, multimodal-gen-test-2-gpu] if: | diff --git a/python/sglang/multimodal_gen/test/server/gpu_cases.py b/python/sglang/multimodal_gen/test/server/gpu_cases.py index 4ea2a0e2f8b9..5a922309e670 100644 --- a/python/sglang/multimodal_gen/test/server/gpu_cases.py +++ b/python/sglang/multimodal_gen/test/server/gpu_cases.py @@ -1515,6 +1515,13 @@ def _discover_unit_tests() -> list[str]: FILE_SUITES = { "unit": _discover_unit_tests(), + # The full unit suite runs on the 5090; these files also hold cases that only + # run on H100 (SM90 kernels, pre-Blackwell fallbacks, FA on SM90, >=48 GiB free). + "unit-h100": [ + "../unit/test_kandinsky6_bf16_dtype.py", + "../unit/test_mxfp8_online_gemm.py", + "../unit/test_subblock_sparse_attention.py", + ], "component-accuracy": [ "../single_test_file/component_accuracy/test_component_accuracy_1_gpu.py", "../single_test_file/component_accuracy/test_component_accuracy_2_gpu.py", @@ -1631,7 +1638,7 @@ def _discover_unit_tests() -> list[str]: }, } -STRICT_SUITES = {"unit", "bcg-diffusion"} +STRICT_SUITES = {"unit", "unit-h100", "bcg-diffusion"} COMPONENT_ACCURACY_SUITES = { "component-accuracy", "component-accuracy-1-gpu", diff --git a/test/registered/e2e/diffusion/test_diffusion_unit.py b/test/registered/e2e/diffusion/test_diffusion_unit.py index 9c46e22efcdb..dadbb6901bdf 100644 --- a/test/registered/e2e/diffusion/test_diffusion_unit.py +++ b/test/registered/e2e/diffusion/test_diffusion_unit.py @@ -5,7 +5,7 @@ register_cuda_ci( est_time=3600, stage="base-b", - runner_config="diffusion-unit-1-gpu-h100", + runner_config="diffusion-unit-1-gpu-5090", ) diff --git a/test/registered/e2e/diffusion/test_diffusion_unit_h100.py b/test/registered/e2e/diffusion/test_diffusion_unit_h100.py new file mode 100644 index 000000000000..d712f028823c --- /dev/null +++ b/test/registered/e2e/diffusion/test_diffusion_unit_h100.py @@ -0,0 +1,11 @@ +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.ci.diffusion_suite_bridge import run_diffusion_suite + +register_cuda_ci( + est_time=300, + stage="base-b", + runner_config="diffusion-unit-1-gpu-h100", +) + +if __name__ == "__main__": + run_diffusion_suite("unit-h100") diff --git a/test/run_suite.py b/test/run_suite.py index 07ab9371a3fb..65f2c733a794 100644 --- a/test/run_suite.py +++ b/test/run_suite.py @@ -89,6 +89,7 @@ "base-b-test-diffusion-2-gpu-h100", "base-b-test-diffusion-bcg-1-gpu-5090", "base-b-test-diffusion-component-2-gpu-h100", + "base-b-test-diffusion-unit-1-gpu-5090", "base-b-test-diffusion-unit-1-gpu-h100", "base-c-test-4-gpu-h100", "base-c-test-4-gpu-b200", From c639bffe52d3ed82b092b6ae2e86906b10ea3382 Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Fri, 9 Oct 2026 21:37:26 +0000 Subject: [PATCH 07/10] [CI] Split H100-only classes out of 1-gpu-large files so the rest runs on 1-gpu-small Move the portable classes of 9 base-b / extra-a files to 1-gpu-small and keep the classes that need an H100 in sibling files on 1-gpu-large: - FA3 (SM80/SM90 only): test_deterministic_fa3, test_spec_standalone_extra_fa3, test_page_major_gpt_oss_fa3, test_page_major_qwen_hybrid_fa3 - Weights over 32 GB: test_hicache_mla (DeepSeek-Coder-V2-Lite), test_spec_mixed_chunk_dspark (Qwen3-14B, FA3 off SM100), test_vision_openai_server_a_large (Qwen3-VL-30B-A3B, Kimi-VL-A3B), test_vlm_input_format_kimi (Kimi-VL-A3B) - H100-tuned throughput floor: test_srt_engine_throughput (the 5090 reaches ~2970 of 3000 tok/s) The 1-gpu-small halves pass on an RTX 5090 with the CI env; their only skips are static (XPU-only, temporarily disabled MiniCPM, GLM-4.1V nightly-only). est_time of those halves is the measured 5090 time. 1-gpu-large est_time: base-b 89.2 -> 65.5 min, extra-a 136.1 -> 128.5 min. --- .../attention/test_deterministic.py | 19 +-- .../attention/test_deterministic_fa3.py | 21 +++ test/registered/core/test_srt_engine.py | 13 +- .../core/test_srt_engine_throughput.py | 26 ++++ test/registered/hicache/test_hicache_mla.py | 34 +++++ .../hicache/test_hicache_variants.py | 26 +--- .../page_major/test_page_major_gpt_oss.py | 10 +- .../page_major/test_page_major_gpt_oss_fa3.py | 20 +++ .../page_major/test_page_major_qwen_hybrid.py | 9 +- .../test_page_major_qwen_hybrid_fa3.py | 19 +++ test/registered/spec/test_spec_mixed_chunk.py | 56 +------- .../spec/test_spec_mixed_chunk_dspark.py | 68 ++++++++++ .../spec/test_spec_standalone_extra.py | 21 +-- .../spec/test_spec_standalone_extra_fa3.py | 20 +++ .../vlm/test_vision_openai_server_a.py | 102 +------------- .../vlm/test_vision_openai_server_a_large.py | 127 ++++++++++++++++++ test/registered/vlm/test_vlm_input_format.py | 59 +------- .../vlm/test_vlm_input_format_kimi.py | 71 ++++++++++ 18 files changed, 435 insertions(+), 286 deletions(-) create mode 100644 test/registered/attention/test_deterministic_fa3.py create mode 100644 test/registered/core/test_srt_engine_throughput.py create mode 100644 test/registered/hicache/test_hicache_mla.py create mode 100644 test/registered/page_major/test_page_major_gpt_oss_fa3.py create mode 100644 test/registered/page_major/test_page_major_qwen_hybrid_fa3.py create mode 100644 test/registered/spec/test_spec_mixed_chunk_dspark.py create mode 100644 test/registered/spec/test_spec_standalone_extra_fa3.py create mode 100644 test/registered/vlm/test_vision_openai_server_a_large.py create mode 100644 test/registered/vlm/test_vlm_input_format_kimi.py diff --git a/test/registered/attention/test_deterministic.py b/test/registered/attention/test_deterministic.py index 6873c47bf992..f1f8d693b198 100644 --- a/test/registered/attention/test_deterministic.py +++ b/test/registered/attention/test_deterministic.py @@ -24,7 +24,8 @@ is_in_amd_ci, ) -register_cuda_ci(est_time=280, stage="base-b", runner_config="1-gpu-large") +# The FA3 class (SM80/SM90 only) is in test_deterministic_fa3.py. +register_cuda_ci(est_time=300, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=278, suite="stage-b-test-1-gpu-small-amd") register_xpu_ci(est_time=207, suite="stage-b-test-1-gpu-xpu") @@ -47,22 +48,6 @@ def get_server_args(cls): return args -@unittest.skipIf(_is_xpu, "CUDA runner only") -@unittest.skipIf(is_in_amd_ci(), "Skip for AMD CI.") -class TestFa3Deterministic(TestDeterministicBase): - # Test with fa3 attention backend - @classmethod - def get_server_args(cls): - args = COMMON_SERVER_ARGS - args.extend( - [ - "--attention-backend", - "fa3", - ] - ) - return args - - @unittest.skipIf(_is_xpu, "CUDA/AMD runner only") class TestTritonDeterministic(TestDeterministicBase): # Test with triton attention backend diff --git a/test/registered/attention/test_deterministic_fa3.py b/test/registered/attention/test_deterministic_fa3.py new file mode 100644 index 000000000000..41b20f455b0f --- /dev/null +++ b/test/registered/attention/test_deterministic_fa3.py @@ -0,0 +1,21 @@ +"""FA3 class of test_deterministic.py, kept on H100 because FA3 is SM80/SM90 only.""" + +import unittest + +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.test_deterministic_utils import ( + COMMON_SERVER_ARGS, + TestDeterministicBase, +) + +register_cuda_ci(est_time=95, stage="base-b", runner_config="1-gpu-large") + + +class TestFa3Deterministic(TestDeterministicBase): + @classmethod + def get_server_args(cls): + return COMMON_SERVER_ARGS + ["--attention-backend", "fa3"] + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/core/test_srt_engine.py b/test/registered/core/test_srt_engine.py index 1559ad0e1a1f..98015d8e6743 100644 --- a/test/registered/core/test_srt_engine.py +++ b/test/registered/core/test_srt_engine.py @@ -10,8 +10,6 @@ import torch import sglang as sgl -from sglang.benchmark.offline_throughput import BenchArgs, throughput_test -from sglang.srt.server_args import ServerArgs from sglang.srt.utils.hf_transformers_utils import get_tokenizer from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import ( @@ -20,7 +18,8 @@ CustomTestCase, ) -register_cuda_ci(est_time=276, stage="base-b", runner_config="1-gpu-large") +# The offline throughput floor is H100-tuned; it is in test_srt_engine_throughput.py. +register_cuda_ci(est_time=220, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=261, suite="stage-b-test-1-gpu-small-amd") @@ -116,14 +115,6 @@ def test_6_engine_cpu_offload(self): print(out2) self.assertEqual(out1, out2) - def test_7_engine_offline_throughput(self): - server_args = ServerArgs( - model_path=DEFAULT_SMALL_MODEL_NAME_FOR_TEST, - ) - bench_args = BenchArgs(num_prompts=10) - result = throughput_test(server_args=server_args, bench_args=bench_args) - self.assertGreater(result["total_throughput"], 3000) - def test_8_engine_async_encode_consistency(self): prompt = "Today is a sunny day and I like" model_path = DEFAULT_SMALL_EMBEDDING_MODEL_NAME_FOR_TEST diff --git a/test/registered/core/test_srt_engine_throughput.py b/test/registered/core/test_srt_engine_throughput.py new file mode 100644 index 000000000000..f7671f25dd73 --- /dev/null +++ b/test/registered/core/test_srt_engine_throughput.py @@ -0,0 +1,26 @@ +"""Offline throughput floor of the SRT engine, kept on H100: the floor is +H100-tuned (the RTX 5090 reaches ~2970 tok/s).""" + +import unittest + +from sglang.benchmark.offline_throughput import BenchArgs, throughput_test +from sglang.srt.server_args import ServerArgs +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.test_utils import DEFAULT_SMALL_MODEL_NAME_FOR_TEST, CustomTestCase + +register_cuda_ci(est_time=60, stage="base-b", runner_config="1-gpu-large") +register_amd_ci(est_time=40, suite="stage-b-test-1-gpu-small-amd") + + +class TestSRTEngineThroughput(CustomTestCase): + def test_engine_offline_throughput(self): + server_args = ServerArgs( + model_path=DEFAULT_SMALL_MODEL_NAME_FOR_TEST, + ) + bench_args = BenchArgs(num_prompts=10) + result = throughput_test(server_args=server_args, bench_args=bench_args) + self.assertGreater(result["total_throughput"], 3000) + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/hicache/test_hicache_mla.py b/test/registered/hicache/test_hicache_mla.py new file mode 100644 index 000000000000..61b56a2f41a4 --- /dev/null +++ b/test/registered/hicache/test_hicache_mla.py @@ -0,0 +1,34 @@ +"""MLA class of test_hicache_variants.py, kept on H100: DeepSeek-Coder-V2-Lite +(~31 GB bf16) does not fit a 32 GB card.""" + +import unittest + +import test_hicache_variants as base + +from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci +from sglang.test.kits.eval_accuracy_kit import MGSMEnMixin, MMLUMixin +from sglang.test.test_utils import DEFAULT_MLA_MODEL_NAME_FOR_TEST + +register_cuda_ci(est_time=150, stage="base-b", runner_config="1-gpu-large") +register_amd_ci(est_time=150, suite="stage-b-test-1-gpu-small-amd") + + +class TestHiCacheMLA(base.HiCacheBaseServer, MMLUMixin, MGSMEnMixin): + """HiCache with MLA model tests""" + + model_name = DEFAULT_MLA_MODEL_NAME_FOR_TEST + server_env = {"SGLANG_ENABLE_RANK_CONSENSUS_CHECKER": "1"} + hicache_args = [ + "--trust-remote-code", + "--enable-hierarchical-cache", + ] + (["--hicache-size", 200] if base._is_hip else ["--hicache-ratio", 2]) + mmlu_score_threshold = 0.54 + mmlu_num_examples = 256 + mmlu_num_threads = 32 + mgsm_en_score_threshold = 0.8 + if base._is_hip: + mgsm_en_num_threads = 32 + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/hicache/test_hicache_variants.py b/test/registered/hicache/test_hicache_variants.py index 8828e80d4ad3..2c5919408d42 100644 --- a/test/registered/hicache/test_hicache_variants.py +++ b/test/registered/hicache/test_hicache_variants.py @@ -1,20 +1,21 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=534, stage="base-b", runner_config="1-gpu-large") +# The MLA class (DeepSeek-Coder-V2-Lite, ~31 GB bf16) does not fit a 32 GB card; +# it is in test_hicache_mla.py. +register_cuda_ci(est_time=400, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=524, suite="stage-b-test-1-gpu-small-amd") """ Consolidated HiCache variant tests. -Tests HiCache with different configurations: standard, MLA, EAGLE, and page size variants. +Tests HiCache with different configurations: standard, EAGLE, and page size variants. """ import unittest from sglang.benchmark.utils import get_tokenizer from sglang.srt.utils import is_hip -from sglang.test.kits.eval_accuracy_kit import MGSMEnMixin, MMLUMixin +from sglang.test.kits.eval_accuracy_kit import MMLUMixin from sglang.test.test_utils import ( DEFAULT_DRAFT_MODEL_EAGLE3, - DEFAULT_MLA_MODEL_NAME_FOR_TEST, DEFAULT_MODEL_NAME_FOR_TEST, DEFAULT_TARGET_MODEL_EAGLE3, DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, @@ -73,23 +74,6 @@ class TestHiCacheStandard(HiCacheBaseServer, MMLUMixin): mmlu_num_threads = 32 -class TestHiCacheMLA(HiCacheBaseServer, MMLUMixin, MGSMEnMixin): - """HiCache with MLA model tests""" - - model_name = DEFAULT_MLA_MODEL_NAME_FOR_TEST - server_env = {"SGLANG_ENABLE_RANK_CONSENSUS_CHECKER": "1"} - hicache_args = [ - "--trust-remote-code", - "--enable-hierarchical-cache", - ] + (["--hicache-size", 200] if _is_hip else ["--hicache-ratio", 2]) - mmlu_score_threshold = 0.54 - mmlu_num_examples = 256 - mmlu_num_threads = 32 - mgsm_en_score_threshold = 0.8 - if _is_hip: - mgsm_en_num_threads = 32 - - @unittest.skipIf(is_hip(), "Disabled for AMD-aiter") class TestHiCacheEagle(HiCacheBaseServer, MMLUMixin): """HiCache with EAGLE speculative decoding tests""" diff --git a/test/registered/page_major/test_page_major_gpt_oss.py b/test/registered/page_major/test_page_major_gpt_oss.py index b7b0efdd47e3..59aff3fbe2f1 100644 --- a/test/registered/page_major/test_page_major_gpt_oss.py +++ b/test/registered/page_major/test_page_major_gpt_oss.py @@ -14,7 +14,8 @@ from sglang.test.server_fixtures.default_fixture import DefaultServerBase from sglang.test.test_utils import DEFAULT_MODEL_NAME_FOR_TEST_MXFP4_WITH_MOE -register_cuda_ci(est_time=276, stage="extra-a", runner_config="1-gpu-large") +# The fa3 cell (SM80/SM90 only) is in test_page_major_gpt_oss_fa3.py. +register_cuda_ci(est_time=185, stage="extra-a", runner_config="1-gpu-small") _UNIFIED_COMMON_ARGS = [ "--enable-unified-memory", @@ -58,12 +59,5 @@ def test_gsm8k(self): self.assertGreaterEqual(metrics["accuracy"], self.gsm8k_threshold) -class TestUnifiedGptOssFa3(TestUnifiedGptOssTriton): - """fa3 pinned: the per-layer views read through the translator's read - tables (eager direct-bind + captured fused copy).""" - - other_args = _UNIFIED_COMMON_ARGS + ["--attention-backend", "fa3"] - - if __name__ == "__main__": unittest.main() diff --git a/test/registered/page_major/test_page_major_gpt_oss_fa3.py b/test/registered/page_major/test_page_major_gpt_oss_fa3.py new file mode 100644 index 000000000000..b391d69ef6cb --- /dev/null +++ b/test/registered/page_major/test_page_major_gpt_oss_fa3.py @@ -0,0 +1,20 @@ +"""fa3 cell of test_page_major_gpt_oss.py, kept on H100 because FA3 is SM80/SM90 only.""" + +import unittest + +import test_page_major_gpt_oss as base + +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=140, stage="extra-a", runner_config="1-gpu-large") + + +class TestUnifiedGptOssFa3(base.TestUnifiedGptOssTriton): + """fa3 pinned: the per-layer views read through the translator's read + tables (eager direct-bind + captured fused copy).""" + + other_args = base._UNIFIED_COMMON_ARGS + ["--attention-backend", "fa3"] + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/page_major/test_page_major_qwen_hybrid.py b/test/registered/page_major/test_page_major_qwen_hybrid.py index 4ea75e21a184..7d76caf569dc 100644 --- a/test/registered/page_major/test_page_major_qwen_hybrid.py +++ b/test/registered/page_major/test_page_major_qwen_hybrid.py @@ -15,7 +15,8 @@ from sglang.test.server_fixtures.default_fixture import DefaultServerBase from sglang.test.test_utils import DEFAULT_HYBRID_GDN_SMALL_MODEL_NAME_FOR_TEST -register_cuda_ci(est_time=261, stage="extra-a", runner_config="1-gpu-large") +# The fa3 cell (SM80/SM90 only) is in test_page_major_qwen_hybrid_fa3.py. +register_cuda_ci(est_time=175, stage="extra-a", runner_config="1-gpu-small") _UNIFIED_COMMON_ARGS = [ "--trust-remote-code", @@ -66,12 +67,6 @@ def test_gsm8k(self): self.assertGreaterEqual(metrics["accuracy"], self.gsm8k_threshold) -class TestUnifiedQwenHybridFa3(TestUnifiedQwenHybridTriton): - """fa3 pinned: read tables, eager direct-bind + captured fused copy.""" - - other_args = _UNIFIED_COMMON_ARGS + ["--attention-backend", "fa3"] - - class TestUnifiedQwenHybridFlashinfer(TestUnifiedQwenHybridTriton): """flashinfer pinned: token ids reconstructed from the read table by the ENTRY_PAGE_SIZE CSR builder.""" diff --git a/test/registered/page_major/test_page_major_qwen_hybrid_fa3.py b/test/registered/page_major/test_page_major_qwen_hybrid_fa3.py new file mode 100644 index 000000000000..f305da49a350 --- /dev/null +++ b/test/registered/page_major/test_page_major_qwen_hybrid_fa3.py @@ -0,0 +1,19 @@ +"""fa3 cell of test_page_major_qwen_hybrid.py, kept on H100 because FA3 is SM80/SM90 only.""" + +import unittest + +import test_page_major_qwen_hybrid as base + +from sglang.test.ci.ci_register import register_cuda_ci + +register_cuda_ci(est_time=90, stage="extra-a", runner_config="1-gpu-large") + + +class TestUnifiedQwenHybridFa3(base.TestUnifiedQwenHybridTriton): + """fa3 pinned: read tables, eager direct-bind + captured fused copy.""" + + other_args = base._UNIFIED_COMMON_ARGS + ["--attention-backend", "fa3"] + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/spec/test_spec_mixed_chunk.py b/test/registered/spec/test_spec_mixed_chunk.py index 8cd1b09165a9..f88c294b04ef 100644 --- a/test/registered/spec/test_spec_mixed_chunk.py +++ b/test/registered/spec/test_spec_mixed_chunk.py @@ -1,6 +1,7 @@ """Mixed chunk prefill x speculative decoding, overlap scheduler. -One cell per supported algorithm (EAGLE3, DFLASH, DSPARK). Inside a mixed +One cell per supported algorithm (EAGLE3, DFLASH; DSPARK is in +test_spec_mixed_chunk_dspark.py). Inside a mixed step every running request degrades to a 1-token extend of its pending bonus token and drafting resumes the next decode step; under overlap the tail state is late-bound at forward entry. Regression guards for the @@ -15,7 +16,7 @@ import unittest from sglang.srt.environ import envs -from sglang.srt.utils import is_sm100_supported, kill_process_tree +from sglang.srt.utils import kill_process_tree from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.kits.eval_accuracy_kit import GSM8KMixin from sglang.test.kits.spec_server_kits import ( @@ -32,7 +33,7 @@ popen_launch_server, ) -register_cuda_ci(est_time=196, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=170, stage="base-b", runner_config="1-gpu-small") class TestEagle3MixedChunk( @@ -86,54 +87,5 @@ def tearDownClass(cls): kill_process_tree(cls.process.pid) -DSPARK_TARGET_MODEL = "Qwen/Qwen3-14B" -DSPARK_DRAFT_MODEL = "deepseek-ai/dspark_qwen3_14b_block7" - - -class TestDSparkMixedChunk(GSM8KMixin, CustomTestCase): - model = DSPARK_TARGET_MODEL - - gsm8k_num_questions = 200 - gsm8k_accuracy_thres = 0.80 - gsm8k_accept_length_thres = 2.0 - - process = None - - @classmethod - def setUpClass(cls): - cls.base_url = DEFAULT_URL_FOR_TEST - cls.process = popen_launch_server( - cls.model, - cls.base_url, - timeout=DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, - other_args=[ - "--trust-remote-code", - "--attention-backend", - "trtllm_mha" if is_sm100_supported() else "fa3", - "--speculative-draft-attention-backend", - "fa4" if is_sm100_supported() else "fa3", - "--speculative-algorithm", - "DSPARK", - "--speculative-draft-model-path", - DSPARK_DRAFT_MODEL, - "--enable-mixed-chunk", - "--chunked-prefill-size", - "128", - "--cuda-graph-max-bs-decode", - "4", - "--mem-fraction-static", - "0.7", - "--page-size", - "1", - "--cuda-graph-backend-prefill=disabled", - ], - ) - - @classmethod - def tearDownClass(cls): - if cls.process is not None: - kill_process_tree(cls.process.pid) - - if __name__ == "__main__": unittest.main() diff --git a/test/registered/spec/test_spec_mixed_chunk_dspark.py b/test/registered/spec/test_spec_mixed_chunk_dspark.py new file mode 100644 index 000000000000..22731e4a2f2a --- /dev/null +++ b/test/registered/spec/test_spec_mixed_chunk_dspark.py @@ -0,0 +1,68 @@ +"""DSPARK cell of test_spec_mixed_chunk.py, kept on H100: the Qwen3-14B target +does not fit a 32 GB card, and off SM100 it runs on FA3 (SM80/SM90 only).""" + +import unittest + +from sglang.srt.utils import is_sm100_supported, kill_process_tree +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.kits.eval_accuracy_kit import GSM8KMixin +from sglang.test.test_utils import ( + DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, + DEFAULT_URL_FOR_TEST, + CustomTestCase, + popen_launch_server, +) + +register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-large") + +DSPARK_TARGET_MODEL = "Qwen/Qwen3-14B" +DSPARK_DRAFT_MODEL = "deepseek-ai/dspark_qwen3_14b_block7" + + +class TestDSparkMixedChunk(GSM8KMixin, CustomTestCase): + model = DSPARK_TARGET_MODEL + + gsm8k_num_questions = 200 + gsm8k_accuracy_thres = 0.80 + gsm8k_accept_length_thres = 2.0 + + process = None + + @classmethod + def setUpClass(cls): + cls.base_url = DEFAULT_URL_FOR_TEST + cls.process = popen_launch_server( + cls.model, + cls.base_url, + timeout=DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, + other_args=[ + "--trust-remote-code", + "--attention-backend", + "trtllm_mha" if is_sm100_supported() else "fa3", + "--speculative-draft-attention-backend", + "fa4" if is_sm100_supported() else "fa3", + "--speculative-algorithm", + "DSPARK", + "--speculative-draft-model-path", + DSPARK_DRAFT_MODEL, + "--enable-mixed-chunk", + "--chunked-prefill-size", + "128", + "--cuda-graph-max-bs-decode", + "4", + "--mem-fraction-static", + "0.7", + "--page-size", + "1", + "--cuda-graph-backend-prefill=disabled", + ], + ) + + @classmethod + def tearDownClass(cls): + if cls.process is not None: + kill_process_tree(cls.process.pid) + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/spec/test_spec_standalone_extra.py b/test/registered/spec/test_spec_standalone_extra.py index c43dbeb02302..aff0b31e12f5 100644 --- a/test/registered/spec/test_spec_standalone_extra.py +++ b/test/registered/spec/test_spec_standalone_extra.py @@ -5,25 +5,18 @@ from sglang.test.server_fixtures.standalone_fixture import StandaloneServerBase from sglang.test.test_utils import CustomTestCase -# Non-V2 standalone speculative decoding tests (FA3, Triton, FlashInfer -# backends). Sibling V2 classes stay per-commit in test_spec_standalone.py. -register_cuda_ci(est_time=225, stage="extra-a", runner_config="1-gpu-large") -# AMD: fa3 / flashinfer attention backends are not built in the ROCm -# sgl_kernel, so only the triton-backend class runs on ROCm (the fa3 and -# flashinfer classes are skipped on ROCm below). +# Non-V2 standalone speculative decoding tests (Triton, FlashInfer backends); +# the FA3 class is in test_spec_standalone_extra_fa3.py. Sibling V2 classes +# stay per-commit in test_spec_standalone.py. +register_cuda_ci(est_time=175, stage="extra-a", runner_config="1-gpu-small") +# AMD: the flashinfer attention backend is not built in the ROCm sgl_kernel, +# so only the triton-backend class runs on ROCm (the flashinfer class is +# skipped on ROCm below). register_amd_ci(est_time=103, suite="extra-a-test-1-gpu-large-amd") _AMD_SKIP_BACKEND = "fa3 / flashinfer attention backends are CUDA-only (not in the ROCm sgl_kernel build)" -@unittest.skipIf(is_hip(), _AMD_SKIP_BACKEND) -class TestStandaloneSpeculativeDecodingBase(StandaloneServerBase, CustomTestCase): - attention_backend = "fa3" - speculative_eagle_topk = 2 - speculative_num_draft_tokens = 7 - disable_overlap = True - - class TestStandaloneSpeculativeDecodingTriton(StandaloneServerBase, CustomTestCase): attention_backend = "triton" speculative_eagle_topk = 2 diff --git a/test/registered/spec/test_spec_standalone_extra_fa3.py b/test/registered/spec/test_spec_standalone_extra_fa3.py new file mode 100644 index 000000000000..9848f466b806 --- /dev/null +++ b/test/registered/spec/test_spec_standalone_extra_fa3.py @@ -0,0 +1,20 @@ +"""FA3 class of test_spec_standalone_extra.py, kept on H100 because FA3 is SM80/SM90 only.""" + +import unittest + +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.server_fixtures.standalone_fixture import StandaloneServerBase +from sglang.test.test_utils import CustomTestCase + +register_cuda_ci(est_time=80, stage="extra-a", runner_config="1-gpu-large") + + +class TestStandaloneSpeculativeDecodingBase(StandaloneServerBase, CustomTestCase): + attention_backend = "fa3" + speculative_eagle_topk = 2 + speculative_num_draft_tokens = 7 + disable_overlap = True + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/vlm/test_vision_openai_server_a.py b/test/registered/vlm/test_vision_openai_server_a.py index 6910b7aed8c0..3d3ea3ef544b 100644 --- a/test/registered/vlm/test_vision_openai_server_a.py +++ b/test/registered/vlm/test_vision_openai_server_a.py @@ -4,15 +4,10 @@ python3 -m unittest test_vision_openai_server.TestOpenAIVisionServer.test_multi_images_chat_completion """ -import base64 -import io -import re import unittest import openai -from PIL import Image, ImageDraw -from sglang.srt.environ import envs from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import ( DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, @@ -29,35 +24,9 @@ terminate_and_kill_process_tree, ) -register_cuda_ci(est_time=578, stage="base-b", runner_config="1-gpu-large") - - -# --- Qwen3-VL grounding regression (deepstack fusion) -------------------------- -# Guards Qwen3MoeLLMModel.forward: deepstack (multi-scale ViT features) injection -# must keep its original inference order. PR #14636 rerouted it through -# post_residual_addition (for RL on-policy / FSDP), which is FP-order-sensitive -# and regresses FP8 visual grounding (the predicted point drifts by ~150+ px). -_GROUNDING_IMG_SIZE = 1000 -# Target box in pixels; on a 1000x1000 canvas this equals the 0-1000 normalized -# coordinate, so the check is robust to normalized-vs-pixel conventions. -_GROUNDING_BOX = (620, 180, 880, 360) # (x0, y0, x1, y1), center (750, 270) -_GROUNDING_MARGIN = 60 -_GROUNDING_SYSTEM = ( - "You are a UI grounding model. Treat the image as a 1000x1000 normalized " - "coordinate system with the top-left at (0,0) and the bottom-right at " - "(1000,1000). Return the geometric center of the requested element. " - "Output ONLY one coordinate in the form (x, y) and nothing else." -) -_GROUNDING_COORD_RE = re.compile(r"\(?\s*(\d{1,4})\s*,\s*(\d{1,4})\s*\)?") - - -def _make_grounding_image() -> str: - """White canvas with a single red box at _GROUNDING_BOX; base64 data URI.""" - img = Image.new("RGB", (_GROUNDING_IMG_SIZE, _GROUNDING_IMG_SIZE), (255, 255, 255)) - ImageDraw.Draw(img).rectangle(_GROUNDING_BOX, fill=(220, 30, 30)) - buf = io.BytesIO() - img.save(buf, format="PNG") - return "data:image/png;base64," + base64.b64encode(buf.getvalue()).decode("utf-8") +# Qwen3-VL-30B-A3B and Kimi-VL-A3B do not fit a 32 GB card; they are in +# test_vision_openai_server_a_large.py. +register_cuda_ci(est_time=465, stage="base-b", runner_config="1-gpu-small") class TestLlavaServer(ImageOpenAITestMixin): @@ -75,57 +44,6 @@ class TestQwen25VLServer(ImageOpenAITestMixin, VideoOpenAITestMixin): ] -class TestQwen3VLServer(ImageOpenAITestMixin, VideoOpenAITestMixin): - model = "Qwen/Qwen3-VL-30B-A3B-Instruct" - extra_args = ["--cuda-graph-max-bs-decode=4"] - - @classmethod - def setUpClass(cls): - with envs.SGLANG_MM_FEATURE_CACHE_MB.override(512): - super().setUpClass() - - def test_deepstack_grounding_hits_target_box(self): - # Regression guard for the Qwen3-VL MoE deepstack fusion order: the - # predicted point must land inside the target box; a deepstack corruption - # drifts it out (see PR #14636). - client = openai.Client(api_key=self.api_key, base_url=self.base_url) - response = client.chat.completions.create( - model="default", - messages=[ - {"role": "system", "content": _GROUNDING_SYSTEM}, - { - "role": "user", - "content": [ - { - "type": "image_url", - "image_url": {"url": _make_grounding_image()}, - }, - { - "type": "text", - "text": "Point at the center of the red rectangle.", - }, - ], - }, - ], - temperature=0, - **(self.get_vision_request_kwargs()), - ) - out = response.choices[0].message.content - match = _GROUNDING_COORD_RE.search(out or "") - self.assertIsNotNone(match, f"could not parse a coordinate from: {out!r}") - x, y = int(match.group(1)), int(match.group(2)) - x0, y0, x1, y1 = _GROUNDING_BOX - inside = (x0 - _GROUNDING_MARGIN <= x <= x1 + _GROUNDING_MARGIN) and ( - y0 - _GROUNDING_MARGIN <= y <= y1 + _GROUNDING_MARGIN - ) - self.assertTrue( - inside, - f"grounding output {out!r} -> ({x}, {y}) fell outside target box " - f"{_GROUNDING_BOX} (margin {_GROUNDING_MARGIN}); deepstack fusion " - f"likely regressed grounding.", - ) - - class TestQwen2VLContextLengthServer(CustomTestCase): # --context-length 300 is calibrated to this model's mm-token expansion: # it must sit above the warmup image's expanded length but below the test @@ -220,20 +138,6 @@ class TestGemma3itServer(ImageOpenAITestMixin): ] -class TestKimiVLServer(ImageOpenAITestMixin): - model = "moonshotai/Kimi-VL-A3B-Instruct" - extra_args = [ - "--context-length=8192", - "--dtype=bfloat16", - # Weights alone need ~0.39; 0.40 left <0.001 headroom and flaked at load. - "--mem-fraction-static=0.42", - ] - - @unittest.skip("model context length exceeded") - def test_video_images_chat_completion(self): - pass - - @unittest.skip( "Disabling this test to speed up CI. Prefer to test it within nightly test." ) diff --git a/test/registered/vlm/test_vision_openai_server_a_large.py b/test/registered/vlm/test_vision_openai_server_a_large.py new file mode 100644 index 000000000000..4d4bf3685c44 --- /dev/null +++ b/test/registered/vlm/test_vision_openai_server_a_large.py @@ -0,0 +1,127 @@ +"""Vision OpenAI server cases that stay on H100. + +Qwen3-VL-30B-A3B and Kimi-VL-A3B do not fit a 32 GB card. Split from +test_vision_openai_server_a.py so the rest of that file runs on 1-gpu-small. +""" + +import base64 +import io +import re +import unittest + +import openai +from PIL import Image, ImageDraw + +from sglang.srt.environ import envs +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.vlm_utils import ( + ImageOpenAITestMixin, + TestOpenAIMLLMServerBase, + VideoOpenAITestMixin, +) + +register_cuda_ci(est_time=200, stage="base-b", runner_config="1-gpu-large") + +# --- Qwen3-VL grounding regression (deepstack fusion) -------------------------- +# Guards Qwen3MoeLLMModel.forward: deepstack (multi-scale ViT features) injection +# must keep its original inference order. PR #14636 rerouted it through +# post_residual_addition (for RL on-policy / FSDP), which is FP-order-sensitive +# and regresses FP8 visual grounding (the predicted point drifts by ~150+ px). +_GROUNDING_IMG_SIZE = 1000 +# Target box in pixels; on a 1000x1000 canvas this equals the 0-1000 normalized +# coordinate, so the check is robust to normalized-vs-pixel conventions. +_GROUNDING_BOX = (620, 180, 880, 360) # (x0, y0, x1, y1), center (750, 270) +_GROUNDING_MARGIN = 60 +_GROUNDING_SYSTEM = ( + "You are a UI grounding model. Treat the image as a 1000x1000 normalized " + "coordinate system with the top-left at (0,0) and the bottom-right at " + "(1000,1000). Return the geometric center of the requested element. " + "Output ONLY one coordinate in the form (x, y) and nothing else." +) +_GROUNDING_COORD_RE = re.compile(r"\(?\s*(\d{1,4})\s*,\s*(\d{1,4})\s*\)?") + + +def _make_grounding_image() -> str: + """White canvas with a single red box at _GROUNDING_BOX; base64 data URI.""" + img = Image.new("RGB", (_GROUNDING_IMG_SIZE, _GROUNDING_IMG_SIZE), (255, 255, 255)) + ImageDraw.Draw(img).rectangle(_GROUNDING_BOX, fill=(220, 30, 30)) + buf = io.BytesIO() + img.save(buf, format="PNG") + return "data:image/png;base64," + base64.b64encode(buf.getvalue()).decode("utf-8") + + +class TestQwen3VLServer(ImageOpenAITestMixin, VideoOpenAITestMixin): + model = "Qwen/Qwen3-VL-30B-A3B-Instruct" + extra_args = ["--cuda-graph-max-bs-decode=4"] + + @classmethod + def setUpClass(cls): + with envs.SGLANG_MM_FEATURE_CACHE_MB.override(512): + super().setUpClass() + + def test_deepstack_grounding_hits_target_box(self): + # Regression guard for the Qwen3-VL MoE deepstack fusion order: the + # predicted point must land inside the target box; a deepstack corruption + # drifts it out (see PR #14636). + client = openai.Client(api_key=self.api_key, base_url=self.base_url) + response = client.chat.completions.create( + model="default", + messages=[ + {"role": "system", "content": _GROUNDING_SYSTEM}, + { + "role": "user", + "content": [ + { + "type": "image_url", + "image_url": {"url": _make_grounding_image()}, + }, + { + "type": "text", + "text": "Point at the center of the red rectangle.", + }, + ], + }, + ], + temperature=0, + **(self.get_vision_request_kwargs()), + ) + out = response.choices[0].message.content + match = _GROUNDING_COORD_RE.search(out or "") + self.assertIsNotNone(match, f"could not parse a coordinate from: {out!r}") + x, y = int(match.group(1)), int(match.group(2)) + x0, y0, x1, y1 = _GROUNDING_BOX + inside = (x0 - _GROUNDING_MARGIN <= x <= x1 + _GROUNDING_MARGIN) and ( + y0 - _GROUNDING_MARGIN <= y <= y1 + _GROUNDING_MARGIN + ) + self.assertTrue( + inside, + f"grounding output {out!r} -> ({x}, {y}) fell outside target box " + f"{_GROUNDING_BOX} (margin {_GROUNDING_MARGIN}); deepstack fusion " + f"likely regressed grounding.", + ) + + +class TestKimiVLServer(ImageOpenAITestMixin): + model = "moonshotai/Kimi-VL-A3B-Instruct" + extra_args = [ + "--context-length=8192", + "--dtype=bfloat16", + # Weights alone need ~0.39; 0.40 left <0.001 headroom and flaked at load. + "--mem-fraction-static=0.42", + ] + + @unittest.skip("model context length exceeded") + def test_video_images_chat_completion(self): + pass + + +# Delete the mixin classes so that they are not collected by pytest +del ( + TestOpenAIMLLMServerBase, + ImageOpenAITestMixin, + VideoOpenAITestMixin, +) + + +if __name__ == "__main__": + unittest.main() diff --git a/test/registered/vlm/test_vlm_input_format.py b/test/registered/vlm/test_vlm_input_format.py index 6b5a94b6c048..39084fe91b18 100644 --- a/test/registered/vlm/test_vlm_input_format.py +++ b/test/registered/vlm/test_vlm_input_format.py @@ -39,7 +39,8 @@ def forward(self, x): from sglang.srt.utils.hf_transformers_utils import _fix_added_tokens_encoding from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=232, stage="base-b", runner_config="1-gpu-large") +# Kimi-VL-A3B does not fit a 32 GB card; it is in test_vlm_input_format_kimi.py. +register_cuda_ci(est_time=170, stage="base-b", runner_config="1-gpu-small") IMAGE_MAN_IRONING_URL = "https://raw.githubusercontent.com/sgl-project/sgl-test-files/refs/heads/main/images/man_ironing_on_back_of_suv.png" IMAGE_SGL_LOGO_URL = "https://raw.githubusercontent.com/sgl-project/sgl-test-files/refs/heads/main/images/sgl_logo.png" @@ -273,62 +274,6 @@ def _processor_output_image_data(self, processor_output): return dict(processor_output, format="processor_output") -# Updated Kimi-VL test to use the new input format. -class TestKimiVLImageUnderstandsImage(VLMInputTestBase, CustomTestCase): - model_path = "moonshotai/Kimi-VL-A3B-Instruct" - chat_template = "kimi-vl" - - @classmethod - def _init_visual(cls): - import inspect - - from transformers import AutoConfig - from transformers.dynamic_module_utils import get_class_from_dynamic_module - - config = AutoConfig.from_pretrained(cls.model_path, trust_remote_code=True) - - # Transformers v5 auto-populates rope_scaling with - # {"rope_theta": ..., "rope_type": "default"} even when the original - # config had rope_scaling: null. The remote KimiVL code branches on - # `if self.config.rope_scaling is None` so we must reset it. - tc = getattr(config, "text_config", None) - if tc is not None: - rs = getattr(tc, "rope_scaling", None) - if isinstance(rs, dict) and rs.get("rope_type") == "default": - tc.rope_scaling = None - - # Transformers v5 calls tie_weights(recompute_mapping=False) in - # post_init, but KimiVL's tie_weights doesn't accept that kwarg. - auto_map = getattr(config, "auto_map", {}) - model_ref = auto_map.get("AutoModel") - if model_ref: - model_cls = get_class_from_dynamic_module(model_ref, cls.model_path) - orig_tie = model_cls.tie_weights - if "recompute_mapping" not in inspect.signature(orig_tie).parameters: - - def _patched_tie(self, **kwargs): - return orig_tie(self) - - model_cls.tie_weights = _patched_tie - - model = AutoModel.from_pretrained( - cls.model_path, config=config, trust_remote_code=True - ) - cls.vision_tower = model.vision_tower.eval().to(cls.device) - cls.mm_projector = model.multi_modal_projector.eval().to(cls.device) - _vt_dtype = next(cls.vision_tower.parameters()).dtype - - cls.visual = lambda tokenizer_output: cls.mm_projector( - cls.vision_tower( - pixel_values=tokenizer_output["pixel_values"].to(_vt_dtype), - grid_hws=tokenizer_output["image_grid_hws"], - ) - ) - - def _processor_output_image_data(self, processor_output): - return dict(processor_output, format="processor_output") - - # not for CI: too large # class TestLlama4ImageUnderstandsImage( # VLMInputTestBase, CustomTestCase diff --git a/test/registered/vlm/test_vlm_input_format_kimi.py b/test/registered/vlm/test_vlm_input_format_kimi.py new file mode 100644 index 000000000000..b80fcaf73bbb --- /dev/null +++ b/test/registered/vlm/test_vlm_input_format_kimi.py @@ -0,0 +1,71 @@ +"""Kimi-VL case of test_vlm_input_format.py, kept on H100: Kimi-VL-A3B plus its +HF vision tower does not fit a 32 GB card.""" + +import unittest + +import test_vlm_input_format as base +from transformers import AutoModel + +from sglang.test.ci.ci_register import register_cuda_ci +from sglang.test.test_utils import CustomTestCase + +register_cuda_ci(est_time=75, stage="base-b", runner_config="1-gpu-large") + + +class TestKimiVLImageUnderstandsImage(base.VLMInputTestBase, CustomTestCase): + model_path = "moonshotai/Kimi-VL-A3B-Instruct" + chat_template = "kimi-vl" + + @classmethod + def _init_visual(cls): + import inspect + + from transformers import AutoConfig + from transformers.dynamic_module_utils import get_class_from_dynamic_module + + config = AutoConfig.from_pretrained(cls.model_path, trust_remote_code=True) + + # Transformers v5 auto-populates rope_scaling with + # {"rope_theta": ..., "rope_type": "default"} even when the original + # config had rope_scaling: null. The remote KimiVL code branches on + # `if self.config.rope_scaling is None` so we must reset it. + tc = getattr(config, "text_config", None) + if tc is not None: + rs = getattr(tc, "rope_scaling", None) + if isinstance(rs, dict) and rs.get("rope_type") == "default": + tc.rope_scaling = None + + # Transformers v5 calls tie_weights(recompute_mapping=False) in + # post_init, but KimiVL's tie_weights doesn't accept that kwarg. + auto_map = getattr(config, "auto_map", {}) + model_ref = auto_map.get("AutoModel") + if model_ref: + model_cls = get_class_from_dynamic_module(model_ref, cls.model_path) + orig_tie = model_cls.tie_weights + if "recompute_mapping" not in inspect.signature(orig_tie).parameters: + + def _patched_tie(self, **kwargs): + return orig_tie(self) + + model_cls.tie_weights = _patched_tie + + model = AutoModel.from_pretrained( + cls.model_path, config=config, trust_remote_code=True + ) + cls.vision_tower = model.vision_tower.eval().to(cls.device) + cls.mm_projector = model.multi_modal_projector.eval().to(cls.device) + _vt_dtype = next(cls.vision_tower.parameters()).dtype + + cls.visual = lambda tokenizer_output: cls.mm_projector( + cls.vision_tower( + pixel_values=tokenizer_output["pixel_values"].to(_vt_dtype), + grid_hws=tokenizer_output["image_grid_hws"], + ) + ) + + def _processor_output_image_data(self, processor_output): + return dict(processor_output, format="processor_output") + + +if __name__ == "__main__": + unittest.main() From ffe5f4394b5e03ee30f6eaf52d510a7a378b3a14 Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Sat, 17 Oct 2026 09:00:59 +0000 Subject: [PATCH 08/10] [CI] Move 1-gpu-large tests added since the last merge to 1-gpu-small Re-run the 16 base-b, base-b-kernel-unit and extra-a files registered on 1-gpu-large since the previous sweep on an RTX 5090 with the CI environment. - base-b-kernel-unit: 13 files move; base-b: test_mamba_extra_buffer_spec_grammar; extra-a: test_mamba_cache_kl. - Raise est_time to the measured 5090 duration where the estimate was more than 1.5x low: test_qknorm_rope_out_of_place 60->136, test_watermark_selector 15->43, test_block_topk_triton 10->17. - test_compressed_tensors_fp4_numeric keeps est_time 30: 14s with a warm JIT cache (407s cold). Its 4 skipped FlashInfer cases need SM100 and also skip on H100. Kept on 1-gpu-large: test_q8kv8_dsa_nope_bf16_prefill (its kpool-width test runs the SM90-only WGMMA sparse prefill kernel). --- .../e2e/radix_cache/test_mamba_extra_buffer_spec_grammar.py | 2 +- .../kernels/ops/attention/test_dsa_kpool_topk_transform.py | 2 +- .../kernels/ops/attention/test_dsv41_request_window_copy.py | 2 +- test/registered/kernels/ops/attention/test_topk_v1.py | 2 +- .../registered/kernels/ops/diffusion/test_conv_bias_epilogue.py | 2 +- test/registered/kernels/ops/diffusion/test_fast_launch.py | 2 +- .../kernels/ops/diffusion/test_qknorm_rope_out_of_place.py | 2 +- test/registered/kernels/ops/mamba/test_mtp_hicache.py | 2 +- test/registered/kernels/ops/moe/test_block_topk_triton.py | 2 +- .../kernels/ops/moe/test_moe_tma_down_padding_rows.py | 2 +- .../ops/quantization/test_compressed_tensors_fp4_numeric.py | 2 +- test/registered/kernels/ops/sampling/test_watermark_selector.py | 2 +- test/registered/kernels/ops/sampling/test_watermark_state.py | 2 +- test/registered/radix_cache/test_mamba_cache_kl.py | 2 +- .../unit/multimodal_gen/test_minimax_h3_step_contract.py | 2 +- 15 files changed, 15 insertions(+), 15 deletions(-) diff --git a/test/registered/e2e/radix_cache/test_mamba_extra_buffer_spec_grammar.py b/test/registered/e2e/radix_cache/test_mamba_extra_buffer_spec_grammar.py index 72186c78873a..db970265e38a 100644 --- a/test/registered/e2e/radix_cache/test_mamba_extra_buffer_spec_grammar.py +++ b/test/registered/e2e/radix_cache/test_mamba_extra_buffer_spec_grammar.py @@ -16,7 +16,7 @@ terminate_and_kill_process_tree, ) -register_cuda_ci(est_time=180, stage="base-b", runner_config="1-gpu-large") +register_cuda_ci(est_time=180, stage="base-b", runner_config="1-gpu-small") MODEL = "Qwen/Qwen3.5-0.8B" # Default --mamba-track-interval: the grid on which decode checkpoints the diff --git a/test/registered/kernels/ops/attention/test_dsa_kpool_topk_transform.py b/test/registered/kernels/ops/attention/test_dsa_kpool_topk_transform.py index 100bb1dcc4a3..05752cdc80a7 100644 --- a/test/registered/kernels/ops/attention/test_dsa_kpool_topk_transform.py +++ b/test/registered/kernels/ops/attention/test_dsa_kpool_topk_transform.py @@ -30,7 +30,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=180, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=180, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=180, stage="jit-kernel-unit", runner_config="amd") # kSmem / (2 * sizeof(int)) in kpool_topk_transform.cuh: entries per stash round. diff --git a/test/registered/kernels/ops/attention/test_dsv41_request_window_copy.py b/test/registered/kernels/ops/attention/test_dsv41_request_window_copy.py index 858d20929ba8..623e7248e9ae 100644 --- a/test/registered/kernels/ops/attention/test_dsv41_request_window_copy.py +++ b/test/registered/kernels/ops/attention/test_dsv41_request_window_copy.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") def reference(src, dst, source, target, page, layout): diff --git a/test/registered/kernels/ops/attention/test_topk_v1.py b/test/registered/kernels/ops/attention/test_topk_v1.py index f9fe40695dda..b135ba4d1cdc 100644 --- a/test/registered/kernels/ops/attention/test_topk_v1.py +++ b/test/registered/kernels/ops/attention/test_topk_v1.py @@ -26,7 +26,7 @@ from sglang.srt.utils import is_hip from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif( not torch.cuda.is_available() or is_hip(), diff --git a/test/registered/kernels/ops/diffusion/test_conv_bias_epilogue.py b/test/registered/kernels/ops/diffusion/test_conv_bias_epilogue.py index 9c7a3c76e56f..4c2f05b6b85f 100644 --- a/test/registered/kernels/ops/diffusion/test_conv_bias_epilogue.py +++ b/test/registered/kernels/ops/diffusion/test_conv_bias_epilogue.py @@ -18,7 +18,7 @@ ) 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") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-small") DEVICE = "cuda" diff --git a/test/registered/kernels/ops/diffusion/test_fast_launch.py b/test/registered/kernels/ops/diffusion/test_fast_launch.py index 964d091b150a..5271513fe181 100644 --- a/test/registered/kernels/ops/diffusion/test_fast_launch.py +++ b/test/registered/kernels/ops/diffusion/test_fast_launch.py @@ -38,7 +38,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/diffusion/test_qknorm_rope_out_of_place.py b/test/registered/kernels/ops/diffusion/test_qknorm_rope_out_of_place.py index b7bb5bf4cb5b..068b71e8445a 100644 --- a/test/registered/kernels/ops/diffusion/test_qknorm_rope_out_of_place.py +++ b/test/registered/kernels/ops/diffusion/test_qknorm_rope_out_of_place.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="4-gpu-b200") -register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=136, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_out_of_place_qknorm_rope_matches_inplace_and_keeps_inputs() -> None: diff --git a/test/registered/kernels/ops/mamba/test_mtp_hicache.py b/test/registered/kernels/ops/mamba/test_mtp_hicache.py index 7866719a7bea..869dec7575f3 100644 --- a/test/registered/kernels/ops/mamba/test_mtp_hicache.py +++ b/test/registered/kernels/ops/mamba/test_mtp_hicache.py @@ -30,7 +30,7 @@ from sglang.srt.server_args import ServerArgs from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") diff --git a/test/registered/kernels/ops/moe/test_block_topk_triton.py b/test/registered/kernels/ops/moe/test_block_topk_triton.py index 2d1094085f5a..e2f10ead7b35 100644 --- a/test/registered/kernels/ops/moe/test_block_topk_triton.py +++ b/test/registered/kernels/ops/moe/test_block_topk_triton.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_amd_ci(est_time=10, suite="jit-kernel-unit-test-amd") diff --git a/test/registered/kernels/ops/moe/test_moe_tma_down_padding_rows.py b/test/registered/kernels/ops/moe/test_moe_tma_down_padding_rows.py index ee5503e2f4ea..1beec3fbcf31 100644 --- a/test/registered/kernels/ops/moe/test_moe_tma_down_padding_rows.py +++ b/test/registered/kernels/ops/moe/test_moe_tma_down_padding_rows.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") E, TOPK, HIDDEN, INTER, BLOCK = 256, 8, 2048, 512, [128, 128] CASES = ((1, 16), (32, 16), (300, 32), (2048, 64)) # (num_tokens, BLOCK_SIZE_M) diff --git a/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py b/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py index 3602118f6c3c..669b9c6c7717 100644 --- a/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py +++ b/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py @@ -36,7 +36,7 @@ ) register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="4-gpu-b200") requires_fp4_marlin = pytest.mark.skipif( diff --git a/test/registered/kernels/ops/sampling/test_watermark_selector.py b/test/registered/kernels/ops/sampling/test_watermark_selector.py index 32ca81070744..0ff9886f6701 100644 --- a/test/registered/kernels/ops/sampling/test_watermark_selector.py +++ b/test/registered/kernels/ops/sampling/test_watermark_selector.py @@ -18,7 +18,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=43, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_selector_matches_torch_across_split_boundaries(): diff --git a/test/registered/kernels/ops/sampling/test_watermark_state.py b/test/registered/kernels/ops/sampling/test_watermark_state.py index ea40c897508b..df287e16b435 100644 --- a/test/registered/kernels/ops/sampling/test_watermark_state.py +++ b/test/registered/kernels/ops/sampling/test_watermark_state.py @@ -12,7 +12,7 @@ from sglang.srt.sampling.watermarking.detector import hash_context from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-small") def test_repeated_context_and_greedy_bypass(): diff --git a/test/registered/radix_cache/test_mamba_cache_kl.py b/test/registered/radix_cache/test_mamba_cache_kl.py index 511c80b0fe04..5c2c27c82b97 100644 --- a/test/registered/radix_cache/test_mamba_cache_kl.py +++ b/test/registered/radix_cache/test_mamba_cache_kl.py @@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.server_fixtures.default_fixture import DefaultServerBase -register_cuda_ci(est_time=120, stage="extra-a", runner_config="1-gpu-large") +register_cuda_ci(est_time=120, stage="extra-a", runner_config="1-gpu-small") class TestMambaLazyInplaceKL(DefaultServerBase): diff --git a/test/registered/unit/multimodal_gen/test_minimax_h3_step_contract.py b/test/registered/unit/multimodal_gen/test_minimax_h3_step_contract.py index 1b15b99d156c..36c119fe9e9c 100644 --- a/test/registered/unit/multimodal_gen/test_minimax_h3_step_contract.py +++ b/test/registered/unit/multimodal_gen/test_minimax_h3_step_contract.py @@ -46,7 +46,7 @@ ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") @pytest.fixture(autouse=True) From f127b7b5aeb0f982f409b18802b94aad548cbedf Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Sat, 17 Oct 2026 09:13:39 +0000 Subject: [PATCH 09/10] [CI] Drop the duplicate 1-gpu-small registration of test_compressed_tensors_fp4_numeric Main already registered it on both 1-gpu-small (est_time 15) and 1-gpu-large; the previous commit switched the 1-gpu-large line to 1-gpu-small, so the file was registered twice in base-b-kernel-unit. Keep main's 1-gpu-small line (14s measured on the 5090 with a warm JIT cache). --- .../ops/quantization/test_compressed_tensors_fp4_numeric.py | 1 - 1 file changed, 1 deletion(-) diff --git a/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py b/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py index 669b9c6c7717..7c02211d66a5 100644 --- a/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py +++ b/test/registered/kernels/ops/quantization/test_compressed_tensors_fp4_numeric.py @@ -36,7 +36,6 @@ ) register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-small") -register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-small") register_cuda_ci(est_time=40, stage="base-b-kernel-unit", runner_config="4-gpu-b200") requires_fp4_marlin = pytest.mark.skipif( From b23e1d74c7ebe021eafa9db17565ec6158c36e01 Mon Sep 17 00:00:00 2001 From: Mohammad Angkad Date: Sat, 17 Oct 2026 19:19:52 +0800 Subject: [PATCH 10/10] Keep test_gdn_replayssm_spec_fold, test_modelopt_fp8 and test_mhc_kernels on 1-gpu-large; set est_time from the 5090 reruns --- test/registered/attention/test_gdn_prefill_layout.py | 2 +- test/registered/attention/unittests/dense/test_fa4.py | 2 +- test/registered/attention/unittests/dense/test_flashinfer.py | 2 +- .../registered/attention/unittests/dense/test_flex_attention.py | 2 +- test/registered/attention/unittests/dense/test_torch_native.py | 2 +- test/registered/attention/unittests/dense/test_triton.py | 2 +- .../attention/unittests/gdn/test_gdn_replayssm_spec_fold.py | 2 +- test/registered/attention/unittests/gdn/test_torch_native.py | 2 +- test/registered/attention/unittests/gdn/test_triton.py | 2 +- test/registered/attention/unittests/kda/test_triton.py | 2 +- test/registered/attention/unittests/mamba/test_mamba2.py | 2 +- test/registered/attention/unittests/swa/test_triton.py | 2 +- test/registered/e2e/quantization/test_modelopt_fp8.py | 2 +- test/registered/kernels/ops/attention/test_dcp_lse_combine.py | 2 +- .../kernels/ops/attention/test_dsv32_indexer_fusion.py | 2 +- .../registered/kernels/ops/attention/test_flash_mla_backends.py | 2 +- test/registered/kernels/ops/gemm/test_nvfp4_marlin.py | 2 +- test/registered/kernels/ops/kvcache/test_hicache.py | 2 +- test/registered/kernels/ops/layernorm/test_mhc_kernels.py | 2 +- test/registered/layers/test_layernorm_fusion.py | 2 +- test/registered/lora/test_fused_moe_lora_kernel.py | 2 +- test/registered/moe/test_fused_append_shared_experts.py | 2 +- test/registered/moe/test_topk_padded_region.py | 2 +- test/registered/moe/test_triton_moe_channel_fp8_kernel.py | 2 +- test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py | 2 +- test/registered/unit/layers/quantization/test_fp8_utils.py | 2 +- test/registered/unit/models/test_kimi_k3_bfa_overlap.py | 2 +- test/registered/unit/multimodal/test_cuda_ipc_transport.py | 2 +- 28 files changed, 28 insertions(+), 28 deletions(-) diff --git a/test/registered/attention/test_gdn_prefill_layout.py b/test/registered/attention/test_gdn_prefill_layout.py index ce579779f9a6..b6151152f84f 100644 --- a/test/registered/attention/test_gdn_prefill_layout.py +++ b/test/registered/attention/test_gdn_prefill_layout.py @@ -17,7 +17,7 @@ from sglang.srt.model_executor.forward_batch_info import ForwardMode from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "Test requires CUDA") diff --git a/test/registered/attention/unittests/dense/test_fa4.py b/test/registered/attention/unittests/dense/test_fa4.py index a52288dd74c1..e0347186124b 100644 --- a/test/registered/attention/unittests/dense/test_fa4.py +++ b/test/registered/attention/unittests/dense/test_fa4.py @@ -30,7 +30,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=27, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=22, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=42, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/dense/test_flashinfer.py b/test/registered/attention/unittests/dense/test_flashinfer.py index 1b761240e399..049725bd7124 100644 --- a/test/registered/attention/unittests/dense/test_flashinfer.py +++ b/test/registered/attention/unittests/dense/test_flashinfer.py @@ -27,7 +27,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=38, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf( diff --git a/test/registered/attention/unittests/dense/test_flex_attention.py b/test/registered/attention/unittests/dense/test_flex_attention.py index f4e1e06ef603..6499303073ab 100644 --- a/test/registered/attention/unittests/dense/test_flex_attention.py +++ b/test/registered/attention/unittests/dense/test_flex_attention.py @@ -15,7 +15,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=46, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/dense/test_torch_native.py b/test/registered/attention/unittests/dense/test_torch_native.py index 2f92ec5b5aed..978bb27f38cf 100644 --- a/test/registered/attention/unittests/dense/test_torch_native.py +++ b/test/registered/attention/unittests/dense/test_torch_native.py @@ -12,7 +12,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=34, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/dense/test_triton.py b/test/registered/attention/unittests/dense/test_triton.py index 016ae11be9ff..41013bc82dd5 100644 --- a/test/registered/attention/unittests/dense/test_triton.py +++ b/test/registered/attention/unittests/dense/test_triton.py @@ -31,7 +31,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=25, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py b/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py index 5af352d8369c..f5e531e2284e 100644 --- a/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py +++ b/test/registered/attention/unittests/gdn/test_gdn_replayssm_spec_fold.py @@ -20,7 +20,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") B, T = 3, 4 H, HV = 4, 8 diff --git a/test/registered/attention/unittests/gdn/test_torch_native.py b/test/registered/attention/unittests/gdn/test_torch_native.py index 3959bfda6e37..bde8c5b4fffc 100644 --- a/test/registered/attention/unittests/gdn/test_torch_native.py +++ b/test/registered/attention/unittests/gdn/test_torch_native.py @@ -16,7 +16,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/gdn/test_triton.py b/test/registered/attention/unittests/gdn/test_triton.py index 03b3fec90435..2cc2d0e11b41 100644 --- a/test/registered/attention/unittests/gdn/test_triton.py +++ b/test/registered/attention/unittests/gdn/test_triton.py @@ -30,7 +30,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/kda/test_triton.py b/test/registered/attention/unittests/kda/test_triton.py index 2a25cad9bba1..003e29994e93 100644 --- a/test/registered/attention/unittests/kda/test_triton.py +++ b/test/registered/attention/unittests/kda/test_triton.py @@ -23,7 +23,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=55, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/attention/unittests/mamba/test_mamba2.py b/test/registered/attention/unittests/mamba/test_mamba2.py index 0623ead82ff3..440213e22c83 100644 --- a/test/registered/attention/unittests/mamba/test_mamba2.py +++ b/test/registered/attention/unittests/mamba/test_mamba2.py @@ -33,7 +33,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=31, stage="base-b", runner_config="1-gpu-small") @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required") diff --git a/test/registered/attention/unittests/swa/test_triton.py b/test/registered/attention/unittests/swa/test_triton.py index 897e0a0349cd..c43cf9e84a13 100644 --- a/test/registered/attention/unittests/swa/test_triton.py +++ b/test/registered/attention/unittests/swa/test_triton.py @@ -24,7 +24,7 @@ from sglang.test.test_utils import CustomTestCase register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200") -register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=25, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd") diff --git a/test/registered/e2e/quantization/test_modelopt_fp8.py b/test/registered/e2e/quantization/test_modelopt_fp8.py index cd9fe495ddd9..3958fbdde51a 100644 --- a/test/registered/e2e/quantization/test_modelopt_fp8.py +++ b/test/registered/e2e/quantization/test_modelopt_fp8.py @@ -15,7 +15,7 @@ # Two classes run from this file: the default server plus the Rust-frontend # variant (when the embedded extension is built), each launches a server + eval. -register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-large") class TestModeloptFP8(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_dcp_lse_combine.py b/test/registered/kernels/ops/attention/test_dcp_lse_combine.py index 7f5ff559fa09..d9e58b520344 100644 --- a/test/registered/kernels/ops/attention/test_dcp_lse_combine.py +++ b/test/registered/kernels/ops/attention/test_dcp_lse_combine.py @@ -16,7 +16,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-small") class TestLSECombineTritonVsCPU(CustomTestCase): diff --git a/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py b/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py index 56edbc97288b..554ae9f98dc0 100644 --- a/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py +++ b/test/registered/kernels/ops/attention/test_dsv32_indexer_fusion.py @@ -31,7 +31,7 @@ _is_hip = is_hip() -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=36, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=45, suite="jit-kernel-unit-test-amd") HEAD_DIM = 128 diff --git a/test/registered/kernels/ops/attention/test_flash_mla_backends.py b/test/registered/kernels/ops/attention/test_flash_mla_backends.py index 6781213eeb16..f48000093c32 100644 --- a/test/registered/kernels/ops/attention/test_flash_mla_backends.py +++ b/test/registered/kernels/ops/attention/test_flash_mla_backends.py @@ -54,7 +54,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=18, stage="base-b", runner_config="1-gpu-small") # Per-token byte layout diff --git a/test/registered/kernels/ops/gemm/test_nvfp4_marlin.py b/test/registered/kernels/ops/gemm/test_nvfp4_marlin.py index 9e7436acfc5e..9efb049fbebe 100644 --- a/test/registered/kernels/ops/gemm/test_nvfp4_marlin.py +++ b/test/registered/kernels/ops/gemm/test_nvfp4_marlin.py @@ -19,7 +19,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_marlin_utils import make_nvfp4_weight_and_ref -register_cuda_ci(est_time=142, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-small") @pytest.mark.skipif( diff --git a/test/registered/kernels/ops/kvcache/test_hicache.py b/test/registered/kernels/ops/kvcache/test_hicache.py index b934e40b2271..904c494719f4 100644 --- a/test/registered/kernels/ops/kvcache/test_hicache.py +++ b/test/registered/kernels/ops/kvcache/test_hicache.py @@ -24,7 +24,7 @@ from sglang.srt.utils import is_cuda, is_hip, is_npu, is_xpu from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=43, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=86, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=12, stage="jit-kernel-unit", runner_config="amd") pytestmark = pytest.mark.skipif( diff --git a/test/registered/kernels/ops/layernorm/test_mhc_kernels.py b/test/registered/kernels/ops/layernorm/test_mhc_kernels.py index 5193ea7722c3..d54456a4c5c5 100644 --- a/test/registered/kernels/ops/layernorm/test_mhc_kernels.py +++ b/test/registered/kernels/ops/layernorm/test_mhc_kernels.py @@ -8,7 +8,7 @@ from sglang.kernels.ops.layernorm.mhc import mhc_fused_post_pre, mhc_post, mhc_pre from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=127, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large") @pytest.fixture diff --git a/test/registered/layers/test_layernorm_fusion.py b/test/registered/layers/test_layernorm_fusion.py index 7145626f727f..3a48c6e033f2 100644 --- a/test/registered/layers/test_layernorm_fusion.py +++ b/test/registered/layers/test_layernorm_fusion.py @@ -7,7 +7,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=23, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=2, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/lora/test_fused_moe_lora_kernel.py b/test/registered/lora/test_fused_moe_lora_kernel.py index 73ad010cd3e0..c38edfff8b03 100644 --- a/test/registered/lora/test_fused_moe_lora_kernel.py +++ b/test/registered/lora/test_fused_moe_lora_kernel.py @@ -17,7 +17,7 @@ # ============================================================================== -register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=33, stage="base-b", runner_config="1-gpu-small") def round_up(x, base): diff --git a/test/registered/moe/test_fused_append_shared_experts.py b/test/registered/moe/test_fused_append_shared_experts.py index e3d2eb0f7d2a..d2dbefa2a220 100644 --- a/test/registered/moe/test_fused_append_shared_experts.py +++ b/test/registered/moe/test_fused_append_shared_experts.py @@ -22,7 +22,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase, empty_gpu_cache -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-small-amd") diff --git a/test/registered/moe/test_topk_padded_region.py b/test/registered/moe/test_topk_padded_region.py index 5c8faed2f20d..601c618db1b6 100644 --- a/test/registered/moe/test_topk_padded_region.py +++ b/test/registered/moe/test_topk_padded_region.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-small") register_amd_ci(est_time=60, stage="stage-b", runner_config="1-gpu-small-amd") _IS_HIP = is_hip() diff --git a/test/registered/moe/test_triton_moe_channel_fp8_kernel.py b/test/registered/moe/test_triton_moe_channel_fp8_kernel.py index 0057660069a3..45c44530e7cd 100644 --- a/test/registered/moe/test_triton_moe_channel_fp8_kernel.py +++ b/test/registered/moe/test_triton_moe_channel_fp8_kernel.py @@ -11,7 +11,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=30, stage="base-b", runner_config="1-gpu-small") def native_w8a8_per_token_matmul(A, B, As, Bs, output_dtype=torch.float16): diff --git a/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py b/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py index a27d43ac4612..9357eb637847 100644 --- a/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py +++ b/test/registered/unit/layers/moe/test_deepep_v2_masked_slab.py @@ -12,7 +12,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=25, stage="base-b", runner_config="1-gpu-small") DEVICE = "cuda" diff --git a/test/registered/unit/layers/quantization/test_fp8_utils.py b/test/registered/unit/layers/quantization/test_fp8_utils.py index 3fa1216a1d07..f3226cdfb2df 100644 --- a/test/registered/unit/layers/quantization/test_fp8_utils.py +++ b/test/registered/unit/layers/quantization/test_fp8_utils.py @@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-small") class TestMxfp8MoeScaleLayout(CustomTestCase): diff --git a/test/registered/unit/models/test_kimi_k3_bfa_overlap.py b/test/registered/unit/models/test_kimi_k3_bfa_overlap.py index 9cf0493bc22a..1f9ceaee08f6 100644 --- a/test/registered/unit/models/test_kimi_k3_bfa_overlap.py +++ b/test/registered/unit/models/test_kimi_k3_bfa_overlap.py @@ -14,7 +14,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=24, stage="base-b", runner_config="1-gpu-small") _H = 7168 _QKVG = 6144 # q,k,v,g slices per rank at TP8 diff --git a/test/registered/unit/multimodal/test_cuda_ipc_transport.py b/test/registered/unit/multimodal/test_cuda_ipc_transport.py index 227a584c3dbe..7e37189f93b8 100644 --- a/test/registered/unit/multimodal/test_cuda_ipc_transport.py +++ b/test/registered/unit/multimodal/test_cuda_ipc_transport.py @@ -28,7 +28,7 @@ from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=37, stage="base-b", runner_config="1-gpu-small") +register_cuda_ci(est_time=57, stage="base-b", runner_config="1-gpu-small") def _produce_pooled_tensor(proxy_queue, consumer_done, result_queue):