diff --git a/.claude/skills/add-jit-kernel/SKILL.md b/.claude/skills/add-jit-kernel/SKILL.md index b2a1fcb25..8095ea5f1 100644 --- a/.claude/skills/add-jit-kernel/SKILL.md +++ b/.claude/skills/add-jit-kernel/SKILL.md @@ -435,31 +435,34 @@ if torch.cuda.get_device_capability()[0] < 9: JIT kernel correctness tests and benchmarks live under `test/registered/jit/` and `test/registered/jit/benchmark/` (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/jit_kernel/` 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-1-gpu-large` on H100 and `base-b-kernel-unit-1-gpu-b200` on B200/SM100 paths (see `.github/workflows/pr-test-jit-kernel.yml`). Multi-GPU JIT tests use `base-b-kernel-unit-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-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`. - **Nightly kernel suite**: `nightly-kernel-1-gpu` with `--nightly` — typically used with `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` in CI for expanded parameter grids (see `python/sglang/jit_kernel/utils.py` → `should_run_full_tests` / `get_ci_test_range`). Wired in `.github/workflows/nightly-test-nvidia.yml` (e.g. `python3 run_suite.py --hw cuda --suite nightly-kernel-1-gpu --nightly --continue-on-error`). -Registration pattern (module level, **literal** `est_time` and `suite` strings — required for AST parsing): +Registration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing): ```python from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") # Optional B200/SM100 registration for tests that cover Blackwell-specific code paths -# register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-b200") +# register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="4-gpu-b200") # Optional second registration: same file also listed under the nightly kernel suite +# (nightly suites use the legacy single-string suite=, not stage/runner_config) # register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) ``` -Keep `est_time` and `suite` as literal values. `run_suite.py` collects them from the file AST, so computed values and helper wrappers can break CI discovery. +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`. The single-string `suite=` form is only for nightly/stress/weekly suites. + +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. Use `register_cuda_ci(..., disabled="reason")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only). **Run like CI** (from repo root): ```bash -(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-1-gpu-large) +(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large) # For B200/SM100-specific coverage: -(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-1-gpu-b200) +(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200) ``` For fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`. @@ -472,7 +475,7 @@ import torch from sglang.jit_kernel.scale import scale from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32]) @@ -517,7 +520,7 @@ if __name__ == "__main__": ## Step 5: Add a benchmark (required) -Benchmarks are `bench_*.py` files under `test/registered/jit/benchmark/`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-1-gpu-large`). +Benchmarks are `bench_*.py` files under `test/registered/jit/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 use the project's own `marker` framework (in `python/sglang/jit_kernel/benchmark/marker.py`) — **do not** use `triton.testing.perf_report` / `triton.testing.do_bench` directly. The marker framework provides (public names: `benchmark`, `parametrize`, `do_bench`, `skip`, `BenchResult`, `BenchSkip`): @@ -544,7 +547,7 @@ from sglang.jit_kernel.benchmark.utils import create_random from sglang.jit_kernel.scale import scale as jit_scale from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large") @torch.compile() @@ -598,14 +601,14 @@ python test/registered/jit/benchmark/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-1-gpu-large +cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large ``` --- ## Troubleshooting -- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time` and `suite` (and optional `nightly=True`); starred args and non-literal values break AST collection +- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time`, `stage`, and `runner_config` (and optional `nightly=True`); starred args and non-literal values break AST collection - **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/jit_kernel/csrc/`; reduce template argument combinations - **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...` - **Unstable benchmark results**: `marker.do_bench` uses CUDA-graph-based timing by default; set `use_cuda_graph=False` only if the kernel can't be captured. Make sure `graph_clone_args` covers every *read* tensor — reusing a single buffer keeps it L2-hot and skews results diff --git a/.claude/skills/ci-workflow-guide/SKILL.md b/.claude/skills/ci-workflow-guide/SKILL.md index e5eb3d6ef..531a42b9b 100644 --- a/.claude/skills/ci-workflow-guide/SKILL.md +++ b/.claude/skills/ci-workflow-guide/SKILL.md @@ -277,10 +277,10 @@ 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-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — | -| `base-b-kernel-unit-1-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — | -| `base-b-kernel-unit-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — | -| `base-b-kernel-benchmark-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — | +| `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-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` | — | | `base-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — | @@ -290,6 +290,8 @@ Large suites are split across matrix jobs using the **LPT (Longest Processing Ti | `base-c-test-8-gpu-b200` | registered only | `8-gpu-b200` | — | | `base-c-test-4-gpu-gb200` | registered only | `4-gpu-gb200` | — | +> **Suite names are generated**, not hand-written: each comes from a test's `register_*_ci(stage=..., runner_config=...)` as `{stage}-test-{runner_config}`, and `runner_config` maps to the `Runner` column via `scripts/ci/runner_configs.yml`. +> > **Note**: Kernel suites (`base-b-kernel-*`) run via `pr-test-jit-kernel.yml` and `pr-test-sgl-kernel.yml`, not the main `pr-test.yml`. `base-c-test-8-gpu-b200` is registered in `test/run_suite.py` but not wired to PR CI. The GB200 job is currently commented out in `pr-test.yml` until a company-owned runner is provisioned. Multimodal diffusion uses `python/sglang/multimodal_gen/test/run_suite.py`, not `test/run_suite.py`. **Workflow usage:** diff --git a/.claude/skills/write-sglang-test/SKILL.md b/.claude/skills/write-sglang-test/SKILL.md index 2f1f832f6..d6bbd86e9 100644 --- a/.claude/skills/write-sglang-test/SKILL.md +++ b/.claude/skills/write-sglang-test/SKILL.md @@ -50,8 +50,13 @@ Defined in `python/sglang/test/test_utils.py`: ### Naming Conventions -- **Suite**: `base-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `base-b-test-1-gpu-small`) -- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`) +A per-commit suite name is **generated** from registration metadata as `{stage}-test-{runner_config}` — you don't hand-write it: + +- **`stage`** — the CI stage (e.g. `base-b`, `base-b-kernel-unit`, `base-c`). +- **`runner_config`** — a runner-pool key from `scripts/ci/runner_configs.yml`, which maps it to the physical runner label (so `1-gpu-large` runs on `1-gpu-h100`). AMD/NPU use their own keys (e.g. `amd`). +- **Suite** — `register_cuda_ci(stage="base-b", runner_config="1-gpu-small")` → `base-b-test-1-gpu-small`, the name you pass to `run_suite.py --suite`. The `-test-` is just the connector; never put it in `register_*_ci`. + +> Legacy single-string `suite=` is only for suites that don't fit that shape — nightly/stress/weekly and some AMD/CPU/NPU pools (e.g. `suite="nightly-kernel-1-gpu", nightly=True`). Per-commit tests always use `stage=` + `runner_config=`. ### All CI Suites @@ -65,10 +70,10 @@ Defined in `python/sglang/test/test_utils.py`: | `base-b-test-1-gpu-large` | `1-gpu-h100` | Tests that need H100-class memory or kernels (e.g. FA3) | | `base-b-test-2-gpu-large` | `2-gpu-h100` | Two-GPU correctness and parallelism (TP/PP) on H100 | | `base-b-test-4-gpu-b200` | `4-gpu-b200` | Early Blackwell coverage (SM100+ paths) on four GPUs | -| `base-b-kernel-unit-1-gpu-large` | `1-gpu-h100` | JIT kernel correctness tests under `test/registered/jit/` | -| `base-b-kernel-unit-1-gpu-b200` | `4-gpu-b200` | JIT kernel correctness tests for Blackwell / SM100-specific paths | -| `base-b-kernel-unit-8-gpu-h200` | `8-gpu-h200` | Multi-GPU JIT kernel correctness tests under `test/registered/jit/` | -| `base-b-kernel-benchmark-1-gpu-large` | `1-gpu-h100` | JIT kernel benchmark files under `test/registered/jit/benchmark/` | +| `base-b-kernel-unit-test-1-gpu-large` | `1-gpu-h100` | JIT kernel correctness tests under `test/registered/jit/` | +| `base-b-kernel-unit-test-4-gpu-b200` | `4-gpu-b200` | JIT kernel correctness tests for Blackwell / SM100-specific paths | +| `base-b-kernel-unit-test-8-gpu-h200` | `8-gpu-h200` | Multi-GPU JIT kernel correctness tests under `test/registered/jit/` | +| `base-b-kernel-benchmark-test-1-gpu-large` | `1-gpu-h100` | JIT kernel benchmark files under `test/registered/jit/benchmark/` | | `base-c-test-4-gpu-h100` | `4-gpu-h100` | Large 4-GPU H100 integration and scaling tests | | `base-c-test-8-gpu-h200` | `8-gpu-h200` | Large 8-GPU H200 runs for big models and parallelism | | `base-c-test-8-gpu-h20` | `8-gpu-h20` | Large 8-GPU H20 runs for big models | @@ -132,9 +137,9 @@ 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-1-gpu-large` -- **JIT kernel correctness for B200 / SM100 paths** → `base-b-kernel-unit-1-gpu-b200` -- **JIT kernel benchmarks** → `base-b-kernel-benchmark-1-gpu-large` +- **JIT kernel correctness** → `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 --- @@ -353,19 +358,19 @@ JIT kernel files live outside `test/registered/` but still use registration: from sglang.test.ci.ci_register import register_cuda_ci # Correctness tests in test/registered/jit/ -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-b200") -register_cuda_ci(est_time=120, suite="base-b-kernel-unit-8-gpu-h200") +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="4-gpu-b200") +register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="8-gpu-h200") # Benchmarks in test/registered/jit/benchmark/ -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large") -# Optional nightly registration +# Optional nightly registration — nightly suites use the legacy single-string suite= register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) register_cuda_ci(est_time=120, suite="nightly-kernel-8-gpu-h200", nightly=True) ``` -Keep `est_time` and `suite` as **literal values** — `run_suite.py` collects them by AST parsing +The `stage` + `runner_config` calls generate suites like `base-b-kernel-unit-test-1-gpu-large`; nightly keeps the legacy `suite=` string. Keep `est_time`, `stage`, `runner_config`, and `suite` as **literal values** — `run_suite.py` collects them by AST parsing. --- diff --git a/.github/workflows/pr-test-jit-kernel.yml b/.github/workflows/pr-test-jit-kernel.yml index d14c4fc3b..b7a923325 100644 --- a/.github/workflows/pr-test-jit-kernel.yml +++ b/.github/workflows/pr-test-jit-kernel.yml @@ -76,7 +76,7 @@ jobs: timeout-minutes: 30 run: | cd test/ - python3 run_suite.py --hw cuda --suite base-b-kernel-unit-1-gpu-large + python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large jit-kernel-multigpu-unit-test: if: | @@ -114,7 +114,7 @@ jobs: timeout-minutes: 45 run: | cd test/ - python3 run_suite.py --hw cuda --suite base-b-kernel-unit-8-gpu-h200 + python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-8-gpu-h200 jit-kernel-benchmark-test: if: | @@ -154,7 +154,7 @@ jobs: timeout-minutes: 45 run: | cd test/ - python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-1-gpu-large + python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large jit-kernel-b200-test: if: | @@ -194,4 +194,4 @@ jobs: timeout-minutes: 30 run: | cd test/ - python3 run_suite.py --hw cuda --suite base-b-kernel-unit-1-gpu-b200 + python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200 diff --git a/test/README.md b/test/README.md index bfd28ef12..cfa439061 100644 --- a/test/README.md +++ b/test/README.md @@ -73,8 +73,8 @@ Parameters: `est_time` (seconds), `stage` + `runner_config` (target stage and ru Keep `est_time`, `stage`, `runner_config` as **literal values** — `run_suite.py` collects them by AST parsing. JIT kernel correctness tests and benchmarks live under `test/registered/jit/`, same as other registered tests (their helpers stay alongside the kernel source under `python/sglang/jit_kernel/` and are imported by absolute path): -- Correctness tests: `test/registered/jit/test_*.py` → `base-b-kernel-unit-1-gpu-large` -- Benchmarks: `test/registered/jit/benchmark/bench_*.py` → `base-b-kernel-benchmark-1-gpu-large` +- Correctness tests: `test/registered/jit/test_*.py` → `base-b-kernel-unit-test-1-gpu-large` +- Benchmarks: `test/registered/jit/benchmark/bench_*.py` → `base-b-kernel-benchmark-test-1-gpu-large` ## Choosing a Suite @@ -85,8 +85,8 @@ 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-1-gpu-large` | -| JIT kernel benchmarks | `base-b-kernel-benchmark-1-gpu-large` | +| JIT kernel correctness | `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_gdn_prefill_cutedsl.py b/test/registered/attention/test_gdn_prefill_cutedsl.py index 807be8a66..cefaa78e5 100644 --- a/test/registered/attention/test_gdn_prefill_cutedsl.py +++ b/test/registered/attention/test_gdn_prefill_cutedsl.py @@ -15,7 +15,7 @@ from sglang.test.ci.ci_register import register_cuda_ci # CuteDSL prefill kernel only exists on Blackwell. Single-GPU kernel-unit # suite is the right slot (matches existing jit_kernel test_*.py pattern). -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-b200") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="4-gpu-b200") if not (torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10): pytest.skip( diff --git a/test/registered/attention/test_kda_prefill_cutedsl.py b/test/registered/attention/test_kda_prefill_cutedsl.py index 2efefdc76..96ffc08bb 100644 --- a/test/registered/attention/test_kda_prefill_cutedsl.py +++ b/test/registered/attention/test_kda_prefill_cutedsl.py @@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci # CuteDSL prefill kernel only exists on Blackwell. Single-GPU kernel-unit suite, # same slot as the GDN prefill test. -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-b200") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="4-gpu-b200") if not (torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10): pytest.skip( diff --git a/test/registered/jit/benchmark/bench_activation.py b/test/registered/jit/benchmark/bench_activation.py index 9277ec132..a46c9a3e9 100644 --- a/test/registered/jit/benchmark/bench_activation.py +++ b/test/registered/jit/benchmark/bench_activation.py @@ -12,7 +12,9 @@ from sglang.jit_kernel.benchmark import marker from sglang.jit_kernel.benchmark.utils import create_random from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) @torch.compile diff --git a/test/registered/jit/benchmark/bench_add_constant.py b/test/registered/jit/benchmark/bench_add_constant.py index fbadf112c..97d0b75df 100644 --- a/test/registered/jit/benchmark/bench_add_constant.py +++ b/test/registered/jit/benchmark/bench_add_constant.py @@ -10,7 +10,9 @@ from sglang.jit_kernel.benchmark.utils import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) CONSTANT = 7 SIZE_LIST = get_benchmark_range( diff --git a/test/registered/jit/benchmark/bench_awq_dequantize.py b/test/registered/jit/benchmark/bench_awq_dequantize.py index 85640b14c..91b9d2dff 100644 --- a/test/registered/jit/benchmark/bench_awq_dequantize.py +++ b/test/registered/jit/benchmark/bench_awq_dequantize.py @@ -9,7 +9,9 @@ from sglang.jit_kernel.benchmark.utils import run_benchmark from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) try: from sgl_kernel import awq_dequantize as aot_awq_dequantize diff --git a/test/registered/jit/benchmark/bench_clamp_position.py b/test/registered/jit/benchmark/bench_clamp_position.py index dfbb5b1e4..c2d74ef5f 100644 --- a/test/registered/jit/benchmark/bench_clamp_position.py +++ b/test/registered/jit/benchmark/bench_clamp_position.py @@ -13,8 +13,10 @@ from sglang.jit_kernel.clamp_position import clamp_position_cuda from sglang.srt.utils import get_compiler_backend from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=13, suite="base-b-kernel-benchmark-1-gpu-large") -register_amd_ci(est_time=16, suite="jit-kernel-unit-test-amd") +register_cuda_ci( + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) +register_amd_ci(est_time=16, stage="jit-kernel-unit", runner_config="amd") SIZE_LIST = get_benchmark_range( full_range=[2**n for n in range(4, 16)], diff --git a/test/registered/jit/benchmark/bench_concat_mla.py b/test/registered/jit/benchmark/bench_concat_mla.py index 0a0d98aa3..34f04c6b4 100644 --- a/test/registered/jit/benchmark/bench_concat_mla.py +++ b/test/registered/jit/benchmark/bench_concat_mla.py @@ -12,7 +12,9 @@ from sglang.jit_kernel.concat_mla import concat_mla_k as jit_k from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) IS_CI = is_in_ci() diff --git a/test/registered/jit/benchmark/bench_custom_all_reduce.py b/test/registered/jit/benchmark/bench_custom_all_reduce.py index f93d04850..a52d443a9 100644 --- a/test/registered/jit/benchmark/bench_custom_all_reduce.py +++ b/test/registered/jit/benchmark/bench_custom_all_reduce.py @@ -33,7 +33,8 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( est_time=120, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="requires multi-GPU, self-skips in CI", ) diff --git a/test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py b/test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py index a593bbfc8..bb7569e72 100644 --- a/test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py +++ b/test/registered/jit/benchmark/bench_dsv4_fp4_indexer.py @@ -10,7 +10,9 @@ from sglang.jit_kernel.benchmark.utils import get_benchmark_range from sglang.srt.utils import is_sm100_supported from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) try: import deep_gemm diff --git a/test/registered/jit/benchmark/bench_fused_qknorm_rope.py b/test/registered/jit/benchmark/bench_fused_qknorm_rope.py index 543be0489..7490a901c 100644 --- a/test/registered/jit/benchmark/bench_fused_qknorm_rope.py +++ b/test/registered/jit/benchmark/bench_fused_qknorm_rope.py @@ -20,7 +20,9 @@ from sglang.jit_kernel.fused_qknorm_rope import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) try: from sgl_kernel import fused_qk_norm_rope as fused_qk_norm_rope_aot diff --git a/test/registered/jit/benchmark/bench_hadamard.py b/test/registered/jit/benchmark/bench_hadamard.py index 704107c2d..121e4f0f3 100644 --- a/test/registered/jit/benchmark/bench_hadamard.py +++ b/test/registered/jit/benchmark/bench_hadamard.py @@ -16,7 +16,9 @@ from sglang.jit_kernel.benchmark.utils import ( from sglang.jit_kernel.hadamard import hadamard_transform from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) # AOT kernel: might not be available in all environments. # This is used for performance baseline comparison. diff --git a/test/registered/jit/benchmark/bench_hicache.py b/test/registered/jit/benchmark/bench_hicache.py index ba897fc12..6811d8b96 100644 --- a/test/registered/jit/benchmark/bench_hicache.py +++ b/test/registered/jit/benchmark/bench_hicache.py @@ -33,7 +33,9 @@ from sglang.jit_kernel.hicache import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=29, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=29, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) DISABLE_TORCH = os.environ.get("DISABLE_TORCH", "0") == "1" PAGE_SIZE = 1 diff --git a/test/registered/jit/benchmark/bench_hisparse.py b/test/registered/jit/benchmark/bench_hisparse.py index 2fe0ef901..18fecab23 100644 --- a/test/registered/jit/benchmark/bench_hisparse.py +++ b/test/registered/jit/benchmark/bench_hisparse.py @@ -9,7 +9,9 @@ from sglang.jit_kernel.benchmark.utils import DEFAULT_DEVICE, DEFAULT_DTYPE from sglang.jit_kernel.hisparse import load_cache_to_device_buffer_mla from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=12, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) DEVICE = DEFAULT_DEVICE DTYPE = DEFAULT_DTYPE diff --git a/test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py b/test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py index 6d0f60b1c..606a1ec56 100644 --- a/test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py +++ b/test/registered/jit/benchmark/bench_mla_kv_pack_quantize_fp8.py @@ -20,7 +20,9 @@ from sglang.jit_kernel.mla_kv_pack_quantize_fp8 import ( from sglang.jit_kernel.utils import is_arch_support_pdl from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) @triton.jit diff --git a/test/registered/jit/benchmark/bench_mxfp8_moe.py b/test/registered/jit/benchmark/bench_mxfp8_moe.py index a4215cc1b..c7c697208 100644 --- a/test/registered/jit/benchmark/bench_mxfp8_moe.py +++ b/test/registered/jit/benchmark/bench_mxfp8_moe.py @@ -13,7 +13,9 @@ from sglang.jit_kernel.mxfp8 import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) def is_sm100_supported(device=None) -> bool: diff --git a/test/registered/jit/benchmark/bench_ngram_compute_decode.py b/test/registered/jit/benchmark/bench_ngram_compute_decode.py index 63ecfaf6a..5e31b200f 100644 --- a/test/registered/jit/benchmark/bench_ngram_compute_decode.py +++ b/test/registered/jit/benchmark/bench_ngram_compute_decode.py @@ -13,7 +13,9 @@ from sglang.jit_kernel.ngram_embedding import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) NE_N = 8 NE_K = 2 diff --git a/test/registered/jit/benchmark/bench_ngram_update_token_table.py b/test/registered/jit/benchmark/bench_ngram_update_token_table.py index 118459df3..e0ab181a9 100644 --- a/test/registered/jit/benchmark/bench_ngram_update_token_table.py +++ b/test/registered/jit/benchmark/bench_ngram_update_token_table.py @@ -13,7 +13,9 @@ from sglang.jit_kernel.ngram_embedding import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=15, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) MAX_CONTEXT_LEN = 4096 BATCH_SIZE_LIST = get_benchmark_range( diff --git a/test/registered/jit/benchmark/bench_norm.py b/test/registered/jit/benchmark/bench_norm.py index 6adc7ac64..790e0c62b 100644 --- a/test/registered/jit/benchmark/bench_norm.py +++ b/test/registered/jit/benchmark/bench_norm.py @@ -11,7 +11,9 @@ from sglang.jit_kernel.norm import fused_add_rmsnorm as jit_fused_add_rmsnorm from sglang.jit_kernel.norm import rmsnorm as jit_rmsnorm from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) DTYPE = torch.bfloat16 diff --git a/test/registered/jit/benchmark/bench_nvfp4_blockwise_moe.py b/test/registered/jit/benchmark/bench_nvfp4_blockwise_moe.py index 0a4a1431e..d0381f35f 100644 --- a/test/registered/jit/benchmark/bench_nvfp4_blockwise_moe.py +++ b/test/registered/jit/benchmark/bench_nvfp4_blockwise_moe.py @@ -15,7 +15,9 @@ from sglang.jit_kernel.nvfp4 import ( from sglang.srt.utils import is_sm100_supported from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) FLOAT4_E2M1_MAX = 6.0 FLOAT8_E4M3_MAX = torch.finfo(torch.float8_e4m3fn).max diff --git a/test/registered/jit/benchmark/bench_nvfp4_quant.py b/test/registered/jit/benchmark/bench_nvfp4_quant.py index 5d516a9ee..cd25e3d90 100644 --- a/test/registered/jit/benchmark/bench_nvfp4_quant.py +++ b/test/registered/jit/benchmark/bench_nvfp4_quant.py @@ -10,7 +10,9 @@ from sglang.jit_kernel.nvfp4 import scaled_fp4_quant from sglang.srt.utils import is_sm100_supported from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) FLOAT4_E2M1_MAX = 6.0 FLOAT8_E4M3_MAX = torch.finfo(torch.float8_e4m3fn).max diff --git a/test/registered/jit/benchmark/bench_nvfp4_scaled_mm.py b/test/registered/jit/benchmark/bench_nvfp4_scaled_mm.py index 251840fe6..feaacc5c0 100644 --- a/test/registered/jit/benchmark/bench_nvfp4_scaled_mm.py +++ b/test/registered/jit/benchmark/bench_nvfp4_scaled_mm.py @@ -10,7 +10,9 @@ from sglang.jit_kernel.nvfp4 import cutlass_scaled_fp4_mm, scaled_fp4_quant from sglang.srt.utils import is_sm100_supported, is_sm120_supported from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) FLOAT4_E2M1_MAX = 6.0 FLOAT8_E4M3_MAX = torch.finfo(torch.float8_e4m3fn).max diff --git a/test/registered/jit/benchmark/bench_online_c128_mtp.py b/test/registered/jit/benchmark/bench_online_c128_mtp.py index 9f7540125..cab1c159b 100644 --- a/test/registered/jit/benchmark/bench_online_c128_mtp.py +++ b/test/registered/jit/benchmark/bench_online_c128_mtp.py @@ -18,7 +18,9 @@ from sglang.jit_kernel.benchmark.utils import ( from sglang.jit_kernel.dsv4.online_c128_mtp import _jit_online_c128_mtp_module from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) HEAD_DIM = 512 STATE_DIM = HEAD_DIM * 3 diff --git a/test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py b/test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py index e5c9f486b..e9e049f3e 100644 --- a/test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py +++ b/test/registered/jit/benchmark/bench_per_tensor_quant_fp8.py @@ -8,7 +8,9 @@ from sglang.jit_kernel.benchmark.utils import get_benchmark_range, run_benchmark from sglang.jit_kernel.per_tensor_quant_fp8 import per_tensor_quant_fp8 from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) try: from vllm import _custom_ops as ops diff --git a/test/registered/jit/benchmark/bench_per_token_group_quant_8bit.py b/test/registered/jit/benchmark/bench_per_token_group_quant_8bit.py index 549ab0fa5..62557c362 100644 --- a/test/registered/jit/benchmark/bench_per_token_group_quant_8bit.py +++ b/test/registered/jit/benchmark/bench_per_token_group_quant_8bit.py @@ -20,7 +20,9 @@ from sglang.srt.utils.bench_utils import bench_kineto from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=13, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) IS_CI = is_in_ci() diff --git a/test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py b/test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py index 983c5b697..fbae57b77 100644 --- a/test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py +++ b/test/registered/jit/benchmark/bench_per_token_group_quant_8bit_v2.py @@ -14,7 +14,9 @@ from sglang.srt.layers.quantization.fp8_kernel import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) G = 128 HIDDEN = 4096 diff --git a/test/registered/jit/benchmark/bench_qknorm.py b/test/registered/jit/benchmark/bench_qknorm.py index bbc6bdf52..381ae2d7a 100644 --- a/test/registered/jit/benchmark/bench_qknorm.py +++ b/test/registered/jit/benchmark/bench_qknorm.py @@ -6,7 +6,9 @@ from sglang.jit_kernel.norm import fused_inplace_qknorm from sglang.srt.utils import get_current_device_stream_fast from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) alt_stream = torch.cuda.Stream() diff --git a/test/registered/jit/benchmark/bench_qknorm_across_heads.py b/test/registered/jit/benchmark/bench_qknorm_across_heads.py index f969db429..8429111ef 100644 --- a/test/registered/jit/benchmark/bench_qknorm_across_heads.py +++ b/test/registered/jit/benchmark/bench_qknorm_across_heads.py @@ -12,7 +12,9 @@ from sglang.srt.utils import get_current_device_stream_fast from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=12, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=12, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) IS_CI = is_in_ci() diff --git a/test/registered/jit/benchmark/bench_renorm.py b/test/registered/jit/benchmark/bench_renorm.py index 500c9b1b1..a84427085 100644 --- a/test/registered/jit/benchmark/bench_renorm.py +++ b/test/registered/jit/benchmark/bench_renorm.py @@ -9,7 +9,9 @@ from sglang.jit_kernel.benchmark.utils import run_benchmark_no_cudagraph from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=5, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) def torch_top_k_renorm_probs(probs, top_k): diff --git a/test/registered/jit/benchmark/bench_resolve_future_token_ids.py b/test/registered/jit/benchmark/bench_resolve_future_token_ids.py index b0334c6d3..c3c48a54f 100644 --- a/test/registered/jit/benchmark/bench_resolve_future_token_ids.py +++ b/test/registered/jit/benchmark/bench_resolve_future_token_ids.py @@ -13,8 +13,10 @@ from sglang.jit_kernel.resolve_future_token_ids import resolve_future_token_ids_ from sglang.srt.utils import get_compiler_backend from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-benchmark-1-gpu-large") -register_amd_ci(est_time=10, suite="jit-kernel-unit-test-amd") +register_cuda_ci( + est_time=10, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) +register_amd_ci(est_time=10, stage="jit-kernel-unit", runner_config="amd") SIZE_LIST = get_benchmark_range( full_range=[2**n for n in range(4, 16)], # 16 … 32K elements diff --git a/test/registered/jit/benchmark/bench_rope.py b/test/registered/jit/benchmark/bench_rope.py index 4d578c239..8e6ce7a0a 100644 --- a/test/registered/jit/benchmark/bench_rope.py +++ b/test/registered/jit/benchmark/bench_rope.py @@ -12,7 +12,9 @@ from sglang.jit_kernel.benchmark.utils import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) MAX_SEQ_LEN = 131072 ROPE_BASE = 10000.0 diff --git a/test/registered/jit/benchmark/bench_set_mla_kv_buffer.py b/test/registered/jit/benchmark/bench_set_mla_kv_buffer.py index 71b4d20f7..28f05e964 100644 --- a/test/registered/jit/benchmark/bench_set_mla_kv_buffer.py +++ b/test/registered/jit/benchmark/bench_set_mla_kv_buffer.py @@ -26,7 +26,9 @@ from sglang.srt.mem_cache.utils import set_mla_kv_buffer_kernel as sglang_triton from sglang.srt.mem_cache.utils import set_mla_kv_buffer_triton as sglang_wrapper from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=9, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=9, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) def _triton_baseline(kv_buffer, loc, cache_k_nope, cache_k_rope): diff --git a/test/registered/jit/benchmark/bench_store_cache.py b/test/registered/jit/benchmark/bench_store_cache.py index 5b3b4dfd2..3e2eb80dd 100644 --- a/test/registered/jit/benchmark/bench_store_cache.py +++ b/test/registered/jit/benchmark/bench_store_cache.py @@ -9,7 +9,9 @@ from sglang.jit_kernel.benchmark.utils import ( from sglang.jit_kernel.kvcache import store_cache from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=9, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=9, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) @torch.compile() diff --git a/test/registered/jit/benchmark/bench_tp_qknorm.py b/test/registered/jit/benchmark/bench_tp_qknorm.py index cc0c56dbd..f4f71bcf4 100644 --- a/test/registered/jit/benchmark/bench_tp_qknorm.py +++ b/test/registered/jit/benchmark/bench_tp_qknorm.py @@ -41,7 +41,8 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( est_time=120, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="requires multi-GPU, self-skips in CI", ) diff --git a/test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py b/test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py index 7ba1ffe87..cc7a865d3 100644 --- a/test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py +++ b/test/registered/jit/benchmark/diffusion/bench_diffusion_nvfp4_scaled_mm.py @@ -18,7 +18,8 @@ from sglang.utils import is_in_ci register_cuda_ci( est_time=120, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="standalone diffusion NVFP4 benchmark", ) diff --git a/test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py b/test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py index 07bf7926d..f96da7f72 100644 --- a/test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py +++ b/test/registered/jit/benchmark/diffusion/bench_fused_norm_scale_shift.py @@ -20,7 +20,8 @@ from sglang.utils import is_in_ci register_cuda_ci( est_time=17, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="Temporarily skipped to unblock flashinfer upgrade. Ref: https://github.com/sgl-project/sglang/actions/runs/23735552939/job/69139238979?pr=21422", ) diff --git a/test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py b/test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py index 86971b525..6875a111c 100644 --- a/test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py +++ b/test/registered/jit/benchmark/diffusion/bench_group_norm_silu.py @@ -16,7 +16,8 @@ from sglang.utils import is_in_ci register_cuda_ci( est_time=45, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="standalone benchmark", ) diff --git a/test/registered/jit/benchmark/diffusion/bench_norm_impls.py b/test/registered/jit/benchmark/diffusion/bench_norm_impls.py index 7f0cb058d..59db976af 100644 --- a/test/registered/jit/benchmark/diffusion/bench_norm_impls.py +++ b/test/registered/jit/benchmark/diffusion/bench_norm_impls.py @@ -24,7 +24,8 @@ from sglang.utils import is_in_ci register_cuda_ci( est_time=120, - suite="base-b-kernel-benchmark-1-gpu-large", + stage="base-b-kernel-benchmark", + runner_config="1-gpu-large", disabled="self-skips in CI, standalone tool", ) diff --git a/test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py b/test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py index 46cea6e8e..17b3ecef3 100644 --- a/test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py +++ b/test/registered/jit/benchmark/diffusion/bench_qknorm_rope.py @@ -13,7 +13,9 @@ from sglang.jit_kernel.benchmark.utils import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=13, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) MAX_SEQ_LEN = 131072 ROPE_BASE = 10000.0 diff --git a/test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py b/test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py index f87f3e3b3..e614624a6 100644 --- a/test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py +++ b/test/registered/jit/benchmark/diffusion/bench_qwen_image_modulation.py @@ -12,7 +12,9 @@ from sglang.jit_kernel.diffusion.triton.scale_shift import ( from sglang.test.ci.ci_register import register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=13, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=13, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) if is_in_ci(): B_RANGE, S_RANGE, D_RANGE = [1], [128], [3072] diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py b/test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py index 7c3722fc8..069b24d48 100644 --- a/test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py +++ b/test/registered/jit/benchmark/minimax/bench_minimax_decode_topk.py @@ -17,7 +17,9 @@ from sglang.srt.layers.attention.minimax_sparse_ops.decode.flash_with_topk_idx i ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=8, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=8, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) BLOCK_SIZE = 128 TOPK = 16 diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py b/test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py index 89b1ad666..586d4106e 100644 --- a/test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py +++ b/test/registered/jit/benchmark/minimax/bench_minimax_qknorm_rope.py @@ -12,7 +12,9 @@ from sglang.jit_kernel.minimax_qknorm_rope import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) HEAD_DIM, ROTARY_DIM, BASE, EPS, MAXPOS = 128, 64, 5_000_000, 1e-6, 131072 NQ, NK = 64, 4 diff --git a/test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py b/test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py index 367a5b9c2..d5f5a6032 100644 --- a/test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py +++ b/test/registered/jit/benchmark/minimax/bench_minimax_store_kv_index.py @@ -7,7 +7,9 @@ from sglang.jit_kernel.benchmark import marker from sglang.jit_kernel.minimax_store_kv_index import store_kv_index from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-benchmark-1-gpu-large") +register_cuda_ci( + est_time=6, stage="base-b-kernel-benchmark", runner_config="1-gpu-large" +) HEAD_DIM = 128 NUM_KV_HEADS = 1 diff --git a/test/registered/jit/deepseek_v4/test_c128_v2.py b/test/registered/jit/deepseek_v4/test_c128_v2.py index 105097029..f8818b501 100644 --- a/test/registered/jit/deepseek_v4/test_c128_v2.py +++ b/test/registered/jit/deepseek_v4/test_c128_v2.py @@ -18,7 +18,7 @@ from sglang.jit_kernel.tests.deepseek_v4.common import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=30, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/deepseek_v4/test_c4_v2.py b/test/registered/jit/deepseek_v4/test_c4_v2.py index 0cdda3387..31233da29 100644 --- a/test/registered/jit/deepseek_v4/test_c4_v2.py +++ b/test/registered/jit/deepseek_v4/test_c4_v2.py @@ -18,7 +18,7 @@ from sglang.jit_kernel.tests.deepseek_v4.common import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=30, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/deepseek_v4/test_fp4_indexer.py b/test/registered/jit/deepseek_v4/test_fp4_indexer.py index 36dbfd811..39de4ef0c 100644 --- a/test/registered/jit/deepseek_v4/test_fp4_indexer.py +++ b/test/registered/jit/deepseek_v4/test_fp4_indexer.py @@ -21,7 +21,7 @@ from sglang.srt.layers.deepseek_v4_rope import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=60, suite="nightly-kernel-1-gpu", nightly=True) HEAD_DIM = 128 diff --git a/test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py b/test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py index 8ca94e371..d3c474984 100644 --- a/test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py +++ b/test/registered/jit/diffusion/test_diffusion_modelopt_fp8_scaled_mm.py @@ -14,7 +14,7 @@ from sglang.srt.layers.quantization.fp8_utils import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=20, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=80, suite="nightly-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py b/test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py index 9c016158b..9dbb19ae7 100644 --- a/test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py +++ b/test/registered/jit/diffusion/test_diffusion_nvfp4_scaled_mm.py @@ -17,7 +17,7 @@ from sglang.srt.layers.quantization.modelopt_quant import pad_nvfp4_weight from sglang.test.ci.ci_register import register_cuda_ci # B200-only correctness coverage for diffusion NVFP4 scaled mm. -register_cuda_ci(est_time=15, suite="base-b-kernel-unit-1-gpu-b200") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="4-gpu-b200") DEVICE = "cuda" DTYPE = torch.bfloat16 diff --git a/test/registered/jit/diffusion/test_flydsl_fused_norm.py b/test/registered/jit/diffusion/test_flydsl_fused_norm.py index bfb67cb43..f41f6daf4 100644 --- a/test/registered/jit/diffusion/test_flydsl_fused_norm.py +++ b/test/registered/jit/diffusion/test_flydsl_fused_norm.py @@ -6,7 +6,7 @@ import torch.nn.functional as F from sglang.test.ci.ci_register import register_amd_ci -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" D = 5120 diff --git a/test/registered/jit/diffusion/test_fused_norm_scale_shift.py b/test/registered/jit/diffusion/test_fused_norm_scale_shift.py index 9a0f87c23..74ff53083 100644 --- a/test/registered/jit/diffusion/test_fused_norm_scale_shift.py +++ b/test/registered/jit/diffusion/test_fused_norm_scale_shift.py @@ -13,7 +13,7 @@ from sglang.jit_kernel.diffusion.cutedsl.scale_residual_norm_scale_shift import ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=28, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/diffusion/test_group_norm_silu.py b/test/registered/jit/diffusion/test_group_norm_silu.py index 261960406..793acae53 100644 --- a/test/registered/jit/diffusion/test_group_norm_silu.py +++ b/test/registered/jit/diffusion/test_group_norm_silu.py @@ -9,7 +9,7 @@ from sglang.jit_kernel.diffusion.group_norm_silu import apply_group_norm_silu from sglang.jit_kernel.diffusion.triton.group_norm_silu import triton_group_norm_silu from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=8, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=8, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=15, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/diffusion/test_qknorm_rope.py b/test/registered/jit/diffusion/test_qknorm_rope.py index 0f3370ffc..76b042095 100644 --- a/test/registered/jit/diffusion/test_qknorm_rope.py +++ b/test/registered/jit/diffusion/test_qknorm_rope.py @@ -8,7 +8,7 @@ import triton from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=44, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=44, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=176, suite="nightly-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/diffusion/test_qwen_image_modulation.py b/test/registered/jit/diffusion/test_qwen_image_modulation.py index 26c700ab2..d16d58916 100644 --- a/test/registered/jit/diffusion/test_qwen_image_modulation.py +++ b/test/registered/jit/diffusion/test_qwen_image_modulation.py @@ -12,7 +12,7 @@ from sglang.jit_kernel.diffusion.triton.scale_shift import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=30, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/diffusion/test_varlen_pack_pad.py b/test/registered/jit/diffusion/test_varlen_pack_pad.py index 645e7da63..fa95364bb 100644 --- a/test/registered/jit/diffusion/test_varlen_pack_pad.py +++ b/test/registered/jit/diffusion/test_varlen_pack_pad.py @@ -15,7 +15,7 @@ from sglang.jit_kernel.diffusion.triton.varlen_pack_pad import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=60, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=15, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py b/test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py index ef45abef5..bd276a76d 100644 --- a/test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py +++ b/test/registered/jit/diffusion/test_varlen_uspattn_equivalence.py @@ -29,7 +29,7 @@ from sglang.multimodal_gen.runtime.layers.attention.layer import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=60, suite="nightly-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/kv_canary/test_const_sync.py b/test/registered/jit/kv_canary/test_const_sync.py index ab83a7e52..a026d7671 100644 --- a/test/registered/jit/kv_canary/test_const_sync.py +++ b/test/registered/jit/kv_canary/test_const_sync.py @@ -7,8 +7,8 @@ import sglang.jit_kernel from sglang.jit_kernel.kv_canary import consts from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=5, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=5, stage="jit-kernel-unit", runner_config="amd") # Resolve the kernel source against the installed jit_kernel package rather diff --git a/test/registered/jit/kv_canary/test_kernel_config.py b/test/registered/jit/kv_canary/test_kernel_config.py index ab6d65882..33d911b3e 100644 --- a/test/registered/jit/kv_canary/test_kernel_config.py +++ b/test/registered/jit/kv_canary/test_kernel_config.py @@ -35,8 +35,8 @@ from sglang.jit_kernel.tests.kv_canary._fixtures import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=60, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=60, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_pipeline_e2e.py b/test/registered/jit/kv_canary/test_pipeline_e2e.py index e0bc7e094..7446ca7f3 100644 --- a/test/registered/jit/kv_canary/test_pipeline_e2e.py +++ b/test/registered/jit/kv_canary/test_pipeline_e2e.py @@ -40,8 +40,8 @@ from sglang.jit_kernel.tests.kv_canary._fixtures import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_plan_fuzz.py b/test/registered/jit/kv_canary/test_plan_fuzz.py index 8a94c9ec8..67385ee66 100644 --- a/test/registered/jit/kv_canary/test_plan_fuzz.py +++ b/test/registered/jit/kv_canary/test_plan_fuzz.py @@ -22,8 +22,8 @@ from sglang.jit_kernel.tests.kv_canary._fuzz_driver import ( from sglang.jit_kernel.tests.kv_canary._invariants import PlanInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_plan_hand.py b/test/registered/jit/kv_canary/test_plan_hand.py index 04b5612d5..fa6385b89 100644 --- a/test/registered/jit/kv_canary/test_plan_hand.py +++ b/test/registered/jit/kv_canary/test_plan_hand.py @@ -22,8 +22,8 @@ from sglang.jit_kernel.tests.kv_canary._fixtures import ( from sglang.jit_kernel.tests.kv_canary._invariants import PlanInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_scatter_req_token_ids.py b/test/registered/jit/kv_canary/test_scatter_req_token_ids.py index f065d3d43..d1cf3fe9a 100644 --- a/test/registered/jit/kv_canary/test_scatter_req_token_ids.py +++ b/test/registered/jit/kv_canary/test_scatter_req_token_ids.py @@ -13,8 +13,8 @@ from sglang.jit_kernel.kv_canary.scatter_req_token_ids import ( 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, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=10, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=10, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_utils.py b/test/registered/jit/kv_canary/test_utils.py index 83023f664..9f74f64d0 100644 --- a/test/registered/jit/kv_canary/test_utils.py +++ b/test/registered/jit/kv_canary/test_utils.py @@ -8,7 +8,7 @@ from sglang.jit_kernel.benchmark.kv_canary.utils import ( ) from sglang.test.ci.ci_register import register_cpu_ci -register_cpu_ci(est_time=3, suite="base-a-test-cpu") +register_cpu_ci(est_time=3, stage="base-a", runner_config="cpu") def test_fast_matrix_cases_include_e2e_decode_and_chunked_prefill_scenarios() -> None: diff --git a/test/registered/jit/kv_canary/test_verify_fuzz.py b/test/registered/jit/kv_canary/test_verify_fuzz.py index 7594860b8..c15697a0f 100644 --- a/test/registered/jit/kv_canary/test_verify_fuzz.py +++ b/test/registered/jit/kv_canary/test_verify_fuzz.py @@ -31,8 +31,8 @@ from sglang.jit_kernel.tests.kv_canary._fuzz_driver import ( from sglang.jit_kernel.tests.kv_canary._invariants import VerifyInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_verify_hand.py b/test/registered/jit/kv_canary/test_verify_hand.py index f3d9f50c0..7a77cc4a4 100644 --- a/test/registered/jit/kv_canary/test_verify_hand.py +++ b/test/registered/jit/kv_canary/test_verify_hand.py @@ -53,8 +53,8 @@ from sglang.jit_kernel.tests.kv_canary._hand_oracle import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_write_fuzz.py b/test/registered/jit/kv_canary/test_write_fuzz.py index 8fc04f0d5..394fa833d 100644 --- a/test/registered/jit/kv_canary/test_write_fuzz.py +++ b/test/registered/jit/kv_canary/test_write_fuzz.py @@ -31,8 +31,8 @@ from sglang.jit_kernel.tests.kv_canary._fuzz_driver import ( from sglang.jit_kernel.tests.kv_canary._invariants import WriteInvariants from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/kv_canary/test_write_hand.py b/test/registered/jit/kv_canary/test_write_hand.py index e12dd4953..f77dcb3ee 100644 --- a/test/registered/jit/kv_canary/test_write_hand.py +++ b/test/registered/jit/kv_canary/test_write_hand.py @@ -51,8 +51,8 @@ from sglang.jit_kernel.tests.kv_canary._hand_oracle import ( ) from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") _DEVICE = torch.device("cuda") diff --git a/test/registered/jit/minimax/test_minimax_decode_topk.py b/test/registered/jit/minimax/test_minimax_decode_topk.py index c4860b7bf..c6d553c66 100644 --- a/test/registered/jit/minimax/test_minimax_decode_topk.py +++ b/test/registered/jit/minimax/test_minimax_decode_topk.py @@ -12,8 +12,8 @@ import torch from sglang.jit_kernel.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, suite="base-b-kernel-unit-1-gpu-large") -register_cuda_ci(est_time=40, suite="base-b-kernel-unit-1-gpu-b200") +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="4-gpu-b200") register_amd_ci(est_time=15, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/minimax/test_minimax_decode_topk_page_table.py b/test/registered/jit/minimax/test_minimax_decode_topk_page_table.py index c4d55cb13..55a825e1a 100644 --- a/test/registered/jit/minimax/test_minimax_decode_topk_page_table.py +++ b/test/registered/jit/minimax/test_minimax_decode_topk_page_table.py @@ -25,7 +25,7 @@ from sglang.srt.layers.attention.minimax_sparse_ops.decode.topk_sparse import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=25, suite="base-b-kernel-unit-1-gpu-b200") +register_cuda_ci(est_time=25, stage="base-b-kernel-unit", runner_config="4-gpu-b200") dev = "cuda" diff --git a/test/registered/jit/minimax/test_minimax_qknorm_rope.py b/test/registered/jit/minimax/test_minimax_qknorm_rope.py index 91d653a6f..274f571bc 100644 --- a/test/registered/jit/minimax/test_minimax_qknorm_rope.py +++ b/test/registered/jit/minimax/test_minimax_qknorm_rope.py @@ -14,8 +14,8 @@ from sglang.jit_kernel.minimax_qknorm_rope import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-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="4-gpu-b200") dev = "cuda" HEAD_DIM, ROTARY_DIM, BASE, EPS = 128, 64, 5_000_000, 1e-6 diff --git a/test/registered/jit/minimax/test_minimax_store_kv_index.py b/test/registered/jit/minimax/test_minimax_store_kv_index.py index d23880bcf..60902df35 100644 --- a/test/registered/jit/minimax/test_minimax_store_kv_index.py +++ b/test/registered/jit/minimax/test_minimax_store_kv_index.py @@ -11,8 +11,8 @@ import torch from sglang.jit_kernel.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, suite="base-b-kernel-unit-1-gpu-large") -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-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="4-gpu-b200") register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) dev = "cuda" diff --git a/test/registered/jit/test_activation.py b/test/registered/jit/test_activation.py index b01660eec..c4683f11b 100644 --- a/test/registered/jit/test_activation.py +++ b/test/registered/jit/test_activation.py @@ -12,9 +12,9 @@ from sglang.jit_kernel.activation import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=20, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=20, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=30, suite="nightly-kernel-1-gpu", nightly=True) -register_amd_ci(est_time=20, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") OPS = SUPPORTED_ACTIVATIONS diff --git a/test/registered/jit/test_add_constant.py b/test/registered/jit/test_add_constant.py index 6a4ecf178..48f2476bf 100644 --- a/test/registered/jit/test_add_constant.py +++ b/test/registered/jit/test_add_constant.py @@ -6,7 +6,7 @@ import torch from sglang.jit_kernel.add_constant import add_constant from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=180, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_amd_deterministic_custom_allreduce.py b/test/registered/jit/test_amd_deterministic_custom_allreduce.py index c5a0525a6..64522f430 100644 --- a/test/registered/jit/test_amd_deterministic_custom_allreduce.py +++ b/test/registered/jit/test_amd_deterministic_custom_allreduce.py @@ -25,7 +25,7 @@ import torch.distributed as dist from sglang.srt.environ import envs from sglang.test.ci.ci_register import register_amd_ci -register_amd_ci(est_time=120, suite="sgl-kernel-unit-test-2-gpu-amd") +register_amd_ci(est_time=120, stage="sgl-kernel-unit", runner_config="2-gpu-amd") def get_open_port(): diff --git a/test/registered/jit/test_amd_nccl_allreduce_determinism.py b/test/registered/jit/test_amd_nccl_allreduce_determinism.py index 480786a8b..a7985eb5f 100644 --- a/test/registered/jit/test_amd_nccl_allreduce_determinism.py +++ b/test/registered/jit/test_amd_nccl_allreduce_determinism.py @@ -22,7 +22,7 @@ import torch.distributed as dist from sglang.test.ci.ci_register import register_amd_ci -register_amd_ci(est_time=120, suite="sgl-kernel-unit-test-2-gpu-amd") +register_amd_ci(est_time=120, stage="sgl-kernel-unit", runner_config="2-gpu-amd") def get_open_port(): diff --git a/test/registered/jit/test_awq_dequantize.py b/test/registered/jit/test_awq_dequantize.py index 19d22e669..b6f5cb37e 100644 --- a/test/registered/jit/test_awq_dequantize.py +++ b/test/registered/jit/test_awq_dequantize.py @@ -7,7 +7,7 @@ import torch from sglang.jit_kernel.awq_dequantize import awq_dequantize as jit_awq_dequantize from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=9, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) try: diff --git a/test/registered/jit/test_awq_marlin_moe_repack.py b/test/registered/jit/test_awq_marlin_moe_repack.py index 7148b057c..d05708d20 100644 --- a/test/registered/jit/test_awq_marlin_moe_repack.py +++ b/test/registered/jit/test_awq_marlin_moe_repack.py @@ -11,7 +11,7 @@ from sglang.jit_kernel.awq_marlin_repack import ( from sglang.srt.layers.quantization.utils import pack_cols, quantize_weights from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_awq_marlin_repack.py b/test/registered/jit/test_awq_marlin_repack.py index 4b0746b85..4034a5ac9 100644 --- a/test/registered/jit/test_awq_marlin_repack.py +++ b/test/registered/jit/test_awq_marlin_repack.py @@ -12,7 +12,7 @@ from sglang.srt.layers.quantization.utils import pack_cols, quantize_weights 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=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_clamp_position.py b/test/registered/jit/test_clamp_position.py index 846d29017..1a7a68b6c 100644 --- a/test/registered/jit/test_clamp_position.py +++ b/test/registered/jit/test_clamp_position.py @@ -6,9 +6,9 @@ import torch from sglang.jit_kernel.clamp_position import clamp_position_cuda from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=12, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=12, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) -register_amd_ci(est_time=12, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=12, stage="jit-kernel-unit", runner_config="amd") def _reference_clamp_position(seq_lens): diff --git a/test/registered/jit/test_concat_mla.py b/test/registered/jit/test_concat_mla.py index ceb70d012..94905bf93 100644 --- a/test/registered/jit/test_concat_mla.py +++ b/test/registered/jit/test_concat_mla.py @@ -7,7 +7,7 @@ import triton from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=17, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=17, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_custom_all_reduce.py b/test/registered/jit/test_custom_all_reduce.py index b70b98bd3..0f3cc84f4 100644 --- a/test/registered/jit/test_custom_all_reduce.py +++ b/test/registered/jit/test_custom_all_reduce.py @@ -46,7 +46,8 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( est_time=300, - suite="base-b-kernel-unit-8-gpu-h200", + stage="base-b-kernel-unit", + runner_config="8-gpu-h200", ) register_cuda_ci( est_time=300, diff --git a/test/registered/jit/test_cutedsl_gdn.py b/test/registered/jit/test_cutedsl_gdn.py index fb6700d5d..d74b89ac8 100644 --- a/test/registered/jit/test_cutedsl_gdn.py +++ b/test/registered/jit/test_cutedsl_gdn.py @@ -29,7 +29,7 @@ try: except ImportError: TRITON_AVAILABLE = False -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py b/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py index a66215863..03dcc776c 100644 --- a/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py +++ b/test/registered/jit/test_deepseek_v4_compress_state_runtime_shapes.py @@ -101,7 +101,7 @@ from sglang.jit_kernel.tests.deepseek_v4.common import ( from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci from sglang.utils import is_in_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=25, suite="nightly-amd-kernel-1-gpu", nightly=True) Mode = Literal["decode", "prefill"] diff --git a/test/registered/jit/test_flash_attention_3_only_qv.py b/test/registered/jit/test_flash_attention_3_only_qv.py index 34221e563..c9b77f91e 100644 --- a/test/registered/jit/test_flash_attention_3_only_qv.py +++ b/test/registered/jit/test_flash_attention_3_only_qv.py @@ -15,7 +15,7 @@ from sglang.jit_kernel.flash_attention import flash_attn_with_kvcache from sglang.test.ci.ci_register import register_cuda_ci from sglang.test.test_utils import CustomTestCase -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") # FA3 only_qv path is SM90 (Hopper) only — skip on pre-Hopper and on # Blackwell+ (sm100+) where FA3 is not built. diff --git a/test/registered/jit/test_flash_attention_4.py b/test/registered/jit/test_flash_attention_4.py index a3d45b6bf..d28af546e 100644 --- a/test/registered/jit/test_flash_attention_4.py +++ b/test/registered/jit/test_flash_attention_4.py @@ -14,8 +14,8 @@ from einops import rearrange, repeat from sglang.jit_kernel.flash_attention import flash_attn_varlen_func from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=120, suite="base-b-kernel-unit-1-gpu-large") -register_cuda_ci(est_time=120, suite="base-b-kernel-unit-1-gpu-b200") +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=900, suite="nightly-kernel-1-gpu", nightly=True) # Skip this test on Hopper machine diff --git a/test/registered/jit/test_fused_add_rmsnorm.py b/test/registered/jit/test_fused_add_rmsnorm.py index 2d8ec7a07..70f13c6f2 100644 --- a/test/registered/jit/test_fused_add_rmsnorm.py +++ b/test/registered/jit/test_fused_add_rmsnorm.py @@ -7,7 +7,7 @@ import torch from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_fused_metadata_copy.py b/test/registered/jit/test_fused_metadata_copy.py index 08be9df14..f2229f197 100644 --- a/test/registered/jit/test_fused_metadata_copy.py +++ b/test/registered/jit/test_fused_metadata_copy.py @@ -16,7 +16,7 @@ import torch from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=100, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=100, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=400, suite="nightly-kernel-1-gpu", nightly=True) # ============================================================================= diff --git a/test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py b/test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py index 94302b148..4a6a149ca 100644 --- a/test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py +++ b/test/registered/jit/test_fused_qk_gemma_rmsnorm_gate.py @@ -7,7 +7,7 @@ import torch from sglang.srt.models.utils import fused_qk_gemma_rmsnorm_with_gate from sglang.test.ci.ci_register import register_amd_ci -register_amd_ci(est_time=20, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=20, stage="jit-kernel-unit", runner_config="amd") def reference_qk_gemma_rmsnorm_with_gate( diff --git a/test/registered/jit/test_fused_store_index_cache.py b/test/registered/jit/test_fused_store_index_cache.py index 30494852f..a7cc1b4f7 100644 --- a/test/registered/jit/test_fused_store_index_cache.py +++ b/test/registered/jit/test_fused_store_index_cache.py @@ -48,7 +48,7 @@ try: except ImportError: _is_fp8_fnuz = False -register_cuda_ci(est_time=24, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=24, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=24, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_fused_verify_triton_gdn.py b/test/registered/jit/test_fused_verify_triton_gdn.py index c6af34475..c417fa962 100644 --- a/test/registered/jit/test_fused_verify_triton_gdn.py +++ b/test/registered/jit/test_fused_verify_triton_gdn.py @@ -26,7 +26,7 @@ try: except ImportError: KERNELS_AVAILABLE = False -register_cuda_ci(est_time=6, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) register_amd_ci(est_time=10, suite="nightly-amd-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_gptq_marlin.py b/test/registered/jit/test_gptq_marlin.py index 32c482c15..8d28ef782 100644 --- a/test/registered/jit/test_gptq_marlin.py +++ b/test/registered/jit/test_gptq_marlin.py @@ -23,7 +23,7 @@ from sglang.test.test_marlin_utils import ( marlin_quantize, ) -register_cuda_ci(est_time=13, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=13, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) MNK_FACTORS = [ diff --git a/test/registered/jit/test_gptq_marlin_repack.py b/test/registered/jit/test_gptq_marlin_repack.py index a94062e17..0bb1075df 100644 --- a/test/registered/jit/test_gptq_marlin_repack.py +++ b/test/registered/jit/test_gptq_marlin_repack.py @@ -13,7 +13,7 @@ from sglang.srt.layers.quantization.utils import ( 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, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) MARLIN_K_CHUNKS = [128] diff --git a/test/registered/jit/test_grouped_topk.py b/test/registered/jit/test_grouped_topk.py index e47d04609..be15b65c9 100644 --- a/test/registered/jit/test_grouped_topk.py +++ b/test/registered/jit/test_grouped_topk.py @@ -9,7 +9,7 @@ from sglang.jit_kernel.utils import get_ci_test_range from sglang.srt.layers.moe.topk import biased_grouped_topk_impl from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_hadamard_jit.py b/test/registered/jit/test_hadamard_jit.py index 208a6b1ae..a424fa30c 100644 --- a/test/registered/jit/test_hadamard_jit.py +++ b/test/registered/jit/test_hadamard_jit.py @@ -16,7 +16,7 @@ from sglang.jit_kernel.hadamard import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=128, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=128, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=512, suite="nightly-kernel-1-gpu", nightly=True) # Exact M×N Hadamard matrices (±1 entries) copied from diff --git a/test/registered/jit/test_hicache.py b/test/registered/jit/test_hicache.py index e26b23330..ecdfcca54 100644 --- a/test/registered/jit/test_hicache.py +++ b/test/registered/jit/test_hicache.py @@ -16,7 +16,7 @@ from sglang.srt.mem_cache.pool_host.common import ( from sglang.srt.utils import is_cuda, is_hip, is_npu, is_xpu from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) pytestmark = pytest.mark.skipif( diff --git a/test/registered/jit/test_hisparse.py b/test/registered/jit/test_hisparse.py index 5fbbd56d4..5f0c5cd87 100644 --- a/test/registered/jit/test_hisparse.py +++ b/test/registered/jit/test_hisparse.py @@ -11,8 +11,8 @@ from sglang.jit_kernel.hisparse import ( 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_amd_ci(est_time=30, suite="stage-b-test-1-gpu-small-amd") -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +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=120, suite="nightly-kernel-1-gpu", nightly=True) pytestmark = pytest.mark.skipif( diff --git a/test/registered/jit/test_kpool_topk_transform.py b/test/registered/jit/test_kpool_topk_transform.py index 8ae8e2b45..3b113f0fb 100644 --- a/test/registered/jit/test_kpool_topk_transform.py +++ b/test/registered/jit/test_kpool_topk_transform.py @@ -17,7 +17,7 @@ import torch from sglang.jit_kernel.kpool_topk_transform import fast_kpool_topk_transform_fused from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") def _ref_torch_kpool_transform_impl( diff --git a/test/registered/jit/test_kvcacheio_asymmetric.py b/test/registered/jit/test_kvcacheio_asymmetric.py index 5667e1b43..205533c07 100644 --- a/test/registered/jit/test_kvcacheio_asymmetric.py +++ b/test/registered/jit/test_kvcacheio_asymmetric.py @@ -7,7 +7,7 @@ import torch from sglang.srt.mem_cache.memory_pool_host import AsymmetricMHATokenToKVPoolHost from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") 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/jit/test_minimax_m3_qk_norm_rope.py b/test/registered/jit/test_minimax_m3_qk_norm_rope.py index 2026cc313..66ccca6db 100644 --- a/test/registered/jit/test_minimax_m3_qk_norm_rope.py +++ b/test/registered/jit/test_minimax_m3_qk_norm_rope.py @@ -22,7 +22,7 @@ from sglang.jit_kernel.minimax_m3.qk_norm_rope import ( # noqa: E402 from sglang.test.ci.ci_register import register_amd_ci # noqa: E402 # ROCm-only fused kernel; runs in the AMD jit-kernel unit suite. -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" EPS = 1e-6 diff --git a/test/registered/jit/test_mla_kv_pack_quantize_fp8.py b/test/registered/jit/test_mla_kv_pack_quantize_fp8.py index b4775d3b2..3266a4715 100644 --- a/test/registered/jit/test_mla_kv_pack_quantize_fp8.py +++ b/test/registered/jit/test_mla_kv_pack_quantize_fp8.py @@ -7,7 +7,7 @@ from sglang.jit_kernel.mla_kv_pack_quantize_fp8 import mla_kv_pack_quantize_fp8 from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=60, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_amd_ci(est_time=60, suite="nightly-amd-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/test_moe_align_block_size.py b/test/registered/jit/test_moe_align_block_size.py index a84ba8e40..41ba65d02 100644 --- a/test/registered/jit/test_moe_align_block_size.py +++ b/test/registered/jit/test_moe_align_block_size.py @@ -9,7 +9,7 @@ import triton.language as tl from sglang.jit_kernel.moe_align import moe_align_block_size from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=28, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_moe_lora_align_block_size.py b/test/registered/jit/test_moe_lora_align_block_size.py index b2e721cc9..ca9915fe9 100644 --- a/test/registered/jit/test_moe_lora_align_block_size.py +++ b/test/registered/jit/test_moe_lora_align_block_size.py @@ -11,7 +11,7 @@ import torch from sglang.jit_kernel.moe_lora_align import moe_lora_align_block_size from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=28, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_moe_wna16_marlin.py b/test/registered/jit/test_moe_wna16_marlin.py index 320cf95aa..e5c883dfc 100644 --- a/test/registered/jit/test_moe_wna16_marlin.py +++ b/test/registered/jit/test_moe_wna16_marlin.py @@ -20,7 +20,7 @@ from sglang.test.test_marlin_utils import ( marlin_quantize, ) -register_cuda_ci(est_time=10, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=10, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_mxfp8_moe.py b/test/registered/jit/test_mxfp8_moe.py index 4ba53351a..c392d7478 100644 --- a/test/registered/jit/test_mxfp8_moe.py +++ b/test/registered/jit/test_mxfp8_moe.py @@ -10,7 +10,7 @@ from sglang.jit_kernel.mxfp8 import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_ngram_embedding.py b/test/registered/jit/test_ngram_embedding.py index 0e0f34ca7..ee8f4def4 100644 --- a/test/registered/jit/test_ngram_embedding.py +++ b/test/registered/jit/test_ngram_embedding.py @@ -11,7 +11,7 @@ from sglang.jit_kernel.ngram_embedding import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") def _make_ngram_params(ne_n: int, ne_k: int, vocab_size: int): diff --git a/test/registered/jit/test_nvfp4_blockwise_moe.py b/test/registered/jit/test_nvfp4_blockwise_moe.py index d188ef268..c022628e3 100644 --- a/test/registered/jit/test_nvfp4_blockwise_moe.py +++ b/test/registered/jit/test_nvfp4_blockwise_moe.py @@ -10,7 +10,7 @@ from sglang.jit_kernel.nvfp4 import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) FLOAT4_E2M1_MAX = 6.0 diff --git a/test/registered/jit/test_nvfp4_gemm.py b/test/registered/jit/test_nvfp4_gemm.py index 4e55fe40b..f0fad1b86 100644 --- a/test/registered/jit/test_nvfp4_gemm.py +++ b/test/registered/jit/test_nvfp4_gemm.py @@ -6,7 +6,7 @@ import torch from sglang.jit_kernel.nvfp4 import cutlass_scaled_fp4_mm, scaled_fp4_quant from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_nvfp4_quant.py b/test/registered/jit/test_nvfp4_quant.py index 293d93ead..1ca9c68aa 100644 --- a/test/registered/jit/test_nvfp4_quant.py +++ b/test/registered/jit/test_nvfp4_quant.py @@ -16,7 +16,7 @@ except Exception: from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=5, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=5, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_per_tensor_quant_fp8.py b/test/registered/jit/test_per_tensor_quant_fp8.py index 2e3acea3a..8bae572ec 100644 --- a/test/registered/jit/test_per_tensor_quant_fp8.py +++ b/test/registered/jit/test_per_tensor_quant_fp8.py @@ -8,7 +8,7 @@ import torch from sglang.jit_kernel.per_tensor_quant_fp8 import per_tensor_quant_fp8 from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=16, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) try: diff --git a/test/registered/jit/test_per_token_group_quant_8bit.py b/test/registered/jit/test_per_token_group_quant_8bit.py index f370cb19a..d3c64bc15 100644 --- a/test/registered/jit/test_per_token_group_quant_8bit.py +++ b/test/registered/jit/test_per_token_group_quant_8bit.py @@ -25,7 +25,7 @@ from sglang.srt.layers.quantization.fp8_kernel import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=16, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) configs = list( diff --git a/test/registered/jit/test_per_token_group_quant_8bit_v2.py b/test/registered/jit/test_per_token_group_quant_8bit_v2.py index ce97a7eee..12442cdce 100644 --- a/test/registered/jit/test_per_token_group_quant_8bit_v2.py +++ b/test/registered/jit/test_per_token_group_quant_8bit_v2.py @@ -14,7 +14,7 @@ from sglang.srt.layers.quantization.fp8_kernel import ( ) from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=90, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=90, stage="base-b-kernel-unit", runner_config="1-gpu-large") G = 128 diff --git a/test/registered/jit/test_pos_enc.py b/test/registered/jit/test_pos_enc.py index cdeb3f85f..e2a4f3c78 100644 --- a/test/registered/jit/test_pos_enc.py +++ b/test/registered/jit/test_pos_enc.py @@ -10,7 +10,7 @@ import triton.language as tl from sglang.jit_kernel.rope import rotary_embedding from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=18, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=18, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_qknorm.py b/test/registered/jit/test_qknorm.py index b0026997c..c707a7fe4 100644 --- a/test/registered/jit/test_qknorm.py +++ b/test/registered/jit/test_qknorm.py @@ -8,7 +8,7 @@ import triton from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=37, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=37, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=148, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_qknorm_across_heads.py b/test/registered/jit/test_qknorm_across_heads.py index a62887a73..70daedac8 100644 --- a/test/registered/jit/test_qknorm_across_heads.py +++ b/test/registered/jit/test_qknorm_across_heads.py @@ -8,7 +8,7 @@ import triton from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=15, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=15, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_renorm.py b/test/registered/jit/test_renorm.py index 119fcd80d..b73513a0f 100644 --- a/test/registered/jit/test_renorm.py +++ b/test/registered/jit/test_renorm.py @@ -9,7 +9,7 @@ import torch from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=6, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=6, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_resolve_future_token_ids.py b/test/registered/jit/test_resolve_future_token_ids.py index 760e9fbd9..bf46b12a5 100644 --- a/test/registered/jit/test_resolve_future_token_ids.py +++ b/test/registered/jit/test_resolve_future_token_ids.py @@ -6,9 +6,9 @@ import torch from sglang.jit_kernel.resolve_future_token_ids import resolve_future_token_ids_cuda from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=9, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=9, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) -register_amd_ci(est_time=9, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=9, stage="jit-kernel-unit", runner_config="amd") def _reference_resolve(input_ids, future_map): diff --git a/test/registered/jit/test_rmsnorm.py b/test/registered/jit/test_rmsnorm.py index 37539ee19..6baac1103 100644 --- a/test/registered/jit/test_rmsnorm.py +++ b/test/registered/jit/test_rmsnorm.py @@ -7,7 +7,7 @@ import torch from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=45, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=45, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=240, suite="nightly-kernel-1-gpu", nightly=True) diff --git a/test/registered/jit/test_rmsnorm_hf.py b/test/registered/jit/test_rmsnorm_hf.py index c26cdf0d8..f0a9eda19 100644 --- a/test/registered/jit/test_rmsnorm_hf.py +++ b/test/registered/jit/test_rmsnorm_hf.py @@ -13,9 +13,9 @@ from sglang.jit_kernel.rmsnorm_hf import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) -register_amd_ci(est_time=30, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=30, stage="jit-kernel-unit", runner_config="amd") EPS = 1e-5 DEVICE = "cuda" diff --git a/test/registered/jit/test_rope.py b/test/registered/jit/test_rope.py index 28f4e35f5..b70eaa1a9 100644 --- a/test/registered/jit/test_rope.py +++ b/test/registered/jit/test_rope.py @@ -7,7 +7,7 @@ import triton from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=64, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=64, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=256, suite="nightly-kernel-1-gpu", nightly=True) DEVICE = "cuda" diff --git a/test/registered/jit/test_set_mla_kv_buffer.py b/test/registered/jit/test_set_mla_kv_buffer.py index d7f967647..4062355ef 100644 --- a/test/registered/jit/test_set_mla_kv_buffer.py +++ b/test/registered/jit/test_set_mla_kv_buffer.py @@ -10,7 +10,7 @@ from sglang.jit_kernel.set_mla_kv_buffer import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=30, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=30, stage="base-b-kernel-unit", runner_config="1-gpu-large") DEVICE = "cuda" CACHE_SIZE = 4096 diff --git a/test/registered/jit/test_sigmoid_gate_mul.py b/test/registered/jit/test_sigmoid_gate_mul.py index 83ca811b8..4f59de093 100644 --- a/test/registered/jit/test_sigmoid_gate_mul.py +++ b/test/registered/jit/test_sigmoid_gate_mul.py @@ -5,8 +5,8 @@ import torch from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=4, suite="base-b-kernel-unit-1-gpu-large") -register_amd_ci(est_time=4, suite="jit-kernel-unit-test-amd") +register_cuda_ci(est_time=4, stage="base-b-kernel-unit", runner_config="1-gpu-large") +register_amd_ci(est_time=4, stage="jit-kernel-unit", runner_config="amd") DEVICE = "cuda" diff --git a/test/registered/jit/test_store_cache.py b/test/registered/jit/test_store_cache.py index 16da71362..b5496c8ce 100644 --- a/test/registered/jit/test_store_cache.py +++ b/test/registered/jit/test_store_cache.py @@ -8,9 +8,9 @@ from sglang.jit_kernel.kvcache import can_use_store_cache, store_cache from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci -register_cuda_ci(est_time=28, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=28, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) -register_amd_ci(est_time=55, suite="jit-kernel-unit-test-amd") +register_amd_ci(est_time=55, stage="jit-kernel-unit", runner_config="amd") BS_LIST = [2**n for n in range(0, 15)] BS_LIST += [x + 1 + i for i, x in enumerate(BS_LIST)] diff --git a/test/registered/jit/test_timestep_embedding.py b/test/registered/jit/test_timestep_embedding.py index 94db3f4f7..38b4e951a 100644 --- a/test/registered/jit/test_timestep_embedding.py +++ b/test/registered/jit/test_timestep_embedding.py @@ -16,7 +16,7 @@ from sglang.jit_kernel.timestep_embedding import ( from sglang.jit_kernel.utils import get_ci_test_range from sglang.test.ci.ci_register import register_cuda_ci -register_cuda_ci(est_time=16, suite="base-b-kernel-unit-1-gpu-large") +register_cuda_ci(est_time=16, stage="base-b-kernel-unit", runner_config="1-gpu-large") register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True) CORRECTNESS_BATCH_SIZES = get_ci_test_range( diff --git a/test/registered/jit/test_tp_qknorm.py b/test/registered/jit/test_tp_qknorm.py index 309db1f8e..03bb5dda8 100644 --- a/test/registered/jit/test_tp_qknorm.py +++ b/test/registered/jit/test_tp_qknorm.py @@ -29,7 +29,8 @@ from sglang.test.ci.ci_register import register_cuda_ci register_cuda_ci( est_time=300, - suite="base-b-kernel-unit-8-gpu-h200", + stage="base-b-kernel-unit", + runner_config="8-gpu-h200", ) register_cuda_ci( est_time=300, diff --git a/test/run_suite.py b/test/run_suite.py index e4c2350c4..f733ccd86 100644 --- a/test/run_suite.py +++ b/test/run_suite.py @@ -70,10 +70,10 @@ PER_COMMIT_SUITES = { "base-b-test-1-gpu-large", "base-b-test-2-gpu-large", "base-b-test-4-gpu-b200", - "base-b-kernel-unit-1-gpu-large", - "base-b-kernel-unit-1-gpu-b200", - "base-b-kernel-unit-8-gpu-h200", - "base-b-kernel-benchmark-1-gpu-large", + "base-b-kernel-unit-test-1-gpu-large", + "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-c-test-4-gpu-h100", "base-c-test-4-gpu-b200", "base-c-test-4-gpu-gb300",