docs: improve CI and testing documentation (#21202)

Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

python/sglang/jit_kernel/tests/test_add_constant.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 
@@ -13,4 +15,4 @@ def test_add_constant(size: int, constant: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_awq_dequantize.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -161,4 +162,4 @@ def test_awq_dequantize_jit_vs_aot(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_awq_marlin_moe_repack.py

@@ -1,3 +1,5 @@
+import sys
+
 import numpy as np
 import pytest
 import torch
@@ -116,4 +118,4 @@ def test_awq_marlin_moe_repack_shape(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_awq_marlin_repack.py

@@ -1,3 +1,5 @@
+import sys
+
 import numpy as np
 import pytest
 import torch
@@ -102,4 +104,4 @@ def test_awq_marlin_repack_correct(num_bits, k_tiles, n_tiles, group_size):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_clamp_position.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 
@@ -37,4 +39,4 @@ class TestClampPosition:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_concat_mla.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -166,4 +167,4 @@ def test_concat_mla_absorb_q_jit_vs_aot(dim_0: int, dim_1: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_cutedsl_gdn.py

@@ -1,5 +1,7 @@
 """Tests for CuTe DSL fused sigmoid gating delta rule kernel (GDN)."""
 
+import sys
+
 import numpy as np
 import pytest
 import torch
@@ -302,4 +304,4 @@ def test_cutedsl_gdn_performance(B: int):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_flash_attention_4.py

@@ -4,6 +4,7 @@
 
 import itertools
 import math
+import sys
 
 import pytest
 import torch
@@ -1501,4 +1502,4 @@ def _generate_block_kvcache(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_fused_add_rmsnorm.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -58,4 +59,4 @@ def test_fused_add_rmsnorm(batch_size: int, hidden_size: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_fused_metadata_copy.py

@@ -8,6 +8,7 @@ This test suite verifies:
 4. Performance benchmarks and speedup measurements
 """
 
+import sys
 import time
 
 import pytest
@@ -1064,4 +1065,4 @@ def test_fused_metadata_copy_multi_large_batch(bs):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_fused_norm_scale_shift.py

@@ -1,3 +1,4 @@
+import sys
 from typing import Optional, Tuple
 
 import pytest
@@ -233,4 +234,4 @@ class TestFusedScaleResidualNormScaleShift:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_fused_store_index_cache.py

@@ -16,6 +16,7 @@ Design Notes:
 
 from __future__ import annotations
 
+import sys
 from typing import Optional, Tuple
 
 import pytest
@@ -450,4 +451,4 @@ def test_reference_writes_nonzero():
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_fused_verify_triton_gdn.py

@@ -6,6 +6,8 @@ the reference two-step implementation:
     2. o = fused_recurrent_gated_delta_rule_update(q, k, v, g, beta, ...)
 """
 
+import sys
+
 import pytest
 import torch
 
@@ -228,4 +230,4 @@ def test_mtp_single_step_decode(N: int):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_gptq_marlin.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 from sgl_kernel.scalar_type import scalar_types
@@ -96,4 +98,4 @@ def test_gptq_marlin_gemm(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_gptq_marlin_repack.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 from sgl_kernel.scalar_type import scalar_types
@@ -87,4 +89,4 @@ def test_gptq_marlin_repack(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_hadamard_jit.py

@@ -1,4 +1,5 @@
 import math
+import sys
 
 import numpy as np
 import pytest
@@ -420,4 +421,4 @@ def test_hadamard_transform_40n(dim, dtype):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))

python/sglang/jit_kernel/tests/test_moe_lora_align_block_size.py

@@ -1,5 +1,6 @@
 # Temporarily adapted from https://github.com/vllm-project/vllm/blob/main/tests/lora/test_moe_lora_align_sum.py, will optimize in future refactor
 import random
+import sys
 
 import pytest
 import torch
@@ -163,4 +164,4 @@ def test_moe_lora_align_block_size(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))

python/sglang/jit_kernel/tests/test_moe_wna16_marlin.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -335,4 +336,4 @@ def test_moe_wna16_marlin_gemm(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_norm_jit.py

@@ -1,5 +1,7 @@
 # Adapted from sgl-kernel/tests/test_norm.py
 
+import sys
+
 import pytest
 import torch
 
@@ -135,4 +137,4 @@ def test_rmsnorm_rejects_unsupported_hidden_size(hidden_size):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))

python/sglang/jit_kernel/tests/test_per_tensor_quant_fp8.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 from typing import Optional, Tuple
 
 import pytest
@@ -83,4 +84,4 @@ def test_jit_per_tensor_quant_supports_3d(shape):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_per_token_group_quant_8bit.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -202,4 +203,4 @@ def test_per_token_group_quant_with_column_major(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_pos_enc.py

@@ -1,3 +1,4 @@
+import sys
 import time
 from typing import Optional, Tuple, Union
 
@@ -486,4 +487,4 @@ def test_performance(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_qknorm.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -93,4 +94,4 @@ def test_qknorm(batch_size: int, n_k: int, n_q: int, head_dim: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_qknorm_across_heads.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -75,4 +76,4 @@ def test_qknorm_across_heads(batch_size: int, hidden_dim: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_qwen_image_modulation.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 import triton
@@ -217,4 +219,4 @@ def test_fused_residual_layernorm_scale_shift_gate_select01(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_renorm.py

@@ -1,6 +1,8 @@
 # Adapted from https://github.com/flashinfer-ai/flashinfer/blob/main/tests/test_sampling.py
 # and /sgl-workspace/sglang/sgl-kernel/tests/test_sampling.py
 
+import sys
+
 import pytest
 import sgl_kernel
 import torch
@@ -115,4 +117,4 @@ def test_top_k_mask_logits(batch_size, vocab_size, k, neginf_input):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))

python/sglang/jit_kernel/tests/test_resolve_future_token_ids.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 
@@ -60,4 +62,4 @@ class TestResolveFutureTokenIds:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_rmsnorm.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -44,4 +45,4 @@ def test_rmsnorm(batch_size: int, hidden_size: int) -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_rope.py

@@ -1,3 +1,5 @@
+import sys
+
 import pytest
 import torch
 import triton
@@ -248,4 +250,4 @@ def test_fused_rope_store(
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_store_cache.py

@@ -1,4 +1,5 @@
 import itertools
+import sys
 
 import pytest
 import torch
@@ -124,4 +125,4 @@ def test_can_use_store_cache() -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/jit_kernel/tests/test_timestep_embedding.py

@@ -1,4 +1,5 @@
 import os
+import sys
 
 import numpy as np
 import pytest
@@ -175,4 +176,4 @@ def test_timestep_embedding_perf():
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/multimodal_gen/.claude/skills/diffusion-kernel/add-cuda-kernel.md

@@ -420,7 +420,8 @@ def test_rmsnorm_correctness(dtype, shape, has_weight):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    import sys
+    sys.exit(pytest.main([__file__, "-v", "-s"]))
 ```
 
 ---

python/sglang/multimodal_gen/.claude/skills/diffusion-kernel/add-triton-kernel.md

@@ -292,7 +292,8 @@ def test_fused_scale_shift_output_dtype_preserved():
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v"])
+    import sys
+    sys.exit(pytest.main([__file__, "-v"]))
 ```
 
 Run:

python/sglang/multimodal_gen/.claude/skills/diffusion-kernel/references/kernel-templates.md

@@ -502,7 +502,8 @@ def test_<op>_cpu_error():
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    import sys
+    sys.exit(pytest.main([__file__, "-v", "-s"]))
 ```
 
 ---

python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/test/test_flux_pipeline.py

@@ -1,6 +1,7 @@
 """Test for ComfyUIFluxPipeline with pass-through scheduler."""
 
 import os
+import sys
 
 import pytest
 import torch
@@ -159,4 +160,4 @@ def test_comfyui_flux_pipeline_direct() -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v"])
+    sys.exit(pytest.main([__file__, "-v"]))

python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/test/test_qwen_image_edit_pipeline.py

@@ -1,6 +1,7 @@
 """Test for ComfyUIQwenImageEditPipeline with pass-through scheduler (I2I/edit mode)."""
 
 import os
+import sys
 
 import pytest
 import torch
@@ -132,4 +133,4 @@ def test_comfyui_qwen_image_edit_pipeline_direct() -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v"])
+    sys.exit(pytest.main([__file__, "-v"]))

python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/test/test_qwen_image_pipeline.py

@@ -1,6 +1,7 @@
 """Test for ComfyUIQwenImagePipeline with pass-through scheduler."""
 
 import os
+import sys
 
 import pytest
 import torch
@@ -116,4 +117,4 @@ def test_comfyui_qwen_image_pipeline_direct() -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v"])
+    sys.exit(pytest.main([__file__, "-v"]))

python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/test/test_zimage_pipeline.py

@@ -1,6 +1,7 @@
 """Test for ComfyUIZImagePipeline with pass-through scheduler."""
 
 import os
+import sys
 
 import pytest
 import torch
@@ -118,4 +119,4 @@ def test_comfyui_zimage_pipeline_direct() -> None:
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v"])
+    sys.exit(pytest.main([__file__, "-v"]))

python/sglang/multimodal_gen/test/server/test_update_weights_from_disk.py

@@ -139,6 +139,7 @@ import functools
 import os
 import random
 import shutil
+import sys
 import tempfile
 import threading
 from collections.abc import Callable
@@ -671,4 +672,4 @@ class TestUpdateWeightsFromDiskWithOffload(_UpdateWeightsApiMixin):
 
 
 if __name__ == "__main__":
-    pytest.main([__file__, "-v", "-s"])
+    sys.exit(pytest.main([__file__, "-v", "-s"]))

python/sglang/srt/configs/model_config.py

@@ -1094,13 +1094,6 @@ class ModelConfig:
                     f"or model type {self.hf_config.model_type}. "
                     "Please upgrade transformers to >= 5.0.0."
                 )
-        elif not needs_tf_v5:
-            logger.warning(
-                f"Transformers version {tf_version_str} is used for model type {self.hf_config.model_type}. "
-                "If you experience issues related to RoPE parameters, "
-                "they may be due to incompatibilities between Transformers >=5.0.0 and some models. "
-                "You can try downgrading to transformers==4.57.1 as a workaround."
-            )
 
     def _get_hf_eos_token_id(self) -> Optional[Set[int]]:
         eos_ids = getattr(self.hf_config, "eos_token_id", None)

python/sglang/srt/layers/quantization/fp8_kernel.py

@@ -42,6 +42,7 @@ from sglang.srt.utils import (
     log_info_on_rank0,
 )
 from sglang.srt.utils.custom_op import register_custom_op
+from sglang.srt.utils.patch_torch import register_fake_if_exists
 
 _is_hip = is_hip()
 _is_cuda = is_cuda()
@@ -2058,7 +2059,7 @@ def triton_scaled_mm(
 if _is_cuda:
     if enable_sgl_per_token_group_quant_8bit:
 
-        @torch.library.register_fake("sgl_kernel::sgl_per_token_group_quant_8bit")
+        @register_fake_if_exists("sgl_kernel::sgl_per_token_group_quant_8bit")
         def _(
             input, output_q, output_s, group_size, eps, fp8_min, fp8_max, scale_ue8m0
         ):
@@ -2066,12 +2067,12 @@ if _is_cuda:
 
     else:
 
-        @torch.library.register_fake("sgl_kernel::sgl_per_token_group_quant_fp8")
+        @register_fake_if_exists("sgl_kernel::sgl_per_token_group_quant_fp8")
         def _(
             input, output_q, output_s, group_size, eps, fp8_min, fp8_max, scale_ue8m0
         ):
             return
 
-    @torch.library.register_fake("sgl_kernel::sgl_per_token_quant_fp8")
+    @register_fake_if_exists("sgl_kernel::sgl_per_token_quant_fp8")
     def _(input, output_q, output_s):
         return

python/sglang/srt/layers/quantization/fp8_utils.py

@@ -92,14 +92,16 @@ if _use_aiter:
 if _is_cuda:
     from sgl_kernel import fp8_blockwise_scaled_mm, fp8_scaled_mm
 
-    @torch.library.register_fake("sgl_kernel::fp8_scaled_mm")
+    from sglang.srt.utils.patch_torch import register_fake_if_exists
+
+    @register_fake_if_exists("sgl_kernel::fp8_scaled_mm")
     def _fp8_scaled_mm_abstract(mat_a, mat_b, scales_a, scales_b, out_dtype, bias=None):
         # mat_a: [M, K], mat_b: [K, N] or [N, K] depending on callsite layout; output is [M, N].
         M = mat_a.shape[-2]
         N = mat_b.shape[-1]
         return mat_a.new_empty((M, N), dtype=out_dtype)
 
-    @torch.library.register_fake("sgl_kernel::fp8_blockwise_scaled_mm")
+    @register_fake_if_exists("sgl_kernel::fp8_blockwise_scaled_mm")
     def _fp8_blockwise_scaled_mm_abstract(mat_a, mat_b, scales_a, scales_b, out_dtype):
         # mat_a: [M, K], mat_b: [K, N] or [N, K] depending on callsite layout; output is [M, N].
         M = mat_a.shape[-2]

python/sglang/srt/mem_cache/hisparse_memory_pool.py

@@ -5,17 +5,26 @@ from typing import Optional
 
 import torch
 
-from sglang.srt.utils import is_npu
-
-if not is_npu():
-    from sgl_kernel.kvcacheio import transfer_kv_all_layer_mla
-
 from sglang.srt.layers.radix_attention import RadixAttention
 from sglang.srt.mem_cache.allocator import (
     BaseTokenToKVPoolAllocator,
     PagedTokenToKVPoolAllocator,
 )
 from sglang.srt.mem_cache.memory_pool import NSATokenToKVPool
+from sglang.srt.utils import is_cuda, is_hip
+
+# sgl_kernel.kvcacheio is only available in CUDA/ROCm sgl-kernel builds (not XPU/MPS/NPU/CPU).
+_is_cuda = is_cuda()
+_is_hip = is_hip()
+if _is_cuda or _is_hip:
+    from sgl_kernel.kvcacheio import transfer_kv_all_layer_mla
+else:
+
+    def transfer_kv_all_layer_mla(*args, **kwargs):
+        raise RuntimeError(
+            "HiSparse device KV transfer requires sgl_kernel.kvcacheio (CUDA/ROCm). "
+            "It is not available on this backend."
+        )
 
 
 class HiSparseNSATokenToKVPool(NSATokenToKVPool):

python/sglang/srt/mem_cache/storage/hf3fs/test_hf3fs_utils.py

@@ -1,4 +1,5 @@
 import multiprocessing.shared_memory
+import sys
 from pathlib import Path
 
 import pytest
@@ -40,4 +41,4 @@ def test_rw_shm():
 
 
 if __name__ == "__main__":
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))

python/sglang/srt/utils/hf_transformers_utils.py

@@ -797,11 +797,6 @@ def get_tokenizer(
     # when trust_remote_code=False and the model requires a custom tokenizer.
     # Detect this and auto-retry with trust_remote_code=True.
     if not trust_remote_code and type(tokenizer).__name__ == "TokenizersBackend":
-        logger.info(
-            "Detected generic TokenizersBackend for %s, "
-            "retrying with trust_remote_code=True",
-            tokenizer_name,
-        )
         tokenizer = AutoTokenizer.from_pretrained(
             tokenizer_name,
             *args,

python/sglang/test/ci/ci_register.py

@@ -31,10 +31,14 @@ class HWBackend(Enum):
 class CIRegistry:
     backend: HWBackend
     filename: str
+    # Estimated time to run the test in seconds.
     est_time: float
+    # The suite this test is registered in.
     suite: str
+    # Whether the test is a nightly test.
     nightly: bool = False
-    disabled: Optional[str] = None  # None = enabled, string = disabled with reason
+    # Reason for disabling the test. None = enabled, string = disabled with reason.
+    disabled: Optional[str] = None
 
 
 def register_cpu_ci(

python/sglang/test/test_custom_ops.py

@@ -1,5 +1,7 @@
 # Adapted from https://github.com/vllm-project/vllm/blob/8ca7a71df787ad711ad3ac70a5bd2eb2bb398938/tests/quantization/test_fp8.py
 
+import sys
+
 import pytest
 import torch
 
@@ -145,4 +147,4 @@ if _is_cuda or _is_hip:
 
 if __name__ == "__main__":
     # Run the specific test function directly
-    pytest.main([__file__])
+    sys.exit(pytest.main([__file__]))