{"owner":"sgl-project","repo":"sglang","hasSkills":true,"totalSkillsCount":87,"totalTokensCount":73919,"categories":["anthropic-skill","root-instruction","project-spec","plugin-manifest"],"hasMcp":false,"mcpConfig":null,"found":[".agents/skills/ci-test-audit/SKILL.md",".agents/skills/add-sgl-kernel/SKILL.md",".agents/skills/clean-startup-log/SKILL.md",".agents/skills/compute-mamba-ratio/SKILL.md",".agents/skills/babysit-pr-to-pass-ci/SKILL.md",".agents/skills/cookbook-add-model/SKILL.md",".agents/skills/ci-workflow-guide/SKILL.md",".agents/skills/add-jit-kernel/SKILL.md",".agents/skills/debug-distributed-hang/SKILL.md",".agents/skills/cookbook-migrate-model/SKILL.md",".agents/skills/cookbook-review-pr/SKILL.md",".agents/skills/debug-cuda-crash/SKILL.md",".agents/skills/env-var-conventions/SKILL.md",".agents/skills/generate-profile/SKILL.md",".agents/skills/kernel-organization/SKILL.md",".agents/skills/kl-consistency-test/SKILL.md",".agents/skills/large-class-style/SKILL.md",".agents/skills/llm-torch-profiler-analysis/SKILL.md",".agents/skills/mechanical-refactor-verify/SKILL.md",".agents/skills/mechanical-refactor-verify/guide-modify-skill.md",".agents/skills/processor-model-parity/SKILL.md",".agents/skills/scripted-runtime-notes/SKILL.md",".agents/skills/sglang-bisect-ci-regression/SKILL.md",".agents/skills/sglang-cherrypick/SKILL.md",".agents/skills/sglang-prod-incident-triage/SKILL.md",".agents/skills/sglang-runtime-context/SKILL.md",".agents/skills/speculative-naming/SKILL.md",".agents/skills/write-sglang-test/SKILL.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md","docs/AGENTS.md","sgl-router/ARCHITECTURE.md",".claude/rules/comment-style.md",".claude/rules/cookbook-first.md",".claude/rules/forward-batch-init-new-purity.md",".claude/rules/general-code-style.md",".claude/rules/kernel-benchmark.md",".claude/rules/modify-component-must-read.md",".claude/rules/no-dataclasses.md",".claude/rules/no-getattr-defensive.md",".claude/rules/prefer-sgl-eval.md",".claude/rules/schedule-batch-out-of-place-mutation.md",".claude/rules/unit-test-admission.md","tools/sglang-simulator/README.md",".agents/skills/ci-test-audit/action-items.md",".agents/skills/mechanical-refactor-verify/guide-construct-proof.md",".agents/skills/mechanical-refactor-verify/guide-split.md",".agents/skills/mechanical-refactor-verify/guide-verify-proof.md",".agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md",".agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md","examples/profiler/nsys_profile_tools/README.md","tools/sglang-simulator/examples/README.md",".agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml",".agents/skills/clean-startup-log/references/noise-sources.md",".agents/skills/cookbook-add-model/references/authoring-reference.md",".agents/skills/cookbook-add-model/references/diffusion-authoring.md",".agents/skills/cookbook-add-model/references/engine-axis.md",".agents/skills/cookbook-add-model/references/mintlify-authoring.md",".agents/skills/cookbook-add-model/references/vendor-logo.md",".agents/skills/cookbook-migrate-model/references/dimension-mapping.md",".agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md",".agents/skills/llm-torch-profiler-analysis/references/heuristics.md",".agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md",".agents/skills/llm-torch-profiler-analysis/references/source-map.md",".agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md",".agents/skills/sglang-prod-incident-triage/references/case-studies.md",".agents/skills/sglang-prod-incident-triage/references/decision-tree.md",".agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md",".agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md","tools/sglang-simulator/examples/assets/replay_table.json","tools/sglang-simulator/examples/sim_configs/aic_silicon.json","tools/sglang-simulator/examples/sim_configs/aic_sol.json","tools/sglang-simulator/examples/sim_configs/ml.json","tools/sglang-simulator/examples/sim_configs/replay.json","tools/sglang-simulator/examples/workloads/sharegpt-example.json","tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json","tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json","tools/sglang-simulator/test/assets/qwen3-8b/config.json","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md","tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md","tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md"],"skills":{".agents/skills/ci-test-audit/SKILL.md":"---\nname: ci-test-audit\ndescription: Audit the existing test tree and CI configuration for improvements, using a catalog of patterns previously applied in this repo. Use when asked to audit tests or CI, shrink CI time or cost, find redundant or misplaced tests, clean up a test group, or review whether a CI change follows established practice.\n---\n\n# CI / Test Audit\n\nAudits what is already in `test/`, `.github/` and `scripts/ci/`. For writing a new test see\n[write-sglang-test](../write-sglang-test/SKILL.md); for how the pipeline dispatches and\ngates work see [ci-workflow-guide](../ci-workflow-guide/SKILL.md).\n\n## How to use it\n\nRead [action-items.md](action-items.md) first. It is a catalog of patterns that have\nbeen applied to this repo, each with a way to spot it and example PRs. The catalog is the substance of this skill; everything below is only scaffolding.\n\nThe user says what to audit -- a test group, a stage, a workflow file, or a complaint\nlike \"this suite is slow\" -- and how thoroughly. Read that scope and decide which\npatterns apply.\n\nAn audit is only useful if it is specific. \"This suite could be trimmed\" is not a\nfinding; \"`TestFooLargePage` is a strict subset of `TestFooRetractLargePage`, drop it\"\nis. Every finding needs the file, the pattern it matches, and what to do.\n\n## Reporting\n\nGroup findings by confidence in the claim, not by pattern id. Lead with the ones where\nthe evidence is in the file you just read; keep the speculative ones separate and say\nwhat you would need to check. Skip anything you are not reasonably sure about -- a long\nlist of maybes costs more to triage than it saves.\n\nPropose, do not apply. Deleting a test, moving a registration to another stage, and\nchanging a threshold are all decisions for the user. Land them as separate changes,\nsince a trim and a threshold change fail for different reasons.\n\n## Keeping the catalog current\n\nExample PRs age. When a pattern shows up in newer work, or a new pattern recurs, update\n`action-items.md`: the pattern text should stay general, the examples are\nreplaceable evidence.\n",".agents/skills/add-sgl-kernel/SKILL.md":"---\nname: add-sgl-kernel\ndescription: Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)\n---\n\n# Tutorial: Adding a New Kernel to `sgl-kernel` (AOT / Heavyweight)\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as an AOT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float)\n- Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `python/sglang/kernels/aot/include/utils.h`)\n\n## Two rules of thumb (must follow)\n\n1. **Prefer `python/sglang/kernels/jit` first** when the kernel does **not** depend on CUTLASS or another large C++ project. This is the default path for lightweight kernels that benefit from rapid iteration.\n2. **Prefer `sgl-kernel`** when the kernel **does** depend on CUTLASS or another large C++ project, or when it should be part of the AOT wheel / torch op registration flow.\n3. **Exception**: if the dependency is `flashinfer`, or CUTLASS that is already provided through `flashinfer`, the kernel can still be implemented as `jit_kernel`.\n\nIn addition, every new kernel must ship with:\n\n- **Tests** (pytest)\n- **A benchmark script** (`marker.do_bench` from `sglang.kernels.jit.benchmark` — it works for any callable, not only JIT kernels)\n\n---\n\n## Repository integration map\n\nYou will typically touch these files/areas:\n\n- Implementation: `python/sglang/kernels/aot/csrc/elementwise/scale.cu` (pick the right subdirectory)\n- Public declarations: `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- Torch extension registration: `python/sglang/kernels/aot/csrc/common_extension.cc`\n- Build: `python/sglang/kernels/aot/CMakeLists.txt` (`set(SOURCES ...)`)\n- Python API: `python/sglang/kernels/aot/python/sgl_kernel/` and `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n- Tests: `python/sglang/kernels/aot/tests/test_scale.py`\n- Benchmarks: `python/sglang/kernels/aot/benchmark/bench_scale.py`\n\n---\n\n## Step 1: Implement the kernel in `csrc/`\n\nPick the right subdirectory:\n\n- `csrc/elementwise/` — for element-wise ops (our example)\n- `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories\n\nCreate `python/sglang/kernels/aot/csrc/elementwise/scale.cu`:\n\n```cpp\n#include \n#include \n#include \n\n#include \"utils.h\" // DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16\n\n// scale_kernel: out[i] = input[i] * factor\n// Supports float, half (__half), __nv_bfloat16 via template T\ntemplate \n__global__ void scale_kernel(T* __restrict__ out,\n const T* __restrict__ input,\n float factor,\n int64_t n) {\n int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x;\n if (idx < n) {\n out[idx] = static_cast(static_cast(input[idx]) * factor);\n }\n}\n\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor) {\n TORCH_CHECK(input.is_cuda(), \"input must be a CUDA tensor\");\n TORCH_CHECK(input.is_contiguous(), \"input must be contiguous\");\n TORCH_CHECK(out.is_cuda(), \"out must be a CUDA tensor\");\n TORCH_CHECK(out.is_contiguous(), \"out must be contiguous\");\n TORCH_CHECK(out.sizes() == input.sizes(), \"out and input must have the same shape\");\n TORCH_CHECK(out.scalar_type() == input.scalar_type(),\n \"out and input must have the same dtype\");\n\n const int64_t n = input.numel();\n const int threads = 256;\n const int blocks = (n + threads - 1) / threads;\n\n const cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n const at::cuda::OptionalCUDAGuard device_guard(device_of(input));\n\n // Dispatches over float, float16, bfloat16\n DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {\n scale_kernel<<>>(\n static_cast(out.data_ptr()),\n static_cast(input.data_ptr()),\n static_cast(factor),\n n);\n cudaError_t status = cudaGetLastError();\n TORCH_CHECK(status == cudaSuccess,\n \"scale_kernel launch failed: \", cudaGetErrorString(status));\n return true;\n });\n}\n```\n\n**Key points:**\n\n- Use `at::Tensor` (PyTorch tensors), `TORCH_CHECK` for validation, `at::cuda::getCurrentCUDAStream()` for stream\n- Keep Python wrappers thin; do shape/dtype/device validation in C++ right around the launch path\n- `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` covers `float`, `half` (FP16), `__nv_bfloat16` (BF16)\n- Add device error checking after every kernel launch\n- If a kernel only works on certain architectures, enforce that with `TORCH_CHECK` and skip logic in tests\n\n---\n\n## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`\n\nEdit `python/sglang/kernels/aot/include/sgl_kernel_ops.h`, add to the elementwise section:\n\n```cpp\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor);\n```\n\n---\n\n## Step 3: Register the op in `csrc/common_extension.cc`\n\nEdit `python/sglang/kernels/aot/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:\n\n```cpp\n// From csrc/elementwise\nm.def(\"scale(Tensor! out, Tensor input, float factor) -> ()\");\nm.impl(\"scale\", torch::kCUDA, &scale);\n```\n\n**Key points:**\n\n- `Tensor!` means in-place / mutable output argument\n- The schema is important for `torch.compile` and for consistent call signatures\n- Keep the torch schema in PyTorch scalar types (`float` here), but note that the C++ launcher signature still needs `double` for scalar arguments accepted by `torch::Library`\n\n---\n\n## Step 4: Add the new source file to `CMakeLists.txt`\n\nEdit `python/sglang/kernels/aot/CMakeLists.txt`, add to `set(SOURCES ...)`:\n\n```cmake\ncsrc/elementwise/scale.cu\n```\n\n**Key points:**\n\n- Keep the list **alphabetically sorted** (the file explicitly requires this)\n- If the kernel has arch constraints, reflect that in tests/benchmarks via skip logic\n\n---\n\n## Step 5: Expose a Python API under `python/sglang/kernels/aot/python/sgl_kernel/`\n\nPrefer following the existing module organization first. For elementwise kernels, the usual pattern is:\n\n- implement the Python wrapper in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`\n- then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n\nFor example, in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`, add:\n\n```python\nimport torch\n\ndef scale(\n input: torch.Tensor,\n factor: float,\n out: torch.Tensor | None = None,\n) -> torch.Tensor:\n \"\"\"\n Element-wise scale: out = input * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n input : CUDA input tensor\n factor : scale factor (float)\n out : optional pre-allocated CUDA output tensor (same shape/dtype as input)\n \"\"\"\n if out is None:\n out = torch.empty_like(input)\n torch.ops.sgl_kernel.scale.default(out, input, factor)\n return out\n```\n\nThen re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.\n\n---\n\n## SGLang integration (required for runtime use)\n\nAfter exposing the AOT wheel symbol, add a lazy wrapper and a `KernelSpec` under\n`python/sglang/kernels/ops//`. Set `backend=KernelBackend.AOT`, record\nactual device/architecture support, and keep `sgl_kernel` imports inside the\nimplementation path. SGLang runtime and integration tests import that wrapper.\nFor this example, use `sglang.kernels.ops.elementwise.scale`.\n\nThe wheel-level tests below validate its standalone API/build. They do not\nreplace CI-registered SGLang correctness tests in\n`test/registered/kernels/ops/elementwise/test_scale.py` and benchmarks in\n`test/registered/kernels/benchmark/elementwise/bench_scale.py`. Follow\n`write-sglang-test` for their registration and CI budget; do not register tests\nunder the shipped `python/sglang/` package.\n\n## Step 6: Write tests (required)\n\nCreate `python/sglang/kernels/aot/tests/test_scale.py`:\n```python\nimport pytest\n\nimport torch\nimport sgl_kernel\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [128, 1024, 4096, 65536])\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0])\ndef test_scale_correctness(dtype, size, factor):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n\n result = sgl_kernel.scale(input, factor, out=out)\n assert result is out\n\n expected = input * factor\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\ndef test_scale_shape_mismatch():\n input = torch.randn(128, dtype=torch.float16, device=\"cuda\")\n out = torch.empty(256, dtype=torch.float16, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"same shape\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\ndef test_scale_cpu_input():\n input = torch.randn(128, dtype=torch.float16) # CPU\n out = torch.empty_like(input)\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-q\"]))\n```\n\n---\n\n## Step 7: Add a benchmark (required)\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nCreate `python/sglang/kernels/aot/benchmark/bench_scale.py`:\n\n```python\nimport torch\n\nimport sgl_kernel\nfrom sglang.kernels.jit.benchmark import marker\n\n\ndef sglang_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n sgl_kernel.scale(input, factor, out=out)\n\n\ndef torch_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n torch.mul(input, factor, out=out)\n\n\nFN_MAP = {\"sglang\": sglang_scale, \"torch\": torch_scale}\n\n\n@marker.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32], [torch.float16])\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096]) # 1K .. 512K\n@marker.benchmark(\"provider\", [\"sglang\", \"torch\"])\ndef benchmark(dtype: torch.dtype, size: int, provider: str):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n return marker.do_bench(\n FN_MAP[provider],\n # Pass every tensor through input_args (not a closure) so marker rotates\n # them across CUDA-graph calls to defeat L2 reuse.\n input_args=(input, 2.0, out),\n # Bandwidth = bytes(input) + bytes(out), both already in input_args.\n memory_output=None,\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n---\n\n## Step 8: Build\n\nBuild:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j16\n```\n\nIf you need to limit host resource usage:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j1 MAX_JOBS=2 CMAKE_ARGS=\"-DSGL_KERNEL_COMPILE_THREADS=1\"\n```\n\n---\n\n## Step 9: Validate\n\nAfter building successfully, run the test and benchmark:\n\n```bash\npytest python/sglang/kernels/aot/tests/test_scale.py -q\npython python/sglang/kernels/aot/benchmark/bench_scale.py\n```\n\nPR CI also runs `pr-test-sgl-kernel.yml`, including the B200 job\n`sgl-kernel-b200-test` when kernel changes are detected. Use that job as the\nBlackwell coverage signal for AOT `sgl-kernel` changes.\n\n---\n\n## Troubleshooting\n\n- **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`\n- **Memory errors**: `compute-sanitizer --tool memcheck python ...`\n- **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`\n- **Binary bloat**: use `python/sglang/kernels/aot/analyze_whl_kernel_sizes.py`\n- **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time\n\n---\n\n## References\n\n- `python/sglang/kernels/aot/README.md`\n- `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- `python/sglang/kernels/aot/csrc/common_extension.cc`\n- `python/sglang/kernels/aot/CMakeLists.txt`\n- `python/sglang/kernels/aot/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends\n- `python/sglang/kernels/aot/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern\n\n## Summary of Files Created/Modified\n\n```\npython/sglang/kernels/aot/csrc/elementwise/scale.cu # NEW: CUDA kernel + launcher\npython/sglang/kernels/aot/include/sgl_kernel_ops.h # MODIFIED: C++ declaration\npython/sglang/kernels/aot/csrc/common_extension.cc # MODIFIED: schema + dispatch registration\npython/sglang/kernels/aot/CMakeLists.txt # MODIFIED: add source file (alphabetical)\npython/sglang/kernels/aot/python/sgl_kernel/elementwise.py # MODIFIED: Python wrapper\npython/sglang/kernels/aot/python/sgl_kernel/__init__.py # MODIFIED: re-export Python API\npython/sglang/kernels/aot/tests/test_scale.py # NEW: tests\npython/sglang/kernels/aot/benchmark/bench_scale.py # NEW: benchmark\n```\n",".agents/skills/clean-startup-log/SKILL.md":"---\nname: clean-startup-log\ndescription: Audit SGLang startup logs, save evidence, and propose cleanup for user review. With no arguments, run Qwen3-8B at TP1 and TP2 plus gpt-oss-20b at TP1.\ndisable-model-invocation: true\n---\n\n# Audit SGLang Startup Logs\n\nThe default outcome is saved logs and a findings report. Apply runtime changes\nonly after the user selects them. A request to edit this skill does not itself\nlaunch servers.\n\n## Default runs\n\nA bare `$clean-startup-log` invocation runs these cases sequentially without\nasking for commands. Explicit commands, models, or TP sizes replace this matrix.\n\n| Case / log filename | Command |\n|---|---|\n| `qwen3-8b-tp1.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B` |\n| `qwen3-8b-tp2.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B --tp 2` |\n| `gpt-oss-20b-tp1.log` | `uv run sglang serve --model-path openai/gpt-oss-20b` |\n\nThese cover dense, tensor-parallel, and MoE/hybrid sliding-window attention\nstartup. Reuse complete captures from the current audit when code and environment\nhave not changed.\n\n## Capture logs\n\n1. Check the checkout, free GPUs, and ports once. Use the requested command when\n resources are free; otherwise select free GPUs with `CUDA_VISIBLE_DEVICES` and\n an unused `--port`, recording the adjustments. Leave existing servers alone.\n2. Create a unique directory with `mktemp -d /tmp/sglang-startup-audit-XXXXXX`.\n Save raw stdout and stderr together in a separate log for each case, for\n example with `set -o pipefail` and `COMMAND 2>&1 | tee LOG_PATH`. Record commands,\n GPU IDs, ports, commit, relevant overrides, and readiness status.\n3. Wait for `The server is fired up and ready to roll!`, then stop that server and\n its workers before the next case. First runs can spend many minutes downloading\n weights or compiling FlashInfer kernels; check download/compiler activity\n before treating a quiet log as a hang. Preserve partial logs for failed or\n stalled starts and report the last stage. Continue independent cases when possible.\n4. Preserve the user's logging configuration, including `NCCL_DEBUG`. If NCCL\n verbosity needs explaining, inspect relevant shell settings, `NCCL_CONF_FILE`,\n and `/etc/nccl.conf`. Do not override intentional diagnostics or recommend\n `NCCL_DEBUG=WARN` solely because the output is long. Avoid full environment dumps.\n\n## Investigate efficiently\n\n- Scan for deprecations, duplicate handler output, unrelated import failures,\n unformatted prints, and unexpected warnings. Read representative excerpts and\n counts instead of repeatedly dumping `server_args`, progress redraws, or NCCL\n diagnostics. Normalize carriage returns for analysis only; preserve raw logs.\n- Trace each candidate to its actual emitter with focused `rg` searches. Inspect\n its log level: SGLang's formatter may omit severity. Group shared signatures\n across cases and distinguish handler duplication from separate GPU/process calls.\n- Repetition, WARNING severity, or a different third-party format alone does not\n establish a cleanup need. Consider whether the message explains configuration,\n progress, resource use, or an operational limitation.\n- Consult [noise-source hints](references/noise-sources.md) only for a matching\n signature or an unresolved emitter. Verify current code rather than trusting\n historical line numbers, fix status, or assumptions about unrelated models.\n\n## Accepted output\n\nPreserve these reviewed messages unless the user requests a different policy:\n\n- NCCL diagnostics enabled by the user's environment or host configuration.\n- NUMA permission warnings, including one check per GPU in TP runs.\n- GPT-OSS MXFP4 backend-selection warnings and default page-size selection warnings.\n- `Init Unified Radix Cache. Components: ... Tree Core: ...`, tree-cache summaries,\n SWA allocation details, and per-rank memory/timing records.\n- Useful progress bars, warmup HTTP access logs, uv synchronization messages,\n isolated NCCL/Gloo startup lines, and one timestamped HF authentication warning.\n\nThese can appear in a clean startup log. Do not repeatedly propose the declined\nNUMA deduplication, backend/page-size level changes, or NCCL verbosity override.\nKeep real operational warnings visible: for example, a Harmony vocabulary failure\ncan disable `/v1/responses` even when server readiness and `/generate` succeed.\n\n## Report before changing code\n\nReturn a compact run table with readiness status and clickable raw-log links.\nFor each actual cleanup candidate, give an exact representative message, affected\ncases/counts, source file/function, and specific proposed behavior. Distinguish\nconfirmed findings from suspicions and operational failures from logging noise.\n\nIf there are no actionable cleanup findings, say the logs are clean and no\nfurther cleanup is needed. Otherwise, ask which numbered changes to adopt and\nwait for the user's selections before editing runtime code or preparing patches.\nHonor existing approvals and declined items without asking again.\n\n## Apply selected changes\n\n- Batch compatible approved edits, then verify affected cases once. Repeat a\n startup only for a new change, failure, or unresolved concern; do not relaunch\n after every one-line edit. Save verification logs separately from baselines.\n- Preserve useful warnings and application handlers. HF can warn during early\n CLI model detection before `configure_logger()`, and spawned processes have\n independent logger state. Keep `configure_hf_hub_logger()` in both\n `suppress_noisy_warnings()` and `configure_logger()`; make repeated setup safe.\n- Keep the legacy compiled-kernel cache migration notice at DEBUG. Avoid broad\n library-level suppression or fd redirection for a narrow logging problem.\n- Run relevant formatting and focused existing checks. Add tests only when they\n verify meaningful behavior, not a log-level spelling. Report changes and\n verification; create branches, commits, and PRs when requested.\n",".agents/skills/compute-mamba-ratio/SKILL.md":"---\nname: compute-mamba-ratio\ndescription: Compute the optimal --mamba-full-memory-ratio (or --max-mamba-cache-size pin) for a hybrid attention + linear-attention (Mamba / GDN / KDA) model's two serving memory pools, from the workload and serving config. Use when a user asks what ratio to set, why concurrency is clamped, or how to size the state vs KV pools for a hybrid model.\n---\n\n# Optimal hybrid dual-pool ratio (`--mamba-full-memory-ratio`)\n\nA hybrid model (attention layers + linear-attention layers — the recurrent-state family: Mamba/SSM, GDN, KDA, etc.) splits serving memory into two independently-budgeted pools, fixed once at startup:\n\n- **state pool** (the linear-attention recurrent state) → caps **concurrency** (hard: whole slots, worst-case reserved, fail-loud)\n- **full-KV pool** (attention KV) → caps **context × concurrency** (soft: paged, over-committable via retraction)\n\n`--mamba-full-memory-ratio r` splits the post-weight budget: `mamba_budget = rest · r/(1+r)`, i.e. `mamba_budget : kv_budget = r`. This skill picks the `r` (or the pin-`--max-mamba-cache-size` alternative) at which **neither pool bottlenecks first** for the user's workload.\n\n## The formula\n\n```\nr* = (S + D) · token_equiv · dcp_size / L\ntoken_equiv = state_bytes_per_slot / kv_bytes_per_token\n```\n\n- `L` = average context length per request (input + output tokens)\n- `token_equiv` = full-KV token-equivalent of one state slot\n- `S` = state slots per running request (cache-strategy dependent, table below)\n- `D` = `--speculative-num-draft-tokens` (0 if NOSPEC); each running req carries `D` extra intermediate states\n- `dcp_size` = `--dcp-size` (1 without DCP). DCP shards the per-rank KV by `dcp_size`, so KV gets ~`dcp_size×` cheaper per request → the balance shifts that much toward the state pool.\n\n`r` is dimensionless (just the split). To also predict the actual concurrency you need `rest` (below).\n\n## Inputs to collect from the user\n\n1. **`L`** — average context (input + output) in tokens.\n2. **The two per-GPU byte constants** — one of:\n - **(a) measured (preferred, exact)** — from one boot log at *any* ratio:\n - `Mamba Cache is allocated. ... ssm_state size X GB` with `max_mamba_cache_size: N` → `state_bytes_per_slot = X / N`\n - `KV Cache is allocated. #tokens: M, KV size: Y GB` → `kv_bytes_per_token = Y / M`\n - **(b) derived** — model arch (linear-layer count + state dims `d_state/d_conv/heads/head_dim`; attention type + KV dims: MLA latent dim, or GQA `kv_heads·head_dim·layers`) **× the dtypes** below.\n3. **`S`** — from `--mamba-radix-cache-strategy`, the overlap scheduler, and `SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK` (table below).\n4. **`D`** — `--speculative-num-draft-tokens` (0 if NOSPEC; **also 0 when ReplaySSM spec-verify is enabled** — see caveats).\n5. **`dcp_size`** — `--dcp-size` (1 if no DCP). If DCP **and** spec with a replicated draft KV, apply the draft caveat (below).\n6. **KV dtype** (bf16 / fp8) and **ssm dtype** (fp32 / bf16) — they set the two byte constants (see the `token_equiv` 2×2).\n7. *(only to also predict the clamp, not just `r`)* **`rest`** = per-GPU memory − weights at the chosen `--mem-fraction-static`. Read `avail mem` after `Load weight end`, or `Memory pool end. avail mem` + pool sizes, from the boot log.\n\n### `S` (state slots per running request), set by `--mamba-radix-cache-strategy`\n\nAll strategies keep prefix caching on; they differ in the track buffer that snapshots chunk-boundary state under the overlap scheduler. `S = base + ping-pong`, where base is 3 (live state + radix retention/COW headroom), and ping-pong = 2 (overlap on, non-lazy) / 1 (lazy, **or** overlap off) / 0 (no track buffer at all). This mirrors `kv_cache_configurator._calculate_mamba_ratio`; read it there if a release moves the constants.\n\n`SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK=1` takes the base to 2, since the decode-time skip frees one resident slot per running request. `no_buffer` is the exception and stays at an effective 3: it adds that slot straight back, because its binding limit is the prefill-to-decode peak, which a decode-time change does not shrink.\n\n| strategy | S | S with decode-lock skip | notes |\n|---|---|---|---|\n| `--disable-radix-cache` | 1 | 1 | prefix cache off entirely; most concurrency, no reuse |\n| `no_buffer` | 3 | 3 | no track buffer; requires overlap **off** (+ `page_size 1`) → lower decode throughput |\n| `extra_buffer` (default), overlap on | 5 | 4 | track buffer reserved per running request |\n| `extra_buffer`, overlap off | 4 | 3 | same buffer, but ping-pong costs 1 slot instead of 2 |\n| `extra_buffer_lazy` (recommended) | 4 | 3 | track buffer allocated lazily at the boundary → 1 fewer slot/req; requires overlap **on** |\n\nThe skip is a real concurrency lever, not a rounding detail: at the default strategy it takes `S` from 5 to 4, so the same state budget admits 25% more requests and `r*` drops by the same factor.\n\nOverlap counts as off when `--disable-overlap-schedule` is passed **or** `--pp-size > 1` (pipeline parallelism turns it off for you). So a PP deployment left on the default strategy sits on the `S = 4` row, not `S = 5` — using 5 there overstates the state cost and starves KV.\n\n### `token_equiv` — compute it from the two byte constants\n\n`token_equiv = state_bytes_per_slot / kv_bytes_per_token`. Both bytes are **per-model, per-GPU** — read them from a boot log (or derive from the arch + dtypes) and divide. That's the whole thing; the dtype effects below are just for a quick mental estimate.\n\nThe recurrent state has two parts: `state_bytes = SSM_bytes + conv_bytes`. The **SSM** part scales with `--mamba-ssm-dtype`; the **conv** part is always bf16 (`SGLANG_MAMBA_CONV_DTYPE`, fixed). How the dtype knobs move `token_equiv`:\n\n- **fp8 KV → ×2** (universal, exact): fp8 is 1 byte vs bf16's 2, so `kv_bytes_per_token` halves.\n- **fp32 ssm vs bf16 ssm → ×≈2**: only the SSM tensor doubles, conv stays fixed, so the exact factor is `(2·SSM + conv)/(SSM + conv)` — just under 2. **conv is usually small relative to SSM, so estimate ≈2** (e.g. one measured model: conv ≈7% of state → factor ≈1.9). Only bother with the exact split if a model's conv is unusually large.\n- **fp32 ssm + fp8 KV → ×≈4**.\n\nFor an actual number, don't apply factors — read `state_bytes_per_slot` and `kv_bytes_per_token` for **your** model+dtype from the boot log and divide.\n\nThe `GB` the boot log prints for both pools is **GiB** (`1024³`). Convert both the same way and `token_equiv` is unaffected either way, since it is a ratio — but get it right before comparing an absolute byte figure against your own log.\n\nWorked example (one measured model, TP8, fp32 ssm + fp8 KV), read off one boot log:\n\n```\nMamba Cache is allocated. max_mamba_cache_size: 257, conv_state size: 0.46GB, ssm_state size: 13.04GB\nKV Cache is allocated. dtype: torch.float8_e4m3fn, #tokens: 1167552, KV size: 15.03 GB\n```\n\n`state_bytes_per_slot = (0.46 + 13.04) GiB / 257 ≈ 56.4 MB`, `kv_bytes_per_token = 15.03 GiB / 1167552 ≈ 13.8 KB` → `token_equiv ≈ **4080**`.\n\n## Compute\n\n```python\ndef optimal_ratio(L, state_bytes_per_slot, kv_bytes_per_token, S, D=0, dcp_size=1):\n token_equiv = state_bytes_per_slot / kv_bytes_per_token\n r = (S + D) * token_equiv * dcp_size / L\n return r # value of --mamba-full-memory-ratio (>1 is legal)\n\ndef predict_clamp(rest_bytes, r, state_bytes_per_slot, S, D=0):\n mamba_budget = rest_bytes * r / (1 + r)\n slots = mamba_budget / state_bytes_per_slot\n # spec: each running req reserves (S+D) worth; non-spec just S\n return int(slots // S) # NOSPEC; with spec the budget joint-solves for (S+D)·per_req per req\n```\n\nThen state the result three ways: the **`r` value**, the **predicted clamp** (if `rest` given), and **which pool binds** (`min(mamba_clamp, KV_cap)`, where `KV_cap = kv_tokens · dcp_size / L`).\n\n## Procedure\n\n1. **Free lever first**: if the boot log shows large idle in `Memory pool end. avail mem` (e.g. 30–40 GB at mem-frac 0.85), raise `--mem-fraction-static` (→0.92) before touching the split — it grows `rest` for both pools at no cost. Validate graph-capture headroom once.\n2. Collect the inputs. Prefer a real boot log for the two byte constants.\n3. Compute `r*`. If `r*` would drive the state pool below one request's worth (`mamba_budget < S · per_req`, happens at very long `L`), **switch to pinning `--max-mamba-cache-size = target_concurrency · S`** and let the rest go to KV — a sub-0.15 `r` is fragile.\n4. Report `r`, predicted clamp, binding pool, and any dtype accuracy gate that applies.\n\n## Worked examples (TP8, B300, validated against measured clamps)\n\nConstants (from boot logs): `state_bytes_per_slot ≈ 56.4 MB` (fp32 ssm), `kv_bytes_per_token ≈ 13.8 KB` (fp8) → `token_equiv ≈ 4080`. Cache strategy `extra_buffer_lazy` → `S = 4`. Workload `L = 9216` (8192 in + 1024 out). NOSPEC → `D = 0`.\n\n- **TP (dcp_size=1)**: `r = 4 · 4080 · 1 / 9216 ≈ 1.8`. Measured: clamp 88, KV cap 87 → balanced. ✅\n- **DCP8 (dcp_size=8)**: `r = 4 · 4080 · 8 / 9216 ≈ 14`. Measured at r=14: clamp 125, KV cap 129 → balanced. ✅ (At the naive r=1.8 the DCP KV pool is ~8× over-provisioned — cap 683 vs clamp 86 — wasting budget that should go to the state pool.)\n\n`predict_clamp` on the same model at the **default** `extra_buffer`, overlap on and the decode-lock skip off (`S = 5`), three boot logs:\n\n| config | `max_mamba_cache_size` | `mmcs // S` | measured `max_running_requests` |\n|---|---|---|---|\n| TP8, no DCP | 257 | 51 | 51 |\n| DCP8, `r = 5.97` | 451 | 90 | 90 |\n| DCP8, `r = 9` | 474 | 94 | 94 |\n\nReference `r` for **this example model** (`token_equiv ≈ 4080`, i.e. fp32 ssm + fp8 KV; S=5 default; multiply by `(S+D)/S` for spec, by `dcp_size` for DCP) — recompute with your own `token_equiv` for a different model:\n\n| L | 2K | 4K | 8K | 32K | 64K | 128K |\n|---|---|---|---|---|---|---|\n| r | 10.0 | 5.0 | 2.5 | 0.62 | 0.31 | 0.16 |\n\n## Caveats\n\n- **DCP + spec**: a replicated (non-DCP-sharded) draft KV does **not** shard by `dcp_size`; at long `L` it dominates per-token cost, so the clean `×dcp_size` overstates DCP's advantage — fall back to a per-token direct-solve (KV term `/dcp` + an un-sharded draft-KV term) when spec is on. (NOSPEC → clean `×dcp_size` holds.)\n- **Spec + ReplaySSM → use `D=0`, not the draft-token count.** When ReplaySSM spec-verify is enabled, the `D` intermediate SSM states move off the per-request slot budget onto a **fixed ring** (a one-time deduction from `rest`, not a per-req term). So the mamba-slot cost per running request drops back to `S` (clamp = `mmcs / S`, not `/(S+D)`), and the balance ratio returns to the NOSPEC value (`r* ≈ S·token_equiv·dcp/L`). Measured example (TP8, D=8, L≈9K): applying `D=8` computes `r*≈2.6` but the true optimum is `r≈1.0` — `D=8` lands KV-bound at ~45% below the achievable peak concurrency. Plain spec (no replayssm) keeps `D` = the draft-token count.\n- **Under DCP the balance point sits beyond any realistic `L`** (the state pool binds first for almost everything), so the practical recommendation is to **pin `--max-mamba-cache-size = target_concurrency · S` directly** rather than dial a large `r`.\n- **`--max-mamba-cache-size` overrides `r`**; bytes beyond the `r` budget come out of the KV pool one-for-one.\n- **Precision changes are behind accuracy gates**: `--kv-cache-dtype fp8_e4m3` (doubles `token_equiv` → doubles `r`) and `--mamba-ssm-dtype bfloat16` (~halves `token_equiv`; also silently switches the linear-attention decode backend on SM100+ — pin `--linear-attn-decode-backend triton`) shift outputs; validate accuracy for the workload before production.\n- **Asymmetry**: the state pool is worst-case-reserved and fail-loud; KV degrades gracefully (retraction). When `L` is spiky/uncertain, bias `r` **up** rather than starve the state pool.\n",".agents/skills/babysit-pr-to-pass-ci/SKILL.md":"---\nname: babysit-pr-to-pass-ci\ndescription: Start and persistently pursue a goal to babysit an SGLang pull request until selected GitHub Actions workflows pass on the latest PR head. Use when asked to monitor, babysit, retry, or fix PR CI for lint.yml, pr-test.yml, pr-test-extra.yml, AMD, or other named workflows; classify failures as PR-related versus flaky or infrastructural, auto-fix and push only small clean fixes, rerun failed jobs only up to 10 times, and ignore unselected workflows.\n---\n\n# Babysit a PR Until Selected CI Passes\n\nDrive only the selected workflows to green on the PR's latest head commit. Keep working across turns by starting a durable goal; do not stop merely because a run was dispatched, a retry began, or a fix was pushed.\n\n## Invocation and arguments\n\nAccept:\n\n```text\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] [WORKFLOW ...]\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] --only WORKFLOW [WORKFLOW ...]\n```\n\n- If the PR is omitted, resolve the PR associated with the current branch.\n- If workflows are omitted, select `lint.yml` and `pr-test.yml`.\n- Without `--only`, add named workflows to the two defaults.\n- With `--only`, monitor exactly the named workflows.\n- Accept a workflow basename such as `pr-test-extra.yml`, a repository path such as `.github/workflows/pr-test-extra.yml`, or an unambiguous workflow display name.\n- Treat `extra` as `pr-test-extra.yml`.\n- Accept natural-language equivalents, such as “PR 12345 plus extra and AMD.” Resolve “AMD” to an exact workflow file from the user's wording or repository context; do not guess when multiple AMD workflows fit.\n\nExamples:\n\n```text\n$babysit-pr-to-pass-ci\n$babysit-pr-to-pass-ci 12345\n$babysit-pr-to-pass-ci https://github.com/sgl-project/sglang/pull/12345 pr-test-extra.yml\n$babysit-pr-to-pass-ci 12345 --only pr-test-amd.yml\n```\n\n## Start or continue the durable goal\n\nTreat explicit invocation of this skill as authorization to create a goal and perform the scoped CI actions below.\n\n1. Resolve the PR and exact workflow set first.\n2. Inspect the current goal.\n3. If no unfinished goal exists, create one immediately using the product's goal mechanism. Include the resolved PR URL, selected workflow files, and this stopping condition: every selected workflow has a successful run for the PR's latest head SHA.\n4. Include these constraints in the goal objective:\n - Ignore every unselected workflow.\n - Diagnose each selected-workflow failure as PR-related, unrelated/infrastructure, or still uncertain.\n - Implement, validate, commit, and push a PR-related fix only when it is clean, non-tricky, and at most 100 changed lines.\n - Stop for user review before editing when the proposed fix is tricky, dirty, or over 100 changed lines.\n - Retry unrelated failures at most 10 times per workflow and head SHA, using failed-jobs-only reruns.\n - Never rerun a full workflow merely to recover failed tests.\n5. If the active goal already represents this same PR and workflow set, continue it instead of creating another goal.\n6. If a different unfinished goal is active, do not replace it silently. Report the conflict and ask the user to pause or clear it.\n\nDo not mark the goal complete until the completion check at the end of this skill succeeds.\n\n## Preflight\n\n1. Verify GitHub authentication and resolve the repository without changing remote configuration.\n2. Resolve the PR and record at least:\n - PR number and URL\n - open/draft state\n - base branch and base SHA\n - head branch and head SHA\n - head repository owner/name, including whether it is a fork\n - labels\n - whether maintainers may modify the branch\n3. Require an open PR. If it is draft and a selected workflow is blocked by the draft gate, stop and ask the user to mark it ready; do not change draft state automatically.\n4. Normalize every selected workflow to an existing file under `.github/workflows/`. Fail fast on ambiguous or nonexistent names.\n5. Read each selected workflow's triggers and gates. Do not assume all optional workflows behave like `pr-test.yml`.\n6. Inspect `git status`. Preserve all user changes. If a fix becomes necessary and the current checkout is dirty, on a different branch, or otherwise unsafe, use an isolated worktree for the PR head.\n\nKnown SGLang gates:\n\n- `pr-test.yml` requires the `run-ci` label for relevant PR changes. If its existing gate job failed because the label was absent, add `run-ci` when authorized and rerun failed jobs in that existing run; adding the label alone does not necessarily create a new `pr-test.yml` run.\n- `pr-test-extra.yml` requires both `run-ci` and `run-ci-extra`. It listens for relevant label events, so adding a missing label may create a fresh run. Prefer that fresh current-head run; rerun failed jobs in an existing current run only when needed.\n- AMD-extra workflows may use the same `run-ci` plus `run-ci-extra` gate. Verify the actual selected file.\n- Treat draft, cooldown, maintenance, and missing-label failures as gate/infrastructure conditions, not code regressions. Fix only the safe gate condition in scope; wait for cooldown or maintenance windows rather than burning retries immediately.\n\nAdding the known CI opt-in labels is within scope. Do not add unrelated labels, mark the PR ready, alter reviewers, merge, or mutate other PR metadata.\n\n## Track only current-head runs\n\nUse both PR checks and workflow-run data so workflow names, run IDs, attempts, jobs, and logs can be correlated. Prefer workflow file IDs over display names when filtering runs.\n\nFor each selected workflow:\n\n1. Select the newest run associated with the current PR head SHA.\n2. Ignore runs for older head SHAs, including green runs made stale by a new push.\n3. Treat `queued`, `requested`, `waiting`, `pending`, and `in_progress` as active; keep monitoring.\n4. Treat a workflow as passing only when its current-head run concludes `success`.\n5. If no current-head run exists, wait briefly for event delivery, then inspect the workflow's path filters, event triggers, and label gates.\n6. Trigger only through a mechanism explicitly supported by that workflow. Never invent `workflow_dispatch` inputs or dispatch a workflow against the default branch when the intent is to test the PR head.\n7. If a selected workflow cannot safely be triggered for the PR head, report the exact blocker and request direction.\n\nAfter any user push or skill-authored push, re-read the PR head SHA, discard stale run selections, and start tracking the new head. Reset retry counters for the new head SHA.\n\n## Monitor persistently\n\n- Maintain a compact ledger containing workflow file, head SHA, run ID, attempt, status, conclusion, skill-initiated retry count, and latest failure signature.\n- Poll at reasonable intervals, normally 30–60 seconds while runs are active. Give concise progress updates during long waits.\n- Re-check the PR head SHA on each monitoring cycle and before every rerun or push.\n- Ignore failures, cancellations, and pending checks from unselected workflows. Do not diagnose, retry, fix, or wait for them.\n- Do not declare success while a selected workflow is absent, stale, pending, skipped, cancelled, neutral, or failed.\n\n## Diagnose a selected-workflow failure\n\n1. Inspect the failed jobs and failed-step logs. In `pr-test.yml`, distinguish the root failure from cascade failures caused by `check-pr-test-health`, wait jobs, or final aggregation.\n2. Record the exact failing test/check, error text, affected runner or hardware, and relevant log URL.\n3. Compare the failure against the PR's changed files and diff.\n4. Classify it:\n\n **PR-related** when evidence ties the failure to changed code, tests, dependencies, configuration, generated files, or behavior introduced by the PR. A reproducible failure on the PR that does not occur on the base is strong evidence.\n\n **Unrelated/infrastructure** when evidence points to a transient network or download error, runner loss, GPU/driver instability, service outage, pre-existing main-branch failure, known flaky threshold, unrelated test area, or another environmental cause.\n\n **Uncertain** when evidence is insufficient. Do not label a failure flaky merely because it might be intermittent. Narrowly reproduce it, compare recent main/scheduled CI, inspect relevant code, or otherwise gather enough evidence to choose a category.\n\n5. When several root failures exist, classify each one. Address PR-related failures before spending retries on unrelated cascades.\n\n## Handle a PR-related failure\n\nEstimate the proposed fix before editing.\n\nA fix is eligible for automatic execution only when all are true:\n\n- The fix is well understood, localized, and clean.\n- The fix is not tricky, hacky, speculative, or a broad workaround.\n- The fix changes at most 100 lines, counting additions plus deletions in the fix patch, not the pre-existing PR diff.\n- The fix does not require a broad API or architecture redesign, risky dependency migration, large generated-file rewrite, or unrelated cleanup.\n- A focused validation is available.\n\nIf any condition fails, stop before editing. Report the diagnosis, proposed approach, expected files, estimated size, risks, and validation plan, then wait for user review.\n\nFor an eligible fix:\n\n1. Work from the latest PR head in a clean checkout or isolated worktree.\n2. Make only the necessary change and add or update focused tests when appropriate.\n3. Run the smallest meaningful local validation. For lint failures, reproduce the specific hook/check before considering broader lint validation.\n4. Review the resulting patch and count changed lines. If it becomes over 100 lines or turns tricky/dirty, do not commit or push it; stop for review and clearly identify the uncommitted skill-authored changes.\n5. Commit only skill-authored files. Never use `git add .`, rewrite user commits, amend without request, or include pre-existing local changes.\n6. Push a new commit non-forcibly to the exact PR head repository and branch. Never force-push. If permission or fork routing is unclear, stop and report the push command/target that needs authorization.\n7. Resolve the new PR head SHA and return to current-head monitoring. A push is progress, not completion.\n\n## Handle an unrelated or infrastructure failure\n\nRetry only the failed jobs in the current selected-workflow run:\n\n```bash\ngh run rerun --failed --repo \n```\n\nRules:\n\n- Never use a full-run rerun as a substitute for `--failed`.\n- Never rerun a stale run after the PR head SHA changes or a newer current-head run exists.\n- Count only successfully dispatched skill-initiated reruns.\n- Allow at most 10 such reruns per selected workflow and head SHA.\n- Wait for each new attempt to finish, inspect its result, and update the failure signature before deciding on another retry.\n- If the failure changes or becomes plausibly PR-related, return to diagnosis instead of blindly consuming all retries.\n- If 10 retries are exhausted, stop. Report the attempts, recurring signatures, job/run links, and evidence that the failure is unrelated or still uncertain. Do not exceed the cap without a new explicit user instruction.\n\nIf GitHub cannot rerun only the failed portion of the selected run, stop and ask the user rather than rerunning the full workflow.\n\n## Completion check\n\nBefore completing the goal:\n\n1. Fetch the PR head SHA again.\n2. For every selected workflow, verify that its newest applicable run is for that exact SHA and concludes `success`.\n3. Verify no selected workflow is still pending and no success was invalidated by a newer push.\n4. Ignore all unselected workflows even if they are failing.\n5. Mark the goal complete only after all selected workflows satisfy these checks.\n\nReport the final PR URL, head SHA, selected workflows, passing run links, any commits pushed, and retry counts. Keep the report scoped to the selected workflows.\n",".agents/skills/cookbook-add-model/SKILL.md":"---\nname: cookbook-add-model\ndescription: Add a new model to the SGLang Cookbook (docs/, Mintlify), config-driven format — instantiate the model-agnostic template into a per-model config (+ benchmarks) JSX under src/snippets/configs/, an MDX page, the docs.json nav entry, NEW-tag hygiene, and the homepage vendor card. Interactive, multi-phase. Run with /cookbook-add-model.\ndisable-model-invocation: true\n---\n\n# Add a model to the SGLang Cookbook\n\n> Migrating an **existing legacy-template page** (one that imports a monolithic\n> `…/autoregressive/-deployment.jsx` generator)? Use the\n> `cookbook-migrate-model` skill instead — same target format, but the legacy\n> page (not the user) is the source of truth.\n\nThe cookbook is **config-driven**: two shared engines contain NO model-specific code —\n`docs/src/snippets/_deployment.jsx` (the 5-dim deploy matrix) and\n`_playground.jsx` (the diff-based override Playground). Adding a model = adding **data**:\na per-model `config` (+ optional `benchmarks`) consumed by both engines, plus an MDX page\nthat imports them. No engine edits.\n\n**Instantiate the model-agnostic template** (NOT a clone of any live cookbook — the\ntemplate is decoupled and covers all hardware + all axes):\n- `templates/config.jsx.tmpl` → `docs/src/snippets/configs//.jsx`\n- `templates/benchmarks.jsx.tmpl` → `…/-benchmarks.jsx` (skip if no numbers)\n- `templates/page.mdx.tmpl` → `docs/cookbook///.mdx`\n\nThe template uses explicit `__TOKEN__` placeholders; you fill them, prune what the model\nlacks, and replace the EXAMPLE cells with verified recipes. DeepSeek-V4 is a populated\n*instance* you can consult, but is not the template.\n\n**Deep references (read on demand, don't inline):**\n- [references/authoring-reference.md](references/authoring-reference.md) — field-by-field config / cells / playground / MDX contract.\n- [references/mintlify-authoring.md](references/mintlify-authoring.md) — MDX rules (forbidden syntax, JSX tables, labeled fences) + invocation-example patterns. Read before writing §1–§3 prose.\n- [references/engine-axis.md](references/engine-axis.md) — adding a new Playground feature axis (rare engine work).\n- [references/vendor-logo.md](references/vendor-logo.md) — new-vendor card logo: ask the user for the brand logo, then generate the icon-only 940×525 RGBA PNG (spec + Pillow recipe + `git add -f`).\n- [references/diffusion-authoring.md](references/diffusion-authoring.md) — required opening, tag, command-picker, and feature-overlay contract for diffusion model pages. Read before editing any `docs/cookbook/diffusion//.mdx` page.\n\n## Architecture at a glance\n\n```\n┌─────────────────────────────────────────────────────────────────┐\n│ cookbook///.mdx │\n│ import { Deployment } from \"/src/snippets/_deployment.jsx\"; │\n│ import { Playground } from \"/src/snippets/_playground.jsx\"; │\n│ import { config } from \"/src/snippets/configs/.../X.jsx\"; │\n│ │\n└─────────────────────────────────────────────────────────────────┘\n │ (config passed as React prop)\n ▼\n┌─────────────────────────────────────────────────────────────────┐\n│ src/snippets/configs//.jsx │\n│ export const config = { │\n│ supportedHardware, variants, quantizations, strategies, ... │\n│ cells: [ { match:{hw,variant,quant,strategy,nodes}, │\n│ env:[...], flags:[...] }, ... ], // 5-dim matrix │\n│ playgroundFeatures: { attention, moe, parsers, ... }, │\n│ }; │\n└─────────────────────────────────────────────────────────────────┘\n │ (consumed by BOTH engines — no model code in engines)\n ▼\n┌──────────────────────────────────┬──────────────────────────────┐\n│ _deployment.jsx │ _playground.jsx │\n│ Renders the verified matrix; │ Renders override chips + │\n│ one cell → its env/flags. │ diff against the cell. │\n└──────────────────────────────────┴──────────────────────────────┘\n```\n\nThe two widgets stay in sync via: the **URL hash** (deploy mirrors its selection;\nplayground reads it), the **`sglang-deploy-sel` custom event** (deploy dispatches on\nevery change; playground listens — `replaceState` doesn't fire `hashchange`), and the\nshared **`sglang-deploy-env` localStorage key** (HOST/PORT placeholders).\n\n> The main template is **autoregressive**. Diffusion pages use\n> `templates/diffusion-page.mdx.tmpl` plus\n> [references/diffusion-authoring.md](references/diffusion-authoring.md); do not force the\n> autoregressive deployment matrix on them. Omni pages follow their own category structure.\n> All categories still obey the Mintlify / NEW-tag / docs.json / category-card / validation\n> rules below.\n\n---\n\n**Interactive, multi-step workflow. Collect inputs incrementally — don't ask for\neverything upfront.** The real work is the verified `cells[]` recipes + measured\nbenchmarks (Phases 2 + 4); everything else is filling the template.\n\n## Phase 1 — Collect inputs\n\n1. **Model card** — HuggingFace repo/URL. **Fetch the page** and extract description,\n param count, architecture, context length, **license**. (Fetching guards factual bugs\n like an off-by-a-few-B param count.) If the model isn't public, ask the user.\n2. **Variants / quantizations** — keep separate: variants are size/mode (e.g. Flash/Pro,\n Instruct/Thinking); quantizations come from the HF card / linked repos (BF16/FP8/FP4/…).\n Default to BF16 when a full-precision repo exists.\n3. **Tested hardware + parallelism** — which platforms are actually tested, and TP/EP/DP\n for each. List only tested hw (unlisted greys out).\n4. **Verified launch recipes** — the full `sglang serve` flags per\n (hw × variant × quant × strategy × nodes) combo → these become `cells[]`. Rewrite any\n `python -m sglang.launch_server` to `sglang serve` form.\n5. **sglang version / image tag** — ask which sglang build the recipes + benchmarks ran on\n (a release like `0.5.x`, or `main`/nightly). **Never guess or hallucinate it.** This one\n tag fills `dockerImages` and the benchmarks' `sglang_version`; when the user is unsure,\n default the image to `lmsysorg/sglang:dev` (nightly) rather than inventing a release.\n6. **Pre-flight**: `gh pr list --repo sgl-project/sglang --search \"\"` (dup check).\n\n**Hardware reference** (the shared `HARDWARE_CATALOG` in `_deployment.jsx`). A GPU **not** in\nthis table (RTX PRO 6000, GH200, future chips) goes in the model's own `config.hardware`\n(`{id,label,vram,vendor}`) — the engine merges it in; don't edit the engine catalog:\n\n| Platform | Vendor | VRAM | Docker image |\n|---|---|---|---|\n| H100 | NVIDIA | 80GB | `lmsysorg/sglang:` |\n| H200 | NVIDIA | 141GB | `lmsysorg/sglang:` |\n| B200 | NVIDIA | 192GB | `lmsysorg/sglang:` |\n| B300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130` when required) |\n| GB200 | NVIDIA | 192GB | `lmsysorg/sglang:` (or `-cu130`) |\n| GB300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130`) |\n| DGX Spark | NVIDIA | 128GB (unified) | `lmsysorg/sglang:` — needs a **CUDA 13** build |\n| MI300X | AMD | 192GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI325X | AMD | 256GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI350X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n| MI355X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n\n- **Image tag (``)**: don't guess — ask the user for the tag the recipes ran on, or\n default to `dev` (nightly). The same tag goes in `dockerImages` and benchmarks'\n `sglang_version`; the engine falls back to `lmsysorg/sglang:dev` for any unmapped hw.\n- **TP sizing** (sanity-check recipes): `weight_GB / gpu_mem`, round up to a power of 2,\n ~20–30% headroom. BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5. MoE → **total** weight, not\n active params. FP4 is Blackwell-only (B200/B300/GB200/GB300/DGX Spark). GB200/GB300\n single-node hosts are typically **4 GPUs** (TP=4 ceiling); a DGX Spark is **1 GPU**, so\n its only multi-GPU topology is TP=2 across 2 nodes.\n- **Platform flags**: Blackwell may need `--attention-backend trtllm_mha`; AMD typically\n needs `--attention-backend triton` + env `SGLANG_USE_AITER=1` /\n `SGLANG_ROCM_FUSED_DECODE_MLA=0` (check AITER TP constraints, e.g. `heads_per_gpu % 16 == 0`).\n- **EP** (MoE): 8-GPU NVIDIA `--tp 8 --ep 8`; AMD `EP = TP`; small NVIDIA (TP≤4) omit\n `--ep` unless benchmarked. (The template's AMD example cell shows these.)\n\n## Phase 2 — Instantiate the template\n\nFor a diffusion model, instantiate `templates/diffusion-page.mdx.tmpl` and keep the\nshared `DiffusionModelTags` component, plus `templates/diffusion-config.jsx.tmpl` for the\nopt-in scoped command builder. Put the compact install command and builder in §1 Quick\nstart. The first two paragraphs in §2 Model capabilities are not generic filler: they must\nstate the model's capability range, strongest differentiator, when to choose it, and at\nleast one real deployment or capability boundary. Put orthogonal runtime features in\n`scope: \"serve\"` or `scope: \"request\"`, not in the base recipe; use the schema from\n`references/diffusion-authoring.md`.\n\n1. **Copy** the three template files to their target paths (above). Note the two\n vendor-folder conventions: under `configs/` the folder is the **HuggingFace org**\n (`deepseek-ai`); under `cookbook/` it's the **display vendor** (`DeepSeek`).\n2. **Replace every `__TOKEN__`**: `__MODEL_DISPLAY__`, `__MODEL_SLUG__`, `__HF_ORG__`,\n `__HF_REPO__`, `__REASONING_PARSER__`, `__TOOLCALL_PARSER__`, `__ONE_LINER__`. Verify\n none remain: `grep -rn '__[A-Z_]*__' `.\n3. **Prune** to what the model supports (delete, don't stub) — using\n [references/authoring-reference.md](references/authoring-reference.md):\n - `supportedHardware` + the EXAMPLE cells: keep your tested families; **delete the\n `mi*` ids + AMD example cell if no AMD recipe**, etc. A GPU not in the shared catalog\n (e.g. RTX PRO 6000) → declare it in `config.hardware` and add its id here.\n - `playgroundFeatures` axes: remove the `megamoe` backend option + the\n `megamoeQuant` block from the `moe` axis (non-Blackwell-MoE), delete `hisparse`\n (non-DSA), `pdDisagg`/`router` (no PD), the `parsers` axis (no parsers), etc.\n - `quantizations` / `variants`: drop what the model doesn't ship; collapse `variants`\n to single `default` if there's no variant axis (then drop the `variant` half of\n `modelNames`/`defaultAccuracy` keys).\n4. **Fill `cells[]`** with the verified recipes from Phase 1 (replace every EXAMPLE cell;\n set `verified: true` only on tested combos), and `modelNames` with real HF slugs,\n `dockerImages` for your hw (use the Phase-1 tag, or default `lmsysorg/sglang:dev` — never\n a guessed release; key by `hw`, or a more specific key — `hw|variant|quant`, `variant|quant`,\n `hw|quant|strategy`, `hw|quant`, resolved in that order before `hw` — when one\n variant/quant/strategy needs its own image), `multiNodeHints` only for fabric-specific hw\n (e.g. gb200).\n\n5. **Diffusion pages: add the ComfyUI section.** Every diffusion cookbook page ends with\n a `## . Run in ComfyUI` section so a reader never has to guess whether the model is\n reachable from ComfyUI. It is one component; the per-model facts live in the component,\n not the page:\n\n ```mdx\n ## . Run in ComfyUI\n\n import { ComfyUISupport } from '/src/snippets/diffusion/comfyui-support.jsx';\n\n \" />\n ```\n\n Pick `model` from the table in\n `docs/src/snippets/diffusion/comfyui-support.jsx`. Use the model's own key when it has\n an entry (its executor or dedicated node differs); otherwise use the generic `image` or\n `video`. A model gets its own key only when the plugin actually treats it specially —\n an entry in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a\n dedicated node. Adding a key without the matching plugin support makes the page lie.\n\n### Site-wiring (do all three)\n\n- **`docs/docs.json`** — add the page under Cookbook → `` → ``, at\n the **top** of that vendor's `pages` (root-relative, no `.mdx`:\n `cookbook///`). New vendor group → insert in the section's\n local ordering.\n- **NEW-tag hygiene** — the new page keeps `tag: NEW` (from the template). Scan the\n vendor dir for existing NEW and strip it from siblings; verify ≤1:\n `grep -rn 'tag: NEW' docs/cookbook///` → at most one result. (Scan\n files; don't assume the first `docs.json` entry holds NEW.)\n- **Homepage card** — `docs/cookbook//intro.mdx`: if the org already has a\n ``, update only its `href` (keep `img`). If the org is **new**, add a ``\n (title = nav-group name; keep card order aligned with `docs.json`) **and create its logo**:\n ask the user for the brand logo, then generate the conforming **icon-only 940×525 RGBA\n transparent** PNG → `docs/cards/logos/.png` per\n [references/vendor-logo.md](references/vendor-logo.md) (track with `git add -f` — `*.png`\n is gitignored repo-wide). Never invent or copy a logo.\n\n## Phase 3 — Validate\n\n```bash\ncd docs\nmint validate # frontmatter, missing nav entries, MDX/JSX errors\nmint broken-links\nmint dev # visual smoke test at http://localhost:3000/cookbook///\n```\n\nSpot-check: cells render sensible commands; URL-hash nav persists across reload; the\nPlayground inherits the Deploy selection live; each axis toggle produces the expected\ndiff; Docker mode wraps in `docker run` with the right image; multi-node cells emit the\nhints + `--nnodes N`; cURL resolves the model name; the NEW badge shows on the new page\nand is gone from same-vendor siblings; the homepage card points to the new model.\n\n## Phase 4 — Interactive testing\n\nThe user deploys each cell, runs the benches, and pastes results; you fill the data:\n- mark each tested `cells[]` entry `verified: true` (absent = yellow/unverified badge);\n- fill the `-benchmarks.jsx` entries (one per cell `match`) with measured\n speed/accuracy + the `sglang_version` the user reports (don't invent one — the template's\n `0.0.0` is a deliberate TODO); set model-level `defaultAccuracy` per variant. Leave a\n cell's entry as a bare `match` stub if it has no numbers yet (the card shows \"pending\").\n\n## Phase 5 — Prose & config tips\n\n**Read [references/mintlify-authoring.md](references/mintlify-authoring.md) first** (it\ncarries the parser-output-shape / thinking-mode / Output-Example / no-hardcoded-sampling\nrules + the Mintlify forbidden-syntax list). Then rewrite the MDX prose from the HF card +\nuser notes: §1 Model Introduction (description, links, params, license, variants table),\n§2 Configuration Tips (hw-specific tuning, caveats), §3 Advanced Usage (Reasoning /\nTool-Calling / HiCache — keep only what applies; match the reasoning example to the\nparser's output shape; each runnable block gets an `**Output Example:**`).\n\nFor diffusion pages, follow the category-specific Quick start and capability contract in\n`references/diffusion-authoring.md`. Run `node docs/scripts/check_cookbook_configs.mjs` to\nverify the tag widget and introduction structure before rendering the page.\n\n## Phase 6 — Review\n\n```\n/cookbook-review-pr \n```\n\n## Git workflow\n\nAlways branch — never commit to main directly.\n\n```bash\ngit checkout -b add--cookbook\ngit add docs/src/snippets/configs//.jsx \\\n docs/src/snippets/configs//-benchmarks.jsx \\\n docs/cookbook///.mdx \\\n docs/docs.json docs/cookbook//intro.mdx\ngit commit -m \"Add cookbook\"\ngit push -u origin add--cookbook\ngh pr create --title \"[Doc] cookbook: Add cookbook\" --body \"...\"\n```\n",".agents/skills/ci-workflow-guide/SKILL.md":"---\nname: ci-workflow-guide\ndescription: Guide to SGLang CI workflow orchestration — stage ordering, fail-fast, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages.\n---\n\n# SGLang CI Workflow Orchestration Guide\n\nThis skill covers the CI **infrastructure** layer — how tests are dispatched, gated, and aborted on failure across stages. For test authoring (templates, fixtures, registration, model selection), see the [write-sglang-test skill](../write-sglang-test/SKILL.md).\n\n---\n\n## Naming Conventions\n\n- **Suite**: `base-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `base-b-test-1-gpu-small`)\n- **Test group**: Directory-level registered test group under `test/registered/` (e.g., `hicache` maps to `test/registered/hicache/test_*.py`)\n- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)\n\n---\n\n## Key Files\n\n| File | Role |\n|------|------|\n| `.github/workflows/pr-test.yml` | Main workflow — all stages, jobs, conditions, matrix definitions |\n| `.github/workflows/pr-test-extra.yml` | Extra workflow — gated by BOTH `run-ci` and `run-ci-extra` labels |\n| `.github/workflows/pr-gate.yml` | PR gating: draft check, `run-ci` label, per-user rate limiting |\n| `.github/actions/check-pr-test-health/action.yml` | Cross-job fail-fast: queries API for any failed job |\n| `.github/actions/wait-for-jobs/action.yml` | Stage gating: polls API until stage jobs complete |\n| `.github/actions/check-maintenance/action.yml` | Maintenance mode check |\n| `.github/scripts/ci-labels.cjs` | Resolves the four CI control labels into dispatch axes |\n| `test/run_suite.py` | Suite runner: collects, filters, partitions, executes tests |\n| `python/sglang/test/ci/ci_register.py` | Test registration (AST-parsed markers), LPT auto-partition |\n| `python/sglang/test/ci/ci_utils.py` | `run_unittest_files()`: execution, retry, continue-on-error |\n| `scripts/ci/utils/slash_command_handler.py` | Handles slash commands from PR comments |\n\n---\n\n## Architecture Overview\n\n```\n ┌──────────────┐\n │ build kernel │\n └──────┬───────┘\n │\n ├─ check-changes ──── detects which packages changed\n │ (main_package, sgl_kernel, jit_kernel, multimodal_gen)\n │\n ├─ call-gate ──────── pr-gate.yml (draft? label? rate limit?)\n │\n ├─────────────────────────────────────────────────────┐\n │ │\n ▼ │\n ┌─────────────────────────────────────┐ │\n │ Base A (~3 min) │ │\n │ pre-flight check │ │\n │ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-1-gpu-small │ │ │\n │ │ (small GPUs) │ │ │\n │ └─────────────────────────────┘ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-cpu │ │ │\n │ │ (CPU) │ │ │\n │ └─────────────────────────────┘ │ │\n └──────┬──────────────────────────────┘ │\n │ │\n ▼ ▼\n ┌─────────────────────────────────────┐ ┌──────────────────────────┐\n │ Base B (~30 min) │ │ kernel test │\n │ base tests │ └──────────────────────────┘\n │ │ ┌──────────────────────────┐\n │ ┌─────────────────────────────┐ │ │ multimodal gen test │\n │ │ base-b-test-1-gpu-small │ │ └──────────────────────────┘\n │ │ (small GPUs, e.g. 5090) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-1-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-2-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ Base C (~30 min) │\n │ advanced tests │\n │ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-h100 │ │\n │ │ (H100 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-8-gpu-h200 │ │\n │ │ (8 x H200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-b200 │ │\n │ │ (4 x B200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ Other advanced tests │ │\n │ │ (DeepEP, PD Disagg, GB300) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ pr-test-finish │\n │ aggregates all results, fails if │\n │ any job failed/cancelled │\n └─────────────────────────────────────┘\n```\n\n**Every stage test job** includes a `check-pr-test-health` step after checkout — if any job in the run has already failed, the job fails fast (red X) with a root cause annotation.\n\n**Scheduled runs** skip `wait-for-base-*` jobs, running all stages in parallel. Fail-fast is also disabled.\n\n---\n\n## Fail-Fast Layers\n\n4 layers of fail-fast, from fine to coarse:\n\n| Layer | Mechanism | Granularity | Disabled on schedule? |\n|-------|-----------|-------------|----------------------|\n| **1. Test method → file** | `unittest -f` (failfast) | One test method fails → entire test file stops immediately | Yes |\n| **2. File → suite** | `run_unittest_files()` default | One test file fails → entire suite stops (`--continue-on-error` off) | Yes |\n| **3. Job → job (same stage)** | `check-pr-test-health` action | One job fails → other waiting jobs in same stage fail-fast (red X) | Yes |\n| **4. Stage → stage (cross-stage)** | `wait-for-base-*` + `needs` | Base A fails → base B/C jobs skip entirely (never get a runner) | Yes (wait jobs skipped) |\n\n- **Layer 1**: `-f` flag appended to all `python3 -m pytest` / `unittest` invocations in `ci_utils.py`\n- **Layer 2**: `--continue-on-error` flag in `run_suite.py` — off for PRs, on for scheduled runs\n- **Layer 3**: `check-pr-test-health` auto-detects `schedule` event and skips; filters out cascade failures to show only root cause jobs\n- **Layer 4**: `wait-for-base-*` jobs are conditioned on `github.event_name == 'pull_request'` — skipped for scheduled runs\n\n---\n\n## Execution Modes\n\n| Aspect | PR (`pull_request`) | Scheduled (`cron`, every 6h) | Manual dispatch (`workflow_dispatch`) |\n|--------|---------------------|------------------------------|--------------------------------------|\n| **Stage ordering** | Sequential: A → B → C via `wait-for-base-*` | Parallel (all at once) | Single target stage only |\n| **Cross-job fail-fast** | Yes (`check-pr-test-health`) | Yes | Yes |\n| **continue-on-error** | No (stop at first failure within suite) | Yes (run all tests) | No |\n| **Retry** | Enabled | Enabled | Enabled |\n| **max_parallel** | 3 (default), 14 if `max-concurrency` label | 14 | 3 (default), 14 if `max-concurrency` |\n| **PR gate** | Yes (draft, label, rate limit) | Skipped | Skipped |\n| **Concurrency** | `cancel-in-progress: true` per PR | Queue (no cancel) | Isolated per stage+SHA |\n\nFour labels relax these limits for one PR: `bypass-fail-fast`, `parallel-stages`,\n`max-concurrency`, and `highest-priority` (all three). `.github/scripts/ci-labels.cjs`\nresolves them; the [contribution guide](https://docs.sglang.io/developer_guide/contribution_guide.html#ci-control-labels)\ndescribes what each one does.\n\n---\n\n## Stage Gating (`wait-for-jobs` action)\n\n`wait-for-base-a` and `wait-for-base-b` are lightweight `ubuntu-latest` jobs that poll the GitHub Actions API.\n\n**How it works:**\n1. Calls `listJobsForWorkflowRun` to list all jobs in the current run\n2. Matches jobs by exact name or prefix (for matrix jobs, e.g., `base-b-test-1-gpu-small (3)`)\n3. If any matched job has `conclusion === 'failure'` → fail immediately (fail-fast)\n4. If all matched jobs are completed and count matches `expected_count` → success\n5. Otherwise → sleep `poll-interval-seconds` (default: 60s) and retry\n6. Timeout after `max-wait-minutes` of wall-clock time (240 min for base-a, 350 min for base-b). Keep budgets under the 6h hard cap on GitHub-hosted runners: a wait killed by the cap ends `cancelled`, not `timeout`, and `timeout-minutes` cannot raise it. The base-b wait sets `fail-on-timeout: 'false'`, so a timeout lets base-c start; `pr-test-finish` fails the run on any base-b failure regardless\n\n**Job specs example** (base-b):\n```json\n[\n {\"prefix\": \"base-b-test-1-gpu-small\", \"expected_count\": 8},\n {\"prefix\": \"base-b-test-1-gpu-large\", \"expected_count\": 14},\n {\"prefix\": \"base-b-test-2-gpu-large\", \"expected_count\": 4},\n {\"prefix\": \"base-b-test-4-gpu-b200\", \"expected_count\": 1}\n]\n```\n\n> **Critical**: `expected_count` must match the matrix size. If you add/remove matrix entries, update the wait job's spec accordingly.\n\n**PR only**: Condition `github.event_name == 'pull_request' && !inputs.target_stage` — scheduled runs and manual dispatches skip these entirely, allowing parallel execution.\n\n---\n\n## Cross-Job Fail-Fast (`check-pr-test-health` action)\n\nComposite action called after checkout in every stage test job (21 jobs total across `pr-test.yml`, `pr-test-multimodal-gen.yml`, `pr-test-sgl-kernel.yml`, `pr-test-jit-kernel.yml`).\n\n**How it works:**\n1. Queries `listJobsForWorkflowRun` for the current workflow run\n2. Filters for **root cause failures only** — jobs with `conclusion === 'failure'` whose failing step is NOT `check-pr-test-health` (excludes cascade failures)\n3. If root cause failures found → calls `core.setFailed()` with the list of root cause job names\n4. If none → does nothing (step succeeds)\n\n**Cascade filtering**: When job A fails fast due to the health check, it also has `conclusion: failure`. Without filtering, job B would list both the original failure AND job A's fail-fast. The filter checks each failed job's `steps` array — if the failing step name contains `check-pr-test-health` or `Check PR test health`, it's excluded from the root cause list.\n\n**Usage pattern:**\n```yaml\nsteps:\n - name: Checkout code\n uses: actions/checkout@v4\n ...\n\n - uses: ./.github/actions/check-pr-test-health\n id: pr-test-health\n\n - name: Install dependencies # skipped automatically if health check failed\n ... # (default if: success() is false)\n\n - name: Run test # also skipped\n ...\n```\n\n**Visual effect**: Job shows **red X** (failure) with error annotation showing root cause job names. Subsequent steps are naturally skipped (default `if: success()` is false after a failed step). No per-step `if` guards needed.\n\n**No stage filtering**: Checks ALL jobs in the run, not just the current stage. Any failure anywhere triggers fail-fast.\n\n**Error message example:**\n```\nFail-fast: skipping — root cause job(s): base-b-test-1-gpu-small (0), base-b-test-1-gpu-small (1)\n```\n\n---\n\n## Within-Suite Failure Handling\n\nControlled by `run_unittest_files()` in `python/sglang/test/ci/ci_utils.py`.\n\n### Flags\n\n| Flag | PR default | Scheduled default | Effect |\n|------|------------|-------------------|--------|\n| `--continue-on-error` | Off | On | Off: stop at first failure. On: run all files, report all failures at end |\n| `--enable-retry` | On | On | Retry retriable failures (accuracy/perf assertions) |\n| `--max-attempts` | 2 | 2 | Max attempts per file including initial run |\n\n### Retry Classification\n\nWhen a test fails and retry is enabled, the output is classified:\n\n**Non-retriable** (checked first — real code errors):\n`SyntaxError`, `ImportError`, `ModuleNotFoundError`, `NameError`, `TypeError`, `AttributeError`, `RuntimeError`, `CUDA out of memory`, `OOM`, `Segmentation fault`, `core dumped`, `ConnectionRefusedError`, `FileNotFoundError`\n\n**Retriable** (accuracy/performance):\n`AssertionError` with comparison patterns (`not greater than`, `not less than`, `not equal to`), `accuracy`, `score`, `latency`, `throughput`, `timeout`\n\n**Default**: Unknown `AssertionError` → retriable. Other unknown failures → not retriable.\n\n### How `continue_on_error` is set\n\nIn `pr-test.yml`'s `check-changes` job:\n- `schedule` runs or `run_all_tests` flag → `continue_on_error = 'true'`\n- PR runs → `continue_on_error = 'false'`\n\nEach test job propagates via:\n```yaml\nenv:\n CONTINUE_ON_ERROR_FLAG: ${{ needs.check-changes.outputs.continue_on_error == 'true' && '--continue-on-error' || '' }}\nrun: |\n python3 run_suite.py --hw cuda --suite $CONTINUE_ON_ERROR_FLAG\n```\n\n---\n\n## Test Partitioning\n\nLarge suites are split across matrix jobs using the **LPT (Longest Processing Time) heuristic** in `ci_register.py:auto_partition()`:\n\n1. Sort tests by `est_time` descending, filename as tie-breaker (deterministic)\n2. Greedily assign each test to the partition with smallest cumulative time\n3. Result: roughly equal total time per partition\n\n**Partition table** (CUDA per-commit suites):\n\n| Suite | Partitions | Runner | max_parallel |\n|-------|-----------|--------|-------------|\n| `base-a-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — |\n| `base-a-test-cpu` | 4 | `ubuntu-latest` | — |\n| `base-b-test-1-gpu-small` | 8 | `1-gpu-5090` | 8 |\n| `base-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) |\n| `base-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — |\n| `base-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-b-kernel-unit-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |\n| `base-b-kernel-benchmark-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — |\n| `base-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — |\n| `base-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — |\n| `base-c-test-deepep-4-gpu-h100` | 1 (no matrix) | `4-gpu-h100` | — |\n| `base-c-test-4-gpu-b200` | 3 | `4-gpu-b200` | — |\n| `base-c-test-4-gpu-b200-small` | 3 | `4-gpu-b200-low-disk` | — |\n| `base-c-test-8-gpu-b200` | registered only | `8-gpu-b200` | — |\n| `base-c-test-4-gpu-gb200` | registered only | `4-gpu-gb200` | — |\n\n> **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`.\n>\n> **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`.\n\n**Workflow usage:**\n```yaml\nstrategy:\n matrix:\n partition: [0, 1, 2, 3, 4, 5, 6, 7]\nsteps:\n - run: python3 run_suite.py --hw cuda --suite base-b-test-1-gpu-small \\\n --auto-partition-id ${{ matrix.partition }} --auto-partition-size 8\n```\n\n---\n\n## check-changes Job\n\nDetermines which test suites to run based on file changes.\n\n### Detection Methods\n\n| Trigger | Method | Details |\n|---------|--------|---------|\n| `pull_request` | `dorny/paths-filter` | Detects changes via GitHub diff |\n| `workflow_dispatch` (with `pr_head_sha`) | GitHub API | `repos/{repo}/compare/main...{sha}` |\n| `schedule` / `run_all_tests` | Force all true | Runs everything |\n\n### Output Flags\n\n| Output | Triggers |\n|--------|----------|\n| `main_package` | Base A/B/C test suites |\n| `sgl_kernel` | Kernel wheel builds + kernel test suites; also switches B200 jobs to kernel-build runner labels outside `target_stage` mode |\n| `jit_kernel` | JIT kernel test workflow |\n| `multimodal_gen` | Multimodal-gen test workflow |\n\n> **Note**: In `target_stage` mode, `sgl_kernel` is only active when `include_wheel_build=true`. Without that opt-in, kernel-change reruns fail validation instead of running a target stage without freshly built wheels. Outside `target_stage`, `sgl_kernel=true` switches B200 jobs from `4-gpu-b200` / `4-gpu-b200-low-disk` to `4-gpu-b200-kernel` / `4-gpu-b200-kernel-low-disk`.\n\n---\n\n## Concurrency Control\n\n```\ngroup: pr-test-{event_name}-{branch}-{pr_sha}-{stage}\n```\n\n| Segment | Source | Purpose |\n|---------|--------|---------|\n| `event_name` | `github.event_name` | Prevents scheduled runs colliding with fork PRs named `main` |\n| `branch` | `github.head_ref \\|\\| github.ref_name` | Per-branch isolation |\n| `pr_sha` | `inputs.pr_head_sha \\|\\| 'current'` | Isolates manual dispatches from main runs |\n| `stage` | `inputs.target_stage \\|\\| 'all'` | Allows parallel stage dispatches |\n\n`cancel-in-progress: true` for `pull_request` events (new push cancels old run), `false` for `workflow_call`.\n\n---\n\n## How To: Add a New Stage Job\n\n1. Define the job in `pr-test.yml` with `needs: [check-changes, call-gate, wait-for-base-X, ...]`\n2. Copy the `if:` condition pattern from an existing same-stage job (handles `target_stage`, `schedule`, `main_package`)\n3. Add `checkout` step\n4. Add `check-pr-test-health` step (after checkout) — if any prior job failed, `core.setFailed()` fires and all subsequent steps auto-skip via default `if: success()`\n5. Add `check-maintenance` step\n6. Add `download-artifact` step if `sgl_kernel` changed\n7. Add `install dependencies` step\n8. Add `run test` step with `$CONTINUE_ON_ERROR_FLAG`\n9. Add `upload-cuda-coredumps` step with `if: always()`\n10. Register the suite name in `PER_COMMIT_SUITES` in `test/run_suite.py`\n11. If using matrix, add `--auto-partition-id` and `--auto-partition-size` to the run command\n12. **Update `wait-for-base-X`** job spec with the new job name and `expected_count` (if matrix)\n13. **Add the job to `pr-test-finish.needs`** list\n\n---\n\n## How To: Debug CI Failures\n\n| Symptom | Likely cause | What to check |\n|---------|-------------|---------------|\n| All stage-B/C jobs green but steps skipped | Earlier job failed, `check-pr-test-health` triggered | Find the actual failed job (red X) |\n| `wait-for-base-b` timeout | `expected_count` doesn't match matrix size | Verify job spec counts match `matrix:` array length |\n| `pr-test-finish` fails but all jobs green | A job was `cancelled` (counts as failure in finish) | Check concurrency cancellation, and whether a `wait-for-base-*` job hit the 6h hosted-runner cap |\n| Tests pass locally but fail in CI | Partition assignment, runner GPU type, or `est_time` inaccuracy | Check which partition the test lands in; verify runner label |\n| Flaky test retried and passed | Retriable failure (accuracy/perf) | Check `[CI Retry]` markers in job logs |\n| Flaky test NOT retried | Matched non-retriable pattern | Check if error matches `NON_RETRIABLE_PATTERNS` in `ci_utils.py` |\n\n---\n\n## Slash Commands\n\n| Command | Effect |\n|---------|--------|\n| `/tag-run-ci-label` | Adds `run-ci` label to PR |\n| `/tag-run-ci-label extra` | Adds both `run-ci` and `run-ci-extra` labels |\n| `/rerun-failed-ci` | Reruns failed jobs in the latest workflow run |\n| `/tag-and-rerun-ci` | Adds `run-ci` label + reruns failed |\n| `/tag-and-rerun-ci extra` | Adds both `run-ci` and `run-ci-extra` labels + reruns failed |\n| `/run-full-ci` | Short form of `/tag-and-rerun-ci extra` (baseline + extra). Alias: `/rerun-full-ci` |\n| `/run-extra-ci` | Adds both labels + reruns **only** `PR Test Extra`, leaving baseline runs alone. Alias: `/rerun-extra-ci` |\n| `/rerun-test [ ...]` | Reruns specific test file(s) via `rerun-test.yml`. A file arg containing a glob metacharacter (`*`, `?`, `[...]`) expands against `test/registered/` and the multimodal test dir to every matching `test_*.py` (e.g. `/rerun-test test_*backend*.py` — wrap in backticks so GitHub doesn't italicize the `*`); matches are deduped, grouped by dispatch shape, and can't carry a `::test` selector. No match → single ⛔ reply, nothing dispatched. |\n| `/rerun-group [ ...]` | Expands registered test groups, then reuses `/rerun-test` |\n| `/run-cpu-ci` | Dispatches `run-cpu-ci.yml` on main: only the `base-a-test-cpu` suite, at the PR head SHA, no caches saved. Same gate as `/rerun-test`. Alias: `/rerun-cpu-ci` |\n\nHandled by `scripts/ci/utils/slash_command_handler.py` → `.github/workflows/slash-command-handler.yml`.\n\n### Label-gated workflow dispatch (pr-test, pr-test-extra)\n\n`pr-test.yml` and `pr-test-extra.yml` both listen for `pull_request.labeled` (in addition to `opened`/`synchronize`/`reopened`). The `check-changes.if` gate has two clauses:\n\n1. **For `labeled` events**: the just-added label must be one of the gating labels (`run-ci` for pr-test, `run-ci` or `run-ci-extra` for pr-test-extra) — otherwise every unrelated label addition would dispatch a full CI run.\n2. **All events**: the PR must currently carry the required labels.\n\nThis is what lets `/tag-run-ci-label` (and the `extra` variant) trigger a fresh CI run without an extra push.\n\n**Caveat — skipped runs cannot be un-skipped by `run.rerun()`:** GitHub's rerun API reuses the original event payload, so rerunning a `pull_request`-event run that was skipped because of missing labels will skip again (label set in the frozen payload doesn't update). The only way to recover a label-skipped run is to add the missing label, which fires a fresh `labeled` event with the current label set. `handle_rerun_failed_ci` in the slash handler is for rerunning failed/non-label-skipped runs; it cannot revive label-skipped ones.\n",".agents/skills/add-jit-kernel/SKILL.md":"---\nname: add-jit-kernel\ndescription: Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups\n---\n\n# Tutorial: Adding a New JIT Kernel to SGLang\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as a JIT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float, passed at runtime)\n- Output: `x * factor` (element-wise), allocated internally\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n\n## When to use JIT vs AOT (`sgl-kernel`)\n\n- **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.\n- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `python/sglang/kernels/aot/` and participate in the wheel build / torch op registration flow.\n- **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.\n\n---\n\n## Conventions\n\nThese hold for every step below.\n\n- **`namespace sglang` is where JIT code lives.** Open it after the include block and close it at the end of the file, with the device kernels, traits and host wrapper inside. The shared `host::` / `device::` helpers are nested in it too, so they resolve unqualified. `load_jit` emits the `TVM_FFI_DLL_EXPORT_TYPED_FUNC` wrapper inside `namespace sglang` as well, so the `kernel_name` you pass from Python needs no `sglang::` prefix.\n- **Tuning policy belongs in Python; C++ gets the answer, not the decision.** For a hyperparameter — split factor, threads per item, block size, vector width, an algorithm variant — prefer (not required) making it a **template parameter** and letting the `@cache_once` module factory choose the value, over a runtime `if`/`switch` in the launcher that picks among pre-instantiated kernels. The kernel then compiles exactly one specialisation and asserts its own preconditions with `static_assert`, while the heuristic that produced the value sits in Python where it is readable, adjustable, and inspectable without recompiling CUDA. `store_cache`'s `num_threads` is the worked example: a `get_kernel(num_split)` ladder guarded by byte-alignment `if constexpr`s became one template argument plus a heuristic in `_jit_kvcache_module`. The exception is a knob that genuinely varies per call with the runtime shape — that has to stay a kernel argument.\n- **Check where the check is cheapest: `static_assert` > C++ host check > cached Python > per-call Python.** Anything fixed at compile time is a `static_assert`. Anything about the tensors is a `TensorMatcher` / `CHECK_HOST` in the C++ launcher, free next to a kernel launch. A check Python cannot delegate goes inside the `@cache_once` module factory, where it runs once per specialisation. What remains in the per-call entry point costs interpreter time on *every* forward, so it should be nothing but picking the module and allocating `out`.\n- **Fixed-width integer types.** Prefer `int32_t` / `int64_t` / `uint32_t` / `size_t` over `int`, `long`, or `long long`, so an index has the same width on both sides of the FFI boundary. Bare `int` is fine only where the width plainly cannot matter — an unrolled loop counter over a `constexpr` bound, a template `int` parameter. Shapes arrive as `int64_t` (`SymbolicSize::unwrap()`); narrowing to `uint32_t` for in-kernel indexing is a deliberate act, so write the `static_cast` explicitly and only where the range is known.\n- **Doxygen comments in C++.** Document exported entities with `///` or `/** ... */` blocks using `\\brief`, `\\param`, `\\tparam`, `\\return`, the way `include/sgl_kernel/` does. `python -m sglang.kernels.jit` writes `CommentFormat: Doxygen` into `.clangd` when clangd is 21 or newer, so these render on hover in the editor. Plain `//` remains fine for implementation notes inside a function body.\n- **ASCII only in C++ and CUDA sources.** Write `--`, `->`, `<=` instead of `—`, `→`, `≤`, including in comments. `grep -nP '[^\\x00-\\x7F]' ` before committing.\n- **`namespace details` is private.** Anything inside one is an implementation detail of its own header, free to change without notice. Do not name `details::` from another module, and never `using namespace details`. If you find yourself wanting something in there, that is the signal to promote it to a documented name instead.\n- **`const T* __restrict__` for read-only pointers.** This is what `csrc/` does throughout, and it lets the compiler emit non-coherent (`LDG`) loads.\n- **Watch the register budget.** For memory-bound kernels, keep to roughly 64 registers per thread so occupancy does not become the limit; prefer recomputing a value over letting it spill. Run `SGLANG_JIT_LOG_RESOURCE_USAGE=1 SGLANG_JIT_FORCE_RECOMPILE=1` to log per-kernel registers, spills and shared memory at INFO. Both halves are load-bearing: a cache hit has no compiler output to report, and passing `-Xptxas -v` by hand does nothing on its own because the build captures compiler output and replays it only on failure.\n- **Alignment is a contract you enforce, not a property you hope for.** A vectorized copy derives its access width from the row *size* (`load_bytes` picks 16B for a 1024B row), but the address it lands on is `base + index * row_stride` — and a stride is not constrained by the size. `TensorMatcher` admits any stride unless you say otherwise, so a padded cache row silently produces `misaligned address`. Every kernel that vectorizes is highly recommended to end its validation with `.ensure_alignment(w)`, where `w` is the width the kernel actually uses.\n\n---\n\n## Common Abstractions in `python/sglang/kernels/jit/include/sgl_kernel/`\n\n**Always prefer these abstractions over raw CUDA primitives.** They provide safety, readability, and consistency with the rest of the codebase. The only reason to drop to raw primitives is performance the abstraction cannot reach — a trade you make deliberately, and justify in a comment.\n\n### `utils.h` — Host-side utilities\n\n```cpp\n#include \n```\n\n- **`CHECK_HOST(cond) << \"msg \" << value`** — **Preferred** runtime check: stream-style, throws `PanicError` with file/line info on failure. Zero overhead on the true path — the message expressions are only evaluated when the check fails.\n- **`host::RuntimeCheck(cond, args...)`** — Function-style alternative to `CHECK_HOST`. Note its message args are always evaluated (even when the check passes), so prefer `CHECK_HOST` — especially on hot paths.\n- **`host::Panic(args...)`** — Unconditionally throw a `PanicError` with a descriptive message.\n- **`host::div_ceil(a, b)`** — Integer ceiling division `(a + b - 1) / b`.\n- **`host::irange(n)`** / **`host::irange(start, end)`** — Range views for cleaner loops.\n- **`host::pointer::offset(ptr, offsets...)`** — Byte-safe pointer arithmetic on `void*`. Use this instead of raw casts.\n\n### `utils.cuh` — Device-side utilities + `LaunchKernel`\n\n```cpp\n#include \n```\n\n- **Type aliases**: `fp16_t`, `bf16_t`, `fp32_t`, `fp8_e4m3_t`, `fp8_e5m2_t` and their packed variants `fp16x2_t`, `bf16x2_t`, `fp32x2_t`, etc.\n- **`SGL_DEVICE`** — Expands to `__forceinline__ __device__`. Use on all device functions.\n- **`SGL_DEVICE_HOST`** — `__forceinline__ __device__ __host__`. Use it when a `constexpr` helper must be callable from both the kernel and its launcher, so the two cannot drift.\n- **`device::kWarpThreads`** — Constant `32` (the codebase's *logical* warp width). `kWarpSize` is an alias. Note this is a group size, not the hardware wave: on gfx950 `warpSize` is 64, and `warp::kFullWidth` reflects that.\n- **`device::get_lane_id()`** — this thread's index within its logical group. **Prefer it over `threadIdx.x % kNumThreads` whenever the value feeds an address.** On CUDA the lane register is one read that folds into the address computation, whereas the modulo is a derived value the compiler re-materialises at *every* address scale — 8 instructions and 2 registers on a two-tile warp copy, measured on sm_100a. It picks the cheaper form per platform, so callers never need to care. Only equals the true in-warp lane when `blockDim.x` is a multiple of `kNumThreads`.\n- **`device::load_as(ptr, offset)`** / **`device::store_as(ptr, val, offset)`** — Type-safe loads/stores from `void*`.\n- **`device::pointer::offset(ptr, offsets...)`** — Pointer arithmetic on device.\n- **`host::LaunchKernel(grid, block, device_or_stream [, smem])`** — RAII kernel launcher that:\n - Resolves the CUDA stream from a `DLDevice` via TVM-FFI automatically.\n - Checks the CUDA error with file/line info after launch via `operator()(kernel, args...)`.\n - Supports `.enable_pdl(bool)` for PDL (Programmatic Dependent Launch, SM90+).\n- **`device::PDLWaitPrimary()`** / **`device::PDLTriggerSecondary()`** — The two halves of PDL, on sm_90+ (no-ops on older archs and ROCm). Their guarantees are **not** symmetric:\n - `PDLTriggerSecondary` (`griddepcontrol.launch_dependents`) only lets the next kernel in the stream *start* early. It carries no memory ordering and publishes nothing — matching that, the header's asm has no `\"memory\"` clobber.\n - `PDLWaitPrimary` (`griddepcontrol.wait`) is the ordering point: it waits until the preceding kernel has fully finished and its writes are visible.\n\n So every read of data the preceding kernel produced must come after `PDLWaitPrimary()`. What overlaps with the primary's tail is whatever you put *before* the wait — loading parameters, computing indices, touching buffers the primary never wrote — so a kernel that waits on its first line gains nothing. Neither call is a barrier: threads may reach or skip them independently. See \"Programmatic Dependent Launch and Synchronization\" in the CUDA C++ Programming Guide.\n- **`CHECK_CUDA(expr) << \"context\"`** — Stream-style CUDA error check; evaluates `expr` once and throws `PanicError` with `cudaGetErrorString` + file/line info if it is not `cudaSuccess`. Extra streamed context is optional.\n- **`host::RuntimeDeviceCheck(cudaError_t)`** — Function-style alternative to `CHECK_CUDA`. It takes no context message, so prefer `CHECK_CUDA`, which builds its error object only on the failure path.\n\n### `tensor.h` — Tensor validation (`TensorMatcher`, Symbolic types)\n\n```cpp\n#include \n```\n\nThis is the **primary validation API** for all kernel launchers. Use it to validate every `tvm::ffi::TensorView` argument.\n\n- **`host::SymbolicSize{\"name\"}`** — A named symbolic dimension. Call `.set_value(n)` to pin it, `.unwrap()` to extract after verification.\n- **`host::SymbolicDType`** — Symbolic dtype. Use `.set_options()` to restrict allowed types.\n- **`host::SymbolicDevice`** — Symbolic device. Use `.set_options()` to restrict to CUDA.\n- **`host::TensorMatcher({dims...})`** — Fluent builder for tensor validation:\n - `.with_dtype()` — require a specific C++ type (e.g. `fp16_t`)\n - `.with_dtype()` — allow a set of types\n - `.with_device(device_sym)` — require CUDA and bind the checked device to a `SymbolicDevice`\n - `.with_strides({strides...})` — validate strides (omit to require contiguous)\n - `.ensure_alignment(bytes)` — require the data pointer **and every stride except the innermost** to be a multiple of `bytes` (a power of two). Size-1 dimensions are skipped, since their strides are arbitrary. This is how a vectorized kernel states its precondition; see the alignment convention above.\n - `.verify(tensor_view)` — execute the check; throws `PanicError` with full context on failure; **chainable** (`verify(a).verify(b)` to check multiple tensors with the same shape)\n- **`host::is_type(dtype)`** — whether a `DLDataType` denotes the C++ type `T` (e.g. `fp16_t`).\n\n**Typical pattern:**\n```cpp\nauto N = SymbolicSize{\"num_elements\"};\nauto device = SymbolicDevice{};\ndevice.set_options();\nTensorMatcher({N}) //\n .with_dtype()\n .with_device(device)\n .ensure_alignment(16) // e.g. for 128-bit vectorized load/store\n .verify(dst)\n .verify(src); // same shape, dtype, device as dst\nconst int64_t n = N.unwrap();\nconst DLDevice dev = device.unwrap();\nconst int64_t last_dim = 128;\nTensorMatcher({N, last_dim}) // a fixed dimension can be a plain integer\n .with_dtype()\n .with_device(device)\n .verify(tensor_2d);\n```\n\n### `ffi.h` — Tensor allocation and blob wrapping (`host::ffi::`)\n\n```cpp\n#include \n```\n\nThe counterpart to `tensor.h`: that one validates what came in, this one produces new `tvm::ffi::Tensor` values. Allocation goes through the environment allocator (`TVMFFIEnvTensorAlloc`), so buffers come from PyTorch's caching allocator rather than a raw `cudaMalloc`.\n\n- **`host::alloc_workspace_tensor(nbytes, device)`** (declared in `utils.cuh`) — **the way to get scratch memory**: a 1-D `uint8` tensor of `nbytes`, or an empty tensor when `nbytes == 0`. Hold the returned `Tensor` in a local across every launch that touches it — it frees on destruction.\n- **`host::ffi::empty(shape, dtype, device)`** — Uninitialized tensor; `shape` accepts a braced list, so `ffi::empty({rows, sizeof(Plan)}, dtype, device)` works for a typed scratch array.\n- **`host::ffi::empty_like(tensor_view)`** — Same shape, dtype, and device as an existing tensor.\n- **`host::ffi::from_blob(data, shape, dtype, device[, deleter, stride, byte_offset])`** / **`from_blob_like(data, tensor_view, ...)`** — View memory you already own as a `Tensor`, no copy. The default deleter does nothing, so ownership stays with the caller; pass one only when the `Tensor` should own the block. Strides default to contiguous.\n\n### `type.cuh` — `DTypeTrait`, `packed_t`, and reduction traits\n\n```cpp\n#include \n```\n\n- **`DTypeTrait`** — Static trait struct, specialized for integral types, `fp32_t`, `fp16_t`, `bf16_t`, `fp8_e4m3_t`, and their packed x2/x4 variants. Provides:\n - `DTypeTrait::from(value)` — convert from another type via the right CUDA intrinsic (e.g. `fp32_t` → `fp16_t`)\n - `DTypeTrait::abs/max/min` — type-dispatched math (fp32, fp16/bf16 scalar and x2, integrals)\n - `DTypeTrait::sqrt/rsqrt/exp/sin/cos(x)` — `fp32_t` only\n - Metadata: `packed_t` / `unpacked_t` / `kVecSize` (packed layout), `kFloatMax` (dtype max as float, e.g. 448.0f for fp8-e4m3), `kZeroBits`\n- **`packed_t`** — Two-element packed alias: `packed_t` = `fp16x2_t`, `packed_t` = `bf16x2_t`, `packed_t` = `fp32x2_t`. Use for vectorized loads/stores.\n- **`device::cast(value)`** — Type-safe cast using `DTypeTrait`, e.g. `cast(v)`.\n- **`device::unpack(value)`** — View a packed value as an `unpacked_t[kVecSize]` array reference (e.g. `fp32x2_t` → `fp32_t[2]`); element writes propagate back to the packed value.\n- **`device::ReductionOp` (`SUM`/`MAX`/`MIN`) and `device::ReductionTrait::reduce(x, y)`** — One binary reduction step, dispatched through `DTypeTrait` (packed types reduce elementwise). This is the engine behind `warp::reduce`; use it directly when writing custom reductions.\n\n### `vec.cuh` — Vectorized memory access (`AlignedVector`)\n\n```cpp\n#include \n```\n\n- **`device::AlignedVector`** — Aligned storage for N elements of type T. N must be a power of two, `sizeof(T)*N <= 32`. Enables vectorized loads/stores for bandwidth efficiency. In terms of API/codegen constraints, the upper bound is 256-bit; in practice, 128-bit is the portable default, while 256-bit vectorization is typically only viable on `SM100+` and should be gated by an architecture check when needed.\n - `.load(ptr, offset)` — vectorized load from `ptr[offset]`\n - `.store(ptr, offset)` — vectorized store to `ptr[offset]`\n - `.fill(value)` — fill all N elements with `value`\n - `operator[](i)` — element access\n- **`device::LoadStoreBytes`** — named widths to pass around instead of bare integers: `MAX_GMEM` (arch-dependent, 32 on Blackwell), `MAX_SMEM` (16), `MAX_PORTABLE` (16, safe on CUDA and HIP), `MIN_COALALESCED` (4), and `RAW_1B` .. `RAW_32B`.\n\n### `tile.cuh` — `tile::Memory` (strided memory access pattern)\n\n```cpp\n#include \n```\n\n- `tile::Memory` is fundamentally a **1D cooperative accessor** over a contiguous region.\n- **`device::tile::Memory::cta(blockDim.x)`** — Creates a tile accessor where each thread handles `tid = threadIdx.x` with stride `tsize` (for `cta(blockDim.x)`, this is `blockDim.x`). Common for loops over a 1D array.\n- **`device::tile::Memory::warp()`** — the 32-lane flavour, and what `warp::load_bytes` is built on. It takes its lane index from `get_lane_id()` rather than `threadIdx.x % 32`, for the addressing reason described under `utils.cuh`. Use `warp(int n)` for a narrower sub-group; that overload keeps the modulo, since a hardware lane index is the wrong grouping below 32.\n- **`device::tile::Memory::thread()`** — single-thread accessor, no cooperation.\n- **`.load(ptr, offset)`** — loads `ptr[tid + offset * tsize]`\n- **`.store(ptr, val, offset)`** — stores to `ptr[tid + offset * tsize]`\n- **`.in_bound(n, offset)`** — boundary check\n\nFor a **2D tile**, either flatten `(row, col)` into a linear tile index first, or compute the address manually with `ptr[row * stride + col]` using your thread/block coordinates.\n\n### `bits.h` — Compile-time bit helpers (`host::`)\n\n```cpp\n#include \n```\n\n`constexpr` wrappers over ``, usable in `static_assert` and in template arguments: `host::is_pow2(x)`, `log2_floor(x)`, `log2_ceil(x)` (both `-1` for `x == 0`), `round_up_pow2(x)`, `round_down_pow2(x)`. They live in `host::` but are constant-expression-only, so device-side `static_assert` can call them.\n\n### `math.cuh` — Device math (`device::math::`)\n\n```cpp\n#include \n```\n\n- `device::math::max/min(a, b)` — type-dispatched binary math via `DTypeTrait`\n- `device::math::abs/sqrt/rsqrt/exp/sin/cos(x)` — type-dispatched unary math via `DTypeTrait`\n\n### `warp.cuh` — Warp-level primitives\n\n```cpp\n#include \n```\n\n**Vectorized row copy — reach for this before hand-rolling.** A warp cooperatively moving one contiguous row is the single most common shape in this tree, and it used to be re-implemented per kernel.\n\n- **`warp::load_bytes(src)`** / **`warp::store_bytes(dst, val)`** — the whole warp moves `kBytes` from/to one contiguous row. `kBytes` is arbitrary, down to 1; the helper picks the widest vector that divides it and handles the ragged tail. The returned value is opaque and encodes the width, so a load and its matching store **must use the same ``**.\n- **`warp::LoadStorePattern`** — the `kPattern` values, and the sign carries the meaning. A **negative** one (`WARP_UNIFORM_16B` and friends: `_GMEM`, `_SMEM`, `_4B`/`_8B`/`_32B`) says *the whole warp splits this row*, so the width is derived from the per-lane share (`gcd(kBytes / 32, |kPattern|)`), falling back to `gcd(kBytes, 4)` when the row is too narrow to give every lane 4 bytes. A **positive** one (a bare `LoadStoreBytes` value) makes no warp-splitting assumption and just caps the per-thread vector at `gcd(kBytes, kPattern)`. `WARP_UNIFORM_16B` is the usual choice for a row copy.\n- **`LoadStorePattern::get_vec_bytes()`** — the width actually chosen. `SGL_DEVICE_HOST`, so **the launcher must call this to feed `.ensure_alignment(...)`**. Pass it the same value the kernel passes `load_bytes` — for a split row that is the *per-warp share*, not the whole row; the full row can resolve to a different width and would under-constrain the strides.\n\n**Reductions.**\n\n- `warp::reduce(value, active_mask)` — `Op` is a `device::ReductionOp` (`SUM`/`MAX`/`MIN`). `kStart` and `kFinish` bound a range of lane-index bits: `` reduces within contiguous groups of `N` lanes, `` reduces *across* groups at the same offset. The operation is **symmetric** — `` and `` reduce the same lane set.\n- `warp::reduce_sum/reduce_max/reduce_min(value)` — wrappers. Work for any type with a `ReductionTrait`: floats, integers, packed x2. This is more widely used in practice than generic `reduce`.\n- `warp::inclusive_reduce(val, lane_id, mask)` and `inclusive_sum` / `inclusive_max` / `inclusive_min` — segmented inclusive scan: every lane keeps its own running total, unlike `reduce`. Forward when `kStart < kFinish`, backward when `kStart > kFinish`. `lane_id` defaults to `get_lane_id()`, which is correct by construction — only pass it if you already have the **segment-relative** index (`threadIdx.x % kWidth`); a warp-relative one silently corrupts every segment past the first when `kWidth < 32`.\n- `warp::broadcast(value, src_lane)` — `src_lane` is **segment-relative**: each `kWidth` segment reads its own lane, not one warp-wide source.\n- `warp::get_lane_id`, `warp::elect_one_lane()` (one elected lane, for gating a single-thread TMA issue; CUDA-only), `warp::kFullMask` and `warp::kFullWidth` (32 on CUDA, 64 on HIP — note this differs from `kWarpThreads`).\n\n### `cta.cuh` — CTA-level primitives\n\n```cpp\n#include \n```\n\n- `device::cta::reduce_max(value, smem, min_value)` — CTA-wide max using shared memory + warp reduction. Caller is responsible for a `__syncthreads()` after if the result in `smem[0]` is needed.\n\n### `atomic.cuh` — Atomic operations\n\n```cpp\n#include \n```\n\n- `device::atomic::max(float* addr, float value)` — float atomic max (handles negative values correctly via bit tricks).\n- **`device::atomic::Event`** — a cross-CTA arrive/wait counter in one 32-bit word. Producers call `arrive()`; consumers call `wait(num_producers)` (exactly one consumer) or `wait_multi(num_producers, num_consumers, n)` (several). The last released consumer resets the word, so one `Event` is reusable across launches with no host re-zero. Preconditions, all of them load-bearing:\n - the word must be **zeroed by the host** before first use, and must live in **global** memory — a shared or local `Event` is an illegal address;\n - `arrive()` is a release and the waits are acquires, so a producer's writes before `arrive()` are visible after the wait;\n - **generations must be explicitly ordered** — the word has no phase bit, so overlapping generation N+1 with a consumer still in generation N is undefined behaviour (a lagging consumer never exits);\n - `wait()` and `wait_multi()` lay the word out incompatibly; mixing them on one `Event` is undefined behaviour.\n\n### `runtime.cuh` — Occupancy and device info\n\n```cpp\n#include \n```\n\n- `host::runtime::get_blocks_per_sm(kernel, block_dim)` — max active blocks per SM (occupancy)\n- `host::runtime::get_sm_count(device_id [, use_cache])` — number of SMs on the device. Memoized per device ordinal, so it is cheap enough to call on a launch path; pass `use_cache=false` to force a driver query.\n- `host::runtime::get_cc_major` / `get_cc_minor` / `get_sm_version` — compute capability, same caching.\n\n**Do not query the architecture at runtime.** The JIT compiles for the exact local GPU, so the arch is a *compile-time* fact: `SGL_CUDA_ARCH` is injected by `load_jit`, and `SGL_ARCH_HOPPER_OR_GREATER` / `SGL_ARCH_BLACKWELL_OR_GREATER` / `device::kMaxVecBytes` are derived from it. Reach for `get_cc_*` only for something the arch genuinely does not determine. SM *count* is the opposite case — it varies within an arch (B200 vs B300 vs a MIG slice), so it has to stay a runtime query.\n\n**Persistent kernel pattern** (cap blocks to SM count × occupancy):\n```cpp\nstatic const uint32_t max_occ = runtime::get_blocks_per_sm(kernel, kBlockSize);\nstatic const uint32_t num_sm = runtime::get_sm_count(device.unwrap().device_id);\nconst auto num_blocks = std::min(num_sm * max_occ, div_ceil(n, kBlockSize));\nLaunchKernel(num_blocks, kBlockSize, device.unwrap())(kernel, params);\n```\n\n---\n\n## Step 0 (optional): Generate a `.clangd` config for better IDE support\n\n```bash\npython -m sglang.kernels.jit -h # for verbose help info about clangd configuration\npython -m sglang.kernels.jit\npython -m sglang.kernels.jit --dep cutlass flashinfer # with cutlass/flashinfer dependency\n```\n\n---\n\n## Step 1: Implement the CUDA kernel in `kernels/jit/csrc/`\n\nCreate `python/sglang/kernels/jit/csrc/elementwise/scale.cuh`.\n\nThe implementation fully uses the project abstractions described above:\n\n```cpp\n// NOTE: Comments for headers are not common in practice.\n// It is only shown here for tutorial purposes to highlight the key abstractions.\n#include // For TensorMatcher, SymbolicSize, SymbolicDevice\n#include // For DTypeTrait, fp16_t, bf16_t, fp32_t\n#include // For CHECK_HOST, div_ceil\n#include // For LaunchKernel, SGL_DEVICE\n#include // For AlignedVector\n\n#include \n#include \n\nnamespace sglang {\n\n/**\n * \\brief Element-wise scale using vectorized 128-bit loads/stores.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kVecN Elements per vector load (e.g. 8 for fp16)\n * \\tparam kUsePDL Whether to emit the PDL wait/trigger pair\n * \\param dst Output buffer, `n_total` elements\n * \\param src Input buffer, `n_total` elements\n * \\param factor Runtime scale factor\n * \\param n_total Number of elements to scale\n */\ntemplate \n__global__ void scale_kernel(T* __restrict__ dst,\n const T* __restrict__ src,\n float factor,\n uint32_t n_total) {\n using vec_t = device::AlignedVector;\n const uint32_t n_vecs = n_total / kVecN;\n\n // If using PDL, wait for primary kernel before any global memory load.\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLWaitPrimary();\n\n // --- vectorised body ---\n const uint32_t vec_stride = blockDim.x * gridDim.x;\n for (uint32_t vi = blockIdx.x * blockDim.x + threadIdx.x;\n vi < n_vecs;\n vi += vec_stride) {\n vec_t v;\n v.load(src, vi);\n#pragma unroll\n for (int i = 0; i < kVecN; ++i) {\n v[i] = static_cast(static_cast(v[i]) * factor);\n }\n v.store(dst, vi);\n }\n\n // --- scalar tail ---\n const uint32_t base = n_vecs * kVecN;\n const uint32_t scalar_stride = blockDim.x * gridDim.x;\n for (uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;\n base + i < n_total;\n i += scalar_stride) {\n dst[base + i] = static_cast(static_cast(src[base + i]) * factor);\n }\n\n // If using PDL, signal for the secondary kernel to start after all threads have finished\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLTriggerSecondary();\n}\n\n/**\n * \\brief Validate the tensors, select the vector width, launch `scale_kernel`.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kUsePDL Whether to launch with PDL enabled\n * \\param dst Output tensor; same shape / dtype / device as `src`\n * \\param src Input tensor on CUDA\n * \\param factor Runtime scale factor\n */\ntemplate \nvoid scale(tvm::ffi::TensorView dst, tvm::ffi::TensorView src, float factor) {\n using namespace host;\n\n // 1. Validate input tensors with TensorMatcher\n SymbolicSize N = {\"num_elements\"};\n SymbolicDevice device_;\n device_.set_options();\n\n TensorMatcher({N}) //\n .with_dtype()\n .with_device(device_)\n .verify(dst)\n .verify(src); // same shape / dtype / device as dst\n\n const uint32_t n = static_cast(N.unwrap());\n const DLDevice device = device_.unwrap();\n\n CHECK_HOST(n > 0) << \"scale: num_elements must be > 0, got \" << n;\n\n // 2. Choose vector width for 128-bit loads (16 bytes)\n // fp16/bf16: 8 elements x 2 bytes = 16 bytes\n // fp32: 4 elements x 4 bytes = 16 bytes\n // We encourage using `device::kMaxVecBytes`, which will change according to\n // the target architecture and can enable 256-bit vectorization on SM100+ if desired.\n // But 128-bit is more commonly adapted for better compatibility,\n // so it's still ok to hardcode 16 here just for simplicity.\n constexpr int kVecN = 16 / sizeof(T);\n const uint32_t n_work_items = div_ceil(n, static_cast(kVecN));\n\n // 3. Launch\n constexpr uint32_t kBlockSize = 256;\n const uint32_t grid = div_ceil(n_work_items, kBlockSize);\n\n // PDL feature is 100% optional. Without `enable_pdl`, the code should still be correct.\n // Try to enable it if profiling shows that it can benefit the performance of this kernel.\n LaunchKernel(grid, kBlockSize, device).enable_pdl(kUsePDL)(\n scale_kernel,\n static_cast(dst.data_ptr()),\n static_cast(src.data_ptr()),\n factor,\n n);\n}\n\n} // namespace sglang\n```\n\n**Key points:**\n\n- Include headers from `sgl_kernel/` — **not** raw CUDA headers for anything already covered\n- Use `TensorMatcher` for all tensor validation; never manually check shape/dtype/device\n- Use `AlignedVector` for vectorised 128-bit loads/stores — significant bandwidth win\n- Use `LaunchKernel` — it resolves the stream and checks errors automatically\n- Use `CHECK_HOST(cond) << ...` for runtime assertions with useful error messages (zero overhead when the check passes)\n- Prefer passing runtime scalars like `factor` directly unless compile-time specialisation is genuinely required\n- `fp16_t` / `bf16_t` / `fp32_t` are the project's type aliases (from `utils.cuh`)\n- `device::cast` or `DTypeTrait::from(val)` for cross-type conversions\n- `device::math::` functions for device math instead of bare `__` intrinsics if possible.\n- Consider PDL — it can help when the kernel has prologue work to overlap. Place `PDLWaitPrimary()` right before the first read of upstream data, not at the top of the kernel\n\n---\n\n## Step 2: Add the Python wrapper in `kernels/ops/`\n\nThe wrapper lives next to its functional group under `python/sglang/kernels/ops/`, not beside the CUDA source — `kernels/jit/` holds only the JIT infrastructure (`csrc/`, `include/`, `utils/`, `benchmark/`). Create `python/sglang/kernels/ops/elementwise/scale.py`:\n\n```python\nfrom __future__ import annotations\n\nfrom typing import TYPE_CHECKING\n\nimport torch\n\nfrom sglang.kernels.jit.utils import (\n cache_once,\n is_arch_support_pdl,\n load_jit,\n make_cpp_args,\n)\n\nif TYPE_CHECKING:\n from tvm_ffi.module import Module\n\n\n@cache_once\ndef _jit_scale_module(dtype: torch.dtype) -> Module:\n \"\"\"Compile and cache the JIT scale module for a given dtype.\"\"\"\n # Checks on the compile key live here, not in `scale`: `cache_once` keys on\n # `dtype`, so this runs once per specialisation instead of once per call.\n if dtype not in (torch.float16, torch.bfloat16, torch.float32):\n raise RuntimeError(\n f\"Unsupported dtype {dtype}. Supported: float16, bfloat16, float32\"\n )\n args = make_cpp_args(dtype, is_arch_support_pdl())\n return load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n )\n\n\ndef scale(src: torch.Tensor, factor: float, out: torch.Tensor | None = None) -> torch.Tensor:\n \"\"\"\n Element-wise scale: dst = src * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n src : CUDA tensor (FP16 / BF16 / FP32)\n factor : scale factor\n out : optional pre-allocated output tensor (same shape/dtype as src)\n\n Returns\n -------\n Scaled tensor (dst = src * factor).\n \"\"\"\n # DO NOT add proactive validation here: every check costs interpreter time\n # on a per-forward path. Tensor invariants belong in the C++ launcher, and\n # anything about the compile key belongs in `_jit_scale_module`.\n if out is None:\n out = torch.empty_like(src)\n\n module = _jit_scale_module(src.dtype)\n module.scale(out, src, factor)\n return out\n```\n\n**Key points:**\n\n- Use `cache_once` — **not** `functools.lru_cache` (incompatible with `torch.compile`)\n- `load_jit` first arg(s) form the unique build marker; same marker = same cached binary\n- Only include compile-time specialisation knobs in the build marker; runtime values like `factor` should stay runtime unless the kernel truly needs templating\n- `cuda_wrappers`: `(export_name, kernel_symbol)` — `export_name` is called from Python\n- `make_cpp_args(dtype, ...)` converts `torch.dtype` to C++ type alias:\n- `is_arch_support_pdl()` checks if the current architecture supports PDL, which is typically passed as a template argument to the kernel.\n- Keep the entry point thin (see Conventions). What Python must still check goes in the `@cache_once` module factory, not in the entry point: `cache_once` keys on its arguments, so a check there costs one evaluation per specialisation instead of one per call — that is where the supported-dtype guard lives. Tensor invariants belong in the C++ launcher; if something here never reaches a `.verify(...)`, close that gap on the C++ side rather than in Python\n\n| `torch.dtype` | C++ type |\n|--------------------|------------|\n| `torch.float16` | `fp16_t` |\n| `torch.bfloat16` | `bf16_t` |\n| `torch.float32` | `fp32_t` |\n\n---\n\n## Register the public entry point (required)\n\nAdd metadata in `python/sglang/kernels/ops/elementwise/__init__.py` so the\noperator appears in the registry without importing its implementation:\n\n```python\nfrom sglang.kernels.registry import register_kernel\nfrom sglang.kernels.spec import CapabilityRequirement, KernelBackend, KernelSpec\n\nregister_kernel(\n KernelSpec(\n op=\"elementwise.scale\",\n backend=KernelBackend.JIT,\n target=\"sglang.kernels.ops.elementwise.scale:scale\",\n capabilities=frozenset({CapabilityRequirement.CUDA}),\n )\n)\n```\n\nRecord the devices the implementation actually supports; the CUDA-only scale\nexample above does not imply that every JIT kernel is CUDA-only. Runtime and\ncorrectness tests import `scale` from `sglang.kernels.ops.elementwise.scale`.\nDo not import the implementation eagerly in the group initializer. When an\nexisting `BaseFusedOp` owns the operation, add the implementation there instead\nof registering a second conflicting op/backend pair.\n\n## Step 3 (optional): Tune JIT build flags\n\nIf your kernel uses some math functions like `expf` or `sinf`, consider enabling `--use_fast_math` for better performance (with a potential precision tradeoff):\n\n```python\nreturn load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n)\n```\n\nIf your kernel requires SM90+, raise a clear Python error before calling `load_jit`. Arch gating is one of the checks that has to live in Python — it decides whether to compile at all, so the C++ launcher never gets to run:\n\n```python\nif torch.cuda.get_device_capability()[0] < 9:\n raise RuntimeError(\"This kernel requires SM90 (Hopper) or later\")\n```\n\n---\n\n## Step 4: Write tests (required)\n\nJIT kernel correctness tests and benchmarks live under `test/registered/kernels/ops//` and `test/registered/kernels/benchmark//`, mirroring the wrapper's group under `python/sglang/kernels/ops/` (NOT inside the `sglang` package -- a `register_*_ci(...)` call anywhere under `python/sglang/` is rejected by the `check-no-registered-tests-in-package` pre-commit hook). Only their test-only helpers (e.g. `benchmark/marker.py`) stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path. **CI does not run `pytest` in those directories directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` and `bench_*.py` under `test/registered/`, collects `register_*_ci(...)` calls by **statically parsing each file's AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check.\n\n- **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`.\n- **Nightly kernel suite**: register with `stage=\"nightly\"` plus the `runner_config` of the machine it needs (e.g. `1-gpu-large`), giving the `nightly-test-1-gpu-large` suite. `.github/workflows/nightly-test-nvidia.yml` sets `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` for the whole nightly run, so the expanded parameter grids apply automatically (see `python/sglang/kernels/jit/utils/common.py` → `should_run_full_tests` / `get_ci_test_range`). There is no separate kernel-only nightly job: every nightly test on one machine type shares that machine's suite.\n\nRegistration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing):\n\n```python\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n# Optional B200/SM100 registration for tests that cover Blackwell-specific code paths\n# register_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"4-gpu-b200\")\n# Optional second registration: same file also runs nightly, same form,\n# stage is just \"nightly\" there (and no `nightly=True`)\n# register_cuda_ci(est_time=120, stage=\"nightly\", runner_config=\"1-gpu-large\")\n```\n\nCI generates the suite name as `{stage}-test-{runner_config}`, so `stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\"` becomes the `base-b-kernel-unit-test-1-gpu-large` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage=\"nightly\"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools.\n\nKeep `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.\n\nUse `register_cuda_ci(..., disabled=\"reason\")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only).\n\n**Run like CI** (from repo root):\n\n```bash\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large)\n# For B200/SM100-specific coverage:\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200)\n```\n\nFor fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`.\n\nCreate `test/registered/kernels/ops/elementwise/test_scale.py`:\n\n```python\nimport pytest\nimport torch\nfrom sglang.kernels.ops.elementwise.scale import scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [1, 127, 128, 1024, 4097]) # cover tail remainder\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0, 3.0])\ndef test_scale_correctness(dtype, size, factor):\n src = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = scale(src, factor)\n expected = src * factor\n\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\ndef test_scale_out_param(dtype):\n src = torch.randn(1024, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(src)\n result = scale(src, 2.0, out=out)\n assert result is out\n torch.testing.assert_close(out, src * 2.0, rtol=1e-2, atol=1e-2)\n\n\ndef test_scale_cpu_error():\n src = torch.randn(128, dtype=torch.float16) # CPU tensor\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n scale(src, 2.0)\n\n\ndef test_scale_unsupported_dtype():\n src = torch.randint(0, 10, (128,), dtype=torch.int32, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"dtype\"):\n scale(src, 2.0)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-v\", \"-s\"]))\n```\n\n---\n\n## Step 5: Add a benchmark (required)\n\nBenchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large`).\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nBenchmarks use the project's own `marker` framework (in `python/sglang/kernels/jit/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`):\n\n- **`@marker.benchmark(line_arg, line_vals, *, unit=\"us\")`** — the **innermost** decorator (bottom of the stack, directly above `def benchmark`). Declares the column axis: each value in `line_vals` becomes a result column, and `line_arg` is the parameter name passed into the benchmark function. `unit` is one of `\"us\" | \"ms\" | \"s\"`.\n- **`@marker.parametrize(names, vals, ci_vals=None)`** — stackable decorator that adds a row axis (pytest-style). Each `@parametrize` adds one (or more, correlated) parameter the benchmark is swept over (Cartesian product across all `parametrize` decorators). `names` may be a single name (`\"size\"`) or a comma-separated correlated tuple axis (`\"h,d\"`, with `vals` then a list of tuples like `[(1, 64), (2, 128)]`). Pass the optional third `ci_vals` for a smaller sweep that is auto-selected under `is_in_ci()` — this is the built-in CI-shrinking mechanism, so you usually don't need `get_benchmark_range` for swept axes.\n- **`marker.do_bench(fn, *, input_args=(), input_kwargs={}, ...)`** — runs `fn` under CUDA graph (default) or a naive loop, returns a `BenchResult`. Key knobs:\n - `memory_args`: defaults to `\"all\"` (footprint derived from all input args/kwargs). Pass an explicit tuple of tensors (e.g. `(k, v, indices)`) to count only the inputs the kernel actually touches.\n - `memory_output`: defaults to `\"out\"` — re-runs `fn` once to capture its **returned** tensor and counts it. For in-place kernels (which return `None`), pass the written tensors explicitly (e.g. `memory_output=(k, v)`); the re-run is then skipped. Set to `None` to count no output.\n - Together `memory_args` + `memory_output` give the GB/s column; with both defaults a function `out = f(src)` already reports `bytes(src) + bytes(out)`.\n - `graph_clone_args` / `graph_clone_kwargs`: which inputs to clone per CUDA-graph iteration to defeat L2 cache reuse. Defaults to `\"all\"`; if you narrow it, keep every read, in-place, and output tensor.\n - `use_cuda_graph=False` for kernels that can't be captured.\n - `metrics=(0.5, \"avg\")` controls reported quantiles (the first metric becomes the table latency column).\n - `disable_log_bandwidth` (defaults from `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1`) skips the bandwidth column entirely.\n- **`utils.create_random(*shape)` / `utils.create_empty(*shape)`** — shorthand for `torch.randn` / `torch.empty` with `DEFAULT_DTYPE` (`bfloat16`) and `DEFAULT_DEVICE` (`\"cuda\"`). Override via the `dtype=` / `device=` kwargs.\n- **`utils.get_benchmark_range(full_range, ci_range)`** — returns the smaller `ci_range` under CI (`is_in_ci()`), the `full_range` locally. Still available for the `benchmark(...)` column axis (which has no `ci_vals`); for `parametrize` row axes prefer the built-in `ci_vals` argument.\n\nCreate `test/registered/kernels/benchmark/elementwise/bench_scale.py`:\n\n```python\nimport torch\n\nfrom sglang.kernels.jit.benchmark import marker\nfrom sglang.kernels.jit.benchmark.utils import create_random\nfrom sglang.kernels.ops.elementwise.scale import scale as jit_scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=6, stage=\"base-b-kernel-benchmark\", runner_config=\"1-gpu-large\")\n\n\n@torch.compile()\ndef torch_impl_scale(src: torch.Tensor, factor: float) -> torch.Tensor:\n return src * factor\n\n\nFN_MAP = {\n \"jit\": jit_scale,\n \"torch\": torch_impl_scale,\n}\n\n\n# `parametrize(name, full_vals, ci_vals)`: the 3rd arg is the smaller sweep\n# auto-selected under CI; the full range runs locally.\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096, 65536]) # 1K .. 512K\n@marker.benchmark(\"impl\", [\"jit\", \"torch\"])\ndef benchmark(size: int, impl: str):\n src = create_random(size)\n factor = 2.0\n return marker.do_bench(\n FN_MAP[impl],\n input_args=(src, factor),\n # `src` is read -> clone it per iter to avoid L2 reuse; factor is a scalar.\n graph_clone_args=(0,),\n # Defaults already report bandwidth: memory_args=\"all\" counts src,\n # memory_output=\"out\" counts the returned tensor -> bytes(src)+bytes(out).\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n**Key points:**\n\n- The `line_arg` name passed to `benchmark` (`\"impl\"` here) must match a parameter on `benchmark(...)`; same for every `parametrize` name (`\"size\"`).\n- Stack `@parametrize` once per swept axis. The required `@marker.benchmark` is the **innermost** decorator (bottom of the stack, directly above the function) — `@parametrize` rows go above it.\n- Prefer `create_random` / `create_empty` from `utils.py` over open-coding `torch.randn(..., dtype=..., device=...)`.\n- The GB/s column appears by default (`memory_args=\"all\"` + `memory_output=\"out\"`). For memory-bound kernels it's the most informative number; scope `memory_args` / `memory_output` to the tensors actually touched if the defaults over- or under-count. For compute-bound kernels where bandwidth is misleading, set `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1` (or `disable_log_bandwidth=True`).\n- For in-place kernels (which return `None`), pass the written tensors via `memory_output=(...)` since the `\"out\"` default would capture nothing.\n- Tune `graph_clone_args` / `graph_clone_kwargs` to all the arguments that might be read by the kernel. We can only skip cloning for write-only args. For in-place modified args, we still need to clone them to get accurate timing (reusing the same buffer keeps it L2-hot and skews results).\n- Call `benchmark.run()` (no `print_data=` kwarg — the marker framework prints directly).\n\nRun locally:\n\n```bash\npython test/registered/kernels/benchmark/elementwise/bench_scale.py\n```\n\nRun the benchmark suite the way CI does:\n\n```bash\ncd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large\n```\n\n---\n\n## Troubleshooting\n\n- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time`, `stage`, and `runner_config`; starred args and non-literal values break AST collection\n- **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/kernels/jit/csrc/`; reduce template argument combinations\n- **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...`\n- **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. `graph_clone_args` defaults to `\"all\"`; if you narrow it, it must still cover every *read* tensor — reusing a single buffer keeps it L2-hot and skews results. Keep *write* tensors in it too: they are what sets the rotation count, and a shared output buffer stays L2-hot the same way.\n- **Missing GB/s column**: the column is on by default; check that `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH` is not `1` and `disable_log_bandwidth` is not `True`. For in-place kernels (return `None`) the `memory_output=\"out\"` default counts nothing — pass the written tensors via `memory_output=(...)`\n\n---\n\n## References\n\n- `docs/docs/developer_guide/development_jit_kernel_guide.mdx`\n- `test/run_suite.py` — suite names, discovery of `test/registered/`, execution entrypoint for CI\n- `python/sglang/test/ci/ci_register.py` — `register_cuda_ci` and AST registration rules\n- `python/sglang/kernels/jit/utils/compile/` — `load_jit` (`loader.py`), `make_cpp_args` (`cpp_args.py`)\n- `python/sglang/kernels/jit/utils/common.py` — `cache_once`, `should_run_full_tests`, `get_ci_test_range`\n- `python/sglang/kernels/jit/include/sgl_kernel/tensor.h` — `TensorMatcher`, `SymbolicSize/DType/Device`, `is_type`\n- `python/sglang/kernels/jit/include/sgl_kernel/ffi.h` — `ffi::empty`, `ffi::empty_like`, `ffi::from_blob`\n- `python/sglang/kernels/jit/include/sgl_kernel/utils.cuh` — type aliases, `LaunchKernel`, `SGL_DEVICE`\n- `python/sglang/kernels/jit/include/sgl_kernel/vec.cuh` — `AlignedVector`\n- `python/sglang/kerne",".agents/skills/debug-distributed-hang/SKILL.md":"---\nname: debug-distributed-hang\ndescription: Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations.\n---\n\n# Debugging Distributed Hangs in SGLang\n\n## Overview\n\nHangs in distributed inference happen when ranks diverge in state, causing collective operations (AllGather, AllReduce, Broadcast, Barrier) to deadlock. Common causes:\n\n- **Size mismatch**: ranks pass different tensor sizes to a collective\n- **Branch divergence**: one rank enters a collective, another skips it\n- **Cascading state drift**: a small non-determinism (e.g., floating-point) propagates into different batch structures\n- **Resource exhaustion**: one rank OOMs or crashes, others wait forever\n\n## Prerequisites\n\n- **py-spy**: `pip install py-spy` or system package. Requires root or `CAP_SYS_PTRACE` to attach to running processes.\n- **cuda-gdb**: Ships with the CUDA toolkit. Ensure it's on your `PATH`.\n\n## Step 1: Confirm and Locate the Hang\n\n### 1a. Watchdog / py-spy\n\nSGLang's watchdog automatically dumps py-spy traces on timeout. Look for:\n\n```\nScheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)\n```\n\nThe py-spy dump shows the stack trace of each thread. The hanging thread is typically blocked in a CUDA synchronize or NCCL collective:\n\n```\nThread (active): \"MainThread\"\n cuStreamSynchronize (libcuda.so)\n ...\n forward_extend (model_runner.py)\n```\n\nSGLang has two watchdog modes (see `python/sglang/srt/utils/watchdog.py`):\n- **Hard watchdog** (`soft=False`, default): dumps py-spy traces then sends `SIGQUIT` to kill the parent process.\n- **Soft watchdog** (`soft=True`): only logs the timeout without killing the process, giving you more time to manually attach debuggers or collect coredumps.\n\nIf the watchdog doesn't trigger, manually dump:\n\n```bash\npy-spy dump --pid \n```\n\n### 1b. NCCL Debug Logging\n\n```bash\nexport NCCL_DEBUG=INFO\nexport NCCL_DEBUG_SUBSYS=COLL\n```\n\nLook for the last collective logged before the hang. Mismatched sizes show up as one rank waiting and another never entering.\n\n### 1c. CUDA Coredump\n\nWhen a process hangs, you can trigger a GPU coredump on demand to see which kernel is stuck. Set these env vars before launching:\n\n```bash\nexport CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1\nexport CUDA_COREDUMP_PIPE=\"/tmp/cuda_pipe_%h_%p\"\nexport CUDA_COREDUMP_FILE=\"/tmp/cuda_coredump_%h_%p\"\nexport CUDA_COREDUMP_SHOW_PROGRESS=1\nexport CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'\n```\n\nWhile the process is hanging, find the pipe via `/proc//fd/` and write to it to trigger the dump:\n\n```bash\nls /proc//fd/ -la 2>/dev/null | grep cuda_pipe\ndd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe__\n```\n\nAlternatively, if you don't need to keep the process alive, `kill -SIGABRT ` also triggers a CUDA coredump (but terminates the process).\n\nThen open with `cuda-gdb --batch -ex \"target cudacore \"`. On load, it immediately shows which kernel is stuck. For example:\n\n```\nOpening GPU coredump: \n[Current focus set to CUDA kernel 0, grid 622721, cluster (4,0,0), block (16,0,0), thread (64,0,0), device 0, sm 0, warp 0, lane 0]\n#0 0x00007f8029b2b040 in ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)<<<(24,1,1),(512,1,1)>>> ()\n```\n\nThis told us the hang was in an NCCL AllGather — not a compute kernel. Combined with the py-spy stack pointing to `LogitsProcessor.forward` → `tensor_model_parallel_all_gather`, we knew it was an AllGather size mismatch between TP ranks.\n\n\n### 1d. Identify the Collective\n\nFrom the stack traces and logs, identify:\n- Which collective hangs (AllGather, AllReduce, Broadcast)\n- Which code path invokes it (e.g., `LogitsProcessor`, `tensor_model_parallel_all_gather`)\n- Whether it's a size mismatch or a missing participant\n\n## Step 2: Per-Rank Logging\n\nThe key technique: each rank writes its own log file so you can diff them.\n\n### Setup Pattern\n\n```python\nimport os\n\n_debug_files = {}\n\ndef get_debug_file(rank):\n key = f\"rank{rank}\"\n if key not in _debug_files:\n _debug_files[key] = open(f\"/tmp/debug_rank{rank}.log\", \"w\")\n return _debug_files[key]\n```\n\nGate logging behind an env var to avoid overhead in production. `SGLANG_DEBUG_HANG` is not a built-in SGLang env var — you need to add this check yourself in the code you're instrumenting:\n\n```python\nif os.environ.get(\"SGLANG_DEBUG_HANG\"):\n f = get_debug_file(rank)\n f.write(f\"EVENT_NAME key1={val1} key2={val2}\\n\")\n f.flush()\n```\n\n### What to Log\n\nLog structured events at key state-mutation points:\n\n```python\nf.write(f\"SCHED_BATCH step={step} num_reqs={n} extend_lens={lens}\\n\")\nf.write(f\"VERIFY predict_hash={hash} accept_len={alen}\\n\")\nf.write(f\"CACHE_INSERT rid={rid} num_tokens={n}\\n\")\n```\n\nUse consistent event names (uppercase prefix) for easy grep/diff.\n\n### Hash Large Tensors\n\nFor tensor values, compute a hash instead of dumping raw data:\n\n```python\nimport hashlib\nh = hashlib.md5(tensor.cpu().numpy().tobytes()).hexdigest()[:8]\nf.write(f\"LOGITS logits_hash={h}\\n\")\n```\n\nFor token ID lists, `str(list).encode()` works:\n\n```python\nh = hashlib.md5(str(tensor.tolist()).encode()).hexdigest()[:8]\n```\n\n### Avoid Implicit Synchronization\n\n`tensor.cpu()`, `tensor.tolist()`, and `tensor.numpy()` all trigger CUDA synchronization. This can:\n- Change timing and mask or move the hang\n- Deadlock if the log point is between two collectives that must run back-to-back\n\nPrefer logging values that are already on CPU (e.g., Python ints, list lengths, request IDs). When you must hash a GPU tensor, do it at a point where the GPU is already idle (e.g., between scheduler steps, not inside a model forward pass).\n\n## Step 3: Diff to Find the Diverge Point\n\n### Basic Diff\n\n```bash\n# Extract specific event type\ngrep \"^VERIFY\" /tmp/debug_rank0.log > /tmp/v_r0.txt\ngrep \"^VERIFY\" /tmp/debug_rank1.log > /tmp/v_r1.txt\ndiff /tmp/v_r0.txt /tmp/v_r1.txt | head -20\n```\n\n### Count Events\n\n```bash\ngrep -c \"^VERIFY\" /tmp/debug_rank*.log\n```\n\nIf counts differ, one rank executed more iterations — that's already a diverge signal.\n\n### Find First Diverge\n\nThe first diff line tells you the exact step where ranks diverge. All lines before it are identical — the root cause is at or before this step.\n\n## Step 4: Binary-Search the Root Cause\n\nOnce you find the diverging event, trace backwards:\n\n### 4a. Identify Inputs\n\nFor the diverging operation, list all its inputs. Add hash logging for each:\n\n```python\nf.write(\n f\"OP_INPUTS input_a_hash={h_a} input_b_hash={h_b} \"\n f\"input_c_hash={h_c} input_d_hash={h_d}\\n\"\n)\n```\n\n### 4b. Diff Inputs Across Ranks\n\nCompare the hashes. Some inputs will match, some won't. The non-matching input is where divergence entered.\n\n### 4c. Recurse\n\nFor the non-matching input, trace where it was produced and repeat: hash its inputs, diff across ranks, find the divergent one. Continue until you reach the root cause.\n\n## Step 5: Common Root Causes and Fixes\n\n### Floating-Point Non-Determinism\n\n**Symptom**: All \"logical\" inputs are identical (same logits after all-gather), but derived floating-point values (softmax, probabilities) differ across GPUs.\n\n**Example**: EAGLE speculative decoding — `F.softmax` → `top_k_renorm_prob` → `top_p_renorm_prob` produces slightly different `target_probs` on each GPU. The sampling kernel then picks different tokens. These flow into `output_ids` → radix cache → different prefix match depths → different `extend_seq_lens` → AllGather size mismatch → hang.\n\n### Random Number Divergence\n\n**Symptom**: Operations using `torch.rand` produce different values on each rank.\n\n**Fix**: Generate on rank 0 and broadcast, or use a shared seed.\n\n### Conditional Code Paths\n\n**Symptom**: A condition (e.g., memory check, queue length) evaluates differently on different ranks, causing one rank to enter a collective while another skips it.\n\n**Fix**: Synchronize the condition value before branching, or restructure to ensure all ranks take the same path.\n\n### Pipeline Parallel (PP) Send/Recv Mismatch\n\n**Symptom**: In PP setups, one stage issues a `send` that the next stage never `recv`s (or vice versa), causing both to block indefinitely. Unlike TP hangs (collective mismatches), PP hangs typically involve point-to-point operations.\n\n**Fix**: Ensure all stages agree on the number of microbatches and the sequence of send/recv calls for each microbatch.\n\n## Step 6: Verify the Fix\n\nRun the failing test multiple times to confirm the fix is stable. Intermittent hangs require many runs. A test that hung ~30% of the time needs at least 10 clean passes to be confident.\n\n## Quick Reference\n\n| Technique | When to Use |\n|-----------|-------------|\n| py-spy dump | First step — see where each rank is stuck |\n| `NCCL_DEBUG=INFO` | Identify which collective and sizes |\n| CUDA coredump + `cuda-gdb` | See which GPU kernel is blocked |\n| Per-rank log files | Compare rank states over time |\n| Hash of tensors | Efficiently compare large tensors across ranks |\n| `diff` on extracted events | Find the exact step of divergence |\n| `broadcast(result, src=0)` | Fix floating-point or sampling non-determinism |\n",".agents/skills/cookbook-migrate-model/SKILL.md":"---\nname: cookbook-migrate-model\ndescription: Migrate a legacy-template SGLang cookbook page (monolithic per-model generator under docs/src/snippets/autoregressive/) onto the config-driven template (shared _deployment.jsx / _playground.jsx engines + per-model config). Use when asked to migrate, convert, or port an existing cookbook page — NOT for brand-new models (use cookbook-add-model for those). Run with /cookbook-migrate-model .\n---\n\n# Cookbook Migrate Model\n\nConvert one legacy cookbook page to the config-driven format, faithfully. The\nlegacy page — its generator widget and its measured benchmark blocks — is the\n**single source of truth**. You are transcribing it into the new data model,\nnot improving it.\n\nReuses the `cookbook-add-model` skill's assets (read them on demand):\n- `../cookbook-add-model/templates/config.jsx.tmpl`, `page.mdx.tmpl`, `benchmarks.jsx.tmpl`\n- `../cookbook-add-model/references/authoring-reference.md` (config/cells/playground contract)\n- `../cookbook-add-model/references/mintlify-authoring.md` (MDX rules)\n\nMigration-specific references in this skill:\n- [references/dimension-mapping.md](references/dimension-mapping.md) — legacy-control → new-dimension mapping rules, command rewrite table, per-family strategy sets, and the Qwen3.5 pilot as a worked example (PR #27848).\n\nThe round's per-model inventory (scope, batch order, quirks, measured-data\nsurvey) is tracked by the migration maintainer outside the repo — expect it in\nyour dispatch prompt, or ask for it.\n\n## Hard rules (non-negotiable)\n\n1. **Never modernize.** Env vars, flags, TP values, docker tags, version strings\n are copied verbatim from the legacy page — even when today's defaults differ\n (e.g. a flag whose behavior has since become the default; keep it anyway). The recipe\n that was verified is the recipe as written. Allowed normalizations are ONLY\n the five alias rewrites in dimension-mapping.md §2 (`launch_server`→`sglang\n serve`, `--model`→`--model-path`, `--tp-size`→`--tp`, abbreviated\n `--speculative-algo`→`--speculative-algorithm`,\n `--expert-parallel-size`→`--ep`). **Accuracy-degrading flags**\n (`--kv-cache-dtype fp8_e4m3`, W4A4-style runtime quant) follow a\n deterministic rule — enforced in migration, no asking: offered as a\n legacy **selectable option** → never select it (cells mirror the\n accuracy-safe side), and the option **survives as a Playground axis** —\n an existing one where it fits, else add one via a separate prior engine\n PR (rule 4); a legacy choice never degrades to a tips mention. Baked\n into the recipe's **default/unconditional command** → keep it verbatim\n (the recipe was measured with it, and fp8 KV halves KV memory —\n stripping could OOM it). See dimension-mapping.md §2 caveats.\n2. **Never invent versions or numbers.** Benchmark numbers only from the\n legacy page's measured blocks, and a result migrates ONLY when its\n `sglang_version` is a **reproducible anchor** — the bar is reproducibility,\n not \"must be a release\":\n - ✅ release tag/version (`v0.5.9` / `0.5.9`), commit hash, OR — for\n **Day-0 support** (the enabling PR isn't merged and no release is cut\n yet) — a specific **PR (`PR #27944`) or commit** you can `gh pr checkout`\n / `git checkout `. Commit is most precise; a PR pin is fine for day-0.\n - ❌ a moving ref — `\"main branch\"`, `\"main (2026-06-11)\"`, open-ended\n `\"0.5.8+\"` — is NOT reproducible: **drop the WHOLE result (speed AND\n accuracy)**, not just speed. Keep `benchmarkCommands` so ⚡Reproduce still\n guides re-measurement against a pinned build.\n **Never inherit cross-model numbers** — measurements the legacy page\n attributes to a *different model* (e.g. a K2.6 page carrying K2.5-measured\n speed) are dropped regardless of version. When kept, `sglang_version` is the\n legacy page's string verbatim. Docker tags only the ones the legacy page\n pinned (unmapped hw falls back to `:dev`).\n3. **Verified policy (strictest tier).** `verified: true` ONLY when (a) the\n legacy page has concrete measured data for that exact 5-dim combo AND\n (b) the cell's flags equal the deployment command used for that measurement\n — modulo `{{HOST_IP}}`/`{{PORT}}`, the five alias rewrites, and **parser\n flags**: `--reasoning-parser`/`--tool-call-parser` are stripped from every\n cell (Playground-only feature; when the measured run had them on, say so in\n the benchmarks file header). When the measured command diverges from the\n generator default, **the verified cell follows the measured command**; the\n generator default stays as the sibling strategy/cell or a tips note. Everything else is unverified (yellow) —\n including combos that look memory-infeasible; keep them verbatim and list\n them in the PR body for the re-verification track.\n4. **Engines are read-only.** `_deployment.jsx` / `_playground.jsx` must not\n change in a migration PR. Model-specific features are config DATA consumed\n by generic axis handlers (MegaMoE precedent), so they need NO engine change.\n A **titled single-select that strips a flag family** — KV Cache DType\n (`--kv-cache-dtype`), mamba (`--mamba-radix-cache-strategy`), … — is already\n covered by the merged generic **`flagSelects`** axis: declare it in the\n config (a list of `{ id, title, stripPrefixes, options }`; see the Qwen3.5\n mamba example), **no engine PR**. Only a genuinely new control *shape* that\n `flagSelects` can't express would need a one-time generic primitive (never a\n model-named handler) on a separate prior PR (engine-axis.md).\n5. **`github.cookbookModel` must be set** (`/`, e.g.\n `qwen/qwen3.5`) and the block never pruned — without it Submit ↗ mislabels\n as deepseek-v4. The issue template itself needs NO edits (free-form input).\n6. **Strategy tiers are signal-driven.** A cell goes under `low-latency` /\n `high-throughput` ONLY on a signal present in the legacy source (an\n explicit performance toggle, a named recipe, or prose stating the\n operating point); **no signal → `balanced`**. Never derive a slant from\n your own hardware intuition — re-tiering on measured evidence is the\n hardware owner's follow-up PR, not part of a migration\n (dimension-mapping.md §4).\n\n## Workflow (one model = one PR)\n\n### 1. Inventory the legacy assets\n- Read the legacy generator (`docs/src/snippets/autoregressive/-deployment.jsx`)\n end-to-end: every option dimension (radio vs checkbox vs dynamic), every\n gate/SUPPORT matrix, the full emitted command per reachable combo (env\n prefixes, `# Error` pseudo-commands included).\n- Read the legacy MDX: §2 install docker tags → `dockerImages` (pinned, not\n upgraded); §3.2 tips → new §2; §4 invocation examples → new §3 (keep real\n Output Examples verbatim); §5 benchmark blocks → transcribe each measured\n block: deploy command used, bench command (dataset/isl/osl/num-prompts/\n concurrency), P50 (median) TTFT/TPOT, output tok/s, hardware, version string.\n- Inbound-anchor sweep: `grep -rn \"\" docs/ --include='*.mdx'` —\n find links/`#fragments` into this page (`mint broken-links` does NOT check\n fragments). Fix referrers or add `` shims in the same PR.\n- Check the maintainer-provided inventory notes for this model's known quirks —\n but treat them (and the §4 family table) as a **survey snapshot**: re-verify\n every dimension against the live legacy files before mapping. Pages keep\n receiving updates (precedent: Kimi-K2.6's live generator has a speculative\n toggle the 2026-06-10 survey notes lack).\n\n### 2. Design the 5-dim mapping\nApply [references/dimension-mapping.md](references/dimension-mapping.md). Key\ndecision — which legacy toggle becomes the `strategies` dimension: a toggle\nthat **changes other parts of the command** (TP, mem) must (the Playground\ncan't do coupled changes), and so does a toggle the legacy page itself labels\nwith operating-point words — e.g. a `dpattention` radio whose options are\nsubtitled \"Low Latency\" / \"High Throughput\" (GLM-5.1 / Kimi-K2.6 pattern) —\neven when its flags are uncoupled (`--attn-dp-size N` is a pure\nflag add). Any other toggle that only adds/removes its own flags becomes a\nPlayground axis with the flags baked into cells when the legacy default was\nON — EXCEPT parsers:\n`--reasoning-parser`/`--tool-call-parser` are NEVER baked into cells, they are\nPlayground-only (DSv4 convention). **Every legacy control survives as an\ninteractive control** — a dimension or a Playground axis, never a tips-only\nmention — and a model-specific control is **config data, not engine code**\n(MegaMoE W4A4 is all DSv4 config on the existing `moe` axis). It's pure\nconfig whenever it fits an existing axis's data schema. A **titled\nsingle-select that strips a flag family** (Nemotron3's \"KV Cache DType\",\nmamba `--mamba-radix-cache-strategy`, …) fits the merged generic **`flagSelects`**\naxis — so it too is config-only (declare a `flagSelects` list). Only a control\nwhose *shape* `flagSelects` still can't express would need a ONE-TIME generic\nprimitive (never a model-named handler) on a separate PRIOR engine PR, keeping\nthe migration PR data-only (hard rule 4, engine-axis.md). The strategy count follows the page's\noperating points: **one recipe → a single `balanced`; two → `low-latency` +\n`high-throughput`; three → the full trio (the ideal)**. The tiers apply per\n(hw × variant × quant) combination — a single-recipe combination on a\nmulti-strategy page parks under its semantically honest tier (no\nlatency/throughput slant → `balanced`; the page's list is the union). When the\nlegacy toggle is MTP / speculative decoding, the direction is a deterministic\ndefault — apply without asking: **MTP on → `low-latency`, MTP off →\n`high-throughput`** (reversed only with maintainer confirmation). Tier\nplacement is signal-driven (hard rule 6). Never invent a recipe just to fill\nstrategy chips (see dimension-mapping.md §4). Record the outcome as a\n**strategy mapping table** for the PR body — one row per group of\ncombinations sharing the same legacy signal (e.g. \"all GPU combos: MTP\ntoggle → low-latency / high-throughput\"; \"xeon: (none) → balanced\"), with a\none-line rationale each; don't enumerate 60 identical rows. The table is what\nhardware owners sign off on at review.\n\n### 3. Generate the config (codegen, then audit)\n- For >~30 cells, port the legacy `generateCommand()` into a throwaway Node\n script that enumerates combos and emits the `cells:[...]` literal (output\n must stay a pure literal — Mintlify forbids runtime spreads/calls). Apply the\n verified-cell override in the script. See the pilot scripts embedded in the\n worked example of dimension-mapping.md §5.\n- **Independent equivalence audit (required):** extract the ORIGINAL generator\n from git (`git show main:` — NOT `HEAD:`, the migration branch deletes\n the file, see dimension-mapping.md §5 item 7), stub React hooks, run it for\n every combo, and diff token-by-token against the new cells. Expected deltas\n only: the appended `--host {{HOST_IP}}`/`--port {{PORT}}`, the\n engine-injected multi-node trio, the §2 alias rewrites (the entrypoint\n rewrite doesn't appear in cell tokens — cells hold flags only; the audit\n script normalizes it on the legacy side), and the intentional verified-cell\n override. Paste the PASS count + the audit script in the PR body\n (collapsed `
`).\n- Hand-author the non-cells fields per authoring-reference.md. Structural\n self-checks: every cell resolves a `modelNames` key; no `--nnodes/--node-rank/\n --dist-init-addr/--host/--port` literals; every `{{KEY}}` declared; every\n `supportedHardware` id has ≥1 cell.\n\n### 4. Benchmarks file\nOne entry per measured block only (cells without entries already render\n\"pending\" — bare `{match}` stubs are unnecessary). `tokens_per_sec_per_gpu` =\ntotal (in+out) tok/s/GPU = `output tok/s ÷ (tp × nnodes) ×\n(isl+osl)/osl` — stored directly (the card shows it as-is). TTFT/TPOT\ntake the P50 (median) rows; set `config.latencyPercentile` (default `\"P50\"`; use\n`\"Mean\"` only for legacy Mean-recorded data — temporary, being migrated to P50; an\nentry-level `latencyPercentile` overrides the page value per cell).\nPut the workload's\n`num_prompts` into `workload`. **`config.accuracyLabels` is required whenever\nthe benchmarks carry accuracy data** — the engine ships no default eval set\n(#27842), so missing labels means the accuracy rows silently don't render;\nextra context (sample counts, suites that don't fit) goes in the entry's\n`notes`. Zero-measured-data pages: skip the file and the `benchmarks` prop\nentirely, but keep `benchmarkCommands` so ⚡Reproduce still guides users.\n\n### 5. Rewrite the MDX\nFrom `page.mdx.tmpl`: keep the original `title` (nav identity), write a fresh\nSEO `description` (top-level — delete any legacy `metatags.description`), **no\n`tag: NEW`** (a migration is not a launch), **no `mode:`**. Install accordion\ncarries the legacy install content + pinned images. Keep the template's\nDSv4-style strategy bullets — serving semantics first (single-user chat /\ntypical multi-user / batch throughput), trimmed to the strategies the page\nships, plus at most a one-line note on what each strategy changes on this\nmodel; do NOT rewrite them as toggle-/migration-centric explanations.\nLegacy §5 benchmark prose is deleted (numbers → benchmark card, commands →\n⚡Reproduce); legacy prose deploy commands are deleted (doc↔config parity —\nfold their unique flags into §2 tips). Invocation examples + real outputs carry\nover verbatim, but **wrapped in Accordions** — §3 commands and outputs are\ncollapsible (required, DeepSeek-V4 pattern): code in an\n``, output in a following\n``; legacy pages kept them inline.\n\n### 6. Delete the legacy generator\nRemove `docs/src/snippets/autoregressive/-deployment.jsx` and its\nimport. `grep -rn \"-deployment\" docs/` must return nothing (config\nprovenance comments must not name the deleted path). Site wiring needs **no\nchanges**: docs.json path/title unchanged, vendor card + logo already exist.\n\n### 7. Validate\n- `grep -rn '__[A-Z_]*__'` on the new files (no template tokens).\n- `cd docs && mint validate && mint broken-links` (pre-existing breaks on\n main are not yours — say so in the PR).\n- `mint dev` browser smoke: initial selection = the verified cell (first in\n `cells[]`) with green badge; multi-node cells show the injected trio +\n header; AMD cells show env prefixes; Docker mode wraps with the pinned image\n and passes cell env as `--env`; condition-hidden combos grey out; benchmark\n card values; NO parser flags in any Deploy command; Playground parser\n toggles ADD the parser flags (green additions) while spec toggles strike\n the baked spec flags (red); Submit ↗ prefills this model. **Probe pitfall:** drive at most\n ONE programmatic click per evaluation and wait for React to settle —\n multiple clicks in one synchronous script batch and read stale DOM.\n- Token-level audit from step 3 passes.\n\n### 8. PR + review\nOne PR per model. PR body: migration framing, verified policy applied, the\nstrategy mapping table (step 2), the audit PASS count + script, any\ninherited-infeasible combos flagged for re-verification. Then run `/cookbook-review-pr ` and fix findings.\nFYI: docs previews only build for in-repo (`sgl-project/sglang`) branches —\na fork-headed PR is perfectly fine but renders no preview; a maintainer can\nre-push the branch in-repo if a preview is wanted for review.\n\n### 9. Keep this skill current\nAny new convention, engine behavior, or pitfall you discover while migrating\n(naming decisions, audit-script gotchas, review-rule conflicts, …) MUST be fed\nback into this skill — same PR if it's skill-file-only, or an immediate\nfollow-up commit on the skill's branch/PR. The next agent runs on what's\nwritten here, not on your session's context.\n",".agents/skills/cookbook-review-pr/SKILL.md":"---\nname: cookbook-review-pr\ndescription: Review a pull request against the SGLang Cookbook (docs/, Mintlify) contribution checklist — the config-driven format (per-model config + benchmarks JSX consumed by the shared _deployment.jsx / _playground.jsx engines). Run with /cookbook-review-pr .\n---\n\n# Cookbook Review PR\n\nFetch the diff, run the checklist, report what you find. The cookbook is **config-driven**:\nshared engines (`_deployment.jsx`, `_playground.jsx`) with NO model-specific code; each\nmodel is a data `config` (+ optional `benchmarks`) under `src/snippets/configs//`\nplus an MDX page. This checklist targets that layout. Field-schema detail lives in\n`.claude/skills/cookbook-add-model/references/authoring-reference.md` — defer to it rather\nthan restating.\n\n## Usage\n\n```\n/cookbook-review-pr \n```\n\n## Steps\n\n1. `gh pr view --repo sgl-project/sglang --json title,body,files,author,baseRefName,headRefName,commits,reviews`\n2. `gh pr diff --repo sgl-project/sglang`\n3. `gh pr list --repo sgl-project/sglang --state open --search \"\"` (duplicate check)\n4. Run every checklist item against the diff.\n5. Output per-file verdicts + overall recommendation.\n\n## Checklist\n\n### 1. File hygiene\n- A cookbook PR should only touch: `docs/src/snippets/configs//*.jsx`\n (config + benchmarks), `docs/cookbook/**/*.mdx`, `docs/docs.json`,\n `docs/cookbook//intro.mdx` (vendor card), `docs/cards/logos/.png`\n (new vendor only). Flag stray files (`settings.local.json`, lockfiles, IDE configs).\n- Pages must be `.mdx`, not `.md`. Files end with a trailing newline. Check commit history\n for unrelated commits accidentally included.\n- **Engines untouched**: `_deployment.jsx` / `_playground.jsx` should NOT change in a\n model-add PR (adding a model is data-only). Engine edits = a separate axis/feature PR\n (see `cookbook-add-model/references/engine-axis.md`); review them against that checklist.\n\n### 2. Config quality (the per-model `config`)\n- Single `export const config = { ... }` literal — **no** function calls, spreads,\n fragment refs, or IIFE (Mintlify re-evals at hydration → `ReferenceError`).\n- No `!(x in y)` anywhere (Mintlify AST walker crashes) — use `obj.key === undefined`.\n- `supportedHardware` ⊆ `HARDWARE_CATALOG` (in `_deployment.jsx`) ∪ `config.hardware`. A\n model-specific GPU the shared catalog lacks must be declared in `config.hardware`\n (`{id,label,vram,vendor}`), **not** added to the engine catalog — unless the platform is\n genuinely shared, which is an engine change the PR body has to argue (and the\n `cookbook-add-model` hardware table moves with it). A new catalog entry is inert for\n other pages: each filters the catalog by its own `supportedHardware`.\n- `placeholders` declares every `{{KEY}}` used in `curl` or any cell.\n- `modelNames` covers every cell (by `hw|variant|quant` triple or `variant|quant` pair).\n- `strategies` count matches the page's operating points — 1 recipe → a single `balanced`;\n 2 → `low-latency` + `high-throughput`; 3 → the full trio. Tiers apply per\n (hw × variant × quant) combination: a single-recipe combination must park under its\n semantically honest tier (clear slant → that tier, e.g. a workstation card under\n `low-latency`; no slant → `balanced`, e.g. a CPU platform) — **flag a no-slant recipe\n parked under low-latency/high-throughput**. Mixed unions\n like [low-latency, balanced, high-throughput] with per-selection greying are fine. Also\n flag model-specific ids (e.g. `mtp`), and flag an INVERTED speculative mapping — the\n deterministic default is MTP/spec-decoding ON → `low-latency`, OFF → `high-throughput`\n (at saturation the draft+verify overhead outweighs the speedup); the reverse needs an\n explicit maintainer-confirmed justification in the PR. The MDX strategy bullets describe serving semantics\n in the DSv4 style (single-user chat / typical multi-user / batch jobs), not internal\n toggles.\n- `dockerImages` covers the hw ids that have cells (else users hit the `:dev` fallback); a\n `hw|quant` key (resolved before the plain `hw`) is valid when one quant on a shared GPU needs\n a different image (e.g. an FP4 dev build) — don't flag those.\n- `multiNodeHints` present ONLY for hw whose fabric needs manual NIC env (e.g. `gb200`\n NVL72) — NOT every `multi-N` hw (standard-IB DeepEP / Marlin multi-node don't need it).\n Hints render above BOTH run modes, so flag one that reads as docker-only (\"add these\n docker run flags\") — `docker run` flags belong in the hardware entry's\n `multiNodeDockerFlags`, which the engine emits into the Docker command itself.\n- `github.cookbookModel` is set to the model's HF id (`/`). The issue\n template's `model` field is a free-form input prefilled from this value; if the config\n omits the `github` block, the engine falls back to `deepseek-ai/deepseek-v4` and the\n page's submissions get mislabeled.\n- `playgroundFeatures` is opt-OUT: the **general axes ship on every cookbook by default**\n (`attention` TP/CP/DP-Attn, `moe` backend+EP for MoE models, `parsers`, `speculative`,\n `pdDisagg`, `hicache`) — flag a missing general axis unless the model genuinely cannot\n use it. Model-specific axes only where applicable (MegaMoE backend + `megamoeQuant`\n only on Blackwell MoE, gated by `requiresHw`; `hisparse` only DSA-style). Knobs that are\n meaningless for a subset of variants/hw are `disable`d with a reason, not silently live\n (e.g. MoE knobs greyed on dense variants). No empty/stub axes.\n- **No leftover `__TOKEN__`** — the config was stamped from the template and every\n placeholder is filled (`grep -rn '__[A-Z_]*__'` on the new config/benchmarks/MDX returns\n nothing).\n- **All-hardware considered**: every `supportedHardware` id (from the catalog or `config.hardware`) has ≥1 cell OR is a deliberate\n greyed \"coming soon\"; AMD was pruned or kept on purpose (not a leftover template family).\n\n### 3. Cells / 5-dim matrix\n- Every cell `match` has EXACTLY the 5 keys (`hw`, `variant`, `quant`, `strategy`, `nodes`).\n- `env` / `flags` are flat literals (only `{{PLACEHOLDER}}` subst) — no shared\n `commonFlags` reference (Mintlify won't inline it).\n- NO `--nnodes` / `--node-rank` / `--dist-init-addr` literals in multi-node cells\n (the renderer injects them from `match.nodes`).\n- NO literal `--host` / `--port` — use `{{HOST_IP}}` / `{{PORT}}`.\n- NO `--reasoning-parser` / `--tool-call-parser` in any cell — parsers are a\n Playground-only feature added on top of the base command (DSv4 convention);\n flag any cell that bakes them in.\n- Accuracy-degrading flags in cells — runtime quant below the checkpoint\n (e.g. MegaMoE W4A4 — DSv4 gates it behind the Playground's `megamoeQuant`)\n and lossy `--kv-cache-dtype` (e.g. `fp8_e4m3` over a higher-precision-KV\n checkpoint): **flag for explicit maintainer confirmation**. Output quality\n should be exactly what the quant chip declares, so absent a recorded\n sign-off in the PR (e.g. carried verbatim from a measured legacy recipe's\n default command), request the flag move to Playground/tips.\n- Flag order: `--model-path` first (an optional `--trust-remote-code` may precede it —\n the DSv4 cells do), then parallelism, then MoE, then tuning, `--host`/`--port`\n last (the playground's insert anchors assume this).\n- TP/memory sanity: `model_weight_GB / (tp × gpu_mem)` fits with ~20–30% headroom\n (BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5; MoE uses **total** weight, not active params).\n\n### 4. Benchmarks\n- Each `benchmarks[]` entry's `match` tuple corresponds to a real cell.\n- `accuracyLabels` is present whenever the benchmarks carry accuracy data — the engine\n ships NO default eval set; without it the accuracy rows silently don't render.\n `defaultAccuracy` / per-cell `accuracy` / `benchmarkCommands.accuracy` keys all\n ∈ `config.accuracyLabels`.\n- A benchmark's quantization must match a variant actually listed — `(BF16)` on a model\n that only released FP8/FP4 is a factual bug.\n- `benchmarkCommands.speed` is `python3 -m sglang.bench_serving` (the workload), separate\n from the `sglang serve` deploy command, and should carry `--flush-cache`: bench_serving's\n `random` prompts are deterministic, so a warm rerun hits the radix cache and inflates\n throughput — speed numbers are measured cache-cold.\n- `sglang_version` is a real build the author ran (a release, or `dev`/nightly) — not a\n guessed/placeholder value (no leftover `0.0.0`).\n- **Latency percentile**: `config.latencyPercentile` (default `\"P50\"`, or `\"Mean\"`) matches the\n percentile the TTFT/TPOT values actually are — the card renders `TTFT ()`. A benchmarks\n entry may carry its own `latencyPercentile` to override the page value per cell\n (entry → config → `\"P50\"`): on a P50 page, kept legacy Mean cells must set it — a\n `sglang_version` tag alone doesn't convey the percentile. (`\"Mean\"` is temporary — legacy\n data is being re-measured to P50.)\n- **Throughput convention**: `tokens_per_sec_per_gpu` is stored as **total (in+out)/GPU**\n = `output tok/s/GPU × (isl+osl)/osl`, shown by the card as-is. Flag output-only values.\n- **Consistent accuracy harness across entries**: every value under one `accuracyLabels`\n column must be produced by the SAME harness — flag a page that, say, measures one\n platform's GSM8K with a historical raw completion harness and another's with\n `sgl-eval run gsm8k --num-examples 1319` and shows both as one \"GSM8K %\"\n (the scores aren't comparable). Either standardize on one harness (matching\n `benchmarkCommands.accuracy`) or require an explicit per-entry note. Common when folding\n a second contributor's measurements (e.g. an AMD/ROCm PR) into the page.\n\n### 5. Doc ↔ config parity (the #1 finding)\n- Any `sglang serve` command shown in MDX prose (config tips, benchmark section) must\n equal what the engine emits from the corresponding cell — same flags, same order. Drift\n here is the most common review miss.\n\n### 5b. ComfyUI section (diffusion pages)\n- A diffusion page ends with `## . Run in ComfyUI` rendering ``. A\n reader must not have to guess whether the model is reachable from ComfyUI.\n- The `model` prop is a key that exists in `docs/src/snippets/diffusion/comfyui-support.jsx`.\n A model-specific key is only correct when the plugin really treats it specially — an entry\n in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a dedicated\n node. Otherwise the generic `image` / `video` key is the honest one; a model-specific key\n without matching plugin support makes the page claim support that does not exist.\n- Prose describing ComfyUI support inline instead of using the component is a finding: the\n facts drift from the plugin.\n\n### 6. Commands / port\n- Launch uses `sglang serve` — flag any `python -m sglang.launch_server` /\n `python3 -m sglang.launch_server` (deprecated). The engine already emits `sglang serve`;\n guard against prose/cells reintroducing the old launcher.\n- Port `30000` everywhere (launch, curl, client `base_url`, bench) — flag `8000`.\n Launch port must match client/curl port on the same page.\n\n### 7. Frontmatter\n- Every new MDX page has `title:` and a **top-level** `description:` (a real one-line value\n prop, not copied from another vendor) — NOT `metatags.description` (non-canonical; the\n top-level field is what renders as the subtitle and SEO meta — see mintlify-authoring).\n- **No `mode: wide` on a model page** — it hides the right-hand \"On this page\" ToC that every\n other model page has. Leave `mode` unset (the Deploy/Playground panels self-cap at 900px, so\n the default column holds them fine). `mode: wide` belongs only on category `intro.mdx` grids.\n- `tag: NEW` only for genuine new launches; when one is added, stale `tag: NEW` on older\n pages should be dropped in the same PR (`grep -RlE \"^tag: NEW\" docs/cookbook/`).\n- MDX imports BOTH `Deployment` and `Playground` from `/src/snippets/...` (absolute).\n- Deploy heading slugs to `deployment` (or `deploy`), Playground to `playground` — so\n \"↑ Switch base\" and \"Open the Playground →\" scroll. No numbered headings for these two.\n\n### 8. Navigation & homepage\n- New page → `docs/docs.json` updated: under the right vendor group inside\n `navigation` → Cookbook → Autoregressive Models, root-relative, **no `.mdx`**:\n `cookbook///`.\n- Homepage `` in `docs/cookbook//intro.mdx` points to the vendor's\n flagship; new vendors get a new `` + a logo at `docs/cards/logos/.png` —\n **940×525 RGBA transparent, icon-only (no wordmark)**, lowercase filename, tracked via\n `git add -f` (`*.png` is gitignored repo-wide). Card order matches the `docs.json` nav order.\n- Don't change `docs/cookbook/intro.mdx` for individual model adds (top-level only).\n\n### 9. Links & factual\n- HuggingFace URLs resolve to a real model. License section matches the actual HF license\n (don't copy from another model). Docker images from `lmsysorg/sglang`; no `sgl-project-dev`.\n The image **tag** is a real build (a release the author ran, or `:dev`/nightly) — not a\n guessed version.\n- Internal links root-relative, no extension (`/cookbook/.../`); flag `.md`/`.mdx`\n or `../`-relative links. `docs.sglang.io` is canonical.\n- No Google-Drive image links (don't render). Shell placeholders are `export VAR=`,\n not `${VAR}` (a bash no-op).\n- **Parser ids must exist in the code registries** on the PR's target branch: every\n `--reasoning-parser X` / `--tool-call-parser Y` named in prose or in\n `playgroundFeatures.parsers` flags is a registered key in\n `python/sglang/srt/parser/reasoning_parser.py` (DetectorMap) /\n `python/sglang/srt/function_call/function_call_parser.py` (ToolCallParserEnum) —\n prose naming a near-miss id (e.g. the reasoning id where the tool id differs) is a\n factual bug. `--…-parser auto` is acceptable ONLY if the template-detection rules\n (`python/sglang/srt/parser/template_detection.py`) actually resolve THIS model's\n chat template to the right parser — no rule match means auto silently disables the\n parser; when in doubt require explicit ids (the DSv4 page pins explicit ids).\n\n### 9b. MDX authoring (Mintlify) — detail in `cookbook-add-model/references/mintlify-authoring.md`\n- **Forbidden syntax**: no Docusaurus admonitions (`:::`), `@site`/`@theme`, GitHub alert\n blocks (`> [!NOTE]`), markdown **pipe tables** (use JSX ``), inline `
`,\n or unknown components. ``/`` only on category `intro.mdx`, not model pages.\n- Code fences are **labeled** (e.g. `python Example` / `bash Command` / `text Output` after\n the opening fence); a fenced block nested inside another uses four backticks outside.\n- §3 commands and outputs are **collapsible** (DeepSeek-V4 pattern): every runnable\n example wrapped in an ``, its real output in a following\n `` (`Pending update...` only with user\n acknowledgement). Flag bare/inline example blocks and `**Output Example:**` headings.\n- Reasoning-parser example matches the parser's **output shape**: separate-field\n (`reasoning_content` + `content`) vs inline `` tags parsed out of `content`.\n- No hardcoded sampling params (`temperature` / `top_p`) in sample code (SGLang uses\n `generation_config.json` defaults); listing them in §1 informationally is fine.\n\n### 10. Quantization rules\n- **NVFP4** checkpoints are Blackwell-only (B200/B300/GB300) — never AMD. An AMD FP4 cell\n is legitimate ONLY when the vendor published an **MXFP4** checkpoint for it (e.g.\n `amd/Qwen3.5-397B-A17B-MXFP4` on MI355X) — verify the HF repo resolves; otherwise the\n AMD FP4 chip must be absent/`disabled`.\n- BF16 / FP8 work on NVIDIA and AMD. `--kv-cache-dtype fp8_e4m3` in a cell is an\n accuracy-degrading flag — see §3 (needs explicit maintainer sign-off; default\n home is Playground/tips).\n\n### 11. Scope\n- Changes match the PR title. Flag global changes hiding behind a platform-specific title\n (e.g. an \"H200 FP8\" PR that adds a flag to ALL cells). Unmentioned side-fixes belong in\n the PR body.\n\n### 12. Duplicate PRs\n- Another open PR for the same model? Flag it; compare completeness; note merge-conflict\n risk on `docs.json` + the vendor card; flag a superseded older PR by the same author.\n\n### 13. Build / validate\n```bash\ncd docs\nmint validate\nmint broken-links\n```\nOptional: `mint dev` for a visual smoke test.\n\n### 14. Reviewer feedback\n- `gh api repos/sgl-project/sglang/pulls//comments` — have prior reviewer requests been\n addressed? Unresolved requested-changes should be flagged.\n\n### 15. Grammar & spelling\n- Check added/changed prose for typos and grammar (e.g. \"recommend\" vs \"recommended\").\n Flag each with the exact wrong text + correction.\n\n## Output\n\nPer file:\n- ✅ PASS\n- ⚠️ ISSUE: \\\n- 🔴 BLOCK: \\\n\nOverall: **APPROVE** / **REQUEST CHANGES** / **BLOCKED**\n",".agents/skills/debug-cuda-crash/SKILL.md":"---\nname: debug-cuda-crash\ndescription: Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging\n---\n\n# Tutorial: Debugging CUDA Crashes with Kernel API Logging\n\nThis tutorial shows you how to debug CUDA crashes and errors in SGLang using the `@debug_kernel_api` logging decorator.\n\n## Goal\n\nWhen your code crashes with CUDA errors such as illegal memory access, device-side assert, out-of-bounds, or NaN/Inf, use kernel API logging to:\n- Capture input tensors BEFORE the crash occurs\n- Understand what data caused the problem\n- Track tensor shapes, dtypes, and values through the call boundary that triggered the crash\n- Detect numerical issues such as NaN, Inf, or obviously wrong shapes\n\n## Why Use Kernel API Logging?\n\n**Problem**: CUDA errors often crash the program before normal debugging output is flushed.\n\n**Solution**: SGLang's `@debug_kernel_api` decorator logs inputs before execution, so you can still see what caused the crash even after the program aborts.\n\n## What Is Covered?\n\nThe current logging coverage focuses on the highest-value kernel boundaries in SGLang:\n- Custom ops registered through `register_custom_op(...)`\n- External custom ops registered through `register_custom_op_from_extern(...)`\n- LLM attention, linear, quantization, and multi-platform wrapper entry points\n- Diffusion attention impl, linear, rotary, and custom-op wrapper entry points\n- Selected direct `torch.ops.sglang.*` hotspots and model-specific bypasses\n\nThis means the logging is useful for both LLM and diffusion kernel debugging, but it does not automatically cover every pure PyTorch call in the repository.\n\n## Step 1: Enable Kernel API Logging\n\n### Basic Logging (Function Names Only)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=1\nexport SGLANG_KERNEL_API_LOGDEST=stdout\n\npython my_script.py\n```\n\nOutput:\n```\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: RMSNorm.forward\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.custom_op.fused_inplace_qknorm\n```\n\nThis is a real level-1 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Detailed Logging (Inputs with Metadata)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nOutput in `debug.log`:\n```\n================================================================================\n[2026-03-19 00:47:30] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[0]=QKVParallelLinear(\n repr=QKVParallelLinear(in_features=1024, output_features=4096, bias=False, tp_size=1, gather_output=False)\n )\n arg[1]=Tensor(\n shape=(1, 1024)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n arg[2]=None\nOutput:\n return=Tensor(\n shape=(1, 4096)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n```\n\nThis is a real level-3 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Full Logging (With Tensor Statistics)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nAdditional output:\n```\n================================================================================\n[2026-03-19 01:00:42] SGLang Kernel API Call: diffusion.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[1]=Tensor(\n shape=(1, 77, 768)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-27.250000\n max=28.500000\n mean=0.011723\n nan_count=0\n inf_count=0\n )\nOutput:\n return=Tensor(\n shape=(1, 77, 2304)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-8.937500\n max=9.375000\n mean=0.009460\n nan_count=0\n inf_count=0\n )\n```\n\nThis is a real level-5 excerpt captured from `black-forest-labs/FLUX.1-dev`.\n\n### Crash-Safe Dumps (Inputs Saved Before Execution)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\n\npython my_script.py\n```\n\nAt level 10, SGLang saves the inputs before execution. If the kernel crashes, the dump directory still contains the inputs and exception metadata.\n\nIf CUDA graph capture is active, tensor dumps are skipped automatically to avoid capture-time CUDA errors. In that case, you still get the kernel API call log, but not `inputs.pt` / `outputs.pt`.\n\nLevel-10 dumps are best understood as crash-safe call snapshots. They always preserve the observed call boundary. They do not guarantee one-click replay for every method, because some methods depend on module state that is not serialized into the dump.\n\nReal level-10 dump layout from `Qwen/Qwen3-0.6B`:\n\n```text\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/inputs.pt\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/metadata.json\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/outputs.pt\n```\n\nReal `metadata.json` excerpt:\n\n```json\n{\n \"function_name\": \"RotaryEmbedding.forward\",\n \"timestamp\": \"20260319_004821_182\",\n \"process_id\": 919286,\n \"execution_status\": \"completed\",\n \"input_tensor_keys\": [\"arg_0\", \"arg_1\", \"arg_2\"],\n \"output_tensor_keys\": [\"result_0\", \"result_1\"]\n}\n```\n\n## Step 2: Reproduce an LLM CUDA Crash\n\nCreate a temporary reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_llm_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.srt.utils.custom_op import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(indices, table):\\\\n\"\n \" return torch.empty((*indices.shape, table.shape[-1]), device=table.device, dtype=table.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_llm_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_llm_cuda_crash(indices, table):\\\\n\"\n \" out = F.embedding(indices, table)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"table = torch.randn(4, 8, device='cuda', dtype=torch.float16)\\\\n\"\n \"indices = torch.tensor([0, 7], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_llm_cuda_crash(indices, table)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level1.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nWhat to expect:\n- The script exits with a CUDA `device-side assert`\n- The log still contains the last API boundary before the crash\n\nTry the same example at level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level3.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow the log shows tensor metadata before the crash.\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_llm_level10_dumps \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow you should see:\n- A log entry for `sglang.custom_op.mock_llm_cuda_crash`\n- A dump directory with `inputs.pt`\n- `metadata.json` showing `execution_status: \"exception\"`\n- No `outputs.pt`, because the kernel crashed before producing output\n\nFor real-model success-path level-10 dumps, it is often easier to temporarily disable CUDA graph and piecewise CUDA graph for the debug run.\n\n## Step 3: Reproduce a Diffusion CUDA Crash\n\nCreate a temporary diffusion-side reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_diffusion_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.multimodal_gen.runtime.layers.utils import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(positions, cache):\\\\n\"\n \" return torch.empty((*positions.shape, cache.shape[-1]), device=cache.device, dtype=cache.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_diffusion_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_diffusion_cuda_crash(positions, cache):\\\\n\"\n \" out = F.embedding(positions, cache)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"cache = torch.randn(4, 64, device='cuda', dtype=torch.float16)\\\\n\"\n \"positions = torch.tensor([0, 9], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_diffusion_cuda_crash(positions, cache)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level1.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level3.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_diffusion_level10_dumps \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nIf your local environment has unrelated FlashInfer import issues, resolve them in the shell before running the example. The example itself does not set any `FLASHINFER_*` environment variable.\n\n## Step 4: Multi-Process Debugging\n\nWhen running with multiple GPUs or worker processes, use `%i` in the log path:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\n\ntorchrun --nproc_per_node=4 my_script.py\n```\n\nThis creates separate logs such as:\n- `debug_rank_12345.log`\n- `debug_rank_12346.log`\n- `debug_rank_12347.log`\n- `debug_rank_12348.log`\n\nReal multi-process example from a 2-GPU `Qwen/Qwen2.5-0.5B-Instruct` run:\n\n```text\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950201.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950349.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950350.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950351.log\n```\n\nYou should usually do the same for level-10 dump directories:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps_%i\n```\n\nThis avoids multiple ranks writing into the same dump directory tree.\n\n## Step 5: Filter Level-10 Dumps\n\nIf level 10 is too noisy, restrict dumps to specific APIs:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\nexport SGLANG_KERNEL_API_DUMP_INCLUDE='sglang.custom_op.*'\nexport SGLANG_KERNEL_API_DUMP_EXCLUDE='*.fake_impl'\n```\n\n`SGLANG_KERNEL_API_DUMP_INCLUDE` and `SGLANG_KERNEL_API_DUMP_EXCLUDE` use shell-style wildcard matching.\n\n## Step 6: Common CUDA Errors and What to Check\n\n### Illegal Memory Access or Device-Side Assert\n\n**Typical errors**:\n```\nRuntimeError: CUDA error: an illegal memory access was encountered\ntorch.AcceleratorError: CUDA error: device-side assert triggered\n```\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck in the logs:\n- ✅ Tensor shapes\n- ✅ Tensor dtypes\n- ✅ CUDA vs CPU device placement\n- ✅ Tensor stride / contiguity\n- ✅ Whether the failing call has inputs logged but no outputs logged\n\nTypical shape-mismatch pattern:\n\n```text\nSGLang Kernel API Call: ...\narg[0]=Tensor(shape=(..., 128), ...) # ✅ expected dimension\narg[1]=Tensor(shape=(..., 64), ...) # ❌ mismatch\n```\n\nThis often points to head-dim, hidden-dim, or cache-layout mismatch rather than a random CUDA failure.\n\n### NaN or Inf\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nCheck:\n- `min`\n- `max`\n- `mean`\n- `nan_count`\n- `inf_count`\n\nTypical bad pattern:\n\n```text\nTensor(\n ...\n min=-1234567.000000 # ❌ suspiciously large\n max=9876543.000000 # ❌ suspiciously large\n mean=nan # ❌ bad\n nan_count=128 # ❌ found NaNs\n inf_count=0 # ✅ no Infs here\n)\n```\n\nThis usually means the bad values were already present before the crashing kernel.\n\n### Out of Memory\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck:\n- Unexpectedly large tensor shapes\n- Batch size\n- Sequence length\n- Frame count or image resolution in diffusion workloads\n\nAlso check whether a supposedly per-token or per-frame tensor accidentally became full-sequence or full-image sized.\n\nTypical bad pattern:\n\n```text\nTensor(\n shape=(1024, 8192, 128, 128) # ❌ way too large\n ...\n)\n```\n\n### Example: Spot a Shape Bug from the Log\n\nSuppose the failing API log looks like this:\n\n```text\n[2026-03-19 00:47:30] SGLang Kernel API Call: RotaryEmbedding.forward\nPositional input arguments:\n arg[0]=Tensor(shape=(1, 8), dtype=torch.int64, ...)\n arg[1]=Tensor(shape=(1, 8, 8, 256), dtype=torch.bfloat16, ...) # ✅ query\n arg[2]=Tensor(shape=(1, 8, 4, 64), dtype=torch.bfloat16, ...) # ❌ key head_dim mismatch\n```\n\nWhat this tells you:\n- ✅ positions look reasonable\n- ✅ query looks plausible\n- ❌ key last dimension is inconsistent with the expected rotary/head dimension\n\nThat usually means the bug is in projection layout, head packing, or cache format rather than in the rotary kernel itself.\n\n## Step 7: Combine with compute-sanitizer\n\nFor harder bugs, combine kernel API logging with CUDA memory checking:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncompute-sanitizer --tool memcheck python3 /tmp/sglang_llm_crash.py\n```\n\nUse `debug.log` to see the exact inputs that reached the crashing API boundary.\n\nTypical `compute-sanitizer` output:\n\n```text\n========= COMPUTE-SANITIZER\n========= Invalid __global__ write of size 4 bytes\n========= at 0x1234 in SomeKernel\n========= by thread (256,0,0) in block (10,0,0)\n========= Address 0x... is out of bounds\n```\n\nUse the sanitizer output to identify the failing kernel and use `debug.log` to identify the exact tensors that reached the API boundary right before it.\n\nIf you need more synchronous host-side error reporting, you can try `CUDA_LAUNCH_BLOCKING=1` as a separate follow-up experiment. It is not part of the default workflow because it changes execution timing and can hide concurrency-related behavior.\n\n## Step 8: Combine with cuda-gdb\n\nFor crashes that need a stack trace instead of only memory diagnostics:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncuda-gdb --args python3 /tmp/sglang_llm_crash.py\n```\n\nInside `cuda-gdb`:\n\n```text\n(cuda-gdb) run\n(cuda-gdb) where\n```\n\nThen correlate the backtrace with `debug.log`.\n\n## Step 9: Kernel-Level Debugging with printf()\n\nWhen you own the CUDA kernel, `printf()` is still useful for narrowing down bad indices, bad launch geometry, or broken state propagation.\n\nBasic pattern:\n\n```cpp\n__global__ void MyKernel(const float* input, float* output, int n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n\n if (threadIdx.x == 0 && blockIdx.x == 0) {\n printf(\"n=%d input0=%f\\n\", n, input[0]);\n }\n\n if (idx < n) {\n output[idx] = input[idx] * 2.0f;\n }\n}\n```\n\nAfter launch, force the output to flush:\n\n```python\nmy_kernel(...)\ntorch.cuda.synchronize()\n```\n\nFor warp-specialized kernels, do not blindly print only on `threadIdx.x == 0`. Pick one representative thread per warp or per specialization group instead.\n\n### Warp-Specialized Kernels: Choosing the Right Print Thread\n\nProblem:\n- `threadIdx.x == 0` only prints from the first warp in the block\n- for warp-specialized kernels, that often misses the warp or group that is actually wrong\n\nBetter pattern:\n\n```cpp\n__global__ void WarpSpecializedKernel(...) {\n // Example: first lane of each warp\n if ((threadIdx.x % 32) == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n }\n}\n```\n\nOr, if the kernel is organized in larger specialization groups, print once per group instead of once per block.\n\nCommon mistake:\n\n```cpp\n// Only warp 0 prints\nif (threadIdx.x == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n}\n```\n\n### Quick Reference\n\n| Kernel Type | Print Condition | Notes |\n|----------|----------|-------------|\n| Simple kernel | `threadIdx.x == 0` | One thread per block is usually enough |\n| Warp-specialized kernel | one representative lane per warp | e.g. `threadIdx.x % 32 == 0` |\n| Group-specialized kernel | one representative lane per group | choose based on the kernel's scheduling layout |\n\n### Other Kernel Debugging Tools\n\n```cpp\nassert(value >= 0.0f && \"value must be non-negative\");\nstatic_assert(BLOCK_SIZE % 32 == 0, \"BLOCK_SIZE must be warp aligned\");\n```\n\n## Environment Variables Reference\n\n| Variable | Values | Description |\n|----------|--------|-------------|\n| `SGLANG_KERNEL_API_LOGLEVEL` | `0` | No logging (default) |\n| | `1` | Function names only |\n| | `3` | Inputs and outputs with metadata |\n| | `5` | Level 3 plus tensor statistics |\n| | `10` | Level 5 plus crash-safe tensor dumps |\n| `SGLANG_KERNEL_API_LOGDEST` | `stdout` | Log to stdout |\n| | `stderr` | Log to stderr |\n| | `` | Log to file |\n| | `log_%i.txt` | `%i` expands to process ID |\n| `SGLANG_KERNEL_API_DUMP_DIR` | `` | Directory for level-10 dumps |\n| `SGLANG_KERNEL_API_DUMP_INCLUDE` | wildcard list | Only dump matching API names |\n| `SGLANG_KERNEL_API_DUMP_EXCLUDE` | wildcard list | Skip matching API names |\n\n## Best Practices\n\n### 1. Start with Level 3\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nLevel 3 is usually enough to catch wrong shapes, wrong dtypes, and wrong devices.\n\n### 2. Use Level 5 for Numerical Issues\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nUse it when you suspect NaN or Inf values.\n\n### 3. Use Level 10 for Crash Reproduction\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\n```\n\nThis is the most useful mode when the process crashes before you can inspect live tensors.\n\nIf you need successful input/output dumps from a real model run, temporarily disable CUDA graph for that debug session.\n\nWhen level 10 is too noisy, pair it with `SGLANG_KERNEL_API_DUMP_INCLUDE` / `SGLANG_KERNEL_API_DUMP_EXCLUDE` instead of dumping every covered API.\n\n### 4. Log to File for Crashes\n\n```bash\nexport SGLANG_KERNEL_API_LOGDEST=crash.log\n```\n\nFile logs are safer than stdout when the process aborts.\n\n### 5. Disable Logging in Production\n\n```bash\nunset SGLANG_KERNEL_API_LOGLEVEL\n```\n\nWhen disabled, the decorator returns the original callable and adds no runtime logging overhead.\n\n## Troubleshooting\n\n### No Logs Appear\n\nCheck:\n1. `echo $SGLANG_KERNEL_API_LOGLEVEL`\n2. `echo $SGLANG_KERNEL_API_LOGDEST`\n3. Whether the failing path goes through a covered API boundary\n\n### Too Much Output\n\nReduce the level:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\n### Statistics Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nstatistics=[skipped: CUDA graph capture in progress]\n```\n\nThat is expected. Level-5 statistics are intentionally skipped during CUDA graph capture to avoid synchronization side effects.\n\n### Tensor Dumps Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nTensor dump skipped: CUDA graph capture in progress\n```\n\nThat is also expected. Level-10 dumps require copying tensors to CPU, which is not allowed during CUDA graph capture.\n",".agents/skills/env-var-conventions/SKILL.md":"# env-var-conventions\n\nPath: `.agents/skills/env-var-conventions/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/env-var-conventions/SKILL.md)",".agents/skills/generate-profile/SKILL.md":"# generate-profile\n\nPath: `.agents/skills/generate-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/generate-profile/SKILL.md)",".agents/skills/kernel-organization/SKILL.md":"# kernel-organization\n\nPath: `.agents/skills/kernel-organization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kernel-organization/SKILL.md)",".agents/skills/kl-consistency-test/SKILL.md":"# kl-consistency-test\n\nPath: `.agents/skills/kl-consistency-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kl-consistency-test/SKILL.md)",".agents/skills/large-class-style/SKILL.md":"# large-class-style\n\nPath: `.agents/skills/large-class-style/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/large-class-style/SKILL.md)",".agents/skills/llm-torch-profiler-analysis/SKILL.md":"# llm-torch-profiler-analysis\n\nPath: `.agents/skills/llm-torch-profiler-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/SKILL.md)",".agents/skills/mechanical-refactor-verify/SKILL.md":"# mechanical-refactor-verify\n\nPath: `.agents/skills/mechanical-refactor-verify/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/SKILL.md)",".agents/skills/mechanical-refactor-verify/guide-modify-skill.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-modify-skill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-modify-skill.md)",".agents/skills/processor-model-parity/SKILL.md":"# processor-model-parity\n\nPath: `.agents/skills/processor-model-parity/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/processor-model-parity/SKILL.md)",".agents/skills/scripted-runtime-notes/SKILL.md":"# scripted-runtime-notes\n\nPath: `.agents/skills/scripted-runtime-notes/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/scripted-runtime-notes/SKILL.md)",".agents/skills/sglang-bisect-ci-regression/SKILL.md":"# sglang-bisect-ci-regression\n\nPath: `.agents/skills/sglang-bisect-ci-regression/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-bisect-ci-regression/SKILL.md)",".agents/skills/sglang-cherrypick/SKILL.md":"# sglang-cherrypick\n\nPath: `.agents/skills/sglang-cherrypick/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-cherrypick/SKILL.md)",".agents/skills/sglang-prod-incident-triage/SKILL.md":"# sglang-prod-incident-triage\n\nPath: `.agents/skills/sglang-prod-incident-triage/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/SKILL.md)",".agents/skills/sglang-runtime-context/SKILL.md":"# sglang-runtime-context\n\nPath: `.agents/skills/sglang-runtime-context/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-runtime-context/SKILL.md)",".agents/skills/speculative-naming/SKILL.md":"# speculative-naming\n\nPath: `.agents/skills/speculative-naming/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/speculative-naming/SKILL.md)",".agents/skills/write-sglang-test/SKILL.md":"# write-sglang-test\n\nPath: `.agents/skills/write-sglang-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/write-sglang-test/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md":"# sglang-diffusion-add-model\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md":"# sglang-diffusion-benchmark-profile\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md":"# sglang-diffusion-cache-dit\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md":"# sglang-diffusion-modelopt-quant\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md":"# sglang-diffusion-performance\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md)","docs/AGENTS.md":"# docs - AI Agent Protocol & Instructions\n\nPath: `docs/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/docs/AGENTS.md)","sgl-router/ARCHITECTURE.md":"# sgl-router - System Architecture & Module Boundaries\n\nPath: `sgl-router/ARCHITECTURE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/sgl-router/ARCHITECTURE.md)",".claude/rules/comment-style.md":"# Rules Skill\n\nPath: `.claude/rules/comment-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/comment-style.md)",".claude/rules/cookbook-first.md":"# Rules Skill\n\nPath: `.claude/rules/cookbook-first.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/cookbook-first.md)",".claude/rules/forward-batch-init-new-purity.md":"# Rules Skill\n\nPath: `.claude/rules/forward-batch-init-new-purity.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/forward-batch-init-new-purity.md)",".claude/rules/general-code-style.md":"# Rules Skill\n\nPath: `.claude/rules/general-code-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/general-code-style.md)",".claude/rules/kernel-benchmark.md":"# Rules Skill\n\nPath: `.claude/rules/kernel-benchmark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/kernel-benchmark.md)",".claude/rules/modify-component-must-read.md":"# Rules Skill\n\nPath: `.claude/rules/modify-component-must-read.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/modify-component-must-read.md)",".claude/rules/no-dataclasses.md":"# Rules Skill\n\nPath: `.claude/rules/no-dataclasses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-dataclasses.md)",".claude/rules/no-getattr-defensive.md":"# Rules Skill\n\nPath: `.claude/rules/no-getattr-defensive.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-getattr-defensive.md)",".claude/rules/prefer-sgl-eval.md":"# Rules Skill\n\nPath: `.claude/rules/prefer-sgl-eval.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/prefer-sgl-eval.md)",".claude/rules/schedule-batch-out-of-place-mutation.md":"# Rules Skill\n\nPath: `.claude/rules/schedule-batch-out-of-place-mutation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/schedule-batch-out-of-place-mutation.md)",".claude/rules/unit-test-admission.md":"# Rules Skill\n\nPath: `.claude/rules/unit-test-admission.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/unit-test-admission.md)","tools/sglang-simulator/README.md":"# sglang-simulator Documentation\n\nPath: `tools/sglang-simulator/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/README.md)",".agents/skills/ci-test-audit/action-items.md":"# Ci-test-audit Skill\n\nPath: `.agents/skills/ci-test-audit/action-items.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/action-items.md)",".agents/skills/mechanical-refactor-verify/guide-construct-proof.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-construct-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-construct-proof.md)",".agents/skills/mechanical-refactor-verify/guide-split.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-split.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-split.md)",".agents/skills/mechanical-refactor-verify/guide-verify-proof.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-verify-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-verify-proof.md)",".agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md)",".agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md)","examples/profiler/nsys_profile_tools/README.md":"# nsys_profile_tools Documentation\n\nPath: `examples/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/examples/profiler/nsys_profile_tools/README.md)","tools/sglang-simulator/examples/README.md":"# examples Documentation\n\nPath: `tools/sglang-simulator/examples/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/README.md)",".agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml":"# Agents Skill\n\nPath: `.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml)",".agents/skills/clean-startup-log/references/noise-sources.md":"# References Skill\n\nPath: `.agents/skills/clean-startup-log/references/noise-sources.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/references/noise-sources.md)",".agents/skills/cookbook-add-model/references/authoring-reference.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/authoring-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/authoring-reference.md)",".agents/skills/cookbook-add-model/references/diffusion-authoring.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/diffusion-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/diffusion-authoring.md)",".agents/skills/cookbook-add-model/references/engine-axis.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/engine-axis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/engine-axis.md)",".agents/skills/cookbook-add-model/references/mintlify-authoring.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/mintlify-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/mintlify-authoring.md)",".agents/skills/cookbook-add-model/references/vendor-logo.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/vendor-logo.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/vendor-logo.md)",".agents/skills/cookbook-migrate-model/references/dimension-mapping.md":"# References Skill\n\nPath: `.agents/skills/cookbook-migrate-model/references/dimension-mapping.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/references/dimension-mapping.md)",".agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md)",".agents/skills/llm-torch-profiler-analysis/references/heuristics.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/heuristics.md)",".agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md)",".agents/skills/llm-torch-profiler-analysis/references/source-map.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/source-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/source-map.md)",".agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md)",".agents/skills/sglang-prod-incident-triage/references/case-studies.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/case-studies.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/case-studies.md)",".agents/skills/sglang-prod-incident-triage/references/decision-tree.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/decision-tree.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/decision-tree.md)",".agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md)",".agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md)","tools/sglang-simulator/examples/assets/replay_table.json":"# Assets Skill\n\nPath: `tools/sglang-simulator/examples/assets/replay_table.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/replay_table.json)","tools/sglang-simulator/examples/sim_configs/aic_silicon.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_silicon.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_silicon.json)","tools/sglang-simulator/examples/sim_configs/aic_sol.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_sol.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_sol.json)","tools/sglang-simulator/examples/sim_configs/ml.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/ml.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/ml.json)","tools/sglang-simulator/examples/sim_configs/replay.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/replay.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/replay.json)","tools/sglang-simulator/examples/workloads/sharegpt-example.json":"# Workloads Skill\n\nPath: `tools/sglang-simulator/examples/workloads/sharegpt-example.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/workloads/sharegpt-example.json)","tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json)","tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json)","tools/sglang-simulator/test/assets/qwen3-8b/config.json":"# Qwen3-8b Skill\n\nPath: `tools/sglang-simulator/test/assets/qwen3-8b/config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/test/assets/qwen3-8b/config.json)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md)","tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md":"# collect_baseline Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md)","tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md":"# train_ml_time_predictor Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md)"},"files":{".agents/skills/ci-test-audit/SKILL.md":"---\nname: ci-test-audit\ndescription: Audit the existing test tree and CI configuration for improvements, using a catalog of patterns previously applied in this repo. Use when asked to audit tests or CI, shrink CI time or cost, find redundant or misplaced tests, clean up a test group, or review whether a CI change follows established practice.\n---\n\n# CI / Test Audit\n\nAudits what is already in `test/`, `.github/` and `scripts/ci/`. For writing a new test see\n[write-sglang-test](../write-sglang-test/SKILL.md); for how the pipeline dispatches and\ngates work see [ci-workflow-guide](../ci-workflow-guide/SKILL.md).\n\n## How to use it\n\nRead [action-items.md](action-items.md) first. It is a catalog of patterns that have\nbeen applied to this repo, each with a way to spot it and example PRs. The catalog is the substance of this skill; everything below is only scaffolding.\n\nThe user says what to audit -- a test group, a stage, a workflow file, or a complaint\nlike \"this suite is slow\" -- and how thoroughly. Read that scope and decide which\npatterns apply.\n\nAn audit is only useful if it is specific. \"This suite could be trimmed\" is not a\nfinding; \"`TestFooLargePage` is a strict subset of `TestFooRetractLargePage`, drop it\"\nis. Every finding needs the file, the pattern it matches, and what to do.\n\n## Reporting\n\nGroup findings by confidence in the claim, not by pattern id. Lead with the ones where\nthe evidence is in the file you just read; keep the speculative ones separate and say\nwhat you would need to check. Skip anything you are not reasonably sure about -- a long\nlist of maybes costs more to triage than it saves.\n\nPropose, do not apply. Deleting a test, moving a registration to another stage, and\nchanging a threshold are all decisions for the user. Land them as separate changes,\nsince a trim and a threshold change fail for different reasons.\n\n## Keeping the catalog current\n\nExample PRs age. When a pattern shows up in newer work, or a new pattern recurs, update\n`action-items.md`: the pattern text should stay general, the examples are\nreplaceable evidence.\n",".agents/skills/add-sgl-kernel/SKILL.md":"---\nname: add-sgl-kernel\ndescription: Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)\n---\n\n# Tutorial: Adding a New Kernel to `sgl-kernel` (AOT / Heavyweight)\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as an AOT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float)\n- Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `python/sglang/kernels/aot/include/utils.h`)\n\n## Two rules of thumb (must follow)\n\n1. **Prefer `python/sglang/kernels/jit` first** when the kernel does **not** depend on CUTLASS or another large C++ project. This is the default path for lightweight kernels that benefit from rapid iteration.\n2. **Prefer `sgl-kernel`** when the kernel **does** depend on CUTLASS or another large C++ project, or when it should be part of the AOT wheel / torch op registration flow.\n3. **Exception**: if the dependency is `flashinfer`, or CUTLASS that is already provided through `flashinfer`, the kernel can still be implemented as `jit_kernel`.\n\nIn addition, every new kernel must ship with:\n\n- **Tests** (pytest)\n- **A benchmark script** (`marker.do_bench` from `sglang.kernels.jit.benchmark` — it works for any callable, not only JIT kernels)\n\n---\n\n## Repository integration map\n\nYou will typically touch these files/areas:\n\n- Implementation: `python/sglang/kernels/aot/csrc/elementwise/scale.cu` (pick the right subdirectory)\n- Public declarations: `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- Torch extension registration: `python/sglang/kernels/aot/csrc/common_extension.cc`\n- Build: `python/sglang/kernels/aot/CMakeLists.txt` (`set(SOURCES ...)`)\n- Python API: `python/sglang/kernels/aot/python/sgl_kernel/` and `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n- Tests: `python/sglang/kernels/aot/tests/test_scale.py`\n- Benchmarks: `python/sglang/kernels/aot/benchmark/bench_scale.py`\n\n---\n\n## Step 1: Implement the kernel in `csrc/`\n\nPick the right subdirectory:\n\n- `csrc/elementwise/` — for element-wise ops (our example)\n- `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories\n\nCreate `python/sglang/kernels/aot/csrc/elementwise/scale.cu`:\n\n```cpp\n#include \n#include \n#include \n\n#include \"utils.h\" // DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16\n\n// scale_kernel: out[i] = input[i] * factor\n// Supports float, half (__half), __nv_bfloat16 via template T\ntemplate \n__global__ void scale_kernel(T* __restrict__ out,\n const T* __restrict__ input,\n float factor,\n int64_t n) {\n int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x;\n if (idx < n) {\n out[idx] = static_cast(static_cast(input[idx]) * factor);\n }\n}\n\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor) {\n TORCH_CHECK(input.is_cuda(), \"input must be a CUDA tensor\");\n TORCH_CHECK(input.is_contiguous(), \"input must be contiguous\");\n TORCH_CHECK(out.is_cuda(), \"out must be a CUDA tensor\");\n TORCH_CHECK(out.is_contiguous(), \"out must be contiguous\");\n TORCH_CHECK(out.sizes() == input.sizes(), \"out and input must have the same shape\");\n TORCH_CHECK(out.scalar_type() == input.scalar_type(),\n \"out and input must have the same dtype\");\n\n const int64_t n = input.numel();\n const int threads = 256;\n const int blocks = (n + threads - 1) / threads;\n\n const cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n const at::cuda::OptionalCUDAGuard device_guard(device_of(input));\n\n // Dispatches over float, float16, bfloat16\n DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {\n scale_kernel<<>>(\n static_cast(out.data_ptr()),\n static_cast(input.data_ptr()),\n static_cast(factor),\n n);\n cudaError_t status = cudaGetLastError();\n TORCH_CHECK(status == cudaSuccess,\n \"scale_kernel launch failed: \", cudaGetErrorString(status));\n return true;\n });\n}\n```\n\n**Key points:**\n\n- Use `at::Tensor` (PyTorch tensors), `TORCH_CHECK` for validation, `at::cuda::getCurrentCUDAStream()` for stream\n- Keep Python wrappers thin; do shape/dtype/device validation in C++ right around the launch path\n- `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` covers `float`, `half` (FP16), `__nv_bfloat16` (BF16)\n- Add device error checking after every kernel launch\n- If a kernel only works on certain architectures, enforce that with `TORCH_CHECK` and skip logic in tests\n\n---\n\n## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`\n\nEdit `python/sglang/kernels/aot/include/sgl_kernel_ops.h`, add to the elementwise section:\n\n```cpp\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor);\n```\n\n---\n\n## Step 3: Register the op in `csrc/common_extension.cc`\n\nEdit `python/sglang/kernels/aot/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:\n\n```cpp\n// From csrc/elementwise\nm.def(\"scale(Tensor! out, Tensor input, float factor) -> ()\");\nm.impl(\"scale\", torch::kCUDA, &scale);\n```\n\n**Key points:**\n\n- `Tensor!` means in-place / mutable output argument\n- The schema is important for `torch.compile` and for consistent call signatures\n- Keep the torch schema in PyTorch scalar types (`float` here), but note that the C++ launcher signature still needs `double` for scalar arguments accepted by `torch::Library`\n\n---\n\n## Step 4: Add the new source file to `CMakeLists.txt`\n\nEdit `python/sglang/kernels/aot/CMakeLists.txt`, add to `set(SOURCES ...)`:\n\n```cmake\ncsrc/elementwise/scale.cu\n```\n\n**Key points:**\n\n- Keep the list **alphabetically sorted** (the file explicitly requires this)\n- If the kernel has arch constraints, reflect that in tests/benchmarks via skip logic\n\n---\n\n## Step 5: Expose a Python API under `python/sglang/kernels/aot/python/sgl_kernel/`\n\nPrefer following the existing module organization first. For elementwise kernels, the usual pattern is:\n\n- implement the Python wrapper in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`\n- then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n\nFor example, in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`, add:\n\n```python\nimport torch\n\ndef scale(\n input: torch.Tensor,\n factor: float,\n out: torch.Tensor | None = None,\n) -> torch.Tensor:\n \"\"\"\n Element-wise scale: out = input * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n input : CUDA input tensor\n factor : scale factor (float)\n out : optional pre-allocated CUDA output tensor (same shape/dtype as input)\n \"\"\"\n if out is None:\n out = torch.empty_like(input)\n torch.ops.sgl_kernel.scale.default(out, input, factor)\n return out\n```\n\nThen re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.\n\n---\n\n## SGLang integration (required for runtime use)\n\nAfter exposing the AOT wheel symbol, add a lazy wrapper and a `KernelSpec` under\n`python/sglang/kernels/ops//`. Set `backend=KernelBackend.AOT`, record\nactual device/architecture support, and keep `sgl_kernel` imports inside the\nimplementation path. SGLang runtime and integration tests import that wrapper.\nFor this example, use `sglang.kernels.ops.elementwise.scale`.\n\nThe wheel-level tests below validate its standalone API/build. They do not\nreplace CI-registered SGLang correctness tests in\n`test/registered/kernels/ops/elementwise/test_scale.py` and benchmarks in\n`test/registered/kernels/benchmark/elementwise/bench_scale.py`. Follow\n`write-sglang-test` for their registration and CI budget; do not register tests\nunder the shipped `python/sglang/` package.\n\n## Step 6: Write tests (required)\n\nCreate `python/sglang/kernels/aot/tests/test_scale.py`:\n```python\nimport pytest\n\nimport torch\nimport sgl_kernel\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [128, 1024, 4096, 65536])\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0])\ndef test_scale_correctness(dtype, size, factor):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n\n result = sgl_kernel.scale(input, factor, out=out)\n assert result is out\n\n expected = input * factor\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\ndef test_scale_shape_mismatch():\n input = torch.randn(128, dtype=torch.float16, device=\"cuda\")\n out = torch.empty(256, dtype=torch.float16, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"same shape\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\ndef test_scale_cpu_input():\n input = torch.randn(128, dtype=torch.float16) # CPU\n out = torch.empty_like(input)\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-q\"]))\n```\n\n---\n\n## Step 7: Add a benchmark (required)\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nCreate `python/sglang/kernels/aot/benchmark/bench_scale.py`:\n\n```python\nimport torch\n\nimport sgl_kernel\nfrom sglang.kernels.jit.benchmark import marker\n\n\ndef sglang_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n sgl_kernel.scale(input, factor, out=out)\n\n\ndef torch_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n torch.mul(input, factor, out=out)\n\n\nFN_MAP = {\"sglang\": sglang_scale, \"torch\": torch_scale}\n\n\n@marker.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32], [torch.float16])\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096]) # 1K .. 512K\n@marker.benchmark(\"provider\", [\"sglang\", \"torch\"])\ndef benchmark(dtype: torch.dtype, size: int, provider: str):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n return marker.do_bench(\n FN_MAP[provider],\n # Pass every tensor through input_args (not a closure) so marker rotates\n # them across CUDA-graph calls to defeat L2 reuse.\n input_args=(input, 2.0, out),\n # Bandwidth = bytes(input) + bytes(out), both already in input_args.\n memory_output=None,\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n---\n\n## Step 8: Build\n\nBuild:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j16\n```\n\nIf you need to limit host resource usage:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j1 MAX_JOBS=2 CMAKE_ARGS=\"-DSGL_KERNEL_COMPILE_THREADS=1\"\n```\n\n---\n\n## Step 9: Validate\n\nAfter building successfully, run the test and benchmark:\n\n```bash\npytest python/sglang/kernels/aot/tests/test_scale.py -q\npython python/sglang/kernels/aot/benchmark/bench_scale.py\n```\n\nPR CI also runs `pr-test-sgl-kernel.yml`, including the B200 job\n`sgl-kernel-b200-test` when kernel changes are detected. Use that job as the\nBlackwell coverage signal for AOT `sgl-kernel` changes.\n\n---\n\n## Troubleshooting\n\n- **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`\n- **Memory errors**: `compute-sanitizer --tool memcheck python ...`\n- **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`\n- **Binary bloat**: use `python/sglang/kernels/aot/analyze_whl_kernel_sizes.py`\n- **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time\n\n---\n\n## References\n\n- `python/sglang/kernels/aot/README.md`\n- `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- `python/sglang/kernels/aot/csrc/common_extension.cc`\n- `python/sglang/kernels/aot/CMakeLists.txt`\n- `python/sglang/kernels/aot/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends\n- `python/sglang/kernels/aot/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern\n\n## Summary of Files Created/Modified\n\n```\npython/sglang/kernels/aot/csrc/elementwise/scale.cu # NEW: CUDA kernel + launcher\npython/sglang/kernels/aot/include/sgl_kernel_ops.h # MODIFIED: C++ declaration\npython/sglang/kernels/aot/csrc/common_extension.cc # MODIFIED: schema + dispatch registration\npython/sglang/kernels/aot/CMakeLists.txt # MODIFIED: add source file (alphabetical)\npython/sglang/kernels/aot/python/sgl_kernel/elementwise.py # MODIFIED: Python wrapper\npython/sglang/kernels/aot/python/sgl_kernel/__init__.py # MODIFIED: re-export Python API\npython/sglang/kernels/aot/tests/test_scale.py # NEW: tests\npython/sglang/kernels/aot/benchmark/bench_scale.py # NEW: benchmark\n```\n",".agents/skills/clean-startup-log/SKILL.md":"---\nname: clean-startup-log\ndescription: Audit SGLang startup logs, save evidence, and propose cleanup for user review. With no arguments, run Qwen3-8B at TP1 and TP2 plus gpt-oss-20b at TP1.\ndisable-model-invocation: true\n---\n\n# Audit SGLang Startup Logs\n\nThe default outcome is saved logs and a findings report. Apply runtime changes\nonly after the user selects them. A request to edit this skill does not itself\nlaunch servers.\n\n## Default runs\n\nA bare `$clean-startup-log` invocation runs these cases sequentially without\nasking for commands. Explicit commands, models, or TP sizes replace this matrix.\n\n| Case / log filename | Command |\n|---|---|\n| `qwen3-8b-tp1.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B` |\n| `qwen3-8b-tp2.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B --tp 2` |\n| `gpt-oss-20b-tp1.log` | `uv run sglang serve --model-path openai/gpt-oss-20b` |\n\nThese cover dense, tensor-parallel, and MoE/hybrid sliding-window attention\nstartup. Reuse complete captures from the current audit when code and environment\nhave not changed.\n\n## Capture logs\n\n1. Check the checkout, free GPUs, and ports once. Use the requested command when\n resources are free; otherwise select free GPUs with `CUDA_VISIBLE_DEVICES` and\n an unused `--port`, recording the adjustments. Leave existing servers alone.\n2. Create a unique directory with `mktemp -d /tmp/sglang-startup-audit-XXXXXX`.\n Save raw stdout and stderr together in a separate log for each case, for\n example with `set -o pipefail` and `COMMAND 2>&1 | tee LOG_PATH`. Record commands,\n GPU IDs, ports, commit, relevant overrides, and readiness status.\n3. Wait for `The server is fired up and ready to roll!`, then stop that server and\n its workers before the next case. First runs can spend many minutes downloading\n weights or compiling FlashInfer kernels; check download/compiler activity\n before treating a quiet log as a hang. Preserve partial logs for failed or\n stalled starts and report the last stage. Continue independent cases when possible.\n4. Preserve the user's logging configuration, including `NCCL_DEBUG`. If NCCL\n verbosity needs explaining, inspect relevant shell settings, `NCCL_CONF_FILE`,\n and `/etc/nccl.conf`. Do not override intentional diagnostics or recommend\n `NCCL_DEBUG=WARN` solely because the output is long. Avoid full environment dumps.\n\n## Investigate efficiently\n\n- Scan for deprecations, duplicate handler output, unrelated import failures,\n unformatted prints, and unexpected warnings. Read representative excerpts and\n counts instead of repeatedly dumping `server_args`, progress redraws, or NCCL\n diagnostics. Normalize carriage returns for analysis only; preserve raw logs.\n- Trace each candidate to its actual emitter with focused `rg` searches. Inspect\n its log level: SGLang's formatter may omit severity. Group shared signatures\n across cases and distinguish handler duplication from separate GPU/process calls.\n- Repetition, WARNING severity, or a different third-party format alone does not\n establish a cleanup need. Consider whether the message explains configuration,\n progress, resource use, or an operational limitation.\n- Consult [noise-source hints](references/noise-sources.md) only for a matching\n signature or an unresolved emitter. Verify current code rather than trusting\n historical line numbers, fix status, or assumptions about unrelated models.\n\n## Accepted output\n\nPreserve these reviewed messages unless the user requests a different policy:\n\n- NCCL diagnostics enabled by the user's environment or host configuration.\n- NUMA permission warnings, including one check per GPU in TP runs.\n- GPT-OSS MXFP4 backend-selection warnings and default page-size selection warnings.\n- `Init Unified Radix Cache. Components: ... Tree Core: ...`, tree-cache summaries,\n SWA allocation details, and per-rank memory/timing records.\n- Useful progress bars, warmup HTTP access logs, uv synchronization messages,\n isolated NCCL/Gloo startup lines, and one timestamped HF authentication warning.\n\nThese can appear in a clean startup log. Do not repeatedly propose the declined\nNUMA deduplication, backend/page-size level changes, or NCCL verbosity override.\nKeep real operational warnings visible: for example, a Harmony vocabulary failure\ncan disable `/v1/responses` even when server readiness and `/generate` succeed.\n\n## Report before changing code\n\nReturn a compact run table with readiness status and clickable raw-log links.\nFor each actual cleanup candidate, give an exact representative message, affected\ncases/counts, source file/function, and specific proposed behavior. Distinguish\nconfirmed findings from suspicions and operational failures from logging noise.\n\nIf there are no actionable cleanup findings, say the logs are clean and no\nfurther cleanup is needed. Otherwise, ask which numbered changes to adopt and\nwait for the user's selections before editing runtime code or preparing patches.\nHonor existing approvals and declined items without asking again.\n\n## Apply selected changes\n\n- Batch compatible approved edits, then verify affected cases once. Repeat a\n startup only for a new change, failure, or unresolved concern; do not relaunch\n after every one-line edit. Save verification logs separately from baselines.\n- Preserve useful warnings and application handlers. HF can warn during early\n CLI model detection before `configure_logger()`, and spawned processes have\n independent logger state. Keep `configure_hf_hub_logger()` in both\n `suppress_noisy_warnings()` and `configure_logger()`; make repeated setup safe.\n- Keep the legacy compiled-kernel cache migration notice at DEBUG. Avoid broad\n library-level suppression or fd redirection for a narrow logging problem.\n- Run relevant formatting and focused existing checks. Add tests only when they\n verify meaningful behavior, not a log-level spelling. Report changes and\n verification; create branches, commits, and PRs when requested.\n",".agents/skills/compute-mamba-ratio/SKILL.md":"---\nname: compute-mamba-ratio\ndescription: Compute the optimal --mamba-full-memory-ratio (or --max-mamba-cache-size pin) for a hybrid attention + linear-attention (Mamba / GDN / KDA) model's two serving memory pools, from the workload and serving config. Use when a user asks what ratio to set, why concurrency is clamped, or how to size the state vs KV pools for a hybrid model.\n---\n\n# Optimal hybrid dual-pool ratio (`--mamba-full-memory-ratio`)\n\nA hybrid model (attention layers + linear-attention layers — the recurrent-state family: Mamba/SSM, GDN, KDA, etc.) splits serving memory into two independently-budgeted pools, fixed once at startup:\n\n- **state pool** (the linear-attention recurrent state) → caps **concurrency** (hard: whole slots, worst-case reserved, fail-loud)\n- **full-KV pool** (attention KV) → caps **context × concurrency** (soft: paged, over-committable via retraction)\n\n`--mamba-full-memory-ratio r` splits the post-weight budget: `mamba_budget = rest · r/(1+r)`, i.e. `mamba_budget : kv_budget = r`. This skill picks the `r` (or the pin-`--max-mamba-cache-size` alternative) at which **neither pool bottlenecks first** for the user's workload.\n\n## The formula\n\n```\nr* = (S + D) · token_equiv · dcp_size / L\ntoken_equiv = state_bytes_per_slot / kv_bytes_per_token\n```\n\n- `L` = average context length per request (input + output tokens)\n- `token_equiv` = full-KV token-equivalent of one state slot\n- `S` = state slots per running request (cache-strategy dependent, table below)\n- `D` = `--speculative-num-draft-tokens` (0 if NOSPEC); each running req carries `D` extra intermediate states\n- `dcp_size` = `--dcp-size` (1 without DCP). DCP shards the per-rank KV by `dcp_size`, so KV gets ~`dcp_size×` cheaper per request → the balance shifts that much toward the state pool.\n\n`r` is dimensionless (just the split). To also predict the actual concurrency you need `rest` (below).\n\n## Inputs to collect from the user\n\n1. **`L`** — average context (input + output) in tokens.\n2. **The two per-GPU byte constants** — one of:\n - **(a) measured (preferred, exact)** — from one boot log at *any* ratio:\n - `Mamba Cache is allocated. ... ssm_state size X GB` with `max_mamba_cache_size: N` → `state_bytes_per_slot = X / N`\n - `KV Cache is allocated. #tokens: M, KV size: Y GB` → `kv_bytes_per_token = Y / M`\n - **(b) derived** — model arch (linear-layer count + state dims `d_state/d_conv/heads/head_dim`; attention type + KV dims: MLA latent dim, or GQA `kv_heads·head_dim·layers`) **× the dtypes** below.\n3. **`S`** — from `--mamba-radix-cache-strategy`, the overlap scheduler, and `SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK` (table below).\n4. **`D`** — `--speculative-num-draft-tokens` (0 if NOSPEC; **also 0 when ReplaySSM spec-verify is enabled** — see caveats).\n5. **`dcp_size`** — `--dcp-size` (1 if no DCP). If DCP **and** spec with a replicated draft KV, apply the draft caveat (below).\n6. **KV dtype** (bf16 / fp8) and **ssm dtype** (fp32 / bf16) — they set the two byte constants (see the `token_equiv` 2×2).\n7. *(only to also predict the clamp, not just `r`)* **`rest`** = per-GPU memory − weights at the chosen `--mem-fraction-static`. Read `avail mem` after `Load weight end`, or `Memory pool end. avail mem` + pool sizes, from the boot log.\n\n### `S` (state slots per running request), set by `--mamba-radix-cache-strategy`\n\nAll strategies keep prefix caching on; they differ in the track buffer that snapshots chunk-boundary state under the overlap scheduler. `S = base + ping-pong`, where base is 3 (live state + radix retention/COW headroom), and ping-pong = 2 (overlap on, non-lazy) / 1 (lazy, **or** overlap off) / 0 (no track buffer at all). This mirrors `kv_cache_configurator._calculate_mamba_ratio`; read it there if a release moves the constants.\n\n`SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK=1` takes the base to 2, since the decode-time skip frees one resident slot per running request. `no_buffer` is the exception and stays at an effective 3: it adds that slot straight back, because its binding limit is the prefill-to-decode peak, which a decode-time change does not shrink.\n\n| strategy | S | S with decode-lock skip | notes |\n|---|---|---|---|\n| `--disable-radix-cache` | 1 | 1 | prefix cache off entirely; most concurrency, no reuse |\n| `no_buffer` | 3 | 3 | no track buffer; requires overlap **off** (+ `page_size 1`) → lower decode throughput |\n| `extra_buffer` (default), overlap on | 5 | 4 | track buffer reserved per running request |\n| `extra_buffer`, overlap off | 4 | 3 | same buffer, but ping-pong costs 1 slot instead of 2 |\n| `extra_buffer_lazy` (recommended) | 4 | 3 | track buffer allocated lazily at the boundary → 1 fewer slot/req; requires overlap **on** |\n\nThe skip is a real concurrency lever, not a rounding detail: at the default strategy it takes `S` from 5 to 4, so the same state budget admits 25% more requests and `r*` drops by the same factor.\n\nOverlap counts as off when `--disable-overlap-schedule` is passed **or** `--pp-size > 1` (pipeline parallelism turns it off for you). So a PP deployment left on the default strategy sits on the `S = 4` row, not `S = 5` — using 5 there overstates the state cost and starves KV.\n\n### `token_equiv` — compute it from the two byte constants\n\n`token_equiv = state_bytes_per_slot / kv_bytes_per_token`. Both bytes are **per-model, per-GPU** — read them from a boot log (or derive from the arch + dtypes) and divide. That's the whole thing; the dtype effects below are just for a quick mental estimate.\n\nThe recurrent state has two parts: `state_bytes = SSM_bytes + conv_bytes`. The **SSM** part scales with `--mamba-ssm-dtype`; the **conv** part is always bf16 (`SGLANG_MAMBA_CONV_DTYPE`, fixed). How the dtype knobs move `token_equiv`:\n\n- **fp8 KV → ×2** (universal, exact): fp8 is 1 byte vs bf16's 2, so `kv_bytes_per_token` halves.\n- **fp32 ssm vs bf16 ssm → ×≈2**: only the SSM tensor doubles, conv stays fixed, so the exact factor is `(2·SSM + conv)/(SSM + conv)` — just under 2. **conv is usually small relative to SSM, so estimate ≈2** (e.g. one measured model: conv ≈7% of state → factor ≈1.9). Only bother with the exact split if a model's conv is unusually large.\n- **fp32 ssm + fp8 KV → ×≈4**.\n\nFor an actual number, don't apply factors — read `state_bytes_per_slot` and `kv_bytes_per_token` for **your** model+dtype from the boot log and divide.\n\nThe `GB` the boot log prints for both pools is **GiB** (`1024³`). Convert both the same way and `token_equiv` is unaffected either way, since it is a ratio — but get it right before comparing an absolute byte figure against your own log.\n\nWorked example (one measured model, TP8, fp32 ssm + fp8 KV), read off one boot log:\n\n```\nMamba Cache is allocated. max_mamba_cache_size: 257, conv_state size: 0.46GB, ssm_state size: 13.04GB\nKV Cache is allocated. dtype: torch.float8_e4m3fn, #tokens: 1167552, KV size: 15.03 GB\n```\n\n`state_bytes_per_slot = (0.46 + 13.04) GiB / 257 ≈ 56.4 MB`, `kv_bytes_per_token = 15.03 GiB / 1167552 ≈ 13.8 KB` → `token_equiv ≈ **4080**`.\n\n## Compute\n\n```python\ndef optimal_ratio(L, state_bytes_per_slot, kv_bytes_per_token, S, D=0, dcp_size=1):\n token_equiv = state_bytes_per_slot / kv_bytes_per_token\n r = (S + D) * token_equiv * dcp_size / L\n return r # value of --mamba-full-memory-ratio (>1 is legal)\n\ndef predict_clamp(rest_bytes, r, state_bytes_per_slot, S, D=0):\n mamba_budget = rest_bytes * r / (1 + r)\n slots = mamba_budget / state_bytes_per_slot\n # spec: each running req reserves (S+D) worth; non-spec just S\n return int(slots // S) # NOSPEC; with spec the budget joint-solves for (S+D)·per_req per req\n```\n\nThen state the result three ways: the **`r` value**, the **predicted clamp** (if `rest` given), and **which pool binds** (`min(mamba_clamp, KV_cap)`, where `KV_cap = kv_tokens · dcp_size / L`).\n\n## Procedure\n\n1. **Free lever first**: if the boot log shows large idle in `Memory pool end. avail mem` (e.g. 30–40 GB at mem-frac 0.85), raise `--mem-fraction-static` (→0.92) before touching the split — it grows `rest` for both pools at no cost. Validate graph-capture headroom once.\n2. Collect the inputs. Prefer a real boot log for the two byte constants.\n3. Compute `r*`. If `r*` would drive the state pool below one request's worth (`mamba_budget < S · per_req`, happens at very long `L`), **switch to pinning `--max-mamba-cache-size = target_concurrency · S`** and let the rest go to KV — a sub-0.15 `r` is fragile.\n4. Report `r`, predicted clamp, binding pool, and any dtype accuracy gate that applies.\n\n## Worked examples (TP8, B300, validated against measured clamps)\n\nConstants (from boot logs): `state_bytes_per_slot ≈ 56.4 MB` (fp32 ssm), `kv_bytes_per_token ≈ 13.8 KB` (fp8) → `token_equiv ≈ 4080`. Cache strategy `extra_buffer_lazy` → `S = 4`. Workload `L = 9216` (8192 in + 1024 out). NOSPEC → `D = 0`.\n\n- **TP (dcp_size=1)**: `r = 4 · 4080 · 1 / 9216 ≈ 1.8`. Measured: clamp 88, KV cap 87 → balanced. ✅\n- **DCP8 (dcp_size=8)**: `r = 4 · 4080 · 8 / 9216 ≈ 14`. Measured at r=14: clamp 125, KV cap 129 → balanced. ✅ (At the naive r=1.8 the DCP KV pool is ~8× over-provisioned — cap 683 vs clamp 86 — wasting budget that should go to the state pool.)\n\n`predict_clamp` on the same model at the **default** `extra_buffer`, overlap on and the decode-lock skip off (`S = 5`), three boot logs:\n\n| config | `max_mamba_cache_size` | `mmcs // S` | measured `max_running_requests` |\n|---|---|---|---|\n| TP8, no DCP | 257 | 51 | 51 |\n| DCP8, `r = 5.97` | 451 | 90 | 90 |\n| DCP8, `r = 9` | 474 | 94 | 94 |\n\nReference `r` for **this example model** (`token_equiv ≈ 4080`, i.e. fp32 ssm + fp8 KV; S=5 default; multiply by `(S+D)/S` for spec, by `dcp_size` for DCP) — recompute with your own `token_equiv` for a different model:\n\n| L | 2K | 4K | 8K | 32K | 64K | 128K |\n|---|---|---|---|---|---|---|\n| r | 10.0 | 5.0 | 2.5 | 0.62 | 0.31 | 0.16 |\n\n## Caveats\n\n- **DCP + spec**: a replicated (non-DCP-sharded) draft KV does **not** shard by `dcp_size`; at long `L` it dominates per-token cost, so the clean `×dcp_size` overstates DCP's advantage — fall back to a per-token direct-solve (KV term `/dcp` + an un-sharded draft-KV term) when spec is on. (NOSPEC → clean `×dcp_size` holds.)\n- **Spec + ReplaySSM → use `D=0`, not the draft-token count.** When ReplaySSM spec-verify is enabled, the `D` intermediate SSM states move off the per-request slot budget onto a **fixed ring** (a one-time deduction from `rest`, not a per-req term). So the mamba-slot cost per running request drops back to `S` (clamp = `mmcs / S`, not `/(S+D)`), and the balance ratio returns to the NOSPEC value (`r* ≈ S·token_equiv·dcp/L`). Measured example (TP8, D=8, L≈9K): applying `D=8` computes `r*≈2.6` but the true optimum is `r≈1.0` — `D=8` lands KV-bound at ~45% below the achievable peak concurrency. Plain spec (no replayssm) keeps `D` = the draft-token count.\n- **Under DCP the balance point sits beyond any realistic `L`** (the state pool binds first for almost everything), so the practical recommendation is to **pin `--max-mamba-cache-size = target_concurrency · S` directly** rather than dial a large `r`.\n- **`--max-mamba-cache-size` overrides `r`**; bytes beyond the `r` budget come out of the KV pool one-for-one.\n- **Precision changes are behind accuracy gates**: `--kv-cache-dtype fp8_e4m3` (doubles `token_equiv` → doubles `r`) and `--mamba-ssm-dtype bfloat16` (~halves `token_equiv`; also silently switches the linear-attention decode backend on SM100+ — pin `--linear-attn-decode-backend triton`) shift outputs; validate accuracy for the workload before production.\n- **Asymmetry**: the state pool is worst-case-reserved and fail-loud; KV degrades gracefully (retraction). When `L` is spiky/uncertain, bias `r` **up** rather than starve the state pool.\n",".agents/skills/babysit-pr-to-pass-ci/SKILL.md":"---\nname: babysit-pr-to-pass-ci\ndescription: Start and persistently pursue a goal to babysit an SGLang pull request until selected GitHub Actions workflows pass on the latest PR head. Use when asked to monitor, babysit, retry, or fix PR CI for lint.yml, pr-test.yml, pr-test-extra.yml, AMD, or other named workflows; classify failures as PR-related versus flaky or infrastructural, auto-fix and push only small clean fixes, rerun failed jobs only up to 10 times, and ignore unselected workflows.\n---\n\n# Babysit a PR Until Selected CI Passes\n\nDrive only the selected workflows to green on the PR's latest head commit. Keep working across turns by starting a durable goal; do not stop merely because a run was dispatched, a retry began, or a fix was pushed.\n\n## Invocation and arguments\n\nAccept:\n\n```text\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] [WORKFLOW ...]\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] --only WORKFLOW [WORKFLOW ...]\n```\n\n- If the PR is omitted, resolve the PR associated with the current branch.\n- If workflows are omitted, select `lint.yml` and `pr-test.yml`.\n- Without `--only`, add named workflows to the two defaults.\n- With `--only`, monitor exactly the named workflows.\n- Accept a workflow basename such as `pr-test-extra.yml`, a repository path such as `.github/workflows/pr-test-extra.yml`, or an unambiguous workflow display name.\n- Treat `extra` as `pr-test-extra.yml`.\n- Accept natural-language equivalents, such as “PR 12345 plus extra and AMD.” Resolve “AMD” to an exact workflow file from the user's wording or repository context; do not guess when multiple AMD workflows fit.\n\nExamples:\n\n```text\n$babysit-pr-to-pass-ci\n$babysit-pr-to-pass-ci 12345\n$babysit-pr-to-pass-ci https://github.com/sgl-project/sglang/pull/12345 pr-test-extra.yml\n$babysit-pr-to-pass-ci 12345 --only pr-test-amd.yml\n```\n\n## Start or continue the durable goal\n\nTreat explicit invocation of this skill as authorization to create a goal and perform the scoped CI actions below.\n\n1. Resolve the PR and exact workflow set first.\n2. Inspect the current goal.\n3. If no unfinished goal exists, create one immediately using the product's goal mechanism. Include the resolved PR URL, selected workflow files, and this stopping condition: every selected workflow has a successful run for the PR's latest head SHA.\n4. Include these constraints in the goal objective:\n - Ignore every unselected workflow.\n - Diagnose each selected-workflow failure as PR-related, unrelated/infrastructure, or still uncertain.\n - Implement, validate, commit, and push a PR-related fix only when it is clean, non-tricky, and at most 100 changed lines.\n - Stop for user review before editing when the proposed fix is tricky, dirty, or over 100 changed lines.\n - Retry unrelated failures at most 10 times per workflow and head SHA, using failed-jobs-only reruns.\n - Never rerun a full workflow merely to recover failed tests.\n5. If the active goal already represents this same PR and workflow set, continue it instead of creating another goal.\n6. If a different unfinished goal is active, do not replace it silently. Report the conflict and ask the user to pause or clear it.\n\nDo not mark the goal complete until the completion check at the end of this skill succeeds.\n\n## Preflight\n\n1. Verify GitHub authentication and resolve the repository without changing remote configuration.\n2. Resolve the PR and record at least:\n - PR number and URL\n - open/draft state\n - base branch and base SHA\n - head branch and head SHA\n - head repository owner/name, including whether it is a fork\n - labels\n - whether maintainers may modify the branch\n3. Require an open PR. If it is draft and a selected workflow is blocked by the draft gate, stop and ask the user to mark it ready; do not change draft state automatically.\n4. Normalize every selected workflow to an existing file under `.github/workflows/`. Fail fast on ambiguous or nonexistent names.\n5. Read each selected workflow's triggers and gates. Do not assume all optional workflows behave like `pr-test.yml`.\n6. Inspect `git status`. Preserve all user changes. If a fix becomes necessary and the current checkout is dirty, on a different branch, or otherwise unsafe, use an isolated worktree for the PR head.\n\nKnown SGLang gates:\n\n- `pr-test.yml` requires the `run-ci` label for relevant PR changes. If its existing gate job failed because the label was absent, add `run-ci` when authorized and rerun failed jobs in that existing run; adding the label alone does not necessarily create a new `pr-test.yml` run.\n- `pr-test-extra.yml` requires both `run-ci` and `run-ci-extra`. It listens for relevant label events, so adding a missing label may create a fresh run. Prefer that fresh current-head run; rerun failed jobs in an existing current run only when needed.\n- AMD-extra workflows may use the same `run-ci` plus `run-ci-extra` gate. Verify the actual selected file.\n- Treat draft, cooldown, maintenance, and missing-label failures as gate/infrastructure conditions, not code regressions. Fix only the safe gate condition in scope; wait for cooldown or maintenance windows rather than burning retries immediately.\n\nAdding the known CI opt-in labels is within scope. Do not add unrelated labels, mark the PR ready, alter reviewers, merge, or mutate other PR metadata.\n\n## Track only current-head runs\n\nUse both PR checks and workflow-run data so workflow names, run IDs, attempts, jobs, and logs can be correlated. Prefer workflow file IDs over display names when filtering runs.\n\nFor each selected workflow:\n\n1. Select the newest run associated with the current PR head SHA.\n2. Ignore runs for older head SHAs, including green runs made stale by a new push.\n3. Treat `queued`, `requested`, `waiting`, `pending`, and `in_progress` as active; keep monitoring.\n4. Treat a workflow as passing only when its current-head run concludes `success`.\n5. If no current-head run exists, wait briefly for event delivery, then inspect the workflow's path filters, event triggers, and label gates.\n6. Trigger only through a mechanism explicitly supported by that workflow. Never invent `workflow_dispatch` inputs or dispatch a workflow against the default branch when the intent is to test the PR head.\n7. If a selected workflow cannot safely be triggered for the PR head, report the exact blocker and request direction.\n\nAfter any user push or skill-authored push, re-read the PR head SHA, discard stale run selections, and start tracking the new head. Reset retry counters for the new head SHA.\n\n## Monitor persistently\n\n- Maintain a compact ledger containing workflow file, head SHA, run ID, attempt, status, conclusion, skill-initiated retry count, and latest failure signature.\n- Poll at reasonable intervals, normally 30–60 seconds while runs are active. Give concise progress updates during long waits.\n- Re-check the PR head SHA on each monitoring cycle and before every rerun or push.\n- Ignore failures, cancellations, and pending checks from unselected workflows. Do not diagnose, retry, fix, or wait for them.\n- Do not declare success while a selected workflow is absent, stale, pending, skipped, cancelled, neutral, or failed.\n\n## Diagnose a selected-workflow failure\n\n1. Inspect the failed jobs and failed-step logs. In `pr-test.yml`, distinguish the root failure from cascade failures caused by `check-pr-test-health`, wait jobs, or final aggregation.\n2. Record the exact failing test/check, error text, affected runner or hardware, and relevant log URL.\n3. Compare the failure against the PR's changed files and diff.\n4. Classify it:\n\n **PR-related** when evidence ties the failure to changed code, tests, dependencies, configuration, generated files, or behavior introduced by the PR. A reproducible failure on the PR that does not occur on the base is strong evidence.\n\n **Unrelated/infrastructure** when evidence points to a transient network or download error, runner loss, GPU/driver instability, service outage, pre-existing main-branch failure, known flaky threshold, unrelated test area, or another environmental cause.\n\n **Uncertain** when evidence is insufficient. Do not label a failure flaky merely because it might be intermittent. Narrowly reproduce it, compare recent main/scheduled CI, inspect relevant code, or otherwise gather enough evidence to choose a category.\n\n5. When several root failures exist, classify each one. Address PR-related failures before spending retries on unrelated cascades.\n\n## Handle a PR-related failure\n\nEstimate the proposed fix before editing.\n\nA fix is eligible for automatic execution only when all are true:\n\n- The fix is well understood, localized, and clean.\n- The fix is not tricky, hacky, speculative, or a broad workaround.\n- The fix changes at most 100 lines, counting additions plus deletions in the fix patch, not the pre-existing PR diff.\n- The fix does not require a broad API or architecture redesign, risky dependency migration, large generated-file rewrite, or unrelated cleanup.\n- A focused validation is available.\n\nIf any condition fails, stop before editing. Report the diagnosis, proposed approach, expected files, estimated size, risks, and validation plan, then wait for user review.\n\nFor an eligible fix:\n\n1. Work from the latest PR head in a clean checkout or isolated worktree.\n2. Make only the necessary change and add or update focused tests when appropriate.\n3. Run the smallest meaningful local validation. For lint failures, reproduce the specific hook/check before considering broader lint validation.\n4. Review the resulting patch and count changed lines. If it becomes over 100 lines or turns tricky/dirty, do not commit or push it; stop for review and clearly identify the uncommitted skill-authored changes.\n5. Commit only skill-authored files. Never use `git add .`, rewrite user commits, amend without request, or include pre-existing local changes.\n6. Push a new commit non-forcibly to the exact PR head repository and branch. Never force-push. If permission or fork routing is unclear, stop and report the push command/target that needs authorization.\n7. Resolve the new PR head SHA and return to current-head monitoring. A push is progress, not completion.\n\n## Handle an unrelated or infrastructure failure\n\nRetry only the failed jobs in the current selected-workflow run:\n\n```bash\ngh run rerun --failed --repo \n```\n\nRules:\n\n- Never use a full-run rerun as a substitute for `--failed`.\n- Never rerun a stale run after the PR head SHA changes or a newer current-head run exists.\n- Count only successfully dispatched skill-initiated reruns.\n- Allow at most 10 such reruns per selected workflow and head SHA.\n- Wait for each new attempt to finish, inspect its result, and update the failure signature before deciding on another retry.\n- If the failure changes or becomes plausibly PR-related, return to diagnosis instead of blindly consuming all retries.\n- If 10 retries are exhausted, stop. Report the attempts, recurring signatures, job/run links, and evidence that the failure is unrelated or still uncertain. Do not exceed the cap without a new explicit user instruction.\n\nIf GitHub cannot rerun only the failed portion of the selected run, stop and ask the user rather than rerunning the full workflow.\n\n## Completion check\n\nBefore completing the goal:\n\n1. Fetch the PR head SHA again.\n2. For every selected workflow, verify that its newest applicable run is for that exact SHA and concludes `success`.\n3. Verify no selected workflow is still pending and no success was invalidated by a newer push.\n4. Ignore all unselected workflows even if they are failing.\n5. Mark the goal complete only after all selected workflows satisfy these checks.\n\nReport the final PR URL, head SHA, selected workflows, passing run links, any commits pushed, and retry counts. Keep the report scoped to the selected workflows.\n",".agents/skills/cookbook-add-model/SKILL.md":"---\nname: cookbook-add-model\ndescription: Add a new model to the SGLang Cookbook (docs/, Mintlify), config-driven format — instantiate the model-agnostic template into a per-model config (+ benchmarks) JSX under src/snippets/configs/, an MDX page, the docs.json nav entry, NEW-tag hygiene, and the homepage vendor card. Interactive, multi-phase. Run with /cookbook-add-model.\ndisable-model-invocation: true\n---\n\n# Add a model to the SGLang Cookbook\n\n> Migrating an **existing legacy-template page** (one that imports a monolithic\n> `…/autoregressive/-deployment.jsx` generator)? Use the\n> `cookbook-migrate-model` skill instead — same target format, but the legacy\n> page (not the user) is the source of truth.\n\nThe cookbook is **config-driven**: two shared engines contain NO model-specific code —\n`docs/src/snippets/_deployment.jsx` (the 5-dim deploy matrix) and\n`_playground.jsx` (the diff-based override Playground). Adding a model = adding **data**:\na per-model `config` (+ optional `benchmarks`) consumed by both engines, plus an MDX page\nthat imports them. No engine edits.\n\n**Instantiate the model-agnostic template** (NOT a clone of any live cookbook — the\ntemplate is decoupled and covers all hardware + all axes):\n- `templates/config.jsx.tmpl` → `docs/src/snippets/configs//.jsx`\n- `templates/benchmarks.jsx.tmpl` → `…/-benchmarks.jsx` (skip if no numbers)\n- `templates/page.mdx.tmpl` → `docs/cookbook///.mdx`\n\nThe template uses explicit `__TOKEN__` placeholders; you fill them, prune what the model\nlacks, and replace the EXAMPLE cells with verified recipes. DeepSeek-V4 is a populated\n*instance* you can consult, but is not the template.\n\n**Deep references (read on demand, don't inline):**\n- [references/authoring-reference.md](references/authoring-reference.md) — field-by-field config / cells / playground / MDX contract.\n- [references/mintlify-authoring.md](references/mintlify-authoring.md) — MDX rules (forbidden syntax, JSX tables, labeled fences) + invocation-example patterns. Read before writing §1–§3 prose.\n- [references/engine-axis.md](references/engine-axis.md) — adding a new Playground feature axis (rare engine work).\n- [references/vendor-logo.md](references/vendor-logo.md) — new-vendor card logo: ask the user for the brand logo, then generate the icon-only 940×525 RGBA PNG (spec + Pillow recipe + `git add -f`).\n- [references/diffusion-authoring.md](references/diffusion-authoring.md) — required opening, tag, command-picker, and feature-overlay contract for diffusion model pages. Read before editing any `docs/cookbook/diffusion//.mdx` page.\n\n## Architecture at a glance\n\n```\n┌─────────────────────────────────────────────────────────────────┐\n│ cookbook///.mdx │\n│ import { Deployment } from \"/src/snippets/_deployment.jsx\"; │\n│ import { Playground } from \"/src/snippets/_playground.jsx\"; │\n│ import { config } from \"/src/snippets/configs/.../X.jsx\"; │\n│ │\n└─────────────────────────────────────────────────────────────────┘\n │ (config passed as React prop)\n ▼\n┌─────────────────────────────────────────────────────────────────┐\n│ src/snippets/configs//.jsx │\n│ export const config = { │\n│ supportedHardware, variants, quantizations, strategies, ... │\n│ cells: [ { match:{hw,variant,quant,strategy,nodes}, │\n│ env:[...], flags:[...] }, ... ], // 5-dim matrix │\n│ playgroundFeatures: { attention, moe, parsers, ... }, │\n│ }; │\n└─────────────────────────────────────────────────────────────────┘\n │ (consumed by BOTH engines — no model code in engines)\n ▼\n┌──────────────────────────────────┬──────────────────────────────┐\n│ _deployment.jsx │ _playground.jsx │\n│ Renders the verified matrix; │ Renders override chips + │\n│ one cell → its env/flags. │ diff against the cell. │\n└──────────────────────────────────┴──────────────────────────────┘\n```\n\nThe two widgets stay in sync via: the **URL hash** (deploy mirrors its selection;\nplayground reads it), the **`sglang-deploy-sel` custom event** (deploy dispatches on\nevery change; playground listens — `replaceState` doesn't fire `hashchange`), and the\nshared **`sglang-deploy-env` localStorage key** (HOST/PORT placeholders).\n\n> The main template is **autoregressive**. Diffusion pages use\n> `templates/diffusion-page.mdx.tmpl` plus\n> [references/diffusion-authoring.md](references/diffusion-authoring.md); do not force the\n> autoregressive deployment matrix on them. Omni pages follow their own category structure.\n> All categories still obey the Mintlify / NEW-tag / docs.json / category-card / validation\n> rules below.\n\n---\n\n**Interactive, multi-step workflow. Collect inputs incrementally — don't ask for\neverything upfront.** The real work is the verified `cells[]` recipes + measured\nbenchmarks (Phases 2 + 4); everything else is filling the template.\n\n## Phase 1 — Collect inputs\n\n1. **Model card** — HuggingFace repo/URL. **Fetch the page** and extract description,\n param count, architecture, context length, **license**. (Fetching guards factual bugs\n like an off-by-a-few-B param count.) If the model isn't public, ask the user.\n2. **Variants / quantizations** — keep separate: variants are size/mode (e.g. Flash/Pro,\n Instruct/Thinking); quantizations come from the HF card / linked repos (BF16/FP8/FP4/…).\n Default to BF16 when a full-precision repo exists.\n3. **Tested hardware + parallelism** — which platforms are actually tested, and TP/EP/DP\n for each. List only tested hw (unlisted greys out).\n4. **Verified launch recipes** — the full `sglang serve` flags per\n (hw × variant × quant × strategy × nodes) combo → these become `cells[]`. Rewrite any\n `python -m sglang.launch_server` to `sglang serve` form.\n5. **sglang version / image tag** — ask which sglang build the recipes + benchmarks ran on\n (a release like `0.5.x`, or `main`/nightly). **Never guess or hallucinate it.** This one\n tag fills `dockerImages` and the benchmarks' `sglang_version`; when the user is unsure,\n default the image to `lmsysorg/sglang:dev` (nightly) rather than inventing a release.\n6. **Pre-flight**: `gh pr list --repo sgl-project/sglang --search \"\"` (dup check).\n\n**Hardware reference** (the shared `HARDWARE_CATALOG` in `_deployment.jsx`). A GPU **not** in\nthis table (RTX PRO 6000, GH200, future chips) goes in the model's own `config.hardware`\n(`{id,label,vram,vendor}`) — the engine merges it in; don't edit the engine catalog:\n\n| Platform | Vendor | VRAM | Docker image |\n|---|---|---|---|\n| H100 | NVIDIA | 80GB | `lmsysorg/sglang:` |\n| H200 | NVIDIA | 141GB | `lmsysorg/sglang:` |\n| B200 | NVIDIA | 192GB | `lmsysorg/sglang:` |\n| B300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130` when required) |\n| GB200 | NVIDIA | 192GB | `lmsysorg/sglang:` (or `-cu130`) |\n| GB300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130`) |\n| DGX Spark | NVIDIA | 128GB (unified) | `lmsysorg/sglang:` — needs a **CUDA 13** build |\n| MI300X | AMD | 192GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI325X | AMD | 256GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI350X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n| MI355X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n\n- **Image tag (``)**: don't guess — ask the user for the tag the recipes ran on, or\n default to `dev` (nightly). The same tag goes in `dockerImages` and benchmarks'\n `sglang_version`; the engine falls back to `lmsysorg/sglang:dev` for any unmapped hw.\n- **TP sizing** (sanity-check recipes): `weight_GB / gpu_mem`, round up to a power of 2,\n ~20–30% headroom. BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5. MoE → **total** weight, not\n active params. FP4 is Blackwell-only (B200/B300/GB200/GB300/DGX Spark). GB200/GB300\n single-node hosts are typically **4 GPUs** (TP=4 ceiling); a DGX Spark is **1 GPU**, so\n its only multi-GPU topology is TP=2 across 2 nodes.\n- **Platform flags**: Blackwell may need `--attention-backend trtllm_mha`; AMD typically\n needs `--attention-backend triton` + env `SGLANG_USE_AITER=1` /\n `SGLANG_ROCM_FUSED_DECODE_MLA=0` (check AITER TP constraints, e.g. `heads_per_gpu % 16 == 0`).\n- **EP** (MoE): 8-GPU NVIDIA `--tp 8 --ep 8`; AMD `EP = TP`; small NVIDIA (TP≤4) omit\n `--ep` unless benchmarked. (The template's AMD example cell shows these.)\n\n## Phase 2 — Instantiate the template\n\nFor a diffusion model, instantiate `templates/diffusion-page.mdx.tmpl` and keep the\nshared `DiffusionModelTags` component, plus `templates/diffusion-config.jsx.tmpl` for the\nopt-in scoped command builder. Put the compact install command and builder in §1 Quick\nstart. The first two paragraphs in §2 Model capabilities are not generic filler: they must\nstate the model's capability range, strongest differentiator, when to choose it, and at\nleast one real deployment or capability boundary. Put orthogonal runtime features in\n`scope: \"serve\"` or `scope: \"request\"`, not in the base recipe; use the schema from\n`references/diffusion-authoring.md`.\n\n1. **Copy** the three template files to their target paths (above). Note the two\n vendor-folder conventions: under `configs/` the folder is the **HuggingFace org**\n (`deepseek-ai`); under `cookbook/` it's the **display vendor** (`DeepSeek`).\n2. **Replace every `__TOKEN__`**: `__MODEL_DISPLAY__`, `__MODEL_SLUG__`, `__HF_ORG__`,\n `__HF_REPO__`, `__REASONING_PARSER__`, `__TOOLCALL_PARSER__`, `__ONE_LINER__`. Verify\n none remain: `grep -rn '__[A-Z_]*__' `.\n3. **Prune** to what the model supports (delete, don't stub) — using\n [references/authoring-reference.md](references/authoring-reference.md):\n - `supportedHardware` + the EXAMPLE cells: keep your tested families; **delete the\n `mi*` ids + AMD example cell if no AMD recipe**, etc. A GPU not in the shared catalog\n (e.g. RTX PRO 6000) → declare it in `config.hardware` and add its id here.\n - `playgroundFeatures` axes: remove the `megamoe` backend option + the\n `megamoeQuant` block from the `moe` axis (non-Blackwell-MoE), delete `hisparse`\n (non-DSA), `pdDisagg`/`router` (no PD), the `parsers` axis (no parsers), etc.\n - `quantizations` / `variants`: drop what the model doesn't ship; collapse `variants`\n to single `default` if there's no variant axis (then drop the `variant` half of\n `modelNames`/`defaultAccuracy` keys).\n4. **Fill `cells[]`** with the verified recipes from Phase 1 (replace every EXAMPLE cell;\n set `verified: true` only on tested combos), and `modelNames` with real HF slugs,\n `dockerImages` for your hw (use the Phase-1 tag, or default `lmsysorg/sglang:dev` — never\n a guessed release; key by `hw`, or a more specific key — `hw|variant|quant`, `variant|quant`,\n `hw|quant|strategy`, `hw|quant`, resolved in that order before `hw` — when one\n variant/quant/strategy needs its own image), `multiNodeHints` only for fabric-specific hw\n (e.g. gb200).\n\n5. **Diffusion pages: add the ComfyUI section.** Every diffusion cookbook page ends with\n a `## . Run in ComfyUI` section so a reader never has to guess whether the model is\n reachable from ComfyUI. It is one component; the per-model facts live in the component,\n not the page:\n\n ```mdx\n ## . Run in ComfyUI\n\n import { ComfyUISupport } from '/src/snippets/diffusion/comfyui-support.jsx';\n\n \" />\n ```\n\n Pick `model` from the table in\n `docs/src/snippets/diffusion/comfyui-support.jsx`. Use the model's own key when it has\n an entry (its executor or dedicated node differs); otherwise use the generic `image` or\n `video`. A model gets its own key only when the plugin actually treats it specially —\n an entry in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a\n dedicated node. Adding a key without the matching plugin support makes the page lie.\n\n### Site-wiring (do all three)\n\n- **`docs/docs.json`** — add the page under Cookbook → `` → ``, at\n the **top** of that vendor's `pages` (root-relative, no `.mdx`:\n `cookbook///`). New vendor group → insert in the section's\n local ordering.\n- **NEW-tag hygiene** — the new page keeps `tag: NEW` (from the template). Scan the\n vendor dir for existing NEW and strip it from siblings; verify ≤1:\n `grep -rn 'tag: NEW' docs/cookbook///` → at most one result. (Scan\n files; don't assume the first `docs.json` entry holds NEW.)\n- **Homepage card** — `docs/cookbook//intro.mdx`: if the org already has a\n ``, update only its `href` (keep `img`). If the org is **new**, add a ``\n (title = nav-group name; keep card order aligned with `docs.json`) **and create its logo**:\n ask the user for the brand logo, then generate the conforming **icon-only 940×525 RGBA\n transparent** PNG → `docs/cards/logos/.png` per\n [references/vendor-logo.md](references/vendor-logo.md) (track with `git add -f` — `*.png`\n is gitignored repo-wide). Never invent or copy a logo.\n\n## Phase 3 — Validate\n\n```bash\ncd docs\nmint validate # frontmatter, missing nav entries, MDX/JSX errors\nmint broken-links\nmint dev # visual smoke test at http://localhost:3000/cookbook///\n```\n\nSpot-check: cells render sensible commands; URL-hash nav persists across reload; the\nPlayground inherits the Deploy selection live; each axis toggle produces the expected\ndiff; Docker mode wraps in `docker run` with the right image; multi-node cells emit the\nhints + `--nnodes N`; cURL resolves the model name; the NEW badge shows on the new page\nand is gone from same-vendor siblings; the homepage card points to the new model.\n\n## Phase 4 — Interactive testing\n\nThe user deploys each cell, runs the benches, and pastes results; you fill the data:\n- mark each tested `cells[]` entry `verified: true` (absent = yellow/unverified badge);\n- fill the `-benchmarks.jsx` entries (one per cell `match`) with measured\n speed/accuracy + the `sglang_version` the user reports (don't invent one — the template's\n `0.0.0` is a deliberate TODO); set model-level `defaultAccuracy` per variant. Leave a\n cell's entry as a bare `match` stub if it has no numbers yet (the card shows \"pending\").\n\n## Phase 5 — Prose & config tips\n\n**Read [references/mintlify-authoring.md](references/mintlify-authoring.md) first** (it\ncarries the parser-output-shape / thinking-mode / Output-Example / no-hardcoded-sampling\nrules + the Mintlify forbidden-syntax list). Then rewrite the MDX prose from the HF card +\nuser notes: §1 Model Introduction (description, links, params, license, variants table),\n§2 Configuration Tips (hw-specific tuning, caveats), §3 Advanced Usage (Reasoning /\nTool-Calling / HiCache — keep only what applies; match the reasoning example to the\nparser's output shape; each runnable block gets an `**Output Example:**`).\n\nFor diffusion pages, follow the category-specific Quick start and capability contract in\n`references/diffusion-authoring.md`. Run `node docs/scripts/check_cookbook_configs.mjs` to\nverify the tag widget and introduction structure before rendering the page.\n\n## Phase 6 — Review\n\n```\n/cookbook-review-pr \n```\n\n## Git workflow\n\nAlways branch — never commit to main directly.\n\n```bash\ngit checkout -b add--cookbook\ngit add docs/src/snippets/configs//.jsx \\\n docs/src/snippets/configs//-benchmarks.jsx \\\n docs/cookbook///.mdx \\\n docs/docs.json docs/cookbook//intro.mdx\ngit commit -m \"Add cookbook\"\ngit push -u origin add--cookbook\ngh pr create --title \"[Doc] cookbook: Add cookbook\" --body \"...\"\n```\n",".agents/skills/ci-workflow-guide/SKILL.md":"---\nname: ci-workflow-guide\ndescription: Guide to SGLang CI workflow orchestration — stage ordering, fail-fast, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages.\n---\n\n# SGLang CI Workflow Orchestration Guide\n\nThis skill covers the CI **infrastructure** layer — how tests are dispatched, gated, and aborted on failure across stages. For test authoring (templates, fixtures, registration, model selection), see the [write-sglang-test skill](../write-sglang-test/SKILL.md).\n\n---\n\n## Naming Conventions\n\n- **Suite**: `base-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `base-b-test-1-gpu-small`)\n- **Test group**: Directory-level registered test group under `test/registered/` (e.g., `hicache` maps to `test/registered/hicache/test_*.py`)\n- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)\n\n---\n\n## Key Files\n\n| File | Role |\n|------|------|\n| `.github/workflows/pr-test.yml` | Main workflow — all stages, jobs, conditions, matrix definitions |\n| `.github/workflows/pr-test-extra.yml` | Extra workflow — gated by BOTH `run-ci` and `run-ci-extra` labels |\n| `.github/workflows/pr-gate.yml` | PR gating: draft check, `run-ci` label, per-user rate limiting |\n| `.github/actions/check-pr-test-health/action.yml` | Cross-job fail-fast: queries API for any failed job |\n| `.github/actions/wait-for-jobs/action.yml` | Stage gating: polls API until stage jobs complete |\n| `.github/actions/check-maintenance/action.yml` | Maintenance mode check |\n| `.github/scripts/ci-labels.cjs` | Resolves the four CI control labels into dispatch axes |\n| `test/run_suite.py` | Suite runner: collects, filters, partitions, executes tests |\n| `python/sglang/test/ci/ci_register.py` | Test registration (AST-parsed markers), LPT auto-partition |\n| `python/sglang/test/ci/ci_utils.py` | `run_unittest_files()`: execution, retry, continue-on-error |\n| `scripts/ci/utils/slash_command_handler.py` | Handles slash commands from PR comments |\n\n---\n\n## Architecture Overview\n\n```\n ┌──────────────┐\n │ build kernel │\n └──────┬───────┘\n │\n ├─ check-changes ──── detects which packages changed\n │ (main_package, sgl_kernel, jit_kernel, multimodal_gen)\n │\n ├─ call-gate ──────── pr-gate.yml (draft? label? rate limit?)\n │\n ├─────────────────────────────────────────────────────┐\n │ │\n ▼ │\n ┌─────────────────────────────────────┐ │\n │ Base A (~3 min) │ │\n │ pre-flight check │ │\n │ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-1-gpu-small │ │ │\n │ │ (small GPUs) │ │ │\n │ └─────────────────────────────┘ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-cpu │ │ │\n │ │ (CPU) │ │ │\n │ └─────────────────────────────┘ │ │\n └──────┬──────────────────────────────┘ │\n │ │\n ▼ ▼\n ┌─────────────────────────────────────┐ ┌──────────────────────────┐\n │ Base B (~30 min) │ │ kernel test │\n │ base tests │ └──────────────────────────┘\n │ │ ┌──────────────────────────┐\n │ ┌─────────────────────────────┐ │ │ multimodal gen test │\n │ │ base-b-test-1-gpu-small │ │ └──────────────────────────┘\n │ │ (small GPUs, e.g. 5090) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-1-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-2-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ Base C (~30 min) │\n │ advanced tests │\n │ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-h100 │ │\n │ │ (H100 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-8-gpu-h200 │ │\n │ │ (8 x H200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-b200 │ │\n │ │ (4 x B200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ Other advanced tests │ │\n │ │ (DeepEP, PD Disagg, GB300) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ pr-test-finish │\n │ aggregates all results, fails if │\n │ any job failed/cancelled │\n └─────────────────────────────────────┘\n```\n\n**Every stage test job** includes a `check-pr-test-health` step after checkout — if any job in the run has already failed, the job fails fast (red X) with a root cause annotation.\n\n**Scheduled runs** skip `wait-for-base-*` jobs, running all stages in parallel. Fail-fast is also disabled.\n\n---\n\n## Fail-Fast Layers\n\n4 layers of fail-fast, from fine to coarse:\n\n| Layer | Mechanism | Granularity | Disabled on schedule? |\n|-------|-----------|-------------|----------------------|\n| **1. Test method → file** | `unittest -f` (failfast) | One test method fails → entire test file stops immediately | Yes |\n| **2. File → suite** | `run_unittest_files()` default | One test file fails → entire suite stops (`--continue-on-error` off) | Yes |\n| **3. Job → job (same stage)** | `check-pr-test-health` action | One job fails → other waiting jobs in same stage fail-fast (red X) | Yes |\n| **4. Stage → stage (cross-stage)** | `wait-for-base-*` + `needs` | Base A fails → base B/C jobs skip entirely (never get a runner) | Yes (wait jobs skipped) |\n\n- **Layer 1**: `-f` flag appended to all `python3 -m pytest` / `unittest` invocations in `ci_utils.py`\n- **Layer 2**: `--continue-on-error` flag in `run_suite.py` — off for PRs, on for scheduled runs\n- **Layer 3**: `check-pr-test-health` auto-detects `schedule` event and skips; filters out cascade failures to show only root cause jobs\n- **Layer 4**: `wait-for-base-*` jobs are conditioned on `github.event_name == 'pull_request'` — skipped for scheduled runs\n\n---\n\n## Execution Modes\n\n| Aspect | PR (`pull_request`) | Scheduled (`cron`, every 6h) | Manual dispatch (`workflow_dispatch`) |\n|--------|---------------------|------------------------------|--------------------------------------|\n| **Stage ordering** | Sequential: A → B → C via `wait-for-base-*` | Parallel (all at once) | Single target stage only |\n| **Cross-job fail-fast** | Yes (`check-pr-test-health`) | Yes | Yes |\n| **continue-on-error** | No (stop at first failure within suite) | Yes (run all tests) | No |\n| **Retry** | Enabled | Enabled | Enabled |\n| **max_parallel** | 3 (default), 14 if `max-concurrency` label | 14 | 3 (default), 14 if `max-concurrency` |\n| **PR gate** | Yes (draft, label, rate limit) | Skipped | Skipped |\n| **Concurrency** | `cancel-in-progress: true` per PR | Queue (no cancel) | Isolated per stage+SHA |\n\nFour labels relax these limits for one PR: `bypass-fail-fast`, `parallel-stages`,\n`max-concurrency`, and `highest-priority` (all three). `.github/scripts/ci-labels.cjs`\nresolves them; the [contribution guide](https://docs.sglang.io/developer_guide/contribution_guide.html#ci-control-labels)\ndescribes what each one does.\n\n---\n\n## Stage Gating (`wait-for-jobs` action)\n\n`wait-for-base-a` and `wait-for-base-b` are lightweight `ubuntu-latest` jobs that poll the GitHub Actions API.\n\n**How it works:**\n1. Calls `listJobsForWorkflowRun` to list all jobs in the current run\n2. Matches jobs by exact name or prefix (for matrix jobs, e.g., `base-b-test-1-gpu-small (3)`)\n3. If any matched job has `conclusion === 'failure'` → fail immediately (fail-fast)\n4. If all matched jobs are completed and count matches `expected_count` → success\n5. Otherwise → sleep `poll-interval-seconds` (default: 60s) and retry\n6. Timeout after `max-wait-minutes` of wall-clock time (240 min for base-a, 350 min for base-b). Keep budgets under the 6h hard cap on GitHub-hosted runners: a wait killed by the cap ends `cancelled`, not `timeout`, and `timeout-minutes` cannot raise it. The base-b wait sets `fail-on-timeout: 'false'`, so a timeout lets base-c start; `pr-test-finish` fails the run on any base-b failure regardless\n\n**Job specs example** (base-b):\n```json\n[\n {\"prefix\": \"base-b-test-1-gpu-small\", \"expected_count\": 8},\n {\"prefix\": \"base-b-test-1-gpu-large\", \"expected_count\": 14},\n {\"prefix\": \"base-b-test-2-gpu-large\", \"expected_count\": 4},\n {\"prefix\": \"base-b-test-4-gpu-b200\", \"expected_count\": 1}\n]\n```\n\n> **Critical**: `expected_count` must match the matrix size. If you add/remove matrix entries, update the wait job's spec accordingly.\n\n**PR only**: Condition `github.event_name == 'pull_request' && !inputs.target_stage` — scheduled runs and manual dispatches skip these entirely, allowing parallel execution.\n\n---\n\n## Cross-Job Fail-Fast (`check-pr-test-health` action)\n\nComposite action called after checkout in every stage test job (21 jobs total across `pr-test.yml`, `pr-test-multimodal-gen.yml`, `pr-test-sgl-kernel.yml`, `pr-test-jit-kernel.yml`).\n\n**How it works:**\n1. Queries `listJobsForWorkflowRun` for the current workflow run\n2. Filters for **root cause failures only** — jobs with `conclusion === 'failure'` whose failing step is NOT `check-pr-test-health` (excludes cascade failures)\n3. If root cause failures found → calls `core.setFailed()` with the list of root cause job names\n4. If none → does nothing (step succeeds)\n\n**Cascade filtering**: When job A fails fast due to the health check, it also has `conclusion: failure`. Without filtering, job B would list both the original failure AND job A's fail-fast. The filter checks each failed job's `steps` array — if the failing step name contains `check-pr-test-health` or `Check PR test health`, it's excluded from the root cause list.\n\n**Usage pattern:**\n```yaml\nsteps:\n - name: Checkout code\n uses: actions/checkout@v4\n ...\n\n - uses: ./.github/actions/check-pr-test-health\n id: pr-test-health\n\n - name: Install dependencies # skipped automatically if health check failed\n ... # (default if: success() is false)\n\n - name: Run test # also skipped\n ...\n```\n\n**Visual effect**: Job shows **red X** (failure) with error annotation showing root cause job names. Subsequent steps are naturally skipped (default `if: success()` is false after a failed step). No per-step `if` guards needed.\n\n**No stage filtering**: Checks ALL jobs in the run, not just the current stage. Any failure anywhere triggers fail-fast.\n\n**Error message example:**\n```\nFail-fast: skipping — root cause job(s): base-b-test-1-gpu-small (0), base-b-test-1-gpu-small (1)\n```\n\n---\n\n## Within-Suite Failure Handling\n\nControlled by `run_unittest_files()` in `python/sglang/test/ci/ci_utils.py`.\n\n### Flags\n\n| Flag | PR default | Scheduled default | Effect |\n|------|------------|-------------------|--------|\n| `--continue-on-error` | Off | On | Off: stop at first failure. On: run all files, report all failures at end |\n| `--enable-retry` | On | On | Retry retriable failures (accuracy/perf assertions) |\n| `--max-attempts` | 2 | 2 | Max attempts per file including initial run |\n\n### Retry Classification\n\nWhen a test fails and retry is enabled, the output is classified:\n\n**Non-retriable** (checked first — real code errors):\n`SyntaxError`, `ImportError`, `ModuleNotFoundError`, `NameError`, `TypeError`, `AttributeError`, `RuntimeError`, `CUDA out of memory`, `OOM`, `Segmentation fault`, `core dumped`, `ConnectionRefusedError`, `FileNotFoundError`\n\n**Retriable** (accuracy/performance):\n`AssertionError` with comparison patterns (`not greater than`, `not less than`, `not equal to`), `accuracy`, `score`, `latency`, `throughput`, `timeout`\n\n**Default**: Unknown `AssertionError` → retriable. Other unknown failures → not retriable.\n\n### How `continue_on_error` is set\n\nIn `pr-test.yml`'s `check-changes` job:\n- `schedule` runs or `run_all_tests` flag → `continue_on_error = 'true'`\n- PR runs → `continue_on_error = 'false'`\n\nEach test job propagates via:\n```yaml\nenv:\n CONTINUE_ON_ERROR_FLAG: ${{ needs.check-changes.outputs.continue_on_error == 'true' && '--continue-on-error' || '' }}\nrun: |\n python3 run_suite.py --hw cuda --suite $CONTINUE_ON_ERROR_FLAG\n```\n\n---\n\n## Test Partitioning\n\nLarge suites are split across matrix jobs using the **LPT (Longest Processing Time) heuristic** in `ci_register.py:auto_partition()`:\n\n1. Sort tests by `est_time` descending, filename as tie-breaker (deterministic)\n2. Greedily assign each test to the partition with smallest cumulative time\n3. Result: roughly equal total time per partition\n\n**Partition table** (CUDA per-commit suites):\n\n| Suite | Partitions | Runner | max_parallel |\n|-------|-----------|--------|-------------|\n| `base-a-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — |\n| `base-a-test-cpu` | 4 | `ubuntu-latest` | — |\n| `base-b-test-1-gpu-small` | 8 | `1-gpu-5090` | 8 |\n| `base-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) |\n| `base-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — |\n| `base-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-b-kernel-unit-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |\n| `base-b-kernel-benchmark-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — |\n| `base-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — |\n| `base-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — |\n| `base-c-test-deepep-4-gpu-h100` | 1 (no matrix) | `4-gpu-h100` | — |\n| `base-c-test-4-gpu-b200` | 3 | `4-gpu-b200` | — |\n| `base-c-test-4-gpu-b200-small` | 3 | `4-gpu-b200-low-disk` | — |\n| `base-c-test-8-gpu-b200` | registered only | `8-gpu-b200` | — |\n| `base-c-test-4-gpu-gb200` | registered only | `4-gpu-gb200` | — |\n\n> **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`.\n>\n> **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`.\n\n**Workflow usage:**\n```yaml\nstrategy:\n matrix:\n partition: [0, 1, 2, 3, 4, 5, 6, 7]\nsteps:\n - run: python3 run_suite.py --hw cuda --suite base-b-test-1-gpu-small \\\n --auto-partition-id ${{ matrix.partition }} --auto-partition-size 8\n```\n\n---\n\n## check-changes Job\n\nDetermines which test suites to run based on file changes.\n\n### Detection Methods\n\n| Trigger | Method | Details |\n|---------|--------|---------|\n| `pull_request` | `dorny/paths-filter` | Detects changes via GitHub diff |\n| `workflow_dispatch` (with `pr_head_sha`) | GitHub API | `repos/{repo}/compare/main...{sha}` |\n| `schedule` / `run_all_tests` | Force all true | Runs everything |\n\n### Output Flags\n\n| Output | Triggers |\n|--------|----------|\n| `main_package` | Base A/B/C test suites |\n| `sgl_kernel` | Kernel wheel builds + kernel test suites; also switches B200 jobs to kernel-build runner labels outside `target_stage` mode |\n| `jit_kernel` | JIT kernel test workflow |\n| `multimodal_gen` | Multimodal-gen test workflow |\n\n> **Note**: In `target_stage` mode, `sgl_kernel` is only active when `include_wheel_build=true`. Without that opt-in, kernel-change reruns fail validation instead of running a target stage without freshly built wheels. Outside `target_stage`, `sgl_kernel=true` switches B200 jobs from `4-gpu-b200` / `4-gpu-b200-low-disk` to `4-gpu-b200-kernel` / `4-gpu-b200-kernel-low-disk`.\n\n---\n\n## Concurrency Control\n\n```\ngroup: pr-test-{event_name}-{branch}-{pr_sha}-{stage}\n```\n\n| Segment | Source | Purpose |\n|---------|--------|---------|\n| `event_name` | `github.event_name` | Prevents scheduled runs colliding with fork PRs named `main` |\n| `branch` | `github.head_ref \\|\\| github.ref_name` | Per-branch isolation |\n| `pr_sha` | `inputs.pr_head_sha \\|\\| 'current'` | Isolates manual dispatches from main runs |\n| `stage` | `inputs.target_stage \\|\\| 'all'` | Allows parallel stage dispatches |\n\n`cancel-in-progress: true` for `pull_request` events (new push cancels old run), `false` for `workflow_call`.\n\n---\n\n## How To: Add a New Stage Job\n\n1. Define the job in `pr-test.yml` with `needs: [check-changes, call-gate, wait-for-base-X, ...]`\n2. Copy the `if:` condition pattern from an existing same-stage job (handles `target_stage`, `schedule`, `main_package`)\n3. Add `checkout` step\n4. Add `check-pr-test-health` step (after checkout) — if any prior job failed, `core.setFailed()` fires and all subsequent steps auto-skip via default `if: success()`\n5. Add `check-maintenance` step\n6. Add `download-artifact` step if `sgl_kernel` changed\n7. Add `install dependencies` step\n8. Add `run test` step with `$CONTINUE_ON_ERROR_FLAG`\n9. Add `upload-cuda-coredumps` step with `if: always()`\n10. Register the suite name in `PER_COMMIT_SUITES` in `test/run_suite.py`\n11. If using matrix, add `--auto-partition-id` and `--auto-partition-size` to the run command\n12. **Update `wait-for-base-X`** job spec with the new job name and `expected_count` (if matrix)\n13. **Add the job to `pr-test-finish.needs`** list\n\n---\n\n## How To: Debug CI Failures\n\n| Symptom | Likely cause | What to check |\n|---------|-------------|---------------|\n| All stage-B/C jobs green but steps skipped | Earlier job failed, `check-pr-test-health` triggered | Find the actual failed job (red X) |\n| `wait-for-base-b` timeout | `expected_count` doesn't match matrix size | Verify job spec counts match `matrix:` array length |\n| `pr-test-finish` fails but all jobs green | A job was `cancelled` (counts as failure in finish) | Check concurrency cancellation, and whether a `wait-for-base-*` job hit the 6h hosted-runner cap |\n| Tests pass locally but fail in CI | Partition assignment, runner GPU type, or `est_time` inaccuracy | Check which partition the test lands in; verify runner label |\n| Flaky test retried and passed | Retriable failure (accuracy/perf) | Check `[CI Retry]` markers in job logs |\n| Flaky test NOT retried | Matched non-retriable pattern | Check if error matches `NON_RETRIABLE_PATTERNS` in `ci_utils.py` |\n\n---\n\n## Slash Commands\n\n| Command | Effect |\n|---------|--------|\n| `/tag-run-ci-label` | Adds `run-ci` label to PR |\n| `/tag-run-ci-label extra` | Adds both `run-ci` and `run-ci-extra` labels |\n| `/rerun-failed-ci` | Reruns failed jobs in the latest workflow run |\n| `/tag-and-rerun-ci` | Adds `run-ci` label + reruns failed |\n| `/tag-and-rerun-ci extra` | Adds both `run-ci` and `run-ci-extra` labels + reruns failed |\n| `/run-full-ci` | Short form of `/tag-and-rerun-ci extra` (baseline + extra). Alias: `/rerun-full-ci` |\n| `/run-extra-ci` | Adds both labels + reruns **only** `PR Test Extra`, leaving baseline runs alone. Alias: `/rerun-extra-ci` |\n| `/rerun-test [ ...]` | Reruns specific test file(s) via `rerun-test.yml`. A file arg containing a glob metacharacter (`*`, `?`, `[...]`) expands against `test/registered/` and the multimodal test dir to every matching `test_*.py` (e.g. `/rerun-test test_*backend*.py` — wrap in backticks so GitHub doesn't italicize the `*`); matches are deduped, grouped by dispatch shape, and can't carry a `::test` selector. No match → single ⛔ reply, nothing dispatched. |\n| `/rerun-group [ ...]` | Expands registered test groups, then reuses `/rerun-test` |\n| `/run-cpu-ci` | Dispatches `run-cpu-ci.yml` on main: only the `base-a-test-cpu` suite, at the PR head SHA, no caches saved. Same gate as `/rerun-test`. Alias: `/rerun-cpu-ci` |\n\nHandled by `scripts/ci/utils/slash_command_handler.py` → `.github/workflows/slash-command-handler.yml`.\n\n### Label-gated workflow dispatch (pr-test, pr-test-extra)\n\n`pr-test.yml` and `pr-test-extra.yml` both listen for `pull_request.labeled` (in addition to `opened`/`synchronize`/`reopened`). The `check-changes.if` gate has two clauses:\n\n1. **For `labeled` events**: the just-added label must be one of the gating labels (`run-ci` for pr-test, `run-ci` or `run-ci-extra` for pr-test-extra) — otherwise every unrelated label addition would dispatch a full CI run.\n2. **All events**: the PR must currently carry the required labels.\n\nThis is what lets `/tag-run-ci-label` (and the `extra` variant) trigger a fresh CI run without an extra push.\n\n**Caveat — skipped runs cannot be un-skipped by `run.rerun()`:** GitHub's rerun API reuses the original event payload, so rerunning a `pull_request`-event run that was skipped because of missing labels will skip again (label set in the frozen payload doesn't update). The only way to recover a label-skipped run is to add the missing label, which fires a fresh `labeled` event with the current label set. `handle_rerun_failed_ci` in the slash handler is for rerunning failed/non-label-skipped runs; it cannot revive label-skipped ones.\n",".agents/skills/add-jit-kernel/SKILL.md":"---\nname: add-jit-kernel\ndescription: Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups\n---\n\n# Tutorial: Adding a New JIT Kernel to SGLang\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as a JIT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float, passed at runtime)\n- Output: `x * factor` (element-wise), allocated internally\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n\n## When to use JIT vs AOT (`sgl-kernel`)\n\n- **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.\n- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `python/sglang/kernels/aot/` and participate in the wheel build / torch op registration flow.\n- **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.\n\n---\n\n## Conventions\n\nThese hold for every step below.\n\n- **`namespace sglang` is where JIT code lives.** Open it after the include block and close it at the end of the file, with the device kernels, traits and host wrapper inside. The shared `host::` / `device::` helpers are nested in it too, so they resolve unqualified. `load_jit` emits the `TVM_FFI_DLL_EXPORT_TYPED_FUNC` wrapper inside `namespace sglang` as well, so the `kernel_name` you pass from Python needs no `sglang::` prefix.\n- **Tuning policy belongs in Python; C++ gets the answer, not the decision.** For a hyperparameter — split factor, threads per item, block size, vector width, an algorithm variant — prefer (not required) making it a **template parameter** and letting the `@cache_once` module factory choose the value, over a runtime `if`/`switch` in the launcher that picks among pre-instantiated kernels. The kernel then compiles exactly one specialisation and asserts its own preconditions with `static_assert`, while the heuristic that produced the value sits in Python where it is readable, adjustable, and inspectable without recompiling CUDA. `store_cache`'s `num_threads` is the worked example: a `get_kernel(num_split)` ladder guarded by byte-alignment `if constexpr`s became one template argument plus a heuristic in `_jit_kvcache_module`. The exception is a knob that genuinely varies per call with the runtime shape — that has to stay a kernel argument.\n- **Check where the check is cheapest: `static_assert` > C++ host check > cached Python > per-call Python.** Anything fixed at compile time is a `static_assert`. Anything about the tensors is a `TensorMatcher` / `CHECK_HOST` in the C++ launcher, free next to a kernel launch. A check Python cannot delegate goes inside the `@cache_once` module factory, where it runs once per specialisation. What remains in the per-call entry point costs interpreter time on *every* forward, so it should be nothing but picking the module and allocating `out`.\n- **Fixed-width integer types.** Prefer `int32_t` / `int64_t` / `uint32_t` / `size_t` over `int`, `long`, or `long long`, so an index has the same width on both sides of the FFI boundary. Bare `int` is fine only where the width plainly cannot matter — an unrolled loop counter over a `constexpr` bound, a template `int` parameter. Shapes arrive as `int64_t` (`SymbolicSize::unwrap()`); narrowing to `uint32_t` for in-kernel indexing is a deliberate act, so write the `static_cast` explicitly and only where the range is known.\n- **Doxygen comments in C++.** Document exported entities with `///` or `/** ... */` blocks using `\\brief`, `\\param`, `\\tparam`, `\\return`, the way `include/sgl_kernel/` does. `python -m sglang.kernels.jit` writes `CommentFormat: Doxygen` into `.clangd` when clangd is 21 or newer, so these render on hover in the editor. Plain `//` remains fine for implementation notes inside a function body.\n- **ASCII only in C++ and CUDA sources.** Write `--`, `->`, `<=` instead of `—`, `→`, `≤`, including in comments. `grep -nP '[^\\x00-\\x7F]' ` before committing.\n- **`namespace details` is private.** Anything inside one is an implementation detail of its own header, free to change without notice. Do not name `details::` from another module, and never `using namespace details`. If you find yourself wanting something in there, that is the signal to promote it to a documented name instead.\n- **`const T* __restrict__` for read-only pointers.** This is what `csrc/` does throughout, and it lets the compiler emit non-coherent (`LDG`) loads.\n- **Watch the register budget.** For memory-bound kernels, keep to roughly 64 registers per thread so occupancy does not become the limit; prefer recomputing a value over letting it spill. Run `SGLANG_JIT_LOG_RESOURCE_USAGE=1 SGLANG_JIT_FORCE_RECOMPILE=1` to log per-kernel registers, spills and shared memory at INFO. Both halves are load-bearing: a cache hit has no compiler output to report, and passing `-Xptxas -v` by hand does nothing on its own because the build captures compiler output and replays it only on failure.\n- **Alignment is a contract you enforce, not a property you hope for.** A vectorized copy derives its access width from the row *size* (`load_bytes` picks 16B for a 1024B row), but the address it lands on is `base + index * row_stride` — and a stride is not constrained by the size. `TensorMatcher` admits any stride unless you say otherwise, so a padded cache row silently produces `misaligned address`. Every kernel that vectorizes is highly recommended to end its validation with `.ensure_alignment(w)`, where `w` is the width the kernel actually uses.\n\n---\n\n## Common Abstractions in `python/sglang/kernels/jit/include/sgl_kernel/`\n\n**Always prefer these abstractions over raw CUDA primitives.** They provide safety, readability, and consistency with the rest of the codebase. The only reason to drop to raw primitives is performance the abstraction cannot reach — a trade you make deliberately, and justify in a comment.\n\n### `utils.h` — Host-side utilities\n\n```cpp\n#include \n```\n\n- **`CHECK_HOST(cond) << \"msg \" << value`** — **Preferred** runtime check: stream-style, throws `PanicError` with file/line info on failure. Zero overhead on the true path — the message expressions are only evaluated when the check fails.\n- **`host::RuntimeCheck(cond, args...)`** — Function-style alternative to `CHECK_HOST`. Note its message args are always evaluated (even when the check passes), so prefer `CHECK_HOST` — especially on hot paths.\n- **`host::Panic(args...)`** — Unconditionally throw a `PanicError` with a descriptive message.\n- **`host::div_ceil(a, b)`** — Integer ceiling division `(a + b - 1) / b`.\n- **`host::irange(n)`** / **`host::irange(start, end)`** — Range views for cleaner loops.\n- **`host::pointer::offset(ptr, offsets...)`** — Byte-safe pointer arithmetic on `void*`. Use this instead of raw casts.\n\n### `utils.cuh` — Device-side utilities + `LaunchKernel`\n\n```cpp\n#include \n```\n\n- **Type aliases**: `fp16_t`, `bf16_t`, `fp32_t`, `fp8_e4m3_t`, `fp8_e5m2_t` and their packed variants `fp16x2_t`, `bf16x2_t`, `fp32x2_t`, etc.\n- **`SGL_DEVICE`** — Expands to `__forceinline__ __device__`. Use on all device functions.\n- **`SGL_DEVICE_HOST`** — `__forceinline__ __device__ __host__`. Use it when a `constexpr` helper must be callable from both the kernel and its launcher, so the two cannot drift.\n- **`device::kWarpThreads`** — Constant `32` (the codebase's *logical* warp width). `kWarpSize` is an alias. Note this is a group size, not the hardware wave: on gfx950 `warpSize` is 64, and `warp::kFullWidth` reflects that.\n- **`device::get_lane_id()`** — this thread's index within its logical group. **Prefer it over `threadIdx.x % kNumThreads` whenever the value feeds an address.** On CUDA the lane register is one read that folds into the address computation, whereas the modulo is a derived value the compiler re-materialises at *every* address scale — 8 instructions and 2 registers on a two-tile warp copy, measured on sm_100a. It picks the cheaper form per platform, so callers never need to care. Only equals the true in-warp lane when `blockDim.x` is a multiple of `kNumThreads`.\n- **`device::load_as(ptr, offset)`** / **`device::store_as(ptr, val, offset)`** — Type-safe loads/stores from `void*`.\n- **`device::pointer::offset(ptr, offsets...)`** — Pointer arithmetic on device.\n- **`host::LaunchKernel(grid, block, device_or_stream [, smem])`** — RAII kernel launcher that:\n - Resolves the CUDA stream from a `DLDevice` via TVM-FFI automatically.\n - Checks the CUDA error with file/line info after launch via `operator()(kernel, args...)`.\n - Supports `.enable_pdl(bool)` for PDL (Programmatic Dependent Launch, SM90+).\n- **`device::PDLWaitPrimary()`** / **`device::PDLTriggerSecondary()`** — The two halves of PDL, on sm_90+ (no-ops on older archs and ROCm). Their guarantees are **not** symmetric:\n - `PDLTriggerSecondary` (`griddepcontrol.launch_dependents`) only lets the next kernel in the stream *start* early. It carries no memory ordering and publishes nothing — matching that, the header's asm has no `\"memory\"` clobber.\n - `PDLWaitPrimary` (`griddepcontrol.wait`) is the ordering point: it waits until the preceding kernel has fully finished and its writes are visible.\n\n So every read of data the preceding kernel produced must come after `PDLWaitPrimary()`. What overlaps with the primary's tail is whatever you put *before* the wait — loading parameters, computing indices, touching buffers the primary never wrote — so a kernel that waits on its first line gains nothing. Neither call is a barrier: threads may reach or skip them independently. See \"Programmatic Dependent Launch and Synchronization\" in the CUDA C++ Programming Guide.\n- **`CHECK_CUDA(expr) << \"context\"`** — Stream-style CUDA error check; evaluates `expr` once and throws `PanicError` with `cudaGetErrorString` + file/line info if it is not `cudaSuccess`. Extra streamed context is optional.\n- **`host::RuntimeDeviceCheck(cudaError_t)`** — Function-style alternative to `CHECK_CUDA`. It takes no context message, so prefer `CHECK_CUDA`, which builds its error object only on the failure path.\n\n### `tensor.h` — Tensor validation (`TensorMatcher`, Symbolic types)\n\n```cpp\n#include \n```\n\nThis is the **primary validation API** for all kernel launchers. Use it to validate every `tvm::ffi::TensorView` argument.\n\n- **`host::SymbolicSize{\"name\"}`** — A named symbolic dimension. Call `.set_value(n)` to pin it, `.unwrap()` to extract after verification.\n- **`host::SymbolicDType`** — Symbolic dtype. Use `.set_options()` to restrict allowed types.\n- **`host::SymbolicDevice`** — Symbolic device. Use `.set_options()` to restrict to CUDA.\n- **`host::TensorMatcher({dims...})`** — Fluent builder for tensor validation:\n - `.with_dtype()` — require a specific C++ type (e.g. `fp16_t`)\n - `.with_dtype()` — allow a set of types\n - `.with_device(device_sym)` — require CUDA and bind the checked device to a `SymbolicDevice`\n - `.with_strides({strides...})` — validate strides (omit to require contiguous)\n - `.ensure_alignment(bytes)` — require the data pointer **and every stride except the innermost** to be a multiple of `bytes` (a power of two). Size-1 dimensions are skipped, since their strides are arbitrary. This is how a vectorized kernel states its precondition; see the alignment convention above.\n - `.verify(tensor_view)` — execute the check; throws `PanicError` with full context on failure; **chainable** (`verify(a).verify(b)` to check multiple tensors with the same shape)\n- **`host::is_type(dtype)`** — whether a `DLDataType` denotes the C++ type `T` (e.g. `fp16_t`).\n\n**Typical pattern:**\n```cpp\nauto N = SymbolicSize{\"num_elements\"};\nauto device = SymbolicDevice{};\ndevice.set_options();\nTensorMatcher({N}) //\n .with_dtype()\n .with_device(device)\n .ensure_alignment(16) // e.g. for 128-bit vectorized load/store\n .verify(dst)\n .verify(src); // same shape, dtype, device as dst\nconst int64_t n = N.unwrap();\nconst DLDevice dev = device.unwrap();\nconst int64_t last_dim = 128;\nTensorMatcher({N, last_dim}) // a fixed dimension can be a plain integer\n .with_dtype()\n .with_device(device)\n .verify(tensor_2d);\n```\n\n### `ffi.h` — Tensor allocation and blob wrapping (`host::ffi::`)\n\n```cpp\n#include \n```\n\nThe counterpart to `tensor.h`: that one validates what came in, this one produces new `tvm::ffi::Tensor` values. Allocation goes through the environment allocator (`TVMFFIEnvTensorAlloc`), so buffers come from PyTorch's caching allocator rather than a raw `cudaMalloc`.\n\n- **`host::alloc_workspace_tensor(nbytes, device)`** (declared in `utils.cuh`) — **the way to get scratch memory**: a 1-D `uint8` tensor of `nbytes`, or an empty tensor when `nbytes == 0`. Hold the returned `Tensor` in a local across every launch that touches it — it frees on destruction.\n- **`host::ffi::empty(shape, dtype, device)`** — Uninitialized tensor; `shape` accepts a braced list, so `ffi::empty({rows, sizeof(Plan)}, dtype, device)` works for a typed scratch array.\n- **`host::ffi::empty_like(tensor_view)`** — Same shape, dtype, and device as an existing tensor.\n- **`host::ffi::from_blob(data, shape, dtype, device[, deleter, stride, byte_offset])`** / **`from_blob_like(data, tensor_view, ...)`** — View memory you already own as a `Tensor`, no copy. The default deleter does nothing, so ownership stays with the caller; pass one only when the `Tensor` should own the block. Strides default to contiguous.\n\n### `type.cuh` — `DTypeTrait`, `packed_t`, and reduction traits\n\n```cpp\n#include \n```\n\n- **`DTypeTrait`** — Static trait struct, specialized for integral types, `fp32_t`, `fp16_t`, `bf16_t`, `fp8_e4m3_t`, and their packed x2/x4 variants. Provides:\n - `DTypeTrait::from(value)` — convert from another type via the right CUDA intrinsic (e.g. `fp32_t` → `fp16_t`)\n - `DTypeTrait::abs/max/min` — type-dispatched math (fp32, fp16/bf16 scalar and x2, integrals)\n - `DTypeTrait::sqrt/rsqrt/exp/sin/cos(x)` — `fp32_t` only\n - Metadata: `packed_t` / `unpacked_t` / `kVecSize` (packed layout), `kFloatMax` (dtype max as float, e.g. 448.0f for fp8-e4m3), `kZeroBits`\n- **`packed_t`** — Two-element packed alias: `packed_t` = `fp16x2_t`, `packed_t` = `bf16x2_t`, `packed_t` = `fp32x2_t`. Use for vectorized loads/stores.\n- **`device::cast(value)`** — Type-safe cast using `DTypeTrait`, e.g. `cast(v)`.\n- **`device::unpack(value)`** — View a packed value as an `unpacked_t[kVecSize]` array reference (e.g. `fp32x2_t` → `fp32_t[2]`); element writes propagate back to the packed value.\n- **`device::ReductionOp` (`SUM`/`MAX`/`MIN`) and `device::ReductionTrait::reduce(x, y)`** — One binary reduction step, dispatched through `DTypeTrait` (packed types reduce elementwise). This is the engine behind `warp::reduce`; use it directly when writing custom reductions.\n\n### `vec.cuh` — Vectorized memory access (`AlignedVector`)\n\n```cpp\n#include \n```\n\n- **`device::AlignedVector`** — Aligned storage for N elements of type T. N must be a power of two, `sizeof(T)*N <= 32`. Enables vectorized loads/stores for bandwidth efficiency. In terms of API/codegen constraints, the upper bound is 256-bit; in practice, 128-bit is the portable default, while 256-bit vectorization is typically only viable on `SM100+` and should be gated by an architecture check when needed.\n - `.load(ptr, offset)` — vectorized load from `ptr[offset]`\n - `.store(ptr, offset)` — vectorized store to `ptr[offset]`\n - `.fill(value)` — fill all N elements with `value`\n - `operator[](i)` — element access\n- **`device::LoadStoreBytes`** — named widths to pass around instead of bare integers: `MAX_GMEM` (arch-dependent, 32 on Blackwell), `MAX_SMEM` (16), `MAX_PORTABLE` (16, safe on CUDA and HIP), `MIN_COALALESCED` (4), and `RAW_1B` .. `RAW_32B`.\n\n### `tile.cuh` — `tile::Memory` (strided memory access pattern)\n\n```cpp\n#include \n```\n\n- `tile::Memory` is fundamentally a **1D cooperative accessor** over a contiguous region.\n- **`device::tile::Memory::cta(blockDim.x)`** — Creates a tile accessor where each thread handles `tid = threadIdx.x` with stride `tsize` (for `cta(blockDim.x)`, this is `blockDim.x`). Common for loops over a 1D array.\n- **`device::tile::Memory::warp()`** — the 32-lane flavour, and what `warp::load_bytes` is built on. It takes its lane index from `get_lane_id()` rather than `threadIdx.x % 32`, for the addressing reason described under `utils.cuh`. Use `warp(int n)` for a narrower sub-group; that overload keeps the modulo, since a hardware lane index is the wrong grouping below 32.\n- **`device::tile::Memory::thread()`** — single-thread accessor, no cooperation.\n- **`.load(ptr, offset)`** — loads `ptr[tid + offset * tsize]`\n- **`.store(ptr, val, offset)`** — stores to `ptr[tid + offset * tsize]`\n- **`.in_bound(n, offset)`** — boundary check\n\nFor a **2D tile**, either flatten `(row, col)` into a linear tile index first, or compute the address manually with `ptr[row * stride + col]` using your thread/block coordinates.\n\n### `bits.h` — Compile-time bit helpers (`host::`)\n\n```cpp\n#include \n```\n\n`constexpr` wrappers over ``, usable in `static_assert` and in template arguments: `host::is_pow2(x)`, `log2_floor(x)`, `log2_ceil(x)` (both `-1` for `x == 0`), `round_up_pow2(x)`, `round_down_pow2(x)`. They live in `host::` but are constant-expression-only, so device-side `static_assert` can call them.\n\n### `math.cuh` — Device math (`device::math::`)\n\n```cpp\n#include \n```\n\n- `device::math::max/min(a, b)` — type-dispatched binary math via `DTypeTrait`\n- `device::math::abs/sqrt/rsqrt/exp/sin/cos(x)` — type-dispatched unary math via `DTypeTrait`\n\n### `warp.cuh` — Warp-level primitives\n\n```cpp\n#include \n```\n\n**Vectorized row copy — reach for this before hand-rolling.** A warp cooperatively moving one contiguous row is the single most common shape in this tree, and it used to be re-implemented per kernel.\n\n- **`warp::load_bytes(src)`** / **`warp::store_bytes(dst, val)`** — the whole warp moves `kBytes` from/to one contiguous row. `kBytes` is arbitrary, down to 1; the helper picks the widest vector that divides it and handles the ragged tail. The returned value is opaque and encodes the width, so a load and its matching store **must use the same ``**.\n- **`warp::LoadStorePattern`** — the `kPattern` values, and the sign carries the meaning. A **negative** one (`WARP_UNIFORM_16B` and friends: `_GMEM`, `_SMEM`, `_4B`/`_8B`/`_32B`) says *the whole warp splits this row*, so the width is derived from the per-lane share (`gcd(kBytes / 32, |kPattern|)`), falling back to `gcd(kBytes, 4)` when the row is too narrow to give every lane 4 bytes. A **positive** one (a bare `LoadStoreBytes` value) makes no warp-splitting assumption and just caps the per-thread vector at `gcd(kBytes, kPattern)`. `WARP_UNIFORM_16B` is the usual choice for a row copy.\n- **`LoadStorePattern::get_vec_bytes()`** — the width actually chosen. `SGL_DEVICE_HOST`, so **the launcher must call this to feed `.ensure_alignment(...)`**. Pass it the same value the kernel passes `load_bytes` — for a split row that is the *per-warp share*, not the whole row; the full row can resolve to a different width and would under-constrain the strides.\n\n**Reductions.**\n\n- `warp::reduce(value, active_mask)` — `Op` is a `device::ReductionOp` (`SUM`/`MAX`/`MIN`). `kStart` and `kFinish` bound a range of lane-index bits: `` reduces within contiguous groups of `N` lanes, `` reduces *across* groups at the same offset. The operation is **symmetric** — `` and `` reduce the same lane set.\n- `warp::reduce_sum/reduce_max/reduce_min(value)` — wrappers. Work for any type with a `ReductionTrait`: floats, integers, packed x2. This is more widely used in practice than generic `reduce`.\n- `warp::inclusive_reduce(val, lane_id, mask)` and `inclusive_sum` / `inclusive_max` / `inclusive_min` — segmented inclusive scan: every lane keeps its own running total, unlike `reduce`. Forward when `kStart < kFinish`, backward when `kStart > kFinish`. `lane_id` defaults to `get_lane_id()`, which is correct by construction — only pass it if you already have the **segment-relative** index (`threadIdx.x % kWidth`); a warp-relative one silently corrupts every segment past the first when `kWidth < 32`.\n- `warp::broadcast(value, src_lane)` — `src_lane` is **segment-relative**: each `kWidth` segment reads its own lane, not one warp-wide source.\n- `warp::get_lane_id`, `warp::elect_one_lane()` (one elected lane, for gating a single-thread TMA issue; CUDA-only), `warp::kFullMask` and `warp::kFullWidth` (32 on CUDA, 64 on HIP — note this differs from `kWarpThreads`).\n\n### `cta.cuh` — CTA-level primitives\n\n```cpp\n#include \n```\n\n- `device::cta::reduce_max(value, smem, min_value)` — CTA-wide max using shared memory + warp reduction. Caller is responsible for a `__syncthreads()` after if the result in `smem[0]` is needed.\n\n### `atomic.cuh` — Atomic operations\n\n```cpp\n#include \n```\n\n- `device::atomic::max(float* addr, float value)` — float atomic max (handles negative values correctly via bit tricks).\n- **`device::atomic::Event`** — a cross-CTA arrive/wait counter in one 32-bit word. Producers call `arrive()`; consumers call `wait(num_producers)` (exactly one consumer) or `wait_multi(num_producers, num_consumers, n)` (several). The last released consumer resets the word, so one `Event` is reusable across launches with no host re-zero. Preconditions, all of them load-bearing:\n - the word must be **zeroed by the host** before first use, and must live in **global** memory — a shared or local `Event` is an illegal address;\n - `arrive()` is a release and the waits are acquires, so a producer's writes before `arrive()` are visible after the wait;\n - **generations must be explicitly ordered** — the word has no phase bit, so overlapping generation N+1 with a consumer still in generation N is undefined behaviour (a lagging consumer never exits);\n - `wait()` and `wait_multi()` lay the word out incompatibly; mixing them on one `Event` is undefined behaviour.\n\n### `runtime.cuh` — Occupancy and device info\n\n```cpp\n#include \n```\n\n- `host::runtime::get_blocks_per_sm(kernel, block_dim)` — max active blocks per SM (occupancy)\n- `host::runtime::get_sm_count(device_id [, use_cache])` — number of SMs on the device. Memoized per device ordinal, so it is cheap enough to call on a launch path; pass `use_cache=false` to force a driver query.\n- `host::runtime::get_cc_major` / `get_cc_minor` / `get_sm_version` — compute capability, same caching.\n\n**Do not query the architecture at runtime.** The JIT compiles for the exact local GPU, so the arch is a *compile-time* fact: `SGL_CUDA_ARCH` is injected by `load_jit`, and `SGL_ARCH_HOPPER_OR_GREATER` / `SGL_ARCH_BLACKWELL_OR_GREATER` / `device::kMaxVecBytes` are derived from it. Reach for `get_cc_*` only for something the arch genuinely does not determine. SM *count* is the opposite case — it varies within an arch (B200 vs B300 vs a MIG slice), so it has to stay a runtime query.\n\n**Persistent kernel pattern** (cap blocks to SM count × occupancy):\n```cpp\nstatic const uint32_t max_occ = runtime::get_blocks_per_sm(kernel, kBlockSize);\nstatic const uint32_t num_sm = runtime::get_sm_count(device.unwrap().device_id);\nconst auto num_blocks = std::min(num_sm * max_occ, div_ceil(n, kBlockSize));\nLaunchKernel(num_blocks, kBlockSize, device.unwrap())(kernel, params);\n```\n\n---\n\n## Step 0 (optional): Generate a `.clangd` config for better IDE support\n\n```bash\npython -m sglang.kernels.jit -h # for verbose help info about clangd configuration\npython -m sglang.kernels.jit\npython -m sglang.kernels.jit --dep cutlass flashinfer # with cutlass/flashinfer dependency\n```\n\n---\n\n## Step 1: Implement the CUDA kernel in `kernels/jit/csrc/`\n\nCreate `python/sglang/kernels/jit/csrc/elementwise/scale.cuh`.\n\nThe implementation fully uses the project abstractions described above:\n\n```cpp\n// NOTE: Comments for headers are not common in practice.\n// It is only shown here for tutorial purposes to highlight the key abstractions.\n#include // For TensorMatcher, SymbolicSize, SymbolicDevice\n#include // For DTypeTrait, fp16_t, bf16_t, fp32_t\n#include // For CHECK_HOST, div_ceil\n#include // For LaunchKernel, SGL_DEVICE\n#include // For AlignedVector\n\n#include \n#include \n\nnamespace sglang {\n\n/**\n * \\brief Element-wise scale using vectorized 128-bit loads/stores.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kVecN Elements per vector load (e.g. 8 for fp16)\n * \\tparam kUsePDL Whether to emit the PDL wait/trigger pair\n * \\param dst Output buffer, `n_total` elements\n * \\param src Input buffer, `n_total` elements\n * \\param factor Runtime scale factor\n * \\param n_total Number of elements to scale\n */\ntemplate \n__global__ void scale_kernel(T* __restrict__ dst,\n const T* __restrict__ src,\n float factor,\n uint32_t n_total) {\n using vec_t = device::AlignedVector;\n const uint32_t n_vecs = n_total / kVecN;\n\n // If using PDL, wait for primary kernel before any global memory load.\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLWaitPrimary();\n\n // --- vectorised body ---\n const uint32_t vec_stride = blockDim.x * gridDim.x;\n for (uint32_t vi = blockIdx.x * blockDim.x + threadIdx.x;\n vi < n_vecs;\n vi += vec_stride) {\n vec_t v;\n v.load(src, vi);\n#pragma unroll\n for (int i = 0; i < kVecN; ++i) {\n v[i] = static_cast(static_cast(v[i]) * factor);\n }\n v.store(dst, vi);\n }\n\n // --- scalar tail ---\n const uint32_t base = n_vecs * kVecN;\n const uint32_t scalar_stride = blockDim.x * gridDim.x;\n for (uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;\n base + i < n_total;\n i += scalar_stride) {\n dst[base + i] = static_cast(static_cast(src[base + i]) * factor);\n }\n\n // If using PDL, signal for the secondary kernel to start after all threads have finished\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLTriggerSecondary();\n}\n\n/**\n * \\brief Validate the tensors, select the vector width, launch `scale_kernel`.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kUsePDL Whether to launch with PDL enabled\n * \\param dst Output tensor; same shape / dtype / device as `src`\n * \\param src Input tensor on CUDA\n * \\param factor Runtime scale factor\n */\ntemplate \nvoid scale(tvm::ffi::TensorView dst, tvm::ffi::TensorView src, float factor) {\n using namespace host;\n\n // 1. Validate input tensors with TensorMatcher\n SymbolicSize N = {\"num_elements\"};\n SymbolicDevice device_;\n device_.set_options();\n\n TensorMatcher({N}) //\n .with_dtype()\n .with_device(device_)\n .verify(dst)\n .verify(src); // same shape / dtype / device as dst\n\n const uint32_t n = static_cast(N.unwrap());\n const DLDevice device = device_.unwrap();\n\n CHECK_HOST(n > 0) << \"scale: num_elements must be > 0, got \" << n;\n\n // 2. Choose vector width for 128-bit loads (16 bytes)\n // fp16/bf16: 8 elements x 2 bytes = 16 bytes\n // fp32: 4 elements x 4 bytes = 16 bytes\n // We encourage using `device::kMaxVecBytes`, which will change according to\n // the target architecture and can enable 256-bit vectorization on SM100+ if desired.\n // But 128-bit is more commonly adapted for better compatibility,\n // so it's still ok to hardcode 16 here just for simplicity.\n constexpr int kVecN = 16 / sizeof(T);\n const uint32_t n_work_items = div_ceil(n, static_cast(kVecN));\n\n // 3. Launch\n constexpr uint32_t kBlockSize = 256;\n const uint32_t grid = div_ceil(n_work_items, kBlockSize);\n\n // PDL feature is 100% optional. Without `enable_pdl`, the code should still be correct.\n // Try to enable it if profiling shows that it can benefit the performance of this kernel.\n LaunchKernel(grid, kBlockSize, device).enable_pdl(kUsePDL)(\n scale_kernel,\n static_cast(dst.data_ptr()),\n static_cast(src.data_ptr()),\n factor,\n n);\n}\n\n} // namespace sglang\n```\n\n**Key points:**\n\n- Include headers from `sgl_kernel/` — **not** raw CUDA headers for anything already covered\n- Use `TensorMatcher` for all tensor validation; never manually check shape/dtype/device\n- Use `AlignedVector` for vectorised 128-bit loads/stores — significant bandwidth win\n- Use `LaunchKernel` — it resolves the stream and checks errors automatically\n- Use `CHECK_HOST(cond) << ...` for runtime assertions with useful error messages (zero overhead when the check passes)\n- Prefer passing runtime scalars like `factor` directly unless compile-time specialisation is genuinely required\n- `fp16_t` / `bf16_t` / `fp32_t` are the project's type aliases (from `utils.cuh`)\n- `device::cast` or `DTypeTrait::from(val)` for cross-type conversions\n- `device::math::` functions for device math instead of bare `__` intrinsics if possible.\n- Consider PDL — it can help when the kernel has prologue work to overlap. Place `PDLWaitPrimary()` right before the first read of upstream data, not at the top of the kernel\n\n---\n\n## Step 2: Add the Python wrapper in `kernels/ops/`\n\nThe wrapper lives next to its functional group under `python/sglang/kernels/ops/`, not beside the CUDA source — `kernels/jit/` holds only the JIT infrastructure (`csrc/`, `include/`, `utils/`, `benchmark/`). Create `python/sglang/kernels/ops/elementwise/scale.py`:\n\n```python\nfrom __future__ import annotations\n\nfrom typing import TYPE_CHECKING\n\nimport torch\n\nfrom sglang.kernels.jit.utils import (\n cache_once,\n is_arch_support_pdl,\n load_jit,\n make_cpp_args,\n)\n\nif TYPE_CHECKING:\n from tvm_ffi.module import Module\n\n\n@cache_once\ndef _jit_scale_module(dtype: torch.dtype) -> Module:\n \"\"\"Compile and cache the JIT scale module for a given dtype.\"\"\"\n # Checks on the compile key live here, not in `scale`: `cache_once` keys on\n # `dtype`, so this runs once per specialisation instead of once per call.\n if dtype not in (torch.float16, torch.bfloat16, torch.float32):\n raise RuntimeError(\n f\"Unsupported dtype {dtype}. Supported: float16, bfloat16, float32\"\n )\n args = make_cpp_args(dtype, is_arch_support_pdl())\n return load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n )\n\n\ndef scale(src: torch.Tensor, factor: float, out: torch.Tensor | None = None) -> torch.Tensor:\n \"\"\"\n Element-wise scale: dst = src * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n src : CUDA tensor (FP16 / BF16 / FP32)\n factor : scale factor\n out : optional pre-allocated output tensor (same shape/dtype as src)\n\n Returns\n -------\n Scaled tensor (dst = src * factor).\n \"\"\"\n # DO NOT add proactive validation here: every check costs interpreter time\n # on a per-forward path. Tensor invariants belong in the C++ launcher, and\n # anything about the compile key belongs in `_jit_scale_module`.\n if out is None:\n out = torch.empty_like(src)\n\n module = _jit_scale_module(src.dtype)\n module.scale(out, src, factor)\n return out\n```\n\n**Key points:**\n\n- Use `cache_once` — **not** `functools.lru_cache` (incompatible with `torch.compile`)\n- `load_jit` first arg(s) form the unique build marker; same marker = same cached binary\n- Only include compile-time specialisation knobs in the build marker; runtime values like `factor` should stay runtime unless the kernel truly needs templating\n- `cuda_wrappers`: `(export_name, kernel_symbol)` — `export_name` is called from Python\n- `make_cpp_args(dtype, ...)` converts `torch.dtype` to C++ type alias:\n- `is_arch_support_pdl()` checks if the current architecture supports PDL, which is typically passed as a template argument to the kernel.\n- Keep the entry point thin (see Conventions). What Python must still check goes in the `@cache_once` module factory, not in the entry point: `cache_once` keys on its arguments, so a check there costs one evaluation per specialisation instead of one per call — that is where the supported-dtype guard lives. Tensor invariants belong in the C++ launcher; if something here never reaches a `.verify(...)`, close that gap on the C++ side rather than in Python\n\n| `torch.dtype` | C++ type |\n|--------------------|------------|\n| `torch.float16` | `fp16_t` |\n| `torch.bfloat16` | `bf16_t` |\n| `torch.float32` | `fp32_t` |\n\n---\n\n## Register the public entry point (required)\n\nAdd metadata in `python/sglang/kernels/ops/elementwise/__init__.py` so the\noperator appears in the registry without importing its implementation:\n\n```python\nfrom sglang.kernels.registry import register_kernel\nfrom sglang.kernels.spec import CapabilityRequirement, KernelBackend, KernelSpec\n\nregister_kernel(\n KernelSpec(\n op=\"elementwise.scale\",\n backend=KernelBackend.JIT,\n target=\"sglang.kernels.ops.elementwise.scale:scale\",\n capabilities=frozenset({CapabilityRequirement.CUDA}),\n )\n)\n```\n\nRecord the devices the implementation actually supports; the CUDA-only scale\nexample above does not imply that every JIT kernel is CUDA-only. Runtime and\ncorrectness tests import `scale` from `sglang.kernels.ops.elementwise.scale`.\nDo not import the implementation eagerly in the group initializer. When an\nexisting `BaseFusedOp` owns the operation, add the implementation there instead\nof registering a second conflicting op/backend pair.\n\n## Step 3 (optional): Tune JIT build flags\n\nIf your kernel uses some math functions like `expf` or `sinf`, consider enabling `--use_fast_math` for better performance (with a potential precision tradeoff):\n\n```python\nreturn load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n)\n```\n\nIf your kernel requires SM90+, raise a clear Python error before calling `load_jit`. Arch gating is one of the checks that has to live in Python — it decides whether to compile at all, so the C++ launcher never gets to run:\n\n```python\nif torch.cuda.get_device_capability()[0] < 9:\n raise RuntimeError(\"This kernel requires SM90 (Hopper) or later\")\n```\n\n---\n\n## Step 4: Write tests (required)\n\nJIT kernel correctness tests and benchmarks live under `test/registered/kernels/ops//` and `test/registered/kernels/benchmark//`, mirroring the wrapper's group under `python/sglang/kernels/ops/` (NOT inside the `sglang` package -- a `register_*_ci(...)` call anywhere under `python/sglang/` is rejected by the `check-no-registered-tests-in-package` pre-commit hook). Only their test-only helpers (e.g. `benchmark/marker.py`) stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path. **CI does not run `pytest` in those directories directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` and `bench_*.py` under `test/registered/`, collects `register_*_ci(...)` calls by **statically parsing each file's AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check.\n\n- **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`.\n- **Nightly kernel suite**: register with `stage=\"nightly\"` plus the `runner_config` of the machine it needs (e.g. `1-gpu-large`), giving the `nightly-test-1-gpu-large` suite. `.github/workflows/nightly-test-nvidia.yml` sets `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` for the whole nightly run, so the expanded parameter grids apply automatically (see `python/sglang/kernels/jit/utils/common.py` → `should_run_full_tests` / `get_ci_test_range`). There is no separate kernel-only nightly job: every nightly test on one machine type shares that machine's suite.\n\nRegistration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing):\n\n```python\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n# Optional B200/SM100 registration for tests that cover Blackwell-specific code paths\n# register_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"4-gpu-b200\")\n# Optional second registration: same file also runs nightly, same form,\n# stage is just \"nightly\" there (and no `nightly=True`)\n# register_cuda_ci(est_time=120, stage=\"nightly\", runner_config=\"1-gpu-large\")\n```\n\nCI generates the suite name as `{stage}-test-{runner_config}`, so `stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\"` becomes the `base-b-kernel-unit-test-1-gpu-large` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage=\"nightly\"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools.\n\nKeep `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.\n\nUse `register_cuda_ci(..., disabled=\"reason\")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only).\n\n**Run like CI** (from repo root):\n\n```bash\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large)\n# For B200/SM100-specific coverage:\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200)\n```\n\nFor fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`.\n\nCreate `test/registered/kernels/ops/elementwise/test_scale.py`:\n\n```python\nimport pytest\nimport torch\nfrom sglang.kernels.ops.elementwise.scale import scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [1, 127, 128, 1024, 4097]) # cover tail remainder\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0, 3.0])\ndef test_scale_correctness(dtype, size, factor):\n src = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = scale(src, factor)\n expected = src * factor\n\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\ndef test_scale_out_param(dtype):\n src = torch.randn(1024, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(src)\n result = scale(src, 2.0, out=out)\n assert result is out\n torch.testing.assert_close(out, src * 2.0, rtol=1e-2, atol=1e-2)\n\n\ndef test_scale_cpu_error():\n src = torch.randn(128, dtype=torch.float16) # CPU tensor\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n scale(src, 2.0)\n\n\ndef test_scale_unsupported_dtype():\n src = torch.randint(0, 10, (128,), dtype=torch.int32, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"dtype\"):\n scale(src, 2.0)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-v\", \"-s\"]))\n```\n\n---\n\n## Step 5: Add a benchmark (required)\n\nBenchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large`).\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nBenchmarks use the project's own `marker` framework (in `python/sglang/kernels/jit/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`):\n\n- **`@marker.benchmark(line_arg, line_vals, *, unit=\"us\")`** — the **innermost** decorator (bottom of the stack, directly above `def benchmark`). Declares the column axis: each value in `line_vals` becomes a result column, and `line_arg` is the parameter name passed into the benchmark function. `unit` is one of `\"us\" | \"ms\" | \"s\"`.\n- **`@marker.parametrize(names, vals, ci_vals=None)`** — stackable decorator that adds a row axis (pytest-style). Each `@parametrize` adds one (or more, correlated) parameter the benchmark is swept over (Cartesian product across all `parametrize` decorators). `names` may be a single name (`\"size\"`) or a comma-separated correlated tuple axis (`\"h,d\"`, with `vals` then a list of tuples like `[(1, 64), (2, 128)]`). Pass the optional third `ci_vals` for a smaller sweep that is auto-selected under `is_in_ci()` — this is the built-in CI-shrinking mechanism, so you usually don't need `get_benchmark_range` for swept axes.\n- **`marker.do_bench(fn, *, input_args=(), input_kwargs={}, ...)`** — runs `fn` under CUDA graph (default) or a naive loop, returns a `BenchResult`. Key knobs:\n - `memory_args`: defaults to `\"all\"` (footprint derived from all input args/kwargs). Pass an explicit tuple of tensors (e.g. `(k, v, indices)`) to count only the inputs the kernel actually touches.\n - `memory_output`: defaults to `\"out\"` — re-runs `fn` once to capture its **returned** tensor and counts it. For in-place kernels (which return `None`), pass the written tensors explicitly (e.g. `memory_output=(k, v)`); the re-run is then skipped. Set to `None` to count no output.\n - Together `memory_args` + `memory_output` give the GB/s column; with both defaults a function `out = f(src)` already reports `bytes(src) + bytes(out)`.\n - `graph_clone_args` / `graph_clone_kwargs`: which inputs to clone per CUDA-graph iteration to defeat L2 cache reuse. Defaults to `\"all\"`; if you narrow it, keep every read, in-place, and output tensor.\n - `use_cuda_graph=False` for kernels that can't be captured.\n - `metrics=(0.5, \"avg\")` controls reported quantiles (the first metric becomes the table latency column).\n - `disable_log_bandwidth` (defaults from `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1`) skips the bandwidth column entirely.\n- **`utils.create_random(*shape)` / `utils.create_empty(*shape)`** — shorthand for `torch.randn` / `torch.empty` with `DEFAULT_DTYPE` (`bfloat16`) and `DEFAULT_DEVICE` (`\"cuda\"`). Override via the `dtype=` / `device=` kwargs.\n- **`utils.get_benchmark_range(full_range, ci_range)`** — returns the smaller `ci_range` under CI (`is_in_ci()`), the `full_range` locally. Still available for the `benchmark(...)` column axis (which has no `ci_vals`); for `parametrize` row axes prefer the built-in `ci_vals` argument.\n\nCreate `test/registered/kernels/benchmark/elementwise/bench_scale.py`:\n\n```python\nimport torch\n\nfrom sglang.kernels.jit.benchmark import marker\nfrom sglang.kernels.jit.benchmark.utils import create_random\nfrom sglang.kernels.ops.elementwise.scale import scale as jit_scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=6, stage=\"base-b-kernel-benchmark\", runner_config=\"1-gpu-large\")\n\n\n@torch.compile()\ndef torch_impl_scale(src: torch.Tensor, factor: float) -> torch.Tensor:\n return src * factor\n\n\nFN_MAP = {\n \"jit\": jit_scale,\n \"torch\": torch_impl_scale,\n}\n\n\n# `parametrize(name, full_vals, ci_vals)`: the 3rd arg is the smaller sweep\n# auto-selected under CI; the full range runs locally.\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096, 65536]) # 1K .. 512K\n@marker.benchmark(\"impl\", [\"jit\", \"torch\"])\ndef benchmark(size: int, impl: str):\n src = create_random(size)\n factor = 2.0\n return marker.do_bench(\n FN_MAP[impl],\n input_args=(src, factor),\n # `src` is read -> clone it per iter to avoid L2 reuse; factor is a scalar.\n graph_clone_args=(0,),\n # Defaults already report bandwidth: memory_args=\"all\" counts src,\n # memory_output=\"out\" counts the returned tensor -> bytes(src)+bytes(out).\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n**Key points:**\n\n- The `line_arg` name passed to `benchmark` (`\"impl\"` here) must match a parameter on `benchmark(...)`; same for every `parametrize` name (`\"size\"`).\n- Stack `@parametrize` once per swept axis. The required `@marker.benchmark` is the **innermost** decorator (bottom of the stack, directly above the function) — `@parametrize` rows go above it.\n- Prefer `create_random` / `create_empty` from `utils.py` over open-coding `torch.randn(..., dtype=..., device=...)`.\n- The GB/s column appears by default (`memory_args=\"all\"` + `memory_output=\"out\"`). For memory-bound kernels it's the most informative number; scope `memory_args` / `memory_output` to the tensors actually touched if the defaults over- or under-count. For compute-bound kernels where bandwidth is misleading, set `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1` (or `disable_log_bandwidth=True`).\n- For in-place kernels (which return `None`), pass the written tensors via `memory_output=(...)` since the `\"out\"` default would capture nothing.\n- Tune `graph_clone_args` / `graph_clone_kwargs` to all the arguments that might be read by the kernel. We can only skip cloning for write-only args. For in-place modified args, we still need to clone them to get accurate timing (reusing the same buffer keeps it L2-hot and skews results).\n- Call `benchmark.run()` (no `print_data=` kwarg — the marker framework prints directly).\n\nRun locally:\n\n```bash\npython test/registered/kernels/benchmark/elementwise/bench_scale.py\n```\n\nRun the benchmark suite the way CI does:\n\n```bash\ncd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large\n```\n\n---\n\n## Troubleshooting\n\n- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time`, `stage`, and `runner_config`; starred args and non-literal values break AST collection\n- **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/kernels/jit/csrc/`; reduce template argument combinations\n- **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...`\n- **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. `graph_clone_args` defaults to `\"all\"`; if you narrow it, it must still cover every *read* tensor — reusing a single buffer keeps it L2-hot and skews results. Keep *write* tensors in it too: they are what sets the rotation count, and a shared output buffer stays L2-hot the same way.\n- **Missing GB/s column**: the column is on by default; check that `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH` is not `1` and `disable_log_bandwidth` is not `True`. For in-place kernels (return `None`) the `memory_output=\"out\"` default counts nothing — pass the written tensors via `memory_output=(...)`\n\n---\n\n## References\n\n- `docs/docs/developer_guide/development_jit_kernel_guide.mdx`\n- `test/run_suite.py` — suite names, discovery of `test/registered/`, execution entrypoint for CI\n- `python/sglang/test/ci/ci_register.py` — `register_cuda_ci` and AST registration rules\n- `python/sglang/kernels/jit/utils/compile/` — `load_jit` (`loader.py`), `make_cpp_args` (`cpp_args.py`)\n- `python/sglang/kernels/jit/utils/common.py` — `cache_once`, `should_run_full_tests`, `get_ci_test_range`\n- `python/sglang/kernels/jit/include/sgl_kernel/tensor.h` — `TensorMatcher`, `SymbolicSize/DType/Device`, `is_type`\n- `python/sglang/kernels/jit/include/sgl_kernel/ffi.h` — `ffi::empty`, `ffi::empty_like`, `ffi::from_blob`\n- `python/sglang/kernels/jit/include/sgl_kernel/utils.cuh` — type aliases, `LaunchKernel`, `SGL_DEVICE`\n- `python/sglang/kernels/jit/include/sgl_kernel/vec.cuh` — `AlignedVector`\n- `python/sglang/kerne",".agents/skills/debug-distributed-hang/SKILL.md":"---\nname: debug-distributed-hang\ndescription: Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations.\n---\n\n# Debugging Distributed Hangs in SGLang\n\n## Overview\n\nHangs in distributed inference happen when ranks diverge in state, causing collective operations (AllGather, AllReduce, Broadcast, Barrier) to deadlock. Common causes:\n\n- **Size mismatch**: ranks pass different tensor sizes to a collective\n- **Branch divergence**: one rank enters a collective, another skips it\n- **Cascading state drift**: a small non-determinism (e.g., floating-point) propagates into different batch structures\n- **Resource exhaustion**: one rank OOMs or crashes, others wait forever\n\n## Prerequisites\n\n- **py-spy**: `pip install py-spy` or system package. Requires root or `CAP_SYS_PTRACE` to attach to running processes.\n- **cuda-gdb**: Ships with the CUDA toolkit. Ensure it's on your `PATH`.\n\n## Step 1: Confirm and Locate the Hang\n\n### 1a. Watchdog / py-spy\n\nSGLang's watchdog automatically dumps py-spy traces on timeout. Look for:\n\n```\nScheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)\n```\n\nThe py-spy dump shows the stack trace of each thread. The hanging thread is typically blocked in a CUDA synchronize or NCCL collective:\n\n```\nThread (active): \"MainThread\"\n cuStreamSynchronize (libcuda.so)\n ...\n forward_extend (model_runner.py)\n```\n\nSGLang has two watchdog modes (see `python/sglang/srt/utils/watchdog.py`):\n- **Hard watchdog** (`soft=False`, default): dumps py-spy traces then sends `SIGQUIT` to kill the parent process.\n- **Soft watchdog** (`soft=True`): only logs the timeout without killing the process, giving you more time to manually attach debuggers or collect coredumps.\n\nIf the watchdog doesn't trigger, manually dump:\n\n```bash\npy-spy dump --pid \n```\n\n### 1b. NCCL Debug Logging\n\n```bash\nexport NCCL_DEBUG=INFO\nexport NCCL_DEBUG_SUBSYS=COLL\n```\n\nLook for the last collective logged before the hang. Mismatched sizes show up as one rank waiting and another never entering.\n\n### 1c. CUDA Coredump\n\nWhen a process hangs, you can trigger a GPU coredump on demand to see which kernel is stuck. Set these env vars before launching:\n\n```bash\nexport CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1\nexport CUDA_COREDUMP_PIPE=\"/tmp/cuda_pipe_%h_%p\"\nexport CUDA_COREDUMP_FILE=\"/tmp/cuda_coredump_%h_%p\"\nexport CUDA_COREDUMP_SHOW_PROGRESS=1\nexport CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'\n```\n\nWhile the process is hanging, find the pipe via `/proc//fd/` and write to it to trigger the dump:\n\n```bash\nls /proc//fd/ -la 2>/dev/null | grep cuda_pipe\ndd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe__\n```\n\nAlternatively, if you don't need to keep the process alive, `kill -SIGABRT ` also triggers a CUDA coredump (but terminates the process).\n\nThen open with `cuda-gdb --batch -ex \"target cudacore \"`. On load, it immediately shows which kernel is stuck. For example:\n\n```\nOpening GPU coredump: \n[Current focus set to CUDA kernel 0, grid 622721, cluster (4,0,0), block (16,0,0), thread (64,0,0), device 0, sm 0, warp 0, lane 0]\n#0 0x00007f8029b2b040 in ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)<<<(24,1,1),(512,1,1)>>> ()\n```\n\nThis told us the hang was in an NCCL AllGather — not a compute kernel. Combined with the py-spy stack pointing to `LogitsProcessor.forward` → `tensor_model_parallel_all_gather`, we knew it was an AllGather size mismatch between TP ranks.\n\n\n### 1d. Identify the Collective\n\nFrom the stack traces and logs, identify:\n- Which collective hangs (AllGather, AllReduce, Broadcast)\n- Which code path invokes it (e.g., `LogitsProcessor`, `tensor_model_parallel_all_gather`)\n- Whether it's a size mismatch or a missing participant\n\n## Step 2: Per-Rank Logging\n\nThe key technique: each rank writes its own log file so you can diff them.\n\n### Setup Pattern\n\n```python\nimport os\n\n_debug_files = {}\n\ndef get_debug_file(rank):\n key = f\"rank{rank}\"\n if key not in _debug_files:\n _debug_files[key] = open(f\"/tmp/debug_rank{rank}.log\", \"w\")\n return _debug_files[key]\n```\n\nGate logging behind an env var to avoid overhead in production. `SGLANG_DEBUG_HANG` is not a built-in SGLang env var — you need to add this check yourself in the code you're instrumenting:\n\n```python\nif os.environ.get(\"SGLANG_DEBUG_HANG\"):\n f = get_debug_file(rank)\n f.write(f\"EVENT_NAME key1={val1} key2={val2}\\n\")\n f.flush()\n```\n\n### What to Log\n\nLog structured events at key state-mutation points:\n\n```python\nf.write(f\"SCHED_BATCH step={step} num_reqs={n} extend_lens={lens}\\n\")\nf.write(f\"VERIFY predict_hash={hash} accept_len={alen}\\n\")\nf.write(f\"CACHE_INSERT rid={rid} num_tokens={n}\\n\")\n```\n\nUse consistent event names (uppercase prefix) for easy grep/diff.\n\n### Hash Large Tensors\n\nFor tensor values, compute a hash instead of dumping raw data:\n\n```python\nimport hashlib\nh = hashlib.md5(tensor.cpu().numpy().tobytes()).hexdigest()[:8]\nf.write(f\"LOGITS logits_hash={h}\\n\")\n```\n\nFor token ID lists, `str(list).encode()` works:\n\n```python\nh = hashlib.md5(str(tensor.tolist()).encode()).hexdigest()[:8]\n```\n\n### Avoid Implicit Synchronization\n\n`tensor.cpu()`, `tensor.tolist()`, and `tensor.numpy()` all trigger CUDA synchronization. This can:\n- Change timing and mask or move the hang\n- Deadlock if the log point is between two collectives that must run back-to-back\n\nPrefer logging values that are already on CPU (e.g., Python ints, list lengths, request IDs). When you must hash a GPU tensor, do it at a point where the GPU is already idle (e.g., between scheduler steps, not inside a model forward pass).\n\n## Step 3: Diff to Find the Diverge Point\n\n### Basic Diff\n\n```bash\n# Extract specific event type\ngrep \"^VERIFY\" /tmp/debug_rank0.log > /tmp/v_r0.txt\ngrep \"^VERIFY\" /tmp/debug_rank1.log > /tmp/v_r1.txt\ndiff /tmp/v_r0.txt /tmp/v_r1.txt | head -20\n```\n\n### Count Events\n\n```bash\ngrep -c \"^VERIFY\" /tmp/debug_rank*.log\n```\n\nIf counts differ, one rank executed more iterations — that's already a diverge signal.\n\n### Find First Diverge\n\nThe first diff line tells you the exact step where ranks diverge. All lines before it are identical — the root cause is at or before this step.\n\n## Step 4: Binary-Search the Root Cause\n\nOnce you find the diverging event, trace backwards:\n\n### 4a. Identify Inputs\n\nFor the diverging operation, list all its inputs. Add hash logging for each:\n\n```python\nf.write(\n f\"OP_INPUTS input_a_hash={h_a} input_b_hash={h_b} \"\n f\"input_c_hash={h_c} input_d_hash={h_d}\\n\"\n)\n```\n\n### 4b. Diff Inputs Across Ranks\n\nCompare the hashes. Some inputs will match, some won't. The non-matching input is where divergence entered.\n\n### 4c. Recurse\n\nFor the non-matching input, trace where it was produced and repeat: hash its inputs, diff across ranks, find the divergent one. Continue until you reach the root cause.\n\n## Step 5: Common Root Causes and Fixes\n\n### Floating-Point Non-Determinism\n\n**Symptom**: All \"logical\" inputs are identical (same logits after all-gather), but derived floating-point values (softmax, probabilities) differ across GPUs.\n\n**Example**: EAGLE speculative decoding — `F.softmax` → `top_k_renorm_prob` → `top_p_renorm_prob` produces slightly different `target_probs` on each GPU. The sampling kernel then picks different tokens. These flow into `output_ids` → radix cache → different prefix match depths → different `extend_seq_lens` → AllGather size mismatch → hang.\n\n### Random Number Divergence\n\n**Symptom**: Operations using `torch.rand` produce different values on each rank.\n\n**Fix**: Generate on rank 0 and broadcast, or use a shared seed.\n\n### Conditional Code Paths\n\n**Symptom**: A condition (e.g., memory check, queue length) evaluates differently on different ranks, causing one rank to enter a collective while another skips it.\n\n**Fix**: Synchronize the condition value before branching, or restructure to ensure all ranks take the same path.\n\n### Pipeline Parallel (PP) Send/Recv Mismatch\n\n**Symptom**: In PP setups, one stage issues a `send` that the next stage never `recv`s (or vice versa), causing both to block indefinitely. Unlike TP hangs (collective mismatches), PP hangs typically involve point-to-point operations.\n\n**Fix**: Ensure all stages agree on the number of microbatches and the sequence of send/recv calls for each microbatch.\n\n## Step 6: Verify the Fix\n\nRun the failing test multiple times to confirm the fix is stable. Intermittent hangs require many runs. A test that hung ~30% of the time needs at least 10 clean passes to be confident.\n\n## Quick Reference\n\n| Technique | When to Use |\n|-----------|-------------|\n| py-spy dump | First step — see where each rank is stuck |\n| `NCCL_DEBUG=INFO` | Identify which collective and sizes |\n| CUDA coredump + `cuda-gdb` | See which GPU kernel is blocked |\n| Per-rank log files | Compare rank states over time |\n| Hash of tensors | Efficiently compare large tensors across ranks |\n| `diff` on extracted events | Find the exact step of divergence |\n| `broadcast(result, src=0)` | Fix floating-point or sampling non-determinism |\n",".agents/skills/cookbook-migrate-model/SKILL.md":"---\nname: cookbook-migrate-model\ndescription: Migrate a legacy-template SGLang cookbook page (monolithic per-model generator under docs/src/snippets/autoregressive/) onto the config-driven template (shared _deployment.jsx / _playground.jsx engines + per-model config). Use when asked to migrate, convert, or port an existing cookbook page — NOT for brand-new models (use cookbook-add-model for those). Run with /cookbook-migrate-model .\n---\n\n# Cookbook Migrate Model\n\nConvert one legacy cookbook page to the config-driven format, faithfully. The\nlegacy page — its generator widget and its measured benchmark blocks — is the\n**single source of truth**. You are transcribing it into the new data model,\nnot improving it.\n\nReuses the `cookbook-add-model` skill's assets (read them on demand):\n- `../cookbook-add-model/templates/config.jsx.tmpl`, `page.mdx.tmpl`, `benchmarks.jsx.tmpl`\n- `../cookbook-add-model/references/authoring-reference.md` (config/cells/playground contract)\n- `../cookbook-add-model/references/mintlify-authoring.md` (MDX rules)\n\nMigration-specific references in this skill:\n- [references/dimension-mapping.md](references/dimension-mapping.md) — legacy-control → new-dimension mapping rules, command rewrite table, per-family strategy sets, and the Qwen3.5 pilot as a worked example (PR #27848).\n\nThe round's per-model inventory (scope, batch order, quirks, measured-data\nsurvey) is tracked by the migration maintainer outside the repo — expect it in\nyour dispatch prompt, or ask for it.\n\n## Hard rules (non-negotiable)\n\n1. **Never modernize.** Env vars, flags, TP values, docker tags, version strings\n are copied verbatim from the legacy page — even when today's defaults differ\n (e.g. a flag whose behavior has since become the default; keep it anyway). The recipe\n that was verified is the recipe as written. Allowed normalizations are ONLY\n the five alias rewrites in dimension-mapping.md §2 (`launch_server`→`sglang\n serve`, `--model`→`--model-path`, `--tp-size`→`--tp`, abbreviated\n `--speculative-algo`→`--speculative-algorithm`,\n `--expert-parallel-size`→`--ep`). **Accuracy-degrading flags**\n (`--kv-cache-dtype fp8_e4m3`, W4A4-style runtime quant) follow a\n deterministic rule — enforced in migration, no asking: offered as a\n legacy **selectable option** → never select it (cells mirror the\n accuracy-safe side), and the option **survives as a Playground axis** —\n an existing one where it fits, else add one via a separate prior engine\n PR (rule 4); a legacy choice never degrades to a tips mention. Baked\n into the recipe's **default/unconditional command** → keep it verbatim\n (the recipe was measured with it, and fp8 KV halves KV memory —\n stripping could OOM it). See dimension-mapping.md §2 caveats.\n2. **Never invent versions or numbers.** Benchmark numbers only from the\n legacy page's measured blocks, and a result migrates ONLY when its\n `sglang_version` is a **reproducible anchor** — the bar is reproducibility,\n not \"must be a release\":\n - ✅ release tag/version (`v0.5.9` / `0.5.9`), commit hash, OR — for\n **Day-0 support** (the enabling PR isn't merged and no release is cut\n yet) — a specific **PR (`PR #27944`) or commit** you can `gh pr checkout`\n / `git checkout `. Commit is most precise; a PR pin is fine for day-0.\n - ❌ a moving ref — `\"main branch\"`, `\"main (2026-06-11)\"`, open-ended\n `\"0.5.8+\"` — is NOT reproducible: **drop the WHOLE result (speed AND\n accuracy)**, not just speed. Keep `benchmarkCommands` so ⚡Reproduce still\n guides re-measurement against a pinned build.\n **Never inherit cross-model numbers** — measurements the legacy page\n attributes to a *different model* (e.g. a K2.6 page carrying K2.5-measured\n speed) are dropped regardless of version. When kept, `sglang_version` is the\n legacy page's string verbatim. Docker tags only the ones the legacy page\n pinned (unmapped hw falls back to `:dev`).\n3. **Verified policy (strictest tier).** `verified: true` ONLY when (a) the\n legacy page has concrete measured data for that exact 5-dim combo AND\n (b) the cell's flags equal the deployment command used for that measurement\n — modulo `{{HOST_IP}}`/`{{PORT}}`, the five alias rewrites, and **parser\n flags**: `--reasoning-parser`/`--tool-call-parser` are stripped from every\n cell (Playground-only feature; when the measured run had them on, say so in\n the benchmarks file header). When the measured command diverges from the\n generator default, **the verified cell follows the measured command**; the\n generator default stays as the sibling strategy/cell or a tips note. Everything else is unverified (yellow) —\n including combos that look memory-infeasible; keep them verbatim and list\n them in the PR body for the re-verification track.\n4. **Engines are read-only.** `_deployment.jsx` / `_playground.jsx` must not\n change in a migration PR. Model-specific features are config DATA consumed\n by generic axis handlers (MegaMoE precedent), so they need NO engine change.\n A **titled single-select that strips a flag family** — KV Cache DType\n (`--kv-cache-dtype`), mamba (`--mamba-radix-cache-strategy`), … — is already\n covered by the merged generic **`flagSelects`** axis: declare it in the\n config (a list of `{ id, title, stripPrefixes, options }`; see the Qwen3.5\n mamba example), **no engine PR**. Only a genuinely new control *shape* that\n `flagSelects` can't express would need a one-time generic primitive (never a\n model-named handler) on a separate prior PR (engine-axis.md).\n5. **`github.cookbookModel` must be set** (`/`, e.g.\n `qwen/qwen3.5`) and the block never pruned — without it Submit ↗ mislabels\n as deepseek-v4. The issue template itself needs NO edits (free-form input).\n6. **Strategy tiers are signal-driven.** A cell goes under `low-latency` /\n `high-throughput` ONLY on a signal present in the legacy source (an\n explicit performance toggle, a named recipe, or prose stating the\n operating point); **no signal → `balanced`**. Never derive a slant from\n your own hardware intuition — re-tiering on measured evidence is the\n hardware owner's follow-up PR, not part of a migration\n (dimension-mapping.md §4).\n\n## Workflow (one model = one PR)\n\n### 1. Inventory the legacy assets\n- Read the legacy generator (`docs/src/snippets/autoregressive/-deployment.jsx`)\n end-to-end: every option dimension (radio vs checkbox vs dynamic), every\n gate/SUPPORT matrix, the full emitted command per reachable combo (env\n prefixes, `# Error` pseudo-commands included).\n- Read the legacy MDX: §2 install docker tags → `dockerImages` (pinned, not\n upgraded); §3.2 tips → new §2; §4 invocation examples → new §3 (keep real\n Output Examples verbatim); §5 benchmark blocks → transcribe each measured\n block: deploy command used, bench command (dataset/isl/osl/num-prompts/\n concurrency), P50 (median) TTFT/TPOT, output tok/s, hardware, version string.\n- Inbound-anchor sweep: `grep -rn \"\" docs/ --include='*.mdx'` —\n find links/`#fragments` into this page (`mint broken-links` does NOT check\n fragments). Fix referrers or add `` shims in the same PR.\n- Check the maintainer-provided inventory notes for this model's known quirks —\n but treat them (and the §4 family table) as a **survey snapshot**: re-verify\n every dimension against the live legacy files before mapping. Pages keep\n receiving updates (precedent: Kimi-K2.6's live generator has a speculative\n toggle the 2026-06-10 survey notes lack).\n\n### 2. Design the 5-dim mapping\nApply [references/dimension-mapping.md](references/dimension-mapping.md). Key\ndecision — which legacy toggle becomes the `strategies` dimension: a toggle\nthat **changes other parts of the command** (TP, mem) must (the Playground\ncan't do coupled changes), and so does a toggle the legacy page itself labels\nwith operating-point words — e.g. a `dpattention` radio whose options are\nsubtitled \"Low Latency\" / \"High Throughput\" (GLM-5.1 / Kimi-K2.6 pattern) —\neven when its flags are uncoupled (`--attn-dp-size N` is a pure\nflag add). Any other toggle that only adds/removes its own flags becomes a\nPlayground axis with the flags baked into cells when the legacy default was\nON — EXCEPT parsers:\n`--reasoning-parser`/`--tool-call-parser` are NEVER baked into cells, they are\nPlayground-only (DSv4 convention). **Every legacy control survives as an\ninteractive control** — a dimension or a Playground axis, never a tips-only\nmention — and a model-specific control is **config data, not engine code**\n(MegaMoE W4A4 is all DSv4 config on the existing `moe` axis). It's pure\nconfig whenever it fits an existing axis's data schema. A **titled\nsingle-select that strips a flag family** (Nemotron3's \"KV Cache DType\",\nmamba `--mamba-radix-cache-strategy`, …) fits the merged generic **`flagSelects`**\naxis — so it too is config-only (declare a `flagSelects` list). Only a control\nwhose *shape* `flagSelects` still can't express would need a ONE-TIME generic\nprimitive (never a model-named handler) on a separate PRIOR engine PR, keeping\nthe migration PR data-only (hard rule 4, engine-axis.md). The strategy count follows the page's\noperating points: **one recipe → a single `balanced`; two → `low-latency` +\n`high-throughput`; three → the full trio (the ideal)**. The tiers apply per\n(hw × variant × quant) combination — a single-recipe combination on a\nmulti-strategy page parks under its semantically honest tier (no\nlatency/throughput slant → `balanced`; the page's list is the union). When the\nlegacy toggle is MTP / speculative decoding, the direction is a deterministic\ndefault — apply without asking: **MTP on → `low-latency`, MTP off →\n`high-throughput`** (reversed only with maintainer confirmation). Tier\nplacement is signal-driven (hard rule 6). Never invent a recipe just to fill\nstrategy chips (see dimension-mapping.md §4). Record the outcome as a\n**strategy mapping table** for the PR body — one row per group of\ncombinations sharing the same legacy signal (e.g. \"all GPU combos: MTP\ntoggle → low-latency / high-throughput\"; \"xeon: (none) → balanced\"), with a\none-line rationale each; don't enumerate 60 identical rows. The table is what\nhardware owners sign off on at review.\n\n### 3. Generate the config (codegen, then audit)\n- For >~30 cells, port the legacy `generateCommand()` into a throwaway Node\n script that enumerates combos and emits the `cells:[...]` literal (output\n must stay a pure literal — Mintlify forbids runtime spreads/calls). Apply the\n verified-cell override in the script. See the pilot scripts embedded in the\n worked example of dimension-mapping.md §5.\n- **Independent equivalence audit (required):** extract the ORIGINAL generator\n from git (`git show main:` — NOT `HEAD:`, the migration branch deletes\n the file, see dimension-mapping.md §5 item 7), stub React hooks, run it for\n every combo, and diff token-by-token against the new cells. Expected deltas\n only: the appended `--host {{HOST_IP}}`/`--port {{PORT}}`, the\n engine-injected multi-node trio, the §2 alias rewrites (the entrypoint\n rewrite doesn't appear in cell tokens — cells hold flags only; the audit\n script normalizes it on the legacy side), and the intentional verified-cell\n override. Paste the PASS count + the audit script in the PR body\n (collapsed `
`).\n- Hand-author the non-cells fields per authoring-reference.md. Structural\n self-checks: every cell resolves a `modelNames` key; no `--nnodes/--node-rank/\n --dist-init-addr/--host/--port` literals; every `{{KEY}}` declared; every\n `supportedHardware` id has ≥1 cell.\n\n### 4. Benchmarks file\nOne entry per measured block only (cells without entries already render\n\"pending\" — bare `{match}` stubs are unnecessary). `tokens_per_sec_per_gpu` =\ntotal (in+out) tok/s/GPU = `output tok/s ÷ (tp × nnodes) ×\n(isl+osl)/osl` — stored directly (the card shows it as-is). TTFT/TPOT\ntake the P50 (median) rows; set `config.latencyPercentile` (default `\"P50\"`; use\n`\"Mean\"` only for legacy Mean-recorded data — temporary, being migrated to P50; an\nentry-level `latencyPercentile` overrides the page value per cell).\nPut the workload's\n`num_prompts` into `workload`. **`config.accuracyLabels` is required whenever\nthe benchmarks carry accuracy data** — the engine ships no default eval set\n(#27842), so missing labels means the accuracy rows silently don't render;\nextra context (sample counts, suites that don't fit) goes in the entry's\n`notes`. Zero-measured-data pages: skip the file and the `benchmarks` prop\nentirely, but keep `benchmarkCommands` so ⚡Reproduce still guides users.\n\n### 5. Rewrite the MDX\nFrom `page.mdx.tmpl`: keep the original `title` (nav identity), write a fresh\nSEO `description` (top-level — delete any legacy `metatags.description`), **no\n`tag: NEW`** (a migration is not a launch), **no `mode:`**. Install accordion\ncarries the legacy install content + pinned images. Keep the template's\nDSv4-style strategy bullets — serving semantics first (single-user chat /\ntypical multi-user / batch throughput), trimmed to the strategies the page\nships, plus at most a one-line note on what each strategy changes on this\nmodel; do NOT rewrite them as toggle-/migration-centric explanations.\nLegacy §5 benchmark prose is deleted (numbers → benchmark card, commands →\n⚡Reproduce); legacy prose deploy commands are deleted (doc↔config parity —\nfold their unique flags into §2 tips). Invocation examples + real outputs carry\nover verbatim, but **wrapped in Accordions** — §3 commands and outputs are\ncollapsible (required, DeepSeek-V4 pattern): code in an\n``, output in a following\n``; legacy pages kept them inline.\n\n### 6. Delete the legacy generator\nRemove `docs/src/snippets/autoregressive/-deployment.jsx` and its\nimport. `grep -rn \"-deployment\" docs/` must return nothing (config\nprovenance comments must not name the deleted path). Site wiring needs **no\nchanges**: docs.json path/title unchanged, vendor card + logo already exist.\n\n### 7. Validate\n- `grep -rn '__[A-Z_]*__'` on the new files (no template tokens).\n- `cd docs && mint validate && mint broken-links` (pre-existing breaks on\n main are not yours — say so in the PR).\n- `mint dev` browser smoke: initial selection = the verified cell (first in\n `cells[]`) with green badge; multi-node cells show the injected trio +\n header; AMD cells show env prefixes; Docker mode wraps with the pinned image\n and passes cell env as `--env`; condition-hidden combos grey out; benchmark\n card values; NO parser flags in any Deploy command; Playground parser\n toggles ADD the parser flags (green additions) while spec toggles strike\n the baked spec flags (red); Submit ↗ prefills this model. **Probe pitfall:** drive at most\n ONE programmatic click per evaluation and wait for React to settle —\n multiple clicks in one synchronous script batch and read stale DOM.\n- Token-level audit from step 3 passes.\n\n### 8. PR + review\nOne PR per model. PR body: migration framing, verified policy applied, the\nstrategy mapping table (step 2), the audit PASS count + script, any\ninherited-infeasible combos flagged for re-verification. Then run `/cookbook-review-pr ` and fix findings.\nFYI: docs previews only build for in-repo (`sgl-project/sglang`) branches —\na fork-headed PR is perfectly fine but renders no preview; a maintainer can\nre-push the branch in-repo if a preview is wanted for review.\n\n### 9. Keep this skill current\nAny new convention, engine behavior, or pitfall you discover while migrating\n(naming decisions, audit-script gotchas, review-rule conflicts, …) MUST be fed\nback into this skill — same PR if it's skill-file-only, or an immediate\nfollow-up commit on the skill's branch/PR. The next agent runs on what's\nwritten here, not on your session's context.\n",".agents/skills/cookbook-review-pr/SKILL.md":"---\nname: cookbook-review-pr\ndescription: Review a pull request against the SGLang Cookbook (docs/, Mintlify) contribution checklist — the config-driven format (per-model config + benchmarks JSX consumed by the shared _deployment.jsx / _playground.jsx engines). Run with /cookbook-review-pr .\n---\n\n# Cookbook Review PR\n\nFetch the diff, run the checklist, report what you find. The cookbook is **config-driven**:\nshared engines (`_deployment.jsx`, `_playground.jsx`) with NO model-specific code; each\nmodel is a data `config` (+ optional `benchmarks`) under `src/snippets/configs//`\nplus an MDX page. This checklist targets that layout. Field-schema detail lives in\n`.claude/skills/cookbook-add-model/references/authoring-reference.md` — defer to it rather\nthan restating.\n\n## Usage\n\n```\n/cookbook-review-pr \n```\n\n## Steps\n\n1. `gh pr view --repo sgl-project/sglang --json title,body,files,author,baseRefName,headRefName,commits,reviews`\n2. `gh pr diff --repo sgl-project/sglang`\n3. `gh pr list --repo sgl-project/sglang --state open --search \"\"` (duplicate check)\n4. Run every checklist item against the diff.\n5. Output per-file verdicts + overall recommendation.\n\n## Checklist\n\n### 1. File hygiene\n- A cookbook PR should only touch: `docs/src/snippets/configs//*.jsx`\n (config + benchmarks), `docs/cookbook/**/*.mdx`, `docs/docs.json`,\n `docs/cookbook//intro.mdx` (vendor card), `docs/cards/logos/.png`\n (new vendor only). Flag stray files (`settings.local.json`, lockfiles, IDE configs).\n- Pages must be `.mdx`, not `.md`. Files end with a trailing newline. Check commit history\n for unrelated commits accidentally included.\n- **Engines untouched**: `_deployment.jsx` / `_playground.jsx` should NOT change in a\n model-add PR (adding a model is data-only). Engine edits = a separate axis/feature PR\n (see `cookbook-add-model/references/engine-axis.md`); review them against that checklist.\n\n### 2. Config quality (the per-model `config`)\n- Single `export const config = { ... }` literal — **no** function calls, spreads,\n fragment refs, or IIFE (Mintlify re-evals at hydration → `ReferenceError`).\n- No `!(x in y)` anywhere (Mintlify AST walker crashes) — use `obj.key === undefined`.\n- `supportedHardware` ⊆ `HARDWARE_CATALOG` (in `_deployment.jsx`) ∪ `config.hardware`. A\n model-specific GPU the shared catalog lacks must be declared in `config.hardware`\n (`{id,label,vram,vendor}`), **not** added to the engine catalog — unless the platform is\n genuinely shared, which is an engine change the PR body has to argue (and the\n `cookbook-add-model` hardware table moves with it). A new catalog entry is inert for\n other pages: each filters the catalog by its own `supportedHardware`.\n- `placeholders` declares every `{{KEY}}` used in `curl` or any cell.\n- `modelNames` covers every cell (by `hw|variant|quant` triple or `variant|quant` pair).\n- `strategies` count matches the page's operating points — 1 recipe → a single `balanced`;\n 2 → `low-latency` + `high-throughput`; 3 → the full trio. Tiers apply per\n (hw × variant × quant) combination: a single-recipe combination must park under its\n semantically honest tier (clear slant → that tier, e.g. a workstation card under\n `low-latency`; no slant → `balanced`, e.g. a CPU platform) — **flag a no-slant recipe\n parked under low-latency/high-throughput**. Mixed unions\n like [low-latency, balanced, high-throughput] with per-selection greying are fine. Also\n flag model-specific ids (e.g. `mtp`), and flag an INVERTED speculative mapping — the\n deterministic default is MTP/spec-decoding ON → `low-latency`, OFF → `high-throughput`\n (at saturation the draft+verify overhead outweighs the speedup); the reverse needs an\n explicit maintainer-confirmed justification in the PR. The MDX strategy bullets describe serving semantics\n in the DSv4 style (single-user chat / typical multi-user / batch jobs), not internal\n toggles.\n- `dockerImages` covers the hw ids that have cells (else users hit the `:dev` fallback); a\n `hw|quant` key (resolved before the plain `hw`) is valid when one quant on a shared GPU needs\n a different image (e.g. an FP4 dev build) — don't flag those.\n- `multiNodeHints` present ONLY for hw whose fabric needs manual NIC env (e.g. `gb200`\n NVL72) — NOT every `multi-N` hw (standard-IB DeepEP / Marlin multi-node don't need it).\n Hints render above BOTH run modes, so flag one that reads as docker-only (\"add these\n docker run flags\") — `docker run` flags belong in the hardware entry's\n `multiNodeDockerFlags`, which the engine emits into the Docker command itself.\n- `github.cookbookModel` is set to the model's HF id (`/`). The issue\n template's `model` field is a free-form input prefilled from this value; if the config\n omits the `github` block, the engine falls back to `deepseek-ai/deepseek-v4` and the\n page's submissions get mislabeled.\n- `playgroundFeatures` is opt-OUT: the **general axes ship on every cookbook by default**\n (`attention` TP/CP/DP-Attn, `moe` backend+EP for MoE models, `parsers`, `speculative`,\n `pdDisagg`, `hicache`) — flag a missing general axis unless the model genuinely cannot\n use it. Model-specific axes only where applicable (MegaMoE backend + `megamoeQuant`\n only on Blackwell MoE, gated by `requiresHw`; `hisparse` only DSA-style). Knobs that are\n meaningless for a subset of variants/hw are `disable`d with a reason, not silently live\n (e.g. MoE knobs greyed on dense variants). No empty/stub axes.\n- **No leftover `__TOKEN__`** — the config was stamped from the template and every\n placeholder is filled (`grep -rn '__[A-Z_]*__'` on the new config/benchmarks/MDX returns\n nothing).\n- **All-hardware considered**: every `supportedHardware` id (from the catalog or `config.hardware`) has ≥1 cell OR is a deliberate\n greyed \"coming soon\"; AMD was pruned or kept on purpose (not a leftover template family).\n\n### 3. Cells / 5-dim matrix\n- Every cell `match` has EXACTLY the 5 keys (`hw`, `variant`, `quant`, `strategy`, `nodes`).\n- `env` / `flags` are flat literals (only `{{PLACEHOLDER}}` subst) — no shared\n `commonFlags` reference (Mintlify won't inline it).\n- NO `--nnodes` / `--node-rank` / `--dist-init-addr` literals in multi-node cells\n (the renderer injects them from `match.nodes`).\n- NO literal `--host` / `--port` — use `{{HOST_IP}}` / `{{PORT}}`.\n- NO `--reasoning-parser` / `--tool-call-parser` in any cell — parsers are a\n Playground-only feature added on top of the base command (DSv4 convention);\n flag any cell that bakes them in.\n- Accuracy-degrading flags in cells — runtime quant below the checkpoint\n (e.g. MegaMoE W4A4 — DSv4 gates it behind the Playground's `megamoeQuant`)\n and lossy `--kv-cache-dtype` (e.g. `fp8_e4m3` over a higher-precision-KV\n checkpoint): **flag for explicit maintainer confirmation**. Output quality\n should be exactly what the quant chip declares, so absent a recorded\n sign-off in the PR (e.g. carried verbatim from a measured legacy recipe's\n default command), request the flag move to Playground/tips.\n- Flag order: `--model-path` first (an optional `--trust-remote-code` may precede it —\n the DSv4 cells do), then parallelism, then MoE, then tuning, `--host`/`--port`\n last (the playground's insert anchors assume this).\n- TP/memory sanity: `model_weight_GB / (tp × gpu_mem)` fits with ~20–30% headroom\n (BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5; MoE uses **total** weight, not active params).\n\n### 4. Benchmarks\n- Each `benchmarks[]` entry's `match` tuple corresponds to a real cell.\n- `accuracyLabels` is present whenever the benchmarks carry accuracy data — the engine\n ships NO default eval set; without it the accuracy rows silently don't render.\n `defaultAccuracy` / per-cell `accuracy` / `benchmarkCommands.accuracy` keys all\n ∈ `config.accuracyLabels`.\n- A benchmark's quantization must match a variant actually listed — `(BF16)` on a model\n that only released FP8/FP4 is a factual bug.\n- `benchmarkCommands.speed` is `python3 -m sglang.bench_serving` (the workload), separate\n from the `sglang serve` deploy command, and should carry `--flush-cache`: bench_serving's\n `random` prompts are deterministic, so a warm rerun hits the radix cache and inflates\n throughput — speed numbers are measured cache-cold.\n- `sglang_version` is a real build the author ran (a release, or `dev`/nightly) — not a\n guessed/placeholder value (no leftover `0.0.0`).\n- **Latency percentile**: `config.latencyPercentile` (default `\"P50\"`, or `\"Mean\"`) matches the\n percentile the TTFT/TPOT values actually are — the card renders `TTFT ()`. A benchmarks\n entry may carry its own `latencyPercentile` to override the page value per cell\n (entry → config → `\"P50\"`): on a P50 page, kept legacy Mean cells must set it — a\n `sglang_version` tag alone doesn't convey the percentile. (`\"Mean\"` is temporary — legacy\n data is being re-measured to P50.)\n- **Throughput convention**: `tokens_per_sec_per_gpu` is stored as **total (in+out)/GPU**\n = `output tok/s/GPU × (isl+osl)/osl`, shown by the card as-is. Flag output-only values.\n- **Consistent accuracy harness across entries**: every value under one `accuracyLabels`\n column must be produced by the SAME harness — flag a page that, say, measures one\n platform's GSM8K with a historical raw completion harness and another's with\n `sgl-eval run gsm8k --num-examples 1319` and shows both as one \"GSM8K %\"\n (the scores aren't comparable). Either standardize on one harness (matching\n `benchmarkCommands.accuracy`) or require an explicit per-entry note. Common when folding\n a second contributor's measurements (e.g. an AMD/ROCm PR) into the page.\n\n### 5. Doc ↔ config parity (the #1 finding)\n- Any `sglang serve` command shown in MDX prose (config tips, benchmark section) must\n equal what the engine emits from the corresponding cell — same flags, same order. Drift\n here is the most common review miss.\n\n### 5b. ComfyUI section (diffusion pages)\n- A diffusion page ends with `## . Run in ComfyUI` rendering ``. A\n reader must not have to guess whether the model is reachable from ComfyUI.\n- The `model` prop is a key that exists in `docs/src/snippets/diffusion/comfyui-support.jsx`.\n A model-specific key is only correct when the plugin really treats it specially — an entry\n in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a dedicated\n node. Otherwise the generic `image` / `video` key is the honest one; a model-specific key\n without matching plugin support makes the page claim support that does not exist.\n- Prose describing ComfyUI support inline instead of using the component is a finding: the\n facts drift from the plugin.\n\n### 6. Commands / port\n- Launch uses `sglang serve` — flag any `python -m sglang.launch_server` /\n `python3 -m sglang.launch_server` (deprecated). The engine already emits `sglang serve`;\n guard against prose/cells reintroducing the old launcher.\n- Port `30000` everywhere (launch, curl, client `base_url`, bench) — flag `8000`.\n Launch port must match client/curl port on the same page.\n\n### 7. Frontmatter\n- Every new MDX page has `title:` and a **top-level** `description:` (a real one-line value\n prop, not copied from another vendor) — NOT `metatags.description` (non-canonical; the\n top-level field is what renders as the subtitle and SEO meta — see mintlify-authoring).\n- **No `mode: wide` on a model page** — it hides the right-hand \"On this page\" ToC that every\n other model page has. Leave `mode` unset (the Deploy/Playground panels self-cap at 900px, so\n the default column holds them fine). `mode: wide` belongs only on category `intro.mdx` grids.\n- `tag: NEW` only for genuine new launches; when one is added, stale `tag: NEW` on older\n pages should be dropped in the same PR (`grep -RlE \"^tag: NEW\" docs/cookbook/`).\n- MDX imports BOTH `Deployment` and `Playground` from `/src/snippets/...` (absolute).\n- Deploy heading slugs to `deployment` (or `deploy`), Playground to `playground` — so\n \"↑ Switch base\" and \"Open the Playground →\" scroll. No numbered headings for these two.\n\n### 8. Navigation & homepage\n- New page → `docs/docs.json` updated: under the right vendor group inside\n `navigation` → Cookbook → Autoregressive Models, root-relative, **no `.mdx`**:\n `cookbook///`.\n- Homepage `` in `docs/cookbook//intro.mdx` points to the vendor's\n flagship; new vendors get a new `` + a logo at `docs/cards/logos/.png` —\n **940×525 RGBA transparent, icon-only (no wordmark)**, lowercase filename, tracked via\n `git add -f` (`*.png` is gitignored repo-wide). Card order matches the `docs.json` nav order.\n- Don't change `docs/cookbook/intro.mdx` for individual model adds (top-level only).\n\n### 9. Links & factual\n- HuggingFace URLs resolve to a real model. License section matches the actual HF license\n (don't copy from another model). Docker images from `lmsysorg/sglang`; no `sgl-project-dev`.\n The image **tag** is a real build (a release the author ran, or `:dev`/nightly) — not a\n guessed version.\n- Internal links root-relative, no extension (`/cookbook/.../`); flag `.md`/`.mdx`\n or `../`-relative links. `docs.sglang.io` is canonical.\n- No Google-Drive image links (don't render). Shell placeholders are `export VAR=`,\n not `${VAR}` (a bash no-op).\n- **Parser ids must exist in the code registries** on the PR's target branch: every\n `--reasoning-parser X` / `--tool-call-parser Y` named in prose or in\n `playgroundFeatures.parsers` flags is a registered key in\n `python/sglang/srt/parser/reasoning_parser.py` (DetectorMap) /\n `python/sglang/srt/function_call/function_call_parser.py` (ToolCallParserEnum) —\n prose naming a near-miss id (e.g. the reasoning id where the tool id differs) is a\n factual bug. `--…-parser auto` is acceptable ONLY if the template-detection rules\n (`python/sglang/srt/parser/template_detection.py`) actually resolve THIS model's\n chat template to the right parser — no rule match means auto silently disables the\n parser; when in doubt require explicit ids (the DSv4 page pins explicit ids).\n\n### 9b. MDX authoring (Mintlify) — detail in `cookbook-add-model/references/mintlify-authoring.md`\n- **Forbidden syntax**: no Docusaurus admonitions (`:::`), `@site`/`@theme`, GitHub alert\n blocks (`> [!NOTE]`), markdown **pipe tables** (use JSX `
`), inline `
`,\n or unknown components. ``/`` only on category `intro.mdx`, not model pages.\n- Code fences are **labeled** (e.g. `python Example` / `bash Command` / `text Output` after\n the opening fence); a fenced block nested inside another uses four backticks outside.\n- §3 commands and outputs are **collapsible** (DeepSeek-V4 pattern): every runnable\n example wrapped in an ``, its real output in a following\n `` (`Pending update...` only with user\n acknowledgement). Flag bare/inline example blocks and `**Output Example:**` headings.\n- Reasoning-parser example matches the parser's **output shape**: separate-field\n (`reasoning_content` + `content`) vs inline `` tags parsed out of `content`.\n- No hardcoded sampling params (`temperature` / `top_p`) in sample code (SGLang uses\n `generation_config.json` defaults); listing them in §1 informationally is fine.\n\n### 10. Quantization rules\n- **NVFP4** checkpoints are Blackwell-only (B200/B300/GB300) — never AMD. An AMD FP4 cell\n is legitimate ONLY when the vendor published an **MXFP4** checkpoint for it (e.g.\n `amd/Qwen3.5-397B-A17B-MXFP4` on MI355X) — verify the HF repo resolves; otherwise the\n AMD FP4 chip must be absent/`disabled`.\n- BF16 / FP8 work on NVIDIA and AMD. `--kv-cache-dtype fp8_e4m3` in a cell is an\n accuracy-degrading flag — see §3 (needs explicit maintainer sign-off; default\n home is Playground/tips).\n\n### 11. Scope\n- Changes match the PR title. Flag global changes hiding behind a platform-specific title\n (e.g. an \"H200 FP8\" PR that adds a flag to ALL cells). Unmentioned side-fixes belong in\n the PR body.\n\n### 12. Duplicate PRs\n- Another open PR for the same model? Flag it; compare completeness; note merge-conflict\n risk on `docs.json` + the vendor card; flag a superseded older PR by the same author.\n\n### 13. Build / validate\n```bash\ncd docs\nmint validate\nmint broken-links\n```\nOptional: `mint dev` for a visual smoke test.\n\n### 14. Reviewer feedback\n- `gh api repos/sgl-project/sglang/pulls//comments` — have prior reviewer requests been\n addressed? Unresolved requested-changes should be flagged.\n\n### 15. Grammar & spelling\n- Check added/changed prose for typos and grammar (e.g. \"recommend\" vs \"recommended\").\n Flag each with the exact wrong text + correction.\n\n## Output\n\nPer file:\n- ✅ PASS\n- ⚠️ ISSUE: \\\n- 🔴 BLOCK: \\\n\nOverall: **APPROVE** / **REQUEST CHANGES** / **BLOCKED**\n",".agents/skills/debug-cuda-crash/SKILL.md":"---\nname: debug-cuda-crash\ndescription: Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging\n---\n\n# Tutorial: Debugging CUDA Crashes with Kernel API Logging\n\nThis tutorial shows you how to debug CUDA crashes and errors in SGLang using the `@debug_kernel_api` logging decorator.\n\n## Goal\n\nWhen your code crashes with CUDA errors such as illegal memory access, device-side assert, out-of-bounds, or NaN/Inf, use kernel API logging to:\n- Capture input tensors BEFORE the crash occurs\n- Understand what data caused the problem\n- Track tensor shapes, dtypes, and values through the call boundary that triggered the crash\n- Detect numerical issues such as NaN, Inf, or obviously wrong shapes\n\n## Why Use Kernel API Logging?\n\n**Problem**: CUDA errors often crash the program before normal debugging output is flushed.\n\n**Solution**: SGLang's `@debug_kernel_api` decorator logs inputs before execution, so you can still see what caused the crash even after the program aborts.\n\n## What Is Covered?\n\nThe current logging coverage focuses on the highest-value kernel boundaries in SGLang:\n- Custom ops registered through `register_custom_op(...)`\n- External custom ops registered through `register_custom_op_from_extern(...)`\n- LLM attention, linear, quantization, and multi-platform wrapper entry points\n- Diffusion attention impl, linear, rotary, and custom-op wrapper entry points\n- Selected direct `torch.ops.sglang.*` hotspots and model-specific bypasses\n\nThis means the logging is useful for both LLM and diffusion kernel debugging, but it does not automatically cover every pure PyTorch call in the repository.\n\n## Step 1: Enable Kernel API Logging\n\n### Basic Logging (Function Names Only)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=1\nexport SGLANG_KERNEL_API_LOGDEST=stdout\n\npython my_script.py\n```\n\nOutput:\n```\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: RMSNorm.forward\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.custom_op.fused_inplace_qknorm\n```\n\nThis is a real level-1 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Detailed Logging (Inputs with Metadata)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nOutput in `debug.log`:\n```\n================================================================================\n[2026-03-19 00:47:30] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[0]=QKVParallelLinear(\n repr=QKVParallelLinear(in_features=1024, output_features=4096, bias=False, tp_size=1, gather_output=False)\n )\n arg[1]=Tensor(\n shape=(1, 1024)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n arg[2]=None\nOutput:\n return=Tensor(\n shape=(1, 4096)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n```\n\nThis is a real level-3 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Full Logging (With Tensor Statistics)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nAdditional output:\n```\n================================================================================\n[2026-03-19 01:00:42] SGLang Kernel API Call: diffusion.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[1]=Tensor(\n shape=(1, 77, 768)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-27.250000\n max=28.500000\n mean=0.011723\n nan_count=0\n inf_count=0\n )\nOutput:\n return=Tensor(\n shape=(1, 77, 2304)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-8.937500\n max=9.375000\n mean=0.009460\n nan_count=0\n inf_count=0\n )\n```\n\nThis is a real level-5 excerpt captured from `black-forest-labs/FLUX.1-dev`.\n\n### Crash-Safe Dumps (Inputs Saved Before Execution)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\n\npython my_script.py\n```\n\nAt level 10, SGLang saves the inputs before execution. If the kernel crashes, the dump directory still contains the inputs and exception metadata.\n\nIf CUDA graph capture is active, tensor dumps are skipped automatically to avoid capture-time CUDA errors. In that case, you still get the kernel API call log, but not `inputs.pt` / `outputs.pt`.\n\nLevel-10 dumps are best understood as crash-safe call snapshots. They always preserve the observed call boundary. They do not guarantee one-click replay for every method, because some methods depend on module state that is not serialized into the dump.\n\nReal level-10 dump layout from `Qwen/Qwen3-0.6B`:\n\n```text\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/inputs.pt\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/metadata.json\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/outputs.pt\n```\n\nReal `metadata.json` excerpt:\n\n```json\n{\n \"function_name\": \"RotaryEmbedding.forward\",\n \"timestamp\": \"20260319_004821_182\",\n \"process_id\": 919286,\n \"execution_status\": \"completed\",\n \"input_tensor_keys\": [\"arg_0\", \"arg_1\", \"arg_2\"],\n \"output_tensor_keys\": [\"result_0\", \"result_1\"]\n}\n```\n\n## Step 2: Reproduce an LLM CUDA Crash\n\nCreate a temporary reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_llm_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.srt.utils.custom_op import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(indices, table):\\\\n\"\n \" return torch.empty((*indices.shape, table.shape[-1]), device=table.device, dtype=table.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_llm_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_llm_cuda_crash(indices, table):\\\\n\"\n \" out = F.embedding(indices, table)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"table = torch.randn(4, 8, device='cuda', dtype=torch.float16)\\\\n\"\n \"indices = torch.tensor([0, 7], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_llm_cuda_crash(indices, table)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level1.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nWhat to expect:\n- The script exits with a CUDA `device-side assert`\n- The log still contains the last API boundary before the crash\n\nTry the same example at level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level3.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow the log shows tensor metadata before the crash.\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_llm_level10_dumps \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow you should see:\n- A log entry for `sglang.custom_op.mock_llm_cuda_crash`\n- A dump directory with `inputs.pt`\n- `metadata.json` showing `execution_status: \"exception\"`\n- No `outputs.pt`, because the kernel crashed before producing output\n\nFor real-model success-path level-10 dumps, it is often easier to temporarily disable CUDA graph and piecewise CUDA graph for the debug run.\n\n## Step 3: Reproduce a Diffusion CUDA Crash\n\nCreate a temporary diffusion-side reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_diffusion_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.multimodal_gen.runtime.layers.utils import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(positions, cache):\\\\n\"\n \" return torch.empty((*positions.shape, cache.shape[-1]), device=cache.device, dtype=cache.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_diffusion_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_diffusion_cuda_crash(positions, cache):\\\\n\"\n \" out = F.embedding(positions, cache)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"cache = torch.randn(4, 64, device='cuda', dtype=torch.float16)\\\\n\"\n \"positions = torch.tensor([0, 9], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_diffusion_cuda_crash(positions, cache)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level1.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level3.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_diffusion_level10_dumps \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nIf your local environment has unrelated FlashInfer import issues, resolve them in the shell before running the example. The example itself does not set any `FLASHINFER_*` environment variable.\n\n## Step 4: Multi-Process Debugging\n\nWhen running with multiple GPUs or worker processes, use `%i` in the log path:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\n\ntorchrun --nproc_per_node=4 my_script.py\n```\n\nThis creates separate logs such as:\n- `debug_rank_12345.log`\n- `debug_rank_12346.log`\n- `debug_rank_12347.log`\n- `debug_rank_12348.log`\n\nReal multi-process example from a 2-GPU `Qwen/Qwen2.5-0.5B-Instruct` run:\n\n```text\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950201.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950349.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950350.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950351.log\n```\n\nYou should usually do the same for level-10 dump directories:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps_%i\n```\n\nThis avoids multiple ranks writing into the same dump directory tree.\n\n## Step 5: Filter Level-10 Dumps\n\nIf level 10 is too noisy, restrict dumps to specific APIs:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\nexport SGLANG_KERNEL_API_DUMP_INCLUDE='sglang.custom_op.*'\nexport SGLANG_KERNEL_API_DUMP_EXCLUDE='*.fake_impl'\n```\n\n`SGLANG_KERNEL_API_DUMP_INCLUDE` and `SGLANG_KERNEL_API_DUMP_EXCLUDE` use shell-style wildcard matching.\n\n## Step 6: Common CUDA Errors and What to Check\n\n### Illegal Memory Access or Device-Side Assert\n\n**Typical errors**:\n```\nRuntimeError: CUDA error: an illegal memory access was encountered\ntorch.AcceleratorError: CUDA error: device-side assert triggered\n```\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck in the logs:\n- ✅ Tensor shapes\n- ✅ Tensor dtypes\n- ✅ CUDA vs CPU device placement\n- ✅ Tensor stride / contiguity\n- ✅ Whether the failing call has inputs logged but no outputs logged\n\nTypical shape-mismatch pattern:\n\n```text\nSGLang Kernel API Call: ...\narg[0]=Tensor(shape=(..., 128), ...) # ✅ expected dimension\narg[1]=Tensor(shape=(..., 64), ...) # ❌ mismatch\n```\n\nThis often points to head-dim, hidden-dim, or cache-layout mismatch rather than a random CUDA failure.\n\n### NaN or Inf\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nCheck:\n- `min`\n- `max`\n- `mean`\n- `nan_count`\n- `inf_count`\n\nTypical bad pattern:\n\n```text\nTensor(\n ...\n min=-1234567.000000 # ❌ suspiciously large\n max=9876543.000000 # ❌ suspiciously large\n mean=nan # ❌ bad\n nan_count=128 # ❌ found NaNs\n inf_count=0 # ✅ no Infs here\n)\n```\n\nThis usually means the bad values were already present before the crashing kernel.\n\n### Out of Memory\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck:\n- Unexpectedly large tensor shapes\n- Batch size\n- Sequence length\n- Frame count or image resolution in diffusion workloads\n\nAlso check whether a supposedly per-token or per-frame tensor accidentally became full-sequence or full-image sized.\n\nTypical bad pattern:\n\n```text\nTensor(\n shape=(1024, 8192, 128, 128) # ❌ way too large\n ...\n)\n```\n\n### Example: Spot a Shape Bug from the Log\n\nSuppose the failing API log looks like this:\n\n```text\n[2026-03-19 00:47:30] SGLang Kernel API Call: RotaryEmbedding.forward\nPositional input arguments:\n arg[0]=Tensor(shape=(1, 8), dtype=torch.int64, ...)\n arg[1]=Tensor(shape=(1, 8, 8, 256), dtype=torch.bfloat16, ...) # ✅ query\n arg[2]=Tensor(shape=(1, 8, 4, 64), dtype=torch.bfloat16, ...) # ❌ key head_dim mismatch\n```\n\nWhat this tells you:\n- ✅ positions look reasonable\n- ✅ query looks plausible\n- ❌ key last dimension is inconsistent with the expected rotary/head dimension\n\nThat usually means the bug is in projection layout, head packing, or cache format rather than in the rotary kernel itself.\n\n## Step 7: Combine with compute-sanitizer\n\nFor harder bugs, combine kernel API logging with CUDA memory checking:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncompute-sanitizer --tool memcheck python3 /tmp/sglang_llm_crash.py\n```\n\nUse `debug.log` to see the exact inputs that reached the crashing API boundary.\n\nTypical `compute-sanitizer` output:\n\n```text\n========= COMPUTE-SANITIZER\n========= Invalid __global__ write of size 4 bytes\n========= at 0x1234 in SomeKernel\n========= by thread (256,0,0) in block (10,0,0)\n========= Address 0x... is out of bounds\n```\n\nUse the sanitizer output to identify the failing kernel and use `debug.log` to identify the exact tensors that reached the API boundary right before it.\n\nIf you need more synchronous host-side error reporting, you can try `CUDA_LAUNCH_BLOCKING=1` as a separate follow-up experiment. It is not part of the default workflow because it changes execution timing and can hide concurrency-related behavior.\n\n## Step 8: Combine with cuda-gdb\n\nFor crashes that need a stack trace instead of only memory diagnostics:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncuda-gdb --args python3 /tmp/sglang_llm_crash.py\n```\n\nInside `cuda-gdb`:\n\n```text\n(cuda-gdb) run\n(cuda-gdb) where\n```\n\nThen correlate the backtrace with `debug.log`.\n\n## Step 9: Kernel-Level Debugging with printf()\n\nWhen you own the CUDA kernel, `printf()` is still useful for narrowing down bad indices, bad launch geometry, or broken state propagation.\n\nBasic pattern:\n\n```cpp\n__global__ void MyKernel(const float* input, float* output, int n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n\n if (threadIdx.x == 0 && blockIdx.x == 0) {\n printf(\"n=%d input0=%f\\n\", n, input[0]);\n }\n\n if (idx < n) {\n output[idx] = input[idx] * 2.0f;\n }\n}\n```\n\nAfter launch, force the output to flush:\n\n```python\nmy_kernel(...)\ntorch.cuda.synchronize()\n```\n\nFor warp-specialized kernels, do not blindly print only on `threadIdx.x == 0`. Pick one representative thread per warp or per specialization group instead.\n\n### Warp-Specialized Kernels: Choosing the Right Print Thread\n\nProblem:\n- `threadIdx.x == 0` only prints from the first warp in the block\n- for warp-specialized kernels, that often misses the warp or group that is actually wrong\n\nBetter pattern:\n\n```cpp\n__global__ void WarpSpecializedKernel(...) {\n // Example: first lane of each warp\n if ((threadIdx.x % 32) == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n }\n}\n```\n\nOr, if the kernel is organized in larger specialization groups, print once per group instead of once per block.\n\nCommon mistake:\n\n```cpp\n// Only warp 0 prints\nif (threadIdx.x == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n}\n```\n\n### Quick Reference\n\n| Kernel Type | Print Condition | Notes |\n|----------|----------|-------------|\n| Simple kernel | `threadIdx.x == 0` | One thread per block is usually enough |\n| Warp-specialized kernel | one representative lane per warp | e.g. `threadIdx.x % 32 == 0` |\n| Group-specialized kernel | one representative lane per group | choose based on the kernel's scheduling layout |\n\n### Other Kernel Debugging Tools\n\n```cpp\nassert(value >= 0.0f && \"value must be non-negative\");\nstatic_assert(BLOCK_SIZE % 32 == 0, \"BLOCK_SIZE must be warp aligned\");\n```\n\n## Environment Variables Reference\n\n| Variable | Values | Description |\n|----------|--------|-------------|\n| `SGLANG_KERNEL_API_LOGLEVEL` | `0` | No logging (default) |\n| | `1` | Function names only |\n| | `3` | Inputs and outputs with metadata |\n| | `5` | Level 3 plus tensor statistics |\n| | `10` | Level 5 plus crash-safe tensor dumps |\n| `SGLANG_KERNEL_API_LOGDEST` | `stdout` | Log to stdout |\n| | `stderr` | Log to stderr |\n| | `` | Log to file |\n| | `log_%i.txt` | `%i` expands to process ID |\n| `SGLANG_KERNEL_API_DUMP_DIR` | `` | Directory for level-10 dumps |\n| `SGLANG_KERNEL_API_DUMP_INCLUDE` | wildcard list | Only dump matching API names |\n| `SGLANG_KERNEL_API_DUMP_EXCLUDE` | wildcard list | Skip matching API names |\n\n## Best Practices\n\n### 1. Start with Level 3\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nLevel 3 is usually enough to catch wrong shapes, wrong dtypes, and wrong devices.\n\n### 2. Use Level 5 for Numerical Issues\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nUse it when you suspect NaN or Inf values.\n\n### 3. Use Level 10 for Crash Reproduction\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\n```\n\nThis is the most useful mode when the process crashes before you can inspect live tensors.\n\nIf you need successful input/output dumps from a real model run, temporarily disable CUDA graph for that debug session.\n\nWhen level 10 is too noisy, pair it with `SGLANG_KERNEL_API_DUMP_INCLUDE` / `SGLANG_KERNEL_API_DUMP_EXCLUDE` instead of dumping every covered API.\n\n### 4. Log to File for Crashes\n\n```bash\nexport SGLANG_KERNEL_API_LOGDEST=crash.log\n```\n\nFile logs are safer than stdout when the process aborts.\n\n### 5. Disable Logging in Production\n\n```bash\nunset SGLANG_KERNEL_API_LOGLEVEL\n```\n\nWhen disabled, the decorator returns the original callable and adds no runtime logging overhead.\n\n## Troubleshooting\n\n### No Logs Appear\n\nCheck:\n1. `echo $SGLANG_KERNEL_API_LOGLEVEL`\n2. `echo $SGLANG_KERNEL_API_LOGDEST`\n3. Whether the failing path goes through a covered API boundary\n\n### Too Much Output\n\nReduce the level:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\n### Statistics Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nstatistics=[skipped: CUDA graph capture in progress]\n```\n\nThat is expected. Level-5 statistics are intentionally skipped during CUDA graph capture to avoid synchronization side effects.\n\n### Tensor Dumps Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nTensor dump skipped: CUDA graph capture in progress\n```\n\nThat is also expected. Level-10 dumps require copying tensors to CPU, which is not allowed during CUDA graph capture.\n",".agents/skills/env-var-conventions/SKILL.md":"# env-var-conventions\n\nPath: `.agents/skills/env-var-conventions/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/env-var-conventions/SKILL.md)",".agents/skills/generate-profile/SKILL.md":"# generate-profile\n\nPath: `.agents/skills/generate-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/generate-profile/SKILL.md)",".agents/skills/kernel-organization/SKILL.md":"# kernel-organization\n\nPath: `.agents/skills/kernel-organization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kernel-organization/SKILL.md)",".agents/skills/kl-consistency-test/SKILL.md":"# kl-consistency-test\n\nPath: `.agents/skills/kl-consistency-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kl-consistency-test/SKILL.md)",".agents/skills/large-class-style/SKILL.md":"# large-class-style\n\nPath: `.agents/skills/large-class-style/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/large-class-style/SKILL.md)",".agents/skills/llm-torch-profiler-analysis/SKILL.md":"# llm-torch-profiler-analysis\n\nPath: `.agents/skills/llm-torch-profiler-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/SKILL.md)",".agents/skills/mechanical-refactor-verify/SKILL.md":"# mechanical-refactor-verify\n\nPath: `.agents/skills/mechanical-refactor-verify/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/SKILL.md)",".agents/skills/mechanical-refactor-verify/guide-modify-skill.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-modify-skill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-modify-skill.md)",".agents/skills/processor-model-parity/SKILL.md":"# processor-model-parity\n\nPath: `.agents/skills/processor-model-parity/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/processor-model-parity/SKILL.md)",".agents/skills/scripted-runtime-notes/SKILL.md":"# scripted-runtime-notes\n\nPath: `.agents/skills/scripted-runtime-notes/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/scripted-runtime-notes/SKILL.md)",".agents/skills/sglang-bisect-ci-regression/SKILL.md":"# sglang-bisect-ci-regression\n\nPath: `.agents/skills/sglang-bisect-ci-regression/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-bisect-ci-regression/SKILL.md)",".agents/skills/sglang-cherrypick/SKILL.md":"# sglang-cherrypick\n\nPath: `.agents/skills/sglang-cherrypick/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-cherrypick/SKILL.md)",".agents/skills/sglang-prod-incident-triage/SKILL.md":"# sglang-prod-incident-triage\n\nPath: `.agents/skills/sglang-prod-incident-triage/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/SKILL.md)",".agents/skills/sglang-runtime-context/SKILL.md":"# sglang-runtime-context\n\nPath: `.agents/skills/sglang-runtime-context/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-runtime-context/SKILL.md)",".agents/skills/speculative-naming/SKILL.md":"# speculative-naming\n\nPath: `.agents/skills/speculative-naming/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/speculative-naming/SKILL.md)",".agents/skills/write-sglang-test/SKILL.md":"# write-sglang-test\n\nPath: `.agents/skills/write-sglang-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/write-sglang-test/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md":"# sglang-diffusion-add-model\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md":"# sglang-diffusion-benchmark-profile\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md":"# sglang-diffusion-cache-dit\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md":"# sglang-diffusion-modelopt-quant\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md":"# sglang-diffusion-performance\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md)","docs/AGENTS.md":"# docs - AI Agent Protocol & Instructions\n\nPath: `docs/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/docs/AGENTS.md)","sgl-router/ARCHITECTURE.md":"# sgl-router - System Architecture & Module Boundaries\n\nPath: `sgl-router/ARCHITECTURE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/sgl-router/ARCHITECTURE.md)",".claude/rules/comment-style.md":"# Rules Skill\n\nPath: `.claude/rules/comment-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/comment-style.md)",".claude/rules/cookbook-first.md":"# Rules Skill\n\nPath: `.claude/rules/cookbook-first.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/cookbook-first.md)",".claude/rules/forward-batch-init-new-purity.md":"# Rules Skill\n\nPath: `.claude/rules/forward-batch-init-new-purity.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/forward-batch-init-new-purity.md)",".claude/rules/general-code-style.md":"# Rules Skill\n\nPath: `.claude/rules/general-code-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/general-code-style.md)",".claude/rules/kernel-benchmark.md":"# Rules Skill\n\nPath: `.claude/rules/kernel-benchmark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/kernel-benchmark.md)",".claude/rules/modify-component-must-read.md":"# Rules Skill\n\nPath: `.claude/rules/modify-component-must-read.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/modify-component-must-read.md)",".claude/rules/no-dataclasses.md":"# Rules Skill\n\nPath: `.claude/rules/no-dataclasses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-dataclasses.md)",".claude/rules/no-getattr-defensive.md":"# Rules Skill\n\nPath: `.claude/rules/no-getattr-defensive.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-getattr-defensive.md)",".claude/rules/prefer-sgl-eval.md":"# Rules Skill\n\nPath: `.claude/rules/prefer-sgl-eval.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/prefer-sgl-eval.md)",".claude/rules/schedule-batch-out-of-place-mutation.md":"# Rules Skill\n\nPath: `.claude/rules/schedule-batch-out-of-place-mutation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/schedule-batch-out-of-place-mutation.md)",".claude/rules/unit-test-admission.md":"# Rules Skill\n\nPath: `.claude/rules/unit-test-admission.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/unit-test-admission.md)","tools/sglang-simulator/README.md":"# sglang-simulator Documentation\n\nPath: `tools/sglang-simulator/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/README.md)",".agents/skills/ci-test-audit/action-items.md":"# Ci-test-audit Skill\n\nPath: `.agents/skills/ci-test-audit/action-items.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/action-items.md)",".agents/skills/mechanical-refactor-verify/guide-construct-proof.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-construct-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-construct-proof.md)",".agents/skills/mechanical-refactor-verify/guide-split.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-split.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-split.md)",".agents/skills/mechanical-refactor-verify/guide-verify-proof.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-verify-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-verify-proof.md)",".agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md)",".agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md)","examples/profiler/nsys_profile_tools/README.md":"# nsys_profile_tools Documentation\n\nPath: `examples/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/examples/profiler/nsys_profile_tools/README.md)","tools/sglang-simulator/examples/README.md":"# examples Documentation\n\nPath: `tools/sglang-simulator/examples/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/README.md)",".agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml":"# Agents Skill\n\nPath: `.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml)",".agents/skills/clean-startup-log/references/noise-sources.md":"# References Skill\n\nPath: `.agents/skills/clean-startup-log/references/noise-sources.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/references/noise-sources.md)",".agents/skills/cookbook-add-model/references/authoring-reference.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/authoring-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/authoring-reference.md)",".agents/skills/cookbook-add-model/references/diffusion-authoring.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/diffusion-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/diffusion-authoring.md)",".agents/skills/cookbook-add-model/references/engine-axis.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/engine-axis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/engine-axis.md)",".agents/skills/cookbook-add-model/references/mintlify-authoring.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/mintlify-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/mintlify-authoring.md)",".agents/skills/cookbook-add-model/references/vendor-logo.md":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/vendor-logo.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/vendor-logo.md)",".agents/skills/cookbook-migrate-model/references/dimension-mapping.md":"# References Skill\n\nPath: `.agents/skills/cookbook-migrate-model/references/dimension-mapping.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/references/dimension-mapping.md)",".agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md)",".agents/skills/llm-torch-profiler-analysis/references/heuristics.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/heuristics.md)",".agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md)",".agents/skills/llm-torch-profiler-analysis/references/source-map.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/source-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/source-map.md)",".agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md)",".agents/skills/sglang-prod-incident-triage/references/case-studies.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/case-studies.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/case-studies.md)",".agents/skills/sglang-prod-incident-triage/references/decision-tree.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/decision-tree.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/decision-tree.md)",".agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md)",".agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md)","tools/sglang-simulator/examples/assets/replay_table.json":"# Assets Skill\n\nPath: `tools/sglang-simulator/examples/assets/replay_table.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/replay_table.json)","tools/sglang-simulator/examples/sim_configs/aic_silicon.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_silicon.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_silicon.json)","tools/sglang-simulator/examples/sim_configs/aic_sol.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_sol.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_sol.json)","tools/sglang-simulator/examples/sim_configs/ml.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/ml.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/ml.json)","tools/sglang-simulator/examples/sim_configs/replay.json":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/replay.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/replay.json)","tools/sglang-simulator/examples/workloads/sharegpt-example.json":"# Workloads Skill\n\nPath: `tools/sglang-simulator/examples/workloads/sharegpt-example.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/workloads/sharegpt-example.json)","tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json)","tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json)","tools/sglang-simulator/test/assets/qwen3-8b/config.json":"# Qwen3-8b Skill\n\nPath: `tools/sglang-simulator/test/assets/qwen3-8b/config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/test/assets/qwen3-8b/config.json)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md)","tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md":"# collect_baseline Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md)","tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md":"# train_ml_time_predictor Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md)","python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md)"},"items":[{"name":"AGENTS.md","path":"docs/AGENTS.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/docs/AGENTS.md","title":"docs - AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# docs - AI Agent Protocol & Instructions\n\nPath: `docs/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/docs/AGENTS.md)","tokens":300,"sizeBytes":169},{"name":"ARCHITECTURE.md","path":"sgl-router/ARCHITECTURE.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/sgl-router/ARCHITECTURE.md","title":"sgl-router - System Architecture & Module Boundaries","category":"project-spec","format":"markdown","content":"# sgl-router - System Architecture & Module Boundaries\n\nPath: `sgl-router/ARCHITECTURE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/sgl-router/ARCHITECTURE.md)","tokens":300,"sizeBytes":206},{"name":"SKILL.md","path":".agents/skills/add-jit-kernel/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/add-jit-kernel/SKILL.md","title":"add-jit-kernel","category":"anthropic-skill","format":"markdown","content":"---\nname: add-jit-kernel\ndescription: Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups\n---\n\n# Tutorial: Adding a New JIT Kernel to SGLang\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as a JIT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float, passed at runtime)\n- Output: `x * factor` (element-wise), allocated internally\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n\n## When to use JIT vs AOT (`sgl-kernel`)\n\n- **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.\n- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `python/sglang/kernels/aot/` and participate in the wheel build / torch op registration flow.\n- **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.\n\n---\n\n## Conventions\n\nThese hold for every step below.\n\n- **`namespace sglang` is where JIT code lives.** Open it after the include block and close it at the end of the file, with the device kernels, traits and host wrapper inside. The shared `host::` / `device::` helpers are nested in it too, so they resolve unqualified. `load_jit` emits the `TVM_FFI_DLL_EXPORT_TYPED_FUNC` wrapper inside `namespace sglang` as well, so the `kernel_name` you pass from Python needs no `sglang::` prefix.\n- **Tuning policy belongs in Python; C++ gets the answer, not the decision.** For a hyperparameter — split factor, threads per item, block size, vector width, an algorithm variant — prefer (not required) making it a **template parameter** and letting the `@cache_once` module factory choose the value, over a runtime `if`/`switch` in the launcher that picks among pre-instantiated kernels. The kernel then compiles exactly one specialisation and asserts its own preconditions with `static_assert`, while the heuristic that produced the value sits in Python where it is readable, adjustable, and inspectable without recompiling CUDA. `store_cache`'s `num_threads` is the worked example: a `get_kernel(num_split)` ladder guarded by byte-alignment `if constexpr`s became one template argument plus a heuristic in `_jit_kvcache_module`. The exception is a knob that genuinely varies per call with the runtime shape — that has to stay a kernel argument.\n- **Check where the check is cheapest: `static_assert` > C++ host check > cached Python > per-call Python.** Anything fixed at compile time is a `static_assert`. Anything about the tensors is a `TensorMatcher` / `CHECK_HOST` in the C++ launcher, free next to a kernel launch. A check Python cannot delegate goes inside the `@cache_once` module factory, where it runs once per specialisation. What remains in the per-call entry point costs interpreter time on *every* forward, so it should be nothing but picking the module and allocating `out`.\n- **Fixed-width integer types.** Prefer `int32_t` / `int64_t` / `uint32_t` / `size_t` over `int`, `long`, or `long long`, so an index has the same width on both sides of the FFI boundary. Bare `int` is fine only where the width plainly cannot matter — an unrolled loop counter over a `constexpr` bound, a template `int` parameter. Shapes arrive as `int64_t` (`SymbolicSize::unwrap()`); narrowing to `uint32_t` for in-kernel indexing is a deliberate act, so write the `static_cast` explicitly and only where the range is known.\n- **Doxygen comments in C++.** Document exported entities with `///` or `/** ... */` blocks using `\\brief`, `\\param`, `\\tparam`, `\\return`, the way `include/sgl_kernel/` does. `python -m sglang.kernels.jit` writes `CommentFormat: Doxygen` into `.clangd` when clangd is 21 or newer, so these render on hover in the editor. Plain `//` remains fine for implementation notes inside a function body.\n- **ASCII only in C++ and CUDA sources.** Write `--`, `->`, `<=` instead of `—`, `→`, `≤`, including in comments. `grep -nP '[^\\x00-\\x7F]' ` before committing.\n- **`namespace details` is private.** Anything inside one is an implementation detail of its own header, free to change without notice. Do not name `details::` from another module, and never `using namespace details`. If you find yourself wanting something in there, that is the signal to promote it to a documented name instead.\n- **`const T* __restrict__` for read-only pointers.** This is what `csrc/` does throughout, and it lets the compiler emit non-coherent (`LDG`) loads.\n- **Watch the register budget.** For memory-bound kernels, keep to roughly 64 registers per thread so occupancy does not become the limit; prefer recomputing a value over letting it spill. Run `SGLANG_JIT_LOG_RESOURCE_USAGE=1 SGLANG_JIT_FORCE_RECOMPILE=1` to log per-kernel registers, spills and shared memory at INFO. Both halves are load-bearing: a cache hit has no compiler output to report, and passing `-Xptxas -v` by hand does nothing on its own because the build captures compiler output and replays it only on failure.\n- **Alignment is a contract you enforce, not a property you hope for.** A vectorized copy derives its access width from the row *size* (`load_bytes` picks 16B for a 1024B row), but the address it lands on is `base + index * row_stride` — and a stride is not constrained by the size. `TensorMatcher` admits any stride unless you say otherwise, so a padded cache row silently produces `misaligned address`. Every kernel that vectorizes is highly recommended to end its validation with `.ensure_alignment(w)`, where `w` is the width the kernel actually uses.\n\n---\n\n## Common Abstractions in `python/sglang/kernels/jit/include/sgl_kernel/`\n\n**Always prefer these abstractions over raw CUDA primitives.** They provide safety, readability, and consistency with the rest of the codebase. The only reason to drop to raw primitives is performance the abstraction cannot reach — a trade you make deliberately, and justify in a comment.\n\n### `utils.h` — Host-side utilities\n\n```cpp\n#include \n```\n\n- **`CHECK_HOST(cond) << \"msg \" << value`** — **Preferred** runtime check: stream-style, throws `PanicError` with file/line info on failure. Zero overhead on the true path — the message expressions are only evaluated when the check fails.\n- **`host::RuntimeCheck(cond, args...)`** — Function-style alternative to `CHECK_HOST`. Note its message args are always evaluated (even when the check passes), so prefer `CHECK_HOST` — especially on hot paths.\n- **`host::Panic(args...)`** — Unconditionally throw a `PanicError` with a descriptive message.\n- **`host::div_ceil(a, b)`** — Integer ceiling division `(a + b - 1) / b`.\n- **`host::irange(n)`** / **`host::irange(start, end)`** — Range views for cleaner loops.\n- **`host::pointer::offset(ptr, offsets...)`** — Byte-safe pointer arithmetic on `void*`. Use this instead of raw casts.\n\n### `utils.cuh` — Device-side utilities + `LaunchKernel`\n\n```cpp\n#include \n```\n\n- **Type aliases**: `fp16_t`, `bf16_t`, `fp32_t`, `fp8_e4m3_t`, `fp8_e5m2_t` and their packed variants `fp16x2_t`, `bf16x2_t`, `fp32x2_t`, etc.\n- **`SGL_DEVICE`** — Expands to `__forceinline__ __device__`. Use on all device functions.\n- **`SGL_DEVICE_HOST`** — `__forceinline__ __device__ __host__`. Use it when a `constexpr` helper must be callable from both the kernel and its launcher, so the two cannot drift.\n- **`device::kWarpThreads`** — Constant `32` (the codebase's *logical* warp width). `kWarpSize` is an alias. Note this is a group size, not the hardware wave: on gfx950 `warpSize` is 64, and `warp::kFullWidth` reflects that.\n- **`device::get_lane_id()`** — this thread's index within its logical group. **Prefer it over `threadIdx.x % kNumThreads` whenever the value feeds an address.** On CUDA the lane register is one read that folds into the address computation, whereas the modulo is a derived value the compiler re-materialises at *every* address scale — 8 instructions and 2 registers on a two-tile warp copy, measured on sm_100a. It picks the cheaper form per platform, so callers never need to care. Only equals the true in-warp lane when `blockDim.x` is a multiple of `kNumThreads`.\n- **`device::load_as(ptr, offset)`** / **`device::store_as(ptr, val, offset)`** — Type-safe loads/stores from `void*`.\n- **`device::pointer::offset(ptr, offsets...)`** — Pointer arithmetic on device.\n- **`host::LaunchKernel(grid, block, device_or_stream [, smem])`** — RAII kernel launcher that:\n - Resolves the CUDA stream from a `DLDevice` via TVM-FFI automatically.\n - Checks the CUDA error with file/line info after launch via `operator()(kernel, args...)`.\n - Supports `.enable_pdl(bool)` for PDL (Programmatic Dependent Launch, SM90+).\n- **`device::PDLWaitPrimary()`** / **`device::PDLTriggerSecondary()`** — The two halves of PDL, on sm_90+ (no-ops on older archs and ROCm). Their guarantees are **not** symmetric:\n - `PDLTriggerSecondary` (`griddepcontrol.launch_dependents`) only lets the next kernel in the stream *start* early. It carries no memory ordering and publishes nothing — matching that, the header's asm has no `\"memory\"` clobber.\n - `PDLWaitPrimary` (`griddepcontrol.wait`) is the ordering point: it waits until the preceding kernel has fully finished and its writes are visible.\n\n So every read of data the preceding kernel produced must come after `PDLWaitPrimary()`. What overlaps with the primary's tail is whatever you put *before* the wait — loading parameters, computing indices, touching buffers the primary never wrote — so a kernel that waits on its first line gains nothing. Neither call is a barrier: threads may reach or skip them independently. See \"Programmatic Dependent Launch and Synchronization\" in the CUDA C++ Programming Guide.\n- **`CHECK_CUDA(expr) << \"context\"`** — Stream-style CUDA error check; evaluates `expr` once and throws `PanicError` with `cudaGetErrorString` + file/line info if it is not `cudaSuccess`. Extra streamed context is optional.\n- **`host::RuntimeDeviceCheck(cudaError_t)`** — Function-style alternative to `CHECK_CUDA`. It takes no context message, so prefer `CHECK_CUDA`, which builds its error object only on the failure path.\n\n### `tensor.h` — Tensor validation (`TensorMatcher`, Symbolic types)\n\n```cpp\n#include \n```\n\nThis is the **primary validation API** for all kernel launchers. Use it to validate every `tvm::ffi::TensorView` argument.\n\n- **`host::SymbolicSize{\"name\"}`** — A named symbolic dimension. Call `.set_value(n)` to pin it, `.unwrap()` to extract after verification.\n- **`host::SymbolicDType`** — Symbolic dtype. Use `.set_options()` to restrict allowed types.\n- **`host::SymbolicDevice`** — Symbolic device. Use `.set_options()` to restrict to CUDA.\n- **`host::TensorMatcher({dims...})`** — Fluent builder for tensor validation:\n - `.with_dtype()` — require a specific C++ type (e.g. `fp16_t`)\n - `.with_dtype()` — allow a set of types\n - `.with_device(device_sym)` — require CUDA and bind the checked device to a `SymbolicDevice`\n - `.with_strides({strides...})` — validate strides (omit to require contiguous)\n - `.ensure_alignment(bytes)` — require the data pointer **and every stride except the innermost** to be a multiple of `bytes` (a power of two). Size-1 dimensions are skipped, since their strides are arbitrary. This is how a vectorized kernel states its precondition; see the alignment convention above.\n - `.verify(tensor_view)` — execute the check; throws `PanicError` with full context on failure; **chainable** (`verify(a).verify(b)` to check multiple tensors with the same shape)\n- **`host::is_type(dtype)`** — whether a `DLDataType` denotes the C++ type `T` (e.g. `fp16_t`).\n\n**Typical pattern:**\n```cpp\nauto N = SymbolicSize{\"num_elements\"};\nauto device = SymbolicDevice{};\ndevice.set_options();\nTensorMatcher({N}) //\n .with_dtype()\n .with_device(device)\n .ensure_alignment(16) // e.g. for 128-bit vectorized load/store\n .verify(dst)\n .verify(src); // same shape, dtype, device as dst\nconst int64_t n = N.unwrap();\nconst DLDevice dev = device.unwrap();\nconst int64_t last_dim = 128;\nTensorMatcher({N, last_dim}) // a fixed dimension can be a plain integer\n .with_dtype()\n .with_device(device)\n .verify(tensor_2d);\n```\n\n### `ffi.h` — Tensor allocation and blob wrapping (`host::ffi::`)\n\n```cpp\n#include \n```\n\nThe counterpart to `tensor.h`: that one validates what came in, this one produces new `tvm::ffi::Tensor` values. Allocation goes through the environment allocator (`TVMFFIEnvTensorAlloc`), so buffers come from PyTorch's caching allocator rather than a raw `cudaMalloc`.\n\n- **`host::alloc_workspace_tensor(nbytes, device)`** (declared in `utils.cuh`) — **the way to get scratch memory**: a 1-D `uint8` tensor of `nbytes`, or an empty tensor when `nbytes == 0`. Hold the returned `Tensor` in a local across every launch that touches it — it frees on destruction.\n- **`host::ffi::empty(shape, dtype, device)`** — Uninitialized tensor; `shape` accepts a braced list, so `ffi::empty({rows, sizeof(Plan)}, dtype, device)` works for a typed scratch array.\n- **`host::ffi::empty_like(tensor_view)`** — Same shape, dtype, and device as an existing tensor.\n- **`host::ffi::from_blob(data, shape, dtype, device[, deleter, stride, byte_offset])`** / **`from_blob_like(data, tensor_view, ...)`** — View memory you already own as a `Tensor`, no copy. The default deleter does nothing, so ownership stays with the caller; pass one only when the `Tensor` should own the block. Strides default to contiguous.\n\n### `type.cuh` — `DTypeTrait`, `packed_t`, and reduction traits\n\n```cpp\n#include \n```\n\n- **`DTypeTrait`** — Static trait struct, specialized for integral types, `fp32_t`, `fp16_t`, `bf16_t`, `fp8_e4m3_t`, and their packed x2/x4 variants. Provides:\n - `DTypeTrait::from(value)` — convert from another type via the right CUDA intrinsic (e.g. `fp32_t` → `fp16_t`)\n - `DTypeTrait::abs/max/min` — type-dispatched math (fp32, fp16/bf16 scalar and x2, integrals)\n - `DTypeTrait::sqrt/rsqrt/exp/sin/cos(x)` — `fp32_t` only\n - Metadata: `packed_t` / `unpacked_t` / `kVecSize` (packed layout), `kFloatMax` (dtype max as float, e.g. 448.0f for fp8-e4m3), `kZeroBits`\n- **`packed_t`** — Two-element packed alias: `packed_t` = `fp16x2_t`, `packed_t` = `bf16x2_t`, `packed_t` = `fp32x2_t`. Use for vectorized loads/stores.\n- **`device::cast(value)`** — Type-safe cast using `DTypeTrait`, e.g. `cast(v)`.\n- **`device::unpack(value)`** — View a packed value as an `unpacked_t[kVecSize]` array reference (e.g. `fp32x2_t` → `fp32_t[2]`); element writes propagate back to the packed value.\n- **`device::ReductionOp` (`SUM`/`MAX`/`MIN`) and `device::ReductionTrait::reduce(x, y)`** — One binary reduction step, dispatched through `DTypeTrait` (packed types reduce elementwise). This is the engine behind `warp::reduce`; use it directly when writing custom reductions.\n\n### `vec.cuh` — Vectorized memory access (`AlignedVector`)\n\n```cpp\n#include \n```\n\n- **`device::AlignedVector`** — Aligned storage for N elements of type T. N must be a power of two, `sizeof(T)*N <= 32`. Enables vectorized loads/stores for bandwidth efficiency. In terms of API/codegen constraints, the upper bound is 256-bit; in practice, 128-bit is the portable default, while 256-bit vectorization is typically only viable on `SM100+` and should be gated by an architecture check when needed.\n - `.load(ptr, offset)` — vectorized load from `ptr[offset]`\n - `.store(ptr, offset)` — vectorized store to `ptr[offset]`\n - `.fill(value)` — fill all N elements with `value`\n - `operator[](i)` — element access\n- **`device::LoadStoreBytes`** — named widths to pass around instead of bare integers: `MAX_GMEM` (arch-dependent, 32 on Blackwell), `MAX_SMEM` (16), `MAX_PORTABLE` (16, safe on CUDA and HIP), `MIN_COALALESCED` (4), and `RAW_1B` .. `RAW_32B`.\n\n### `tile.cuh` — `tile::Memory` (strided memory access pattern)\n\n```cpp\n#include \n```\n\n- `tile::Memory` is fundamentally a **1D cooperative accessor** over a contiguous region.\n- **`device::tile::Memory::cta(blockDim.x)`** — Creates a tile accessor where each thread handles `tid = threadIdx.x` with stride `tsize` (for `cta(blockDim.x)`, this is `blockDim.x`). Common for loops over a 1D array.\n- **`device::tile::Memory::warp()`** — the 32-lane flavour, and what `warp::load_bytes` is built on. It takes its lane index from `get_lane_id()` rather than `threadIdx.x % 32`, for the addressing reason described under `utils.cuh`. Use `warp(int n)` for a narrower sub-group; that overload keeps the modulo, since a hardware lane index is the wrong grouping below 32.\n- **`device::tile::Memory::thread()`** — single-thread accessor, no cooperation.\n- **`.load(ptr, offset)`** — loads `ptr[tid + offset * tsize]`\n- **`.store(ptr, val, offset)`** — stores to `ptr[tid + offset * tsize]`\n- **`.in_bound(n, offset)`** — boundary check\n\nFor a **2D tile**, either flatten `(row, col)` into a linear tile index first, or compute the address manually with `ptr[row * stride + col]` using your thread/block coordinates.\n\n### `bits.h` — Compile-time bit helpers (`host::`)\n\n```cpp\n#include \n```\n\n`constexpr` wrappers over ``, usable in `static_assert` and in template arguments: `host::is_pow2(x)`, `log2_floor(x)`, `log2_ceil(x)` (both `-1` for `x == 0`), `round_up_pow2(x)`, `round_down_pow2(x)`. They live in `host::` but are constant-expression-only, so device-side `static_assert` can call them.\n\n### `math.cuh` — Device math (`device::math::`)\n\n```cpp\n#include \n```\n\n- `device::math::max/min(a, b)` — type-dispatched binary math via `DTypeTrait`\n- `device::math::abs/sqrt/rsqrt/exp/sin/cos(x)` — type-dispatched unary math via `DTypeTrait`\n\n### `warp.cuh` — Warp-level primitives\n\n```cpp\n#include \n```\n\n**Vectorized row copy — reach for this before hand-rolling.** A warp cooperatively moving one contiguous row is the single most common shape in this tree, and it used to be re-implemented per kernel.\n\n- **`warp::load_bytes(src)`** / **`warp::store_bytes(dst, val)`** — the whole warp moves `kBytes` from/to one contiguous row. `kBytes` is arbitrary, down to 1; the helper picks the widest vector that divides it and handles the ragged tail. The returned value is opaque and encodes the width, so a load and its matching store **must use the same ``**.\n- **`warp::LoadStorePattern`** — the `kPattern` values, and the sign carries the meaning. A **negative** one (`WARP_UNIFORM_16B` and friends: `_GMEM`, `_SMEM`, `_4B`/`_8B`/`_32B`) says *the whole warp splits this row*, so the width is derived from the per-lane share (`gcd(kBytes / 32, |kPattern|)`), falling back to `gcd(kBytes, 4)` when the row is too narrow to give every lane 4 bytes. A **positive** one (a bare `LoadStoreBytes` value) makes no warp-splitting assumption and just caps the per-thread vector at `gcd(kBytes, kPattern)`. `WARP_UNIFORM_16B` is the usual choice for a row copy.\n- **`LoadStorePattern::get_vec_bytes()`** — the width actually chosen. `SGL_DEVICE_HOST`, so **the launcher must call this to feed `.ensure_alignment(...)`**. Pass it the same value the kernel passes `load_bytes` — for a split row that is the *per-warp share*, not the whole row; the full row can resolve to a different width and would under-constrain the strides.\n\n**Reductions.**\n\n- `warp::reduce(value, active_mask)` — `Op` is a `device::ReductionOp` (`SUM`/`MAX`/`MIN`). `kStart` and `kFinish` bound a range of lane-index bits: `` reduces within contiguous groups of `N` lanes, `` reduces *across* groups at the same offset. The operation is **symmetric** — `` and `` reduce the same lane set.\n- `warp::reduce_sum/reduce_max/reduce_min(value)` — wrappers. Work for any type with a `ReductionTrait`: floats, integers, packed x2. This is more widely used in practice than generic `reduce`.\n- `warp::inclusive_reduce(val, lane_id, mask)` and `inclusive_sum` / `inclusive_max` / `inclusive_min` — segmented inclusive scan: every lane keeps its own running total, unlike `reduce`. Forward when `kStart < kFinish`, backward when `kStart > kFinish`. `lane_id` defaults to `get_lane_id()`, which is correct by construction — only pass it if you already have the **segment-relative** index (`threadIdx.x % kWidth`); a warp-relative one silently corrupts every segment past the first when `kWidth < 32`.\n- `warp::broadcast(value, src_lane)` — `src_lane` is **segment-relative**: each `kWidth` segment reads its own lane, not one warp-wide source.\n- `warp::get_lane_id`, `warp::elect_one_lane()` (one elected lane, for gating a single-thread TMA issue; CUDA-only), `warp::kFullMask` and `warp::kFullWidth` (32 on CUDA, 64 on HIP — note this differs from `kWarpThreads`).\n\n### `cta.cuh` — CTA-level primitives\n\n```cpp\n#include \n```\n\n- `device::cta::reduce_max(value, smem, min_value)` — CTA-wide max using shared memory + warp reduction. Caller is responsible for a `__syncthreads()` after if the result in `smem[0]` is needed.\n\n### `atomic.cuh` — Atomic operations\n\n```cpp\n#include \n```\n\n- `device::atomic::max(float* addr, float value)` — float atomic max (handles negative values correctly via bit tricks).\n- **`device::atomic::Event`** — a cross-CTA arrive/wait counter in one 32-bit word. Producers call `arrive()`; consumers call `wait(num_producers)` (exactly one consumer) or `wait_multi(num_producers, num_consumers, n)` (several). The last released consumer resets the word, so one `Event` is reusable across launches with no host re-zero. Preconditions, all of them load-bearing:\n - the word must be **zeroed by the host** before first use, and must live in **global** memory — a shared or local `Event` is an illegal address;\n - `arrive()` is a release and the waits are acquires, so a producer's writes before `arrive()` are visible after the wait;\n - **generations must be explicitly ordered** — the word has no phase bit, so overlapping generation N+1 with a consumer still in generation N is undefined behaviour (a lagging consumer never exits);\n - `wait()` and `wait_multi()` lay the word out incompatibly; mixing them on one `Event` is undefined behaviour.\n\n### `runtime.cuh` — Occupancy and device info\n\n```cpp\n#include \n```\n\n- `host::runtime::get_blocks_per_sm(kernel, block_dim)` — max active blocks per SM (occupancy)\n- `host::runtime::get_sm_count(device_id [, use_cache])` — number of SMs on the device. Memoized per device ordinal, so it is cheap enough to call on a launch path; pass `use_cache=false` to force a driver query.\n- `host::runtime::get_cc_major` / `get_cc_minor` / `get_sm_version` — compute capability, same caching.\n\n**Do not query the architecture at runtime.** The JIT compiles for the exact local GPU, so the arch is a *compile-time* fact: `SGL_CUDA_ARCH` is injected by `load_jit`, and `SGL_ARCH_HOPPER_OR_GREATER` / `SGL_ARCH_BLACKWELL_OR_GREATER` / `device::kMaxVecBytes` are derived from it. Reach for `get_cc_*` only for something the arch genuinely does not determine. SM *count* is the opposite case — it varies within an arch (B200 vs B300 vs a MIG slice), so it has to stay a runtime query.\n\n**Persistent kernel pattern** (cap blocks to SM count × occupancy):\n```cpp\nstatic const uint32_t max_occ = runtime::get_blocks_per_sm(kernel, kBlockSize);\nstatic const uint32_t num_sm = runtime::get_sm_count(device.unwrap().device_id);\nconst auto num_blocks = std::min(num_sm * max_occ, div_ceil(n, kBlockSize));\nLaunchKernel(num_blocks, kBlockSize, device.unwrap())(kernel, params);\n```\n\n---\n\n## Step 0 (optional): Generate a `.clangd` config for better IDE support\n\n```bash\npython -m sglang.kernels.jit -h # for verbose help info about clangd configuration\npython -m sglang.kernels.jit\npython -m sglang.kernels.jit --dep cutlass flashinfer # with cutlass/flashinfer dependency\n```\n\n---\n\n## Step 1: Implement the CUDA kernel in `kernels/jit/csrc/`\n\nCreate `python/sglang/kernels/jit/csrc/elementwise/scale.cuh`.\n\nThe implementation fully uses the project abstractions described above:\n\n```cpp\n// NOTE: Comments for headers are not common in practice.\n// It is only shown here for tutorial purposes to highlight the key abstractions.\n#include // For TensorMatcher, SymbolicSize, SymbolicDevice\n#include // For DTypeTrait, fp16_t, bf16_t, fp32_t\n#include // For CHECK_HOST, div_ceil\n#include // For LaunchKernel, SGL_DEVICE\n#include // For AlignedVector\n\n#include \n#include \n\nnamespace sglang {\n\n/**\n * \\brief Element-wise scale using vectorized 128-bit loads/stores.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kVecN Elements per vector load (e.g. 8 for fp16)\n * \\tparam kUsePDL Whether to emit the PDL wait/trigger pair\n * \\param dst Output buffer, `n_total` elements\n * \\param src Input buffer, `n_total` elements\n * \\param factor Runtime scale factor\n * \\param n_total Number of elements to scale\n */\ntemplate \n__global__ void scale_kernel(T* __restrict__ dst,\n const T* __restrict__ src,\n float factor,\n uint32_t n_total) {\n using vec_t = device::AlignedVector;\n const uint32_t n_vecs = n_total / kVecN;\n\n // If using PDL, wait for primary kernel before any global memory load.\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLWaitPrimary();\n\n // --- vectorised body ---\n const uint32_t vec_stride = blockDim.x * gridDim.x;\n for (uint32_t vi = blockIdx.x * blockDim.x + threadIdx.x;\n vi < n_vecs;\n vi += vec_stride) {\n vec_t v;\n v.load(src, vi);\n#pragma unroll\n for (int i = 0; i < kVecN; ++i) {\n v[i] = static_cast(static_cast(v[i]) * factor);\n }\n v.store(dst, vi);\n }\n\n // --- scalar tail ---\n const uint32_t base = n_vecs * kVecN;\n const uint32_t scalar_stride = blockDim.x * gridDim.x;\n for (uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;\n base + i < n_total;\n i += scalar_stride) {\n dst[base + i] = static_cast(static_cast(src[base + i]) * factor);\n }\n\n // If using PDL, signal for the secondary kernel to start after all threads have finished\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLTriggerSecondary();\n}\n\n/**\n * \\brief Validate the tensors, select the vector width, launch `scale_kernel`.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kUsePDL Whether to launch with PDL enabled\n * \\param dst Output tensor; same shape / dtype / device as `src`\n * \\param src Input tensor on CUDA\n * \\param factor Runtime scale factor\n */\ntemplate \nvoid scale(tvm::ffi::TensorView dst, tvm::ffi::TensorView src, float factor) {\n using namespace host;\n\n // 1. Validate input tensors with TensorMatcher\n SymbolicSize N = {\"num_elements\"};\n SymbolicDevice device_;\n device_.set_options();\n\n TensorMatcher({N}) //\n .with_dtype()\n .with_device(device_)\n .verify(dst)\n .verify(src); // same shape / dtype / device as dst\n\n const uint32_t n = static_cast(N.unwrap());\n const DLDevice device = device_.unwrap();\n\n CHECK_HOST(n > 0) << \"scale: num_elements must be > 0, got \" << n;\n\n // 2. Choose vector width for 128-bit loads (16 bytes)\n // fp16/bf16: 8 elements x 2 bytes = 16 bytes\n // fp32: 4 elements x 4 bytes = 16 bytes\n // We encourage using `device::kMaxVecBytes`, which will change according to\n // the target architecture and can enable 256-bit vectorization on SM100+ if desired.\n // But 128-bit is more commonly adapted for better compatibility,\n // so it's still ok to hardcode 16 here just for simplicity.\n constexpr int kVecN = 16 / sizeof(T);\n const uint32_t n_work_items = div_ceil(n, static_cast(kVecN));\n\n // 3. Launch\n constexpr uint32_t kBlockSize = 256;\n const uint32_t grid = div_ceil(n_work_items, kBlockSize);\n\n // PDL feature is 100% optional. Without `enable_pdl`, the code should still be correct.\n // Try to enable it if profiling shows that it can benefit the performance of this kernel.\n LaunchKernel(grid, kBlockSize, device).enable_pdl(kUsePDL)(\n scale_kernel,\n static_cast(dst.data_ptr()),\n static_cast(src.data_ptr()),\n factor,\n n);\n}\n\n} // namespace sglang\n```\n\n**Key points:**\n\n- Include headers from `sgl_kernel/` — **not** raw CUDA headers for anything already covered\n- Use `TensorMatcher` for all tensor validation; never manually check shape/dtype/device\n- Use `AlignedVector` for vectorised 128-bit loads/stores — significant bandwidth win\n- Use `LaunchKernel` — it resolves the stream and checks errors automatically\n- Use `CHECK_HOST(cond) << ...` for runtime assertions with useful error messages (zero overhead when the check passes)\n- Prefer passing runtime scalars like `factor` directly unless compile-time specialisation is genuinely required\n- `fp16_t` / `bf16_t` / `fp32_t` are the project's type aliases (from `utils.cuh`)\n- `device::cast` or `DTypeTrait::from(val)` for cross-type conversions\n- `device::math::` functions for device math instead of bare `__` intrinsics if possible.\n- Consider PDL — it can help when the kernel has prologue work to overlap. Place `PDLWaitPrimary()` right before the first read of upstream data, not at the top of the kernel\n\n---\n\n## Step 2: Add the Python wrapper in `kernels/ops/`\n\nThe wrapper lives next to its functional group under `python/sglang/kernels/ops/`, not beside the CUDA source — `kernels/jit/` holds only the JIT infrastructure (`csrc/`, `include/`, `utils/`, `benchmark/`). Create `python/sglang/kernels/ops/elementwise/scale.py`:\n\n```python\nfrom __future__ import annotations\n\nfrom typing import TYPE_CHECKING\n\nimport torch\n\nfrom sglang.kernels.jit.utils import (\n cache_once,\n is_arch_support_pdl,\n load_jit,\n make_cpp_args,\n)\n\nif TYPE_CHECKING:\n from tvm_ffi.module import Module\n\n\n@cache_once\ndef _jit_scale_module(dtype: torch.dtype) -> Module:\n \"\"\"Compile and cache the JIT scale module for a given dtype.\"\"\"\n # Checks on the compile key live here, not in `scale`: `cache_once` keys on\n # `dtype`, so this runs once per specialisation instead of once per call.\n if dtype not in (torch.float16, torch.bfloat16, torch.float32):\n raise RuntimeError(\n f\"Unsupported dtype {dtype}. Supported: float16, bfloat16, float32\"\n )\n args = make_cpp_args(dtype, is_arch_support_pdl())\n return load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n )\n\n\ndef scale(src: torch.Tensor, factor: float, out: torch.Tensor | None = None) -> torch.Tensor:\n \"\"\"\n Element-wise scale: dst = src * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n src : CUDA tensor (FP16 / BF16 / FP32)\n factor : scale factor\n out : optional pre-allocated output tensor (same shape/dtype as src)\n\n Returns\n -------\n Scaled tensor (dst = src * factor).\n \"\"\"\n # DO NOT add proactive validation here: every check costs interpreter time\n # on a per-forward path. Tensor invariants belong in the C++ launcher, and\n # anything about the compile key belongs in `_jit_scale_module`.\n if out is None:\n out = torch.empty_like(src)\n\n module = _jit_scale_module(src.dtype)\n module.scale(out, src, factor)\n return out\n```\n\n**Key points:**\n\n- Use `cache_once` — **not** `functools.lru_cache` (incompatible with `torch.compile`)\n- `load_jit` first arg(s) form the unique build marker; same marker = same cached binary\n- Only include compile-time specialisation knobs in the build marker; runtime values like `factor` should stay runtime unless the kernel truly needs templating\n- `cuda_wrappers`: `(export_name, kernel_symbol)` — `export_name` is called from Python\n- `make_cpp_args(dtype, ...)` converts `torch.dtype` to C++ type alias:\n- `is_arch_support_pdl()` checks if the current architecture supports PDL, which is typically passed as a template argument to the kernel.\n- Keep the entry point thin (see Conventions). What Python must still check goes in the `@cache_once` module factory, not in the entry point: `cache_once` keys on its arguments, so a check there costs one evaluation per specialisation instead of one per call — that is where the supported-dtype guard lives. Tensor invariants belong in the C++ launcher; if something here never reaches a `.verify(...)`, close that gap on the C++ side rather than in Python\n\n| `torch.dtype` | C++ type |\n|--------------------|------------|\n| `torch.float16` | `fp16_t` |\n| `torch.bfloat16` | `bf16_t` |\n| `torch.float32` | `fp32_t` |\n\n---\n\n## Register the public entry point (required)\n\nAdd metadata in `python/sglang/kernels/ops/elementwise/__init__.py` so the\noperator appears in the registry without importing its implementation:\n\n```python\nfrom sglang.kernels.registry import register_kernel\nfrom sglang.kernels.spec import CapabilityRequirement, KernelBackend, KernelSpec\n\nregister_kernel(\n KernelSpec(\n op=\"elementwise.scale\",\n backend=KernelBackend.JIT,\n target=\"sglang.kernels.ops.elementwise.scale:scale\",\n capabilities=frozenset({CapabilityRequirement.CUDA}),\n )\n)\n```\n\nRecord the devices the implementation actually supports; the CUDA-only scale\nexample above does not imply that every JIT kernel is CUDA-only. Runtime and\ncorrectness tests import `scale` from `sglang.kernels.ops.elementwise.scale`.\nDo not import the implementation eagerly in the group initializer. When an\nexisting `BaseFusedOp` owns the operation, add the implementation there instead\nof registering a second conflicting op/backend pair.\n\n## Step 3 (optional): Tune JIT build flags\n\nIf your kernel uses some math functions like `expf` or `sinf`, consider enabling `--use_fast_math` for better performance (with a potential precision tradeoff):\n\n```python\nreturn load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n)\n```\n\nIf your kernel requires SM90+, raise a clear Python error before calling `load_jit`. Arch gating is one of the checks that has to live in Python — it decides whether to compile at all, so the C++ launcher never gets to run:\n\n```python\nif torch.cuda.get_device_capability()[0] < 9:\n raise RuntimeError(\"This kernel requires SM90 (Hopper) or later\")\n```\n\n---\n\n## Step 4: Write tests (required)\n\nJIT kernel correctness tests and benchmarks live under `test/registered/kernels/ops//` and `test/registered/kernels/benchmark//`, mirroring the wrapper's group under `python/sglang/kernels/ops/` (NOT inside the `sglang` package -- a `register_*_ci(...)` call anywhere under `python/sglang/` is rejected by the `check-no-registered-tests-in-package` pre-commit hook). Only their test-only helpers (e.g. `benchmark/marker.py`) stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path. **CI does not run `pytest` in those directories directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` and `bench_*.py` under `test/registered/`, collects `register_*_ci(...)` calls by **statically parsing each file's AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check.\n\n- **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`.\n- **Nightly kernel suite**: register with `stage=\"nightly\"` plus the `runner_config` of the machine it needs (e.g. `1-gpu-large`), giving the `nightly-test-1-gpu-large` suite. `.github/workflows/nightly-test-nvidia.yml` sets `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` for the whole nightly run, so the expanded parameter grids apply automatically (see `python/sglang/kernels/jit/utils/common.py` → `should_run_full_tests` / `get_ci_test_range`). There is no separate kernel-only nightly job: every nightly test on one machine type shares that machine's suite.\n\nRegistration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing):\n\n```python\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n# Optional B200/SM100 registration for tests that cover Blackwell-specific code paths\n# register_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"4-gpu-b200\")\n# Optional second registration: same file also runs nightly, same form,\n# stage is just \"nightly\" there (and no `nightly=True`)\n# register_cuda_ci(est_time=120, stage=\"nightly\", runner_config=\"1-gpu-large\")\n```\n\nCI generates the suite name as `{stage}-test-{runner_config}`, so `stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\"` becomes the `base-b-kernel-unit-test-1-gpu-large` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage=\"nightly\"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools.\n\nKeep `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.\n\nUse `register_cuda_ci(..., disabled=\"reason\")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only).\n\n**Run like CI** (from repo root):\n\n```bash\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large)\n# For B200/SM100-specific coverage:\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200)\n```\n\nFor fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`.\n\nCreate `test/registered/kernels/ops/elementwise/test_scale.py`:\n\n```python\nimport pytest\nimport torch\nfrom sglang.kernels.ops.elementwise.scale import scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [1, 127, 128, 1024, 4097]) # cover tail remainder\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0, 3.0])\ndef test_scale_correctness(dtype, size, factor):\n src = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = scale(src, factor)\n expected = src * factor\n\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\ndef test_scale_out_param(dtype):\n src = torch.randn(1024, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(src)\n result = scale(src, 2.0, out=out)\n assert result is out\n torch.testing.assert_close(out, src * 2.0, rtol=1e-2, atol=1e-2)\n\n\ndef test_scale_cpu_error():\n src = torch.randn(128, dtype=torch.float16) # CPU tensor\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n scale(src, 2.0)\n\n\ndef test_scale_unsupported_dtype():\n src = torch.randint(0, 10, (128,), dtype=torch.int32, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"dtype\"):\n scale(src, 2.0)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-v\", \"-s\"]))\n```\n\n---\n\n## Step 5: Add a benchmark (required)\n\nBenchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large`).\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nBenchmarks use the project's own `marker` framework (in `python/sglang/kernels/jit/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`):\n\n- **`@marker.benchmark(line_arg, line_vals, *, unit=\"us\")`** — the **innermost** decorator (bottom of the stack, directly above `def benchmark`). Declares the column axis: each value in `line_vals` becomes a result column, and `line_arg` is the parameter name passed into the benchmark function. `unit` is one of `\"us\" | \"ms\" | \"s\"`.\n- **`@marker.parametrize(names, vals, ci_vals=None)`** — stackable decorator that adds a row axis (pytest-style). Each `@parametrize` adds one (or more, correlated) parameter the benchmark is swept over (Cartesian product across all `parametrize` decorators). `names` may be a single name (`\"size\"`) or a comma-separated correlated tuple axis (`\"h,d\"`, with `vals` then a list of tuples like `[(1, 64), (2, 128)]`). Pass the optional third `ci_vals` for a smaller sweep that is auto-selected under `is_in_ci()` — this is the built-in CI-shrinking mechanism, so you usually don't need `get_benchmark_range` for swept axes.\n- **`marker.do_bench(fn, *, input_args=(), input_kwargs={}, ...)`** — runs `fn` under CUDA graph (default) or a naive loop, returns a `BenchResult`. Key knobs:\n - `memory_args`: defaults to `\"all\"` (footprint derived from all input args/kwargs). Pass an explicit tuple of tensors (e.g. `(k, v, indices)`) to count only the inputs the kernel actually touches.\n - `memory_output`: defaults to `\"out\"` — re-runs `fn` once to capture its **returned** tensor and counts it. For in-place kernels (which return `None`), pass the written tensors explicitly (e.g. `memory_output=(k, v)`); the re-run is then skipped. Set to `None` to count no output.\n - Together `memory_args` + `memory_output` give the GB/s column; with both defaults a function `out = f(src)` already reports `bytes(src) + bytes(out)`.\n - `graph_clone_args` / `graph_clone_kwargs`: which inputs to clone per CUDA-graph iteration to defeat L2 cache reuse. Defaults to `\"all\"`; if you narrow it, keep every read, in-place, and output tensor.\n - `use_cuda_graph=False` for kernels that can't be captured.\n - `metrics=(0.5, \"avg\")` controls reported quantiles (the first metric becomes the table latency column).\n - `disable_log_bandwidth` (defaults from `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1`) skips the bandwidth column entirely.\n- **`utils.create_random(*shape)` / `utils.create_empty(*shape)`** — shorthand for `torch.randn` / `torch.empty` with `DEFAULT_DTYPE` (`bfloat16`) and `DEFAULT_DEVICE` (`\"cuda\"`). Override via the `dtype=` / `device=` kwargs.\n- **`utils.get_benchmark_range(full_range, ci_range)`** — returns the smaller `ci_range` under CI (`is_in_ci()`), the `full_range` locally. Still available for the `benchmark(...)` column axis (which has no `ci_vals`); for `parametrize` row axes prefer the built-in `ci_vals` argument.\n\nCreate `test/registered/kernels/benchmark/elementwise/bench_scale.py`:\n\n```python\nimport torch\n\nfrom sglang.kernels.jit.benchmark import marker\nfrom sglang.kernels.jit.benchmark.utils import create_random\nfrom sglang.kernels.ops.elementwise.scale import scale as jit_scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=6, stage=\"base-b-kernel-benchmark\", runner_config=\"1-gpu-large\")\n\n\n@torch.compile()\ndef torch_impl_scale(src: torch.Tensor, factor: float) -> torch.Tensor:\n return src * factor\n\n\nFN_MAP = {\n \"jit\": jit_scale,\n \"torch\": torch_impl_scale,\n}\n\n\n# `parametrize(name, full_vals, ci_vals)`: the 3rd arg is the smaller sweep\n# auto-selected under CI; the full range runs locally.\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096, 65536]) # 1K .. 512K\n@marker.benchmark(\"impl\", [\"jit\", \"torch\"])\ndef benchmark(size: int, impl: str):\n src = create_random(size)\n factor = 2.0\n return marker.do_bench(\n FN_MAP[impl],\n input_args=(src, factor),\n # `src` is read -> clone it per iter to avoid L2 reuse; factor is a scalar.\n graph_clone_args=(0,),\n # Defaults already report bandwidth: memory_args=\"all\" counts src,\n # memory_output=\"out\" counts the returned tensor -> bytes(src)+bytes(out).\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n**Key points:**\n\n- The `line_arg` name passed to `benchmark` (`\"impl\"` here) must match a parameter on `benchmark(...)`; same for every `parametrize` name (`\"size\"`).\n- Stack `@parametrize` once per swept axis. The required `@marker.benchmark` is the **innermost** decorator (bottom of the stack, directly above the function) — `@parametrize` rows go above it.\n- Prefer `create_random` / `create_empty` from `utils.py` over open-coding `torch.randn(..., dtype=..., device=...)`.\n- The GB/s column appears by default (`memory_args=\"all\"` + `memory_output=\"out\"`). For memory-bound kernels it's the most informative number; scope `memory_args` / `memory_output` to the tensors actually touched if the defaults over- or under-count. For compute-bound kernels where bandwidth is misleading, set `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1` (or `disable_log_bandwidth=True`).\n- For in-place kernels (which return `None`), pass the written tensors via `memory_output=(...)` since the `\"out\"` default would capture nothing.\n- Tune `graph_clone_args` / `graph_clone_kwargs` to all the arguments that might be read by the kernel. We can only skip cloning for write-only args. For in-place modified args, we still need to clone them to get accurate timing (reusing the same buffer keeps it L2-hot and skews results).\n- Call `benchmark.run()` (no `print_data=` kwarg — the marker framework prints directly).\n\nRun locally:\n\n```bash\npython test/registered/kernels/benchmark/elementwise/bench_scale.py\n```\n\nRun the benchmark suite the way CI does:\n\n```bash\ncd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large\n```\n\n---\n\n## Troubleshooting\n\n- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time`, `stage`, and `runner_config`; starred args and non-literal values break AST collection\n- **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/kernels/jit/csrc/`; reduce template argument combinations\n- **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...`\n- **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. `graph_clone_args` defaults to `\"all\"`; if you narrow it, it must still cover every *read* tensor — reusing a single buffer keeps it L2-hot and skews results. Keep *write* tensors in it too: they are what sets the rotation count, and a shared output buffer stays L2-hot the same way.\n- **Missing GB/s column**: the column is on by default; check that `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH` is not `1` and `disable_log_bandwidth` is not `True`. For in-place kernels (return `None`) the `memory_output=\"out\"` default counts nothing — pass the written tensors via `memory_output=(...)`\n\n---\n\n## References\n\n- `docs/docs/developer_guide/development_jit_kernel_guide.mdx`\n- `test/run_suite.py` — suite names, discovery of `test/registered/`, execution entrypoint for CI\n- `python/sglang/test/ci/ci_register.py` — `register_cuda_ci` and AST registration rules\n- `python/sglang/kernels/jit/utils/compile/` — `load_jit` (`loader.py`), `make_cpp_args` (`cpp_args.py`)\n- `python/sglang/kernels/jit/utils/common.py` — `cache_once`, `should_run_full_tests`, `get_ci_test_range`\n- `python/sglang/kernels/jit/include/sgl_kernel/tensor.h` — `TensorMatcher`, `SymbolicSize/DType/Device`, `is_type`\n- `python/sglang/kernels/jit/include/sgl_kernel/ffi.h` — `ffi::empty`, `ffi::empty_like`, `ffi::from_blob`\n- `python/sglang/kernels/jit/include/sgl_kernel/utils.cuh` — type aliases, `LaunchKernel`, `SGL_DEVICE`\n- `python/sglang/kernels/jit/include/sgl_kernel/vec.cuh` — `AlignedVector`\n- `python/sglang/kerne","frontmatter":{"name":"add-jit-kernel","description":"Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups"},"isInternal":false,"tokens":13181,"sizeBytes":50279},{"name":"SKILL.md","path":".agents/skills/add-sgl-kernel/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/add-sgl-kernel/SKILL.md","title":"add-sgl-kernel","category":"anthropic-skill","format":"markdown","content":"---\nname: add-sgl-kernel\ndescription: Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)\n---\n\n# Tutorial: Adding a New Kernel to `sgl-kernel` (AOT / Heavyweight)\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as an AOT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float)\n- Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `python/sglang/kernels/aot/include/utils.h`)\n\n## Two rules of thumb (must follow)\n\n1. **Prefer `python/sglang/kernels/jit` first** when the kernel does **not** depend on CUTLASS or another large C++ project. This is the default path for lightweight kernels that benefit from rapid iteration.\n2. **Prefer `sgl-kernel`** when the kernel **does** depend on CUTLASS or another large C++ project, or when it should be part of the AOT wheel / torch op registration flow.\n3. **Exception**: if the dependency is `flashinfer`, or CUTLASS that is already provided through `flashinfer`, the kernel can still be implemented as `jit_kernel`.\n\nIn addition, every new kernel must ship with:\n\n- **Tests** (pytest)\n- **A benchmark script** (`marker.do_bench` from `sglang.kernels.jit.benchmark` — it works for any callable, not only JIT kernels)\n\n---\n\n## Repository integration map\n\nYou will typically touch these files/areas:\n\n- Implementation: `python/sglang/kernels/aot/csrc/elementwise/scale.cu` (pick the right subdirectory)\n- Public declarations: `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- Torch extension registration: `python/sglang/kernels/aot/csrc/common_extension.cc`\n- Build: `python/sglang/kernels/aot/CMakeLists.txt` (`set(SOURCES ...)`)\n- Python API: `python/sglang/kernels/aot/python/sgl_kernel/` and `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n- Tests: `python/sglang/kernels/aot/tests/test_scale.py`\n- Benchmarks: `python/sglang/kernels/aot/benchmark/bench_scale.py`\n\n---\n\n## Step 1: Implement the kernel in `csrc/`\n\nPick the right subdirectory:\n\n- `csrc/elementwise/` — for element-wise ops (our example)\n- `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories\n\nCreate `python/sglang/kernels/aot/csrc/elementwise/scale.cu`:\n\n```cpp\n#include \n#include \n#include \n\n#include \"utils.h\" // DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16\n\n// scale_kernel: out[i] = input[i] * factor\n// Supports float, half (__half), __nv_bfloat16 via template T\ntemplate \n__global__ void scale_kernel(T* __restrict__ out,\n const T* __restrict__ input,\n float factor,\n int64_t n) {\n int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x;\n if (idx < n) {\n out[idx] = static_cast(static_cast(input[idx]) * factor);\n }\n}\n\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor) {\n TORCH_CHECK(input.is_cuda(), \"input must be a CUDA tensor\");\n TORCH_CHECK(input.is_contiguous(), \"input must be contiguous\");\n TORCH_CHECK(out.is_cuda(), \"out must be a CUDA tensor\");\n TORCH_CHECK(out.is_contiguous(), \"out must be contiguous\");\n TORCH_CHECK(out.sizes() == input.sizes(), \"out and input must have the same shape\");\n TORCH_CHECK(out.scalar_type() == input.scalar_type(),\n \"out and input must have the same dtype\");\n\n const int64_t n = input.numel();\n const int threads = 256;\n const int blocks = (n + threads - 1) / threads;\n\n const cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n const at::cuda::OptionalCUDAGuard device_guard(device_of(input));\n\n // Dispatches over float, float16, bfloat16\n DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {\n scale_kernel<<>>(\n static_cast(out.data_ptr()),\n static_cast(input.data_ptr()),\n static_cast(factor),\n n);\n cudaError_t status = cudaGetLastError();\n TORCH_CHECK(status == cudaSuccess,\n \"scale_kernel launch failed: \", cudaGetErrorString(status));\n return true;\n });\n}\n```\n\n**Key points:**\n\n- Use `at::Tensor` (PyTorch tensors), `TORCH_CHECK` for validation, `at::cuda::getCurrentCUDAStream()` for stream\n- Keep Python wrappers thin; do shape/dtype/device validation in C++ right around the launch path\n- `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` covers `float`, `half` (FP16), `__nv_bfloat16` (BF16)\n- Add device error checking after every kernel launch\n- If a kernel only works on certain architectures, enforce that with `TORCH_CHECK` and skip logic in tests\n\n---\n\n## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`\n\nEdit `python/sglang/kernels/aot/include/sgl_kernel_ops.h`, add to the elementwise section:\n\n```cpp\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor);\n```\n\n---\n\n## Step 3: Register the op in `csrc/common_extension.cc`\n\nEdit `python/sglang/kernels/aot/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:\n\n```cpp\n// From csrc/elementwise\nm.def(\"scale(Tensor! out, Tensor input, float factor) -> ()\");\nm.impl(\"scale\", torch::kCUDA, &scale);\n```\n\n**Key points:**\n\n- `Tensor!` means in-place / mutable output argument\n- The schema is important for `torch.compile` and for consistent call signatures\n- Keep the torch schema in PyTorch scalar types (`float` here), but note that the C++ launcher signature still needs `double` for scalar arguments accepted by `torch::Library`\n\n---\n\n## Step 4: Add the new source file to `CMakeLists.txt`\n\nEdit `python/sglang/kernels/aot/CMakeLists.txt`, add to `set(SOURCES ...)`:\n\n```cmake\ncsrc/elementwise/scale.cu\n```\n\n**Key points:**\n\n- Keep the list **alphabetically sorted** (the file explicitly requires this)\n- If the kernel has arch constraints, reflect that in tests/benchmarks via skip logic\n\n---\n\n## Step 5: Expose a Python API under `python/sglang/kernels/aot/python/sgl_kernel/`\n\nPrefer following the existing module organization first. For elementwise kernels, the usual pattern is:\n\n- implement the Python wrapper in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`\n- then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n\nFor example, in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`, add:\n\n```python\nimport torch\n\ndef scale(\n input: torch.Tensor,\n factor: float,\n out: torch.Tensor | None = None,\n) -> torch.Tensor:\n \"\"\"\n Element-wise scale: out = input * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n input : CUDA input tensor\n factor : scale factor (float)\n out : optional pre-allocated CUDA output tensor (same shape/dtype as input)\n \"\"\"\n if out is None:\n out = torch.empty_like(input)\n torch.ops.sgl_kernel.scale.default(out, input, factor)\n return out\n```\n\nThen re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.\n\n---\n\n## SGLang integration (required for runtime use)\n\nAfter exposing the AOT wheel symbol, add a lazy wrapper and a `KernelSpec` under\n`python/sglang/kernels/ops//`. Set `backend=KernelBackend.AOT`, record\nactual device/architecture support, and keep `sgl_kernel` imports inside the\nimplementation path. SGLang runtime and integration tests import that wrapper.\nFor this example, use `sglang.kernels.ops.elementwise.scale`.\n\nThe wheel-level tests below validate its standalone API/build. They do not\nreplace CI-registered SGLang correctness tests in\n`test/registered/kernels/ops/elementwise/test_scale.py` and benchmarks in\n`test/registered/kernels/benchmark/elementwise/bench_scale.py`. Follow\n`write-sglang-test` for their registration and CI budget; do not register tests\nunder the shipped `python/sglang/` package.\n\n## Step 6: Write tests (required)\n\nCreate `python/sglang/kernels/aot/tests/test_scale.py`:\n```python\nimport pytest\n\nimport torch\nimport sgl_kernel\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [128, 1024, 4096, 65536])\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0])\ndef test_scale_correctness(dtype, size, factor):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n\n result = sgl_kernel.scale(input, factor, out=out)\n assert result is out\n\n expected = input * factor\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\ndef test_scale_shape_mismatch():\n input = torch.randn(128, dtype=torch.float16, device=\"cuda\")\n out = torch.empty(256, dtype=torch.float16, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"same shape\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\ndef test_scale_cpu_input():\n input = torch.randn(128, dtype=torch.float16) # CPU\n out = torch.empty_like(input)\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-q\"]))\n```\n\n---\n\n## Step 7: Add a benchmark (required)\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nCreate `python/sglang/kernels/aot/benchmark/bench_scale.py`:\n\n```python\nimport torch\n\nimport sgl_kernel\nfrom sglang.kernels.jit.benchmark import marker\n\n\ndef sglang_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n sgl_kernel.scale(input, factor, out=out)\n\n\ndef torch_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n torch.mul(input, factor, out=out)\n\n\nFN_MAP = {\"sglang\": sglang_scale, \"torch\": torch_scale}\n\n\n@marker.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32], [torch.float16])\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096]) # 1K .. 512K\n@marker.benchmark(\"provider\", [\"sglang\", \"torch\"])\ndef benchmark(dtype: torch.dtype, size: int, provider: str):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n return marker.do_bench(\n FN_MAP[provider],\n # Pass every tensor through input_args (not a closure) so marker rotates\n # them across CUDA-graph calls to defeat L2 reuse.\n input_args=(input, 2.0, out),\n # Bandwidth = bytes(input) + bytes(out), both already in input_args.\n memory_output=None,\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n---\n\n## Step 8: Build\n\nBuild:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j16\n```\n\nIf you need to limit host resource usage:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j1 MAX_JOBS=2 CMAKE_ARGS=\"-DSGL_KERNEL_COMPILE_THREADS=1\"\n```\n\n---\n\n## Step 9: Validate\n\nAfter building successfully, run the test and benchmark:\n\n```bash\npytest python/sglang/kernels/aot/tests/test_scale.py -q\npython python/sglang/kernels/aot/benchmark/bench_scale.py\n```\n\nPR CI also runs `pr-test-sgl-kernel.yml`, including the B200 job\n`sgl-kernel-b200-test` when kernel changes are detected. Use that job as the\nBlackwell coverage signal for AOT `sgl-kernel` changes.\n\n---\n\n## Troubleshooting\n\n- **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`\n- **Memory errors**: `compute-sanitizer --tool memcheck python ...`\n- **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`\n- **Binary bloat**: use `python/sglang/kernels/aot/analyze_whl_kernel_sizes.py`\n- **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time\n\n---\n\n## References\n\n- `python/sglang/kernels/aot/README.md`\n- `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- `python/sglang/kernels/aot/csrc/common_extension.cc`\n- `python/sglang/kernels/aot/CMakeLists.txt`\n- `python/sglang/kernels/aot/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends\n- `python/sglang/kernels/aot/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern\n\n## Summary of Files Created/Modified\n\n```\npython/sglang/kernels/aot/csrc/elementwise/scale.cu # NEW: CUDA kernel + launcher\npython/sglang/kernels/aot/include/sgl_kernel_ops.h # MODIFIED: C++ declaration\npython/sglang/kernels/aot/csrc/common_extension.cc # MODIFIED: schema + dispatch registration\npython/sglang/kernels/aot/CMakeLists.txt # MODIFIED: add source file (alphabetical)\npython/sglang/kernels/aot/python/sgl_kernel/elementwise.py # MODIFIED: Python wrapper\npython/sglang/kernels/aot/python/sgl_kernel/__init__.py # MODIFIED: re-export Python API\npython/sglang/kernels/aot/tests/test_scale.py # NEW: tests\npython/sglang/kernels/aot/benchmark/bench_scale.py # NEW: benchmark\n```\n","frontmatter":{"name":"add-sgl-kernel","description":"Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)"},"isInternal":false,"tokens":3565,"sizeBytes":13519},{"name":"openai.yaml","path":".agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml","title":"Agents Skill","category":"anthropic-skill","format":"yaml","content":"# Agents Skill\n\nPath: `.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml)","tokens":300,"sizeBytes":224},{"name":"SKILL.md","path":".agents/skills/babysit-pr-to-pass-ci/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/SKILL.md","title":"babysit-pr-to-pass-ci","category":"anthropic-skill","format":"markdown","content":"---\nname: babysit-pr-to-pass-ci\ndescription: Start and persistently pursue a goal to babysit an SGLang pull request until selected GitHub Actions workflows pass on the latest PR head. Use when asked to monitor, babysit, retry, or fix PR CI for lint.yml, pr-test.yml, pr-test-extra.yml, AMD, or other named workflows; classify failures as PR-related versus flaky or infrastructural, auto-fix and push only small clean fixes, rerun failed jobs only up to 10 times, and ignore unselected workflows.\n---\n\n# Babysit a PR Until Selected CI Passes\n\nDrive only the selected workflows to green on the PR's latest head commit. Keep working across turns by starting a durable goal; do not stop merely because a run was dispatched, a retry began, or a fix was pushed.\n\n## Invocation and arguments\n\nAccept:\n\n```text\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] [WORKFLOW ...]\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] --only WORKFLOW [WORKFLOW ...]\n```\n\n- If the PR is omitted, resolve the PR associated with the current branch.\n- If workflows are omitted, select `lint.yml` and `pr-test.yml`.\n- Without `--only`, add named workflows to the two defaults.\n- With `--only`, monitor exactly the named workflows.\n- Accept a workflow basename such as `pr-test-extra.yml`, a repository path such as `.github/workflows/pr-test-extra.yml`, or an unambiguous workflow display name.\n- Treat `extra` as `pr-test-extra.yml`.\n- Accept natural-language equivalents, such as “PR 12345 plus extra and AMD.” Resolve “AMD” to an exact workflow file from the user's wording or repository context; do not guess when multiple AMD workflows fit.\n\nExamples:\n\n```text\n$babysit-pr-to-pass-ci\n$babysit-pr-to-pass-ci 12345\n$babysit-pr-to-pass-ci https://github.com/sgl-project/sglang/pull/12345 pr-test-extra.yml\n$babysit-pr-to-pass-ci 12345 --only pr-test-amd.yml\n```\n\n## Start or continue the durable goal\n\nTreat explicit invocation of this skill as authorization to create a goal and perform the scoped CI actions below.\n\n1. Resolve the PR and exact workflow set first.\n2. Inspect the current goal.\n3. If no unfinished goal exists, create one immediately using the product's goal mechanism. Include the resolved PR URL, selected workflow files, and this stopping condition: every selected workflow has a successful run for the PR's latest head SHA.\n4. Include these constraints in the goal objective:\n - Ignore every unselected workflow.\n - Diagnose each selected-workflow failure as PR-related, unrelated/infrastructure, or still uncertain.\n - Implement, validate, commit, and push a PR-related fix only when it is clean, non-tricky, and at most 100 changed lines.\n - Stop for user review before editing when the proposed fix is tricky, dirty, or over 100 changed lines.\n - Retry unrelated failures at most 10 times per workflow and head SHA, using failed-jobs-only reruns.\n - Never rerun a full workflow merely to recover failed tests.\n5. If the active goal already represents this same PR and workflow set, continue it instead of creating another goal.\n6. If a different unfinished goal is active, do not replace it silently. Report the conflict and ask the user to pause or clear it.\n\nDo not mark the goal complete until the completion check at the end of this skill succeeds.\n\n## Preflight\n\n1. Verify GitHub authentication and resolve the repository without changing remote configuration.\n2. Resolve the PR and record at least:\n - PR number and URL\n - open/draft state\n - base branch and base SHA\n - head branch and head SHA\n - head repository owner/name, including whether it is a fork\n - labels\n - whether maintainers may modify the branch\n3. Require an open PR. If it is draft and a selected workflow is blocked by the draft gate, stop and ask the user to mark it ready; do not change draft state automatically.\n4. Normalize every selected workflow to an existing file under `.github/workflows/`. Fail fast on ambiguous or nonexistent names.\n5. Read each selected workflow's triggers and gates. Do not assume all optional workflows behave like `pr-test.yml`.\n6. Inspect `git status`. Preserve all user changes. If a fix becomes necessary and the current checkout is dirty, on a different branch, or otherwise unsafe, use an isolated worktree for the PR head.\n\nKnown SGLang gates:\n\n- `pr-test.yml` requires the `run-ci` label for relevant PR changes. If its existing gate job failed because the label was absent, add `run-ci` when authorized and rerun failed jobs in that existing run; adding the label alone does not necessarily create a new `pr-test.yml` run.\n- `pr-test-extra.yml` requires both `run-ci` and `run-ci-extra`. It listens for relevant label events, so adding a missing label may create a fresh run. Prefer that fresh current-head run; rerun failed jobs in an existing current run only when needed.\n- AMD-extra workflows may use the same `run-ci` plus `run-ci-extra` gate. Verify the actual selected file.\n- Treat draft, cooldown, maintenance, and missing-label failures as gate/infrastructure conditions, not code regressions. Fix only the safe gate condition in scope; wait for cooldown or maintenance windows rather than burning retries immediately.\n\nAdding the known CI opt-in labels is within scope. Do not add unrelated labels, mark the PR ready, alter reviewers, merge, or mutate other PR metadata.\n\n## Track only current-head runs\n\nUse both PR checks and workflow-run data so workflow names, run IDs, attempts, jobs, and logs can be correlated. Prefer workflow file IDs over display names when filtering runs.\n\nFor each selected workflow:\n\n1. Select the newest run associated with the current PR head SHA.\n2. Ignore runs for older head SHAs, including green runs made stale by a new push.\n3. Treat `queued`, `requested`, `waiting`, `pending`, and `in_progress` as active; keep monitoring.\n4. Treat a workflow as passing only when its current-head run concludes `success`.\n5. If no current-head run exists, wait briefly for event delivery, then inspect the workflow's path filters, event triggers, and label gates.\n6. Trigger only through a mechanism explicitly supported by that workflow. Never invent `workflow_dispatch` inputs or dispatch a workflow against the default branch when the intent is to test the PR head.\n7. If a selected workflow cannot safely be triggered for the PR head, report the exact blocker and request direction.\n\nAfter any user push or skill-authored push, re-read the PR head SHA, discard stale run selections, and start tracking the new head. Reset retry counters for the new head SHA.\n\n## Monitor persistently\n\n- Maintain a compact ledger containing workflow file, head SHA, run ID, attempt, status, conclusion, skill-initiated retry count, and latest failure signature.\n- Poll at reasonable intervals, normally 30–60 seconds while runs are active. Give concise progress updates during long waits.\n- Re-check the PR head SHA on each monitoring cycle and before every rerun or push.\n- Ignore failures, cancellations, and pending checks from unselected workflows. Do not diagnose, retry, fix, or wait for them.\n- Do not declare success while a selected workflow is absent, stale, pending, skipped, cancelled, neutral, or failed.\n\n## Diagnose a selected-workflow failure\n\n1. Inspect the failed jobs and failed-step logs. In `pr-test.yml`, distinguish the root failure from cascade failures caused by `check-pr-test-health`, wait jobs, or final aggregation.\n2. Record the exact failing test/check, error text, affected runner or hardware, and relevant log URL.\n3. Compare the failure against the PR's changed files and diff.\n4. Classify it:\n\n **PR-related** when evidence ties the failure to changed code, tests, dependencies, configuration, generated files, or behavior introduced by the PR. A reproducible failure on the PR that does not occur on the base is strong evidence.\n\n **Unrelated/infrastructure** when evidence points to a transient network or download error, runner loss, GPU/driver instability, service outage, pre-existing main-branch failure, known flaky threshold, unrelated test area, or another environmental cause.\n\n **Uncertain** when evidence is insufficient. Do not label a failure flaky merely because it might be intermittent. Narrowly reproduce it, compare recent main/scheduled CI, inspect relevant code, or otherwise gather enough evidence to choose a category.\n\n5. When several root failures exist, classify each one. Address PR-related failures before spending retries on unrelated cascades.\n\n## Handle a PR-related failure\n\nEstimate the proposed fix before editing.\n\nA fix is eligible for automatic execution only when all are true:\n\n- The fix is well understood, localized, and clean.\n- The fix is not tricky, hacky, speculative, or a broad workaround.\n- The fix changes at most 100 lines, counting additions plus deletions in the fix patch, not the pre-existing PR diff.\n- The fix does not require a broad API or architecture redesign, risky dependency migration, large generated-file rewrite, or unrelated cleanup.\n- A focused validation is available.\n\nIf any condition fails, stop before editing. Report the diagnosis, proposed approach, expected files, estimated size, risks, and validation plan, then wait for user review.\n\nFor an eligible fix:\n\n1. Work from the latest PR head in a clean checkout or isolated worktree.\n2. Make only the necessary change and add or update focused tests when appropriate.\n3. Run the smallest meaningful local validation. For lint failures, reproduce the specific hook/check before considering broader lint validation.\n4. Review the resulting patch and count changed lines. If it becomes over 100 lines or turns tricky/dirty, do not commit or push it; stop for review and clearly identify the uncommitted skill-authored changes.\n5. Commit only skill-authored files. Never use `git add .`, rewrite user commits, amend without request, or include pre-existing local changes.\n6. Push a new commit non-forcibly to the exact PR head repository and branch. Never force-push. If permission or fork routing is unclear, stop and report the push command/target that needs authorization.\n7. Resolve the new PR head SHA and return to current-head monitoring. A push is progress, not completion.\n\n## Handle an unrelated or infrastructure failure\n\nRetry only the failed jobs in the current selected-workflow run:\n\n```bash\ngh run rerun --failed --repo \n```\n\nRules:\n\n- Never use a full-run rerun as a substitute for `--failed`.\n- Never rerun a stale run after the PR head SHA changes or a newer current-head run exists.\n- Count only successfully dispatched skill-initiated reruns.\n- Allow at most 10 such reruns per selected workflow and head SHA.\n- Wait for each new attempt to finish, inspect its result, and update the failure signature before deciding on another retry.\n- If the failure changes or becomes plausibly PR-related, return to diagnosis instead of blindly consuming all retries.\n- If 10 retries are exhausted, stop. Report the attempts, recurring signatures, job/run links, and evidence that the failure is unrelated or still uncertain. Do not exceed the cap without a new explicit user instruction.\n\nIf GitHub cannot rerun only the failed portion of the selected run, stop and ask the user rather than rerunning the full workflow.\n\n## Completion check\n\nBefore completing the goal:\n\n1. Fetch the PR head SHA again.\n2. For every selected workflow, verify that its newest applicable run is for that exact SHA and concludes `success`.\n3. Verify no selected workflow is still pending and no success was invalidated by a newer push.\n4. Ignore all unselected workflows even if they are failing.\n5. Mark the goal complete only after all selected workflows satisfy these checks.\n\nReport the final PR URL, head SHA, selected workflows, passing run links, any commits pushed, and retry counts. Keep the report scoped to the selected workflows.\n","frontmatter":{"name":"babysit-pr-to-pass-ci","description":"Start and persistently pursue a goal to babysit an SGLang pull request until selected GitHub Actions workflows pass on the latest PR head. Use when asked to monitor, babysit, retry, or fix PR CI for lint.yml, pr-test.yml, pr-test-extra.yml, AMD, or other named workflows; classify failures as PR-related versus flaky or infrastructural, auto-fix and push only small clean fixes, rerun failed jobs only up to 10 times, and ignore unselected workflows."},"isInternal":false,"tokens":2557,"sizeBytes":11878},{"name":"action-items.md","path":".agents/skills/ci-test-audit/action-items.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/action-items.md","title":"Ci-test-audit Skill","category":"anthropic-skill","format":"markdown","content":"# Ci-test-audit Skill\n\nPath: `.agents/skills/ci-test-audit/action-items.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/action-items.md)","tokens":300,"sizeBytes":209},{"name":"SKILL.md","path":".agents/skills/ci-test-audit/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/SKILL.md","title":"ci-test-audit","category":"anthropic-skill","format":"markdown","content":"---\nname: ci-test-audit\ndescription: Audit the existing test tree and CI configuration for improvements, using a catalog of patterns previously applied in this repo. Use when asked to audit tests or CI, shrink CI time or cost, find redundant or misplaced tests, clean up a test group, or review whether a CI change follows established practice.\n---\n\n# CI / Test Audit\n\nAudits what is already in `test/`, `.github/` and `scripts/ci/`. For writing a new test see\n[write-sglang-test](../write-sglang-test/SKILL.md); for how the pipeline dispatches and\ngates work see [ci-workflow-guide](../ci-workflow-guide/SKILL.md).\n\n## How to use it\n\nRead [action-items.md](action-items.md) first. It is a catalog of patterns that have\nbeen applied to this repo, each with a way to spot it and example PRs. The catalog is the substance of this skill; everything below is only scaffolding.\n\nThe user says what to audit -- a test group, a stage, a workflow file, or a complaint\nlike \"this suite is slow\" -- and how thoroughly. Read that scope and decide which\npatterns apply.\n\nAn audit is only useful if it is specific. \"This suite could be trimmed\" is not a\nfinding; \"`TestFooLargePage` is a strict subset of `TestFooRetractLargePage`, drop it\"\nis. Every finding needs the file, the pattern it matches, and what to do.\n\n## Reporting\n\nGroup findings by confidence in the claim, not by pattern id. Lead with the ones where\nthe evidence is in the file you just read; keep the speculative ones separate and say\nwhat you would need to check. Skip anything you are not reasonably sure about -- a long\nlist of maybes costs more to triage than it saves.\n\nPropose, do not apply. Deleting a test, moving a registration to another stage, and\nchanging a threshold are all decisions for the user. Land them as separate changes,\nsince a trim and a threshold change fail for different reasons.\n\n## Keeping the catalog current\n\nExample PRs age. When a pattern shows up in newer work, or a new pattern recurs, update\n`action-items.md`: the pattern text should stay general, the examples are\nreplaceable evidence.\n","frontmatter":{"name":"ci-test-audit","description":"Audit the existing test tree and CI configuration for improvements, using a catalog of patterns previously applied in this repo. Use when asked to audit tests or CI, shrink CI time or cost, find redundant or misplaced tests, clean up a test group, or review whether a CI change follows established practice."},"isInternal":false,"tokens":487,"sizeBytes":2079},{"name":"SKILL.md","path":".agents/skills/ci-workflow-guide/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-workflow-guide/SKILL.md","title":"ci-workflow-guide","category":"anthropic-skill","format":"markdown","content":"---\nname: ci-workflow-guide\ndescription: Guide to SGLang CI workflow orchestration — stage ordering, fail-fast, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages.\n---\n\n# SGLang CI Workflow Orchestration Guide\n\nThis skill covers the CI **infrastructure** layer — how tests are dispatched, gated, and aborted on failure across stages. For test authoring (templates, fixtures, registration, model selection), see the [write-sglang-test skill](../write-sglang-test/SKILL.md).\n\n---\n\n## Naming Conventions\n\n- **Suite**: `base-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `base-b-test-1-gpu-small`)\n- **Test group**: Directory-level registered test group under `test/registered/` (e.g., `hicache` maps to `test/registered/hicache/test_*.py`)\n- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)\n\n---\n\n## Key Files\n\n| File | Role |\n|------|------|\n| `.github/workflows/pr-test.yml` | Main workflow — all stages, jobs, conditions, matrix definitions |\n| `.github/workflows/pr-test-extra.yml` | Extra workflow — gated by BOTH `run-ci` and `run-ci-extra` labels |\n| `.github/workflows/pr-gate.yml` | PR gating: draft check, `run-ci` label, per-user rate limiting |\n| `.github/actions/check-pr-test-health/action.yml` | Cross-job fail-fast: queries API for any failed job |\n| `.github/actions/wait-for-jobs/action.yml` | Stage gating: polls API until stage jobs complete |\n| `.github/actions/check-maintenance/action.yml` | Maintenance mode check |\n| `.github/scripts/ci-labels.cjs` | Resolves the four CI control labels into dispatch axes |\n| `test/run_suite.py` | Suite runner: collects, filters, partitions, executes tests |\n| `python/sglang/test/ci/ci_register.py` | Test registration (AST-parsed markers), LPT auto-partition |\n| `python/sglang/test/ci/ci_utils.py` | `run_unittest_files()`: execution, retry, continue-on-error |\n| `scripts/ci/utils/slash_command_handler.py` | Handles slash commands from PR comments |\n\n---\n\n## Architecture Overview\n\n```\n ┌──────────────┐\n │ build kernel │\n └──────┬───────┘\n │\n ├─ check-changes ──── detects which packages changed\n │ (main_package, sgl_kernel, jit_kernel, multimodal_gen)\n │\n ├─ call-gate ──────── pr-gate.yml (draft? label? rate limit?)\n │\n ├─────────────────────────────────────────────────────┐\n │ │\n ▼ │\n ┌─────────────────────────────────────┐ │\n │ Base A (~3 min) │ │\n │ pre-flight check │ │\n │ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-1-gpu-small │ │ │\n │ │ (small GPUs) │ │ │\n │ └─────────────────────────────┘ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-cpu │ │ │\n │ │ (CPU) │ │ │\n │ └─────────────────────────────┘ │ │\n └──────┬──────────────────────────────┘ │\n │ │\n ▼ ▼\n ┌─────────────────────────────────────┐ ┌──────────────────────────┐\n │ Base B (~30 min) │ │ kernel test │\n │ base tests │ └──────────────────────────┘\n │ │ ┌──────────────────────────┐\n │ ┌─────────────────────────────┐ │ │ multimodal gen test │\n │ │ base-b-test-1-gpu-small │ │ └──────────────────────────┘\n │ │ (small GPUs, e.g. 5090) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-1-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-2-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ Base C (~30 min) │\n │ advanced tests │\n │ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-h100 │ │\n │ │ (H100 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-8-gpu-h200 │ │\n │ │ (8 x H200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-b200 │ │\n │ │ (4 x B200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ Other advanced tests │ │\n │ │ (DeepEP, PD Disagg, GB300) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ pr-test-finish │\n │ aggregates all results, fails if │\n │ any job failed/cancelled │\n └─────────────────────────────────────┘\n```\n\n**Every stage test job** includes a `check-pr-test-health` step after checkout — if any job in the run has already failed, the job fails fast (red X) with a root cause annotation.\n\n**Scheduled runs** skip `wait-for-base-*` jobs, running all stages in parallel. Fail-fast is also disabled.\n\n---\n\n## Fail-Fast Layers\n\n4 layers of fail-fast, from fine to coarse:\n\n| Layer | Mechanism | Granularity | Disabled on schedule? |\n|-------|-----------|-------------|----------------------|\n| **1. Test method → file** | `unittest -f` (failfast) | One test method fails → entire test file stops immediately | Yes |\n| **2. File → suite** | `run_unittest_files()` default | One test file fails → entire suite stops (`--continue-on-error` off) | Yes |\n| **3. Job → job (same stage)** | `check-pr-test-health` action | One job fails → other waiting jobs in same stage fail-fast (red X) | Yes |\n| **4. Stage → stage (cross-stage)** | `wait-for-base-*` + `needs` | Base A fails → base B/C jobs skip entirely (never get a runner) | Yes (wait jobs skipped) |\n\n- **Layer 1**: `-f` flag appended to all `python3 -m pytest` / `unittest` invocations in `ci_utils.py`\n- **Layer 2**: `--continue-on-error` flag in `run_suite.py` — off for PRs, on for scheduled runs\n- **Layer 3**: `check-pr-test-health` auto-detects `schedule` event and skips; filters out cascade failures to show only root cause jobs\n- **Layer 4**: `wait-for-base-*` jobs are conditioned on `github.event_name == 'pull_request'` — skipped for scheduled runs\n\n---\n\n## Execution Modes\n\n| Aspect | PR (`pull_request`) | Scheduled (`cron`, every 6h) | Manual dispatch (`workflow_dispatch`) |\n|--------|---------------------|------------------------------|--------------------------------------|\n| **Stage ordering** | Sequential: A → B → C via `wait-for-base-*` | Parallel (all at once) | Single target stage only |\n| **Cross-job fail-fast** | Yes (`check-pr-test-health`) | Yes | Yes |\n| **continue-on-error** | No (stop at first failure within suite) | Yes (run all tests) | No |\n| **Retry** | Enabled | Enabled | Enabled |\n| **max_parallel** | 3 (default), 14 if `max-concurrency` label | 14 | 3 (default), 14 if `max-concurrency` |\n| **PR gate** | Yes (draft, label, rate limit) | Skipped | Skipped |\n| **Concurrency** | `cancel-in-progress: true` per PR | Queue (no cancel) | Isolated per stage+SHA |\n\nFour labels relax these limits for one PR: `bypass-fail-fast`, `parallel-stages`,\n`max-concurrency`, and `highest-priority` (all three). `.github/scripts/ci-labels.cjs`\nresolves them; the [contribution guide](https://docs.sglang.io/developer_guide/contribution_guide.html#ci-control-labels)\ndescribes what each one does.\n\n---\n\n## Stage Gating (`wait-for-jobs` action)\n\n`wait-for-base-a` and `wait-for-base-b` are lightweight `ubuntu-latest` jobs that poll the GitHub Actions API.\n\n**How it works:**\n1. Calls `listJobsForWorkflowRun` to list all jobs in the current run\n2. Matches jobs by exact name or prefix (for matrix jobs, e.g., `base-b-test-1-gpu-small (3)`)\n3. If any matched job has `conclusion === 'failure'` → fail immediately (fail-fast)\n4. If all matched jobs are completed and count matches `expected_count` → success\n5. Otherwise → sleep `poll-interval-seconds` (default: 60s) and retry\n6. Timeout after `max-wait-minutes` of wall-clock time (240 min for base-a, 350 min for base-b). Keep budgets under the 6h hard cap on GitHub-hosted runners: a wait killed by the cap ends `cancelled`, not `timeout`, and `timeout-minutes` cannot raise it. The base-b wait sets `fail-on-timeout: 'false'`, so a timeout lets base-c start; `pr-test-finish` fails the run on any base-b failure regardless\n\n**Job specs example** (base-b):\n```json\n[\n {\"prefix\": \"base-b-test-1-gpu-small\", \"expected_count\": 8},\n {\"prefix\": \"base-b-test-1-gpu-large\", \"expected_count\": 14},\n {\"prefix\": \"base-b-test-2-gpu-large\", \"expected_count\": 4},\n {\"prefix\": \"base-b-test-4-gpu-b200\", \"expected_count\": 1}\n]\n```\n\n> **Critical**: `expected_count` must match the matrix size. If you add/remove matrix entries, update the wait job's spec accordingly.\n\n**PR only**: Condition `github.event_name == 'pull_request' && !inputs.target_stage` — scheduled runs and manual dispatches skip these entirely, allowing parallel execution.\n\n---\n\n## Cross-Job Fail-Fast (`check-pr-test-health` action)\n\nComposite action called after checkout in every stage test job (21 jobs total across `pr-test.yml`, `pr-test-multimodal-gen.yml`, `pr-test-sgl-kernel.yml`, `pr-test-jit-kernel.yml`).\n\n**How it works:**\n1. Queries `listJobsForWorkflowRun` for the current workflow run\n2. Filters for **root cause failures only** — jobs with `conclusion === 'failure'` whose failing step is NOT `check-pr-test-health` (excludes cascade failures)\n3. If root cause failures found → calls `core.setFailed()` with the list of root cause job names\n4. If none → does nothing (step succeeds)\n\n**Cascade filtering**: When job A fails fast due to the health check, it also has `conclusion: failure`. Without filtering, job B would list both the original failure AND job A's fail-fast. The filter checks each failed job's `steps` array — if the failing step name contains `check-pr-test-health` or `Check PR test health`, it's excluded from the root cause list.\n\n**Usage pattern:**\n```yaml\nsteps:\n - name: Checkout code\n uses: actions/checkout@v4\n ...\n\n - uses: ./.github/actions/check-pr-test-health\n id: pr-test-health\n\n - name: Install dependencies # skipped automatically if health check failed\n ... # (default if: success() is false)\n\n - name: Run test # also skipped\n ...\n```\n\n**Visual effect**: Job shows **red X** (failure) with error annotation showing root cause job names. Subsequent steps are naturally skipped (default `if: success()` is false after a failed step). No per-step `if` guards needed.\n\n**No stage filtering**: Checks ALL jobs in the run, not just the current stage. Any failure anywhere triggers fail-fast.\n\n**Error message example:**\n```\nFail-fast: skipping — root cause job(s): base-b-test-1-gpu-small (0), base-b-test-1-gpu-small (1)\n```\n\n---\n\n## Within-Suite Failure Handling\n\nControlled by `run_unittest_files()` in `python/sglang/test/ci/ci_utils.py`.\n\n### Flags\n\n| Flag | PR default | Scheduled default | Effect |\n|------|------------|-------------------|--------|\n| `--continue-on-error` | Off | On | Off: stop at first failure. On: run all files, report all failures at end |\n| `--enable-retry` | On | On | Retry retriable failures (accuracy/perf assertions) |\n| `--max-attempts` | 2 | 2 | Max attempts per file including initial run |\n\n### Retry Classification\n\nWhen a test fails and retry is enabled, the output is classified:\n\n**Non-retriable** (checked first — real code errors):\n`SyntaxError`, `ImportError`, `ModuleNotFoundError`, `NameError`, `TypeError`, `AttributeError`, `RuntimeError`, `CUDA out of memory`, `OOM`, `Segmentation fault`, `core dumped`, `ConnectionRefusedError`, `FileNotFoundError`\n\n**Retriable** (accuracy/performance):\n`AssertionError` with comparison patterns (`not greater than`, `not less than`, `not equal to`), `accuracy`, `score`, `latency`, `throughput`, `timeout`\n\n**Default**: Unknown `AssertionError` → retriable. Other unknown failures → not retriable.\n\n### How `continue_on_error` is set\n\nIn `pr-test.yml`'s `check-changes` job:\n- `schedule` runs or `run_all_tests` flag → `continue_on_error = 'true'`\n- PR runs → `continue_on_error = 'false'`\n\nEach test job propagates via:\n```yaml\nenv:\n CONTINUE_ON_ERROR_FLAG: ${{ needs.check-changes.outputs.continue_on_error == 'true' && '--continue-on-error' || '' }}\nrun: |\n python3 run_suite.py --hw cuda --suite $CONTINUE_ON_ERROR_FLAG\n```\n\n---\n\n## Test Partitioning\n\nLarge suites are split across matrix jobs using the **LPT (Longest Processing Time) heuristic** in `ci_register.py:auto_partition()`:\n\n1. Sort tests by `est_time` descending, filename as tie-breaker (deterministic)\n2. Greedily assign each test to the partition with smallest cumulative time\n3. Result: roughly equal total time per partition\n\n**Partition table** (CUDA per-commit suites):\n\n| Suite | Partitions | Runner | max_parallel |\n|-------|-----------|--------|-------------|\n| `base-a-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — |\n| `base-a-test-cpu` | 4 | `ubuntu-latest` | — |\n| `base-b-test-1-gpu-small` | 8 | `1-gpu-5090` | 8 |\n| `base-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) |\n| `base-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — |\n| `base-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-b-kernel-unit-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |\n| `base-b-kernel-benchmark-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — |\n| `base-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — |\n| `base-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — |\n| `base-c-test-deepep-4-gpu-h100` | 1 (no matrix) | `4-gpu-h100` | — |\n| `base-c-test-4-gpu-b200` | 3 | `4-gpu-b200` | — |\n| `base-c-test-4-gpu-b200-small` | 3 | `4-gpu-b200-low-disk` | — |\n| `base-c-test-8-gpu-b200` | registered only | `8-gpu-b200` | — |\n| `base-c-test-4-gpu-gb200` | registered only | `4-gpu-gb200` | — |\n\n> **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`.\n>\n> **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`.\n\n**Workflow usage:**\n```yaml\nstrategy:\n matrix:\n partition: [0, 1, 2, 3, 4, 5, 6, 7]\nsteps:\n - run: python3 run_suite.py --hw cuda --suite base-b-test-1-gpu-small \\\n --auto-partition-id ${{ matrix.partition }} --auto-partition-size 8\n```\n\n---\n\n## check-changes Job\n\nDetermines which test suites to run based on file changes.\n\n### Detection Methods\n\n| Trigger | Method | Details |\n|---------|--------|---------|\n| `pull_request` | `dorny/paths-filter` | Detects changes via GitHub diff |\n| `workflow_dispatch` (with `pr_head_sha`) | GitHub API | `repos/{repo}/compare/main...{sha}` |\n| `schedule` / `run_all_tests` | Force all true | Runs everything |\n\n### Output Flags\n\n| Output | Triggers |\n|--------|----------|\n| `main_package` | Base A/B/C test suites |\n| `sgl_kernel` | Kernel wheel builds + kernel test suites; also switches B200 jobs to kernel-build runner labels outside `target_stage` mode |\n| `jit_kernel` | JIT kernel test workflow |\n| `multimodal_gen` | Multimodal-gen test workflow |\n\n> **Note**: In `target_stage` mode, `sgl_kernel` is only active when `include_wheel_build=true`. Without that opt-in, kernel-change reruns fail validation instead of running a target stage without freshly built wheels. Outside `target_stage`, `sgl_kernel=true` switches B200 jobs from `4-gpu-b200` / `4-gpu-b200-low-disk` to `4-gpu-b200-kernel` / `4-gpu-b200-kernel-low-disk`.\n\n---\n\n## Concurrency Control\n\n```\ngroup: pr-test-{event_name}-{branch}-{pr_sha}-{stage}\n```\n\n| Segment | Source | Purpose |\n|---------|--------|---------|\n| `event_name` | `github.event_name` | Prevents scheduled runs colliding with fork PRs named `main` |\n| `branch` | `github.head_ref \\|\\| github.ref_name` | Per-branch isolation |\n| `pr_sha` | `inputs.pr_head_sha \\|\\| 'current'` | Isolates manual dispatches from main runs |\n| `stage` | `inputs.target_stage \\|\\| 'all'` | Allows parallel stage dispatches |\n\n`cancel-in-progress: true` for `pull_request` events (new push cancels old run), `false` for `workflow_call`.\n\n---\n\n## How To: Add a New Stage Job\n\n1. Define the job in `pr-test.yml` with `needs: [check-changes, call-gate, wait-for-base-X, ...]`\n2. Copy the `if:` condition pattern from an existing same-stage job (handles `target_stage`, `schedule`, `main_package`)\n3. Add `checkout` step\n4. Add `check-pr-test-health` step (after checkout) — if any prior job failed, `core.setFailed()` fires and all subsequent steps auto-skip via default `if: success()`\n5. Add `check-maintenance` step\n6. Add `download-artifact` step if `sgl_kernel` changed\n7. Add `install dependencies` step\n8. Add `run test` step with `$CONTINUE_ON_ERROR_FLAG`\n9. Add `upload-cuda-coredumps` step with `if: always()`\n10. Register the suite name in `PER_COMMIT_SUITES` in `test/run_suite.py`\n11. If using matrix, add `--auto-partition-id` and `--auto-partition-size` to the run command\n12. **Update `wait-for-base-X`** job spec with the new job name and `expected_count` (if matrix)\n13. **Add the job to `pr-test-finish.needs`** list\n\n---\n\n## How To: Debug CI Failures\n\n| Symptom | Likely cause | What to check |\n|---------|-------------|---------------|\n| All stage-B/C jobs green but steps skipped | Earlier job failed, `check-pr-test-health` triggered | Find the actual failed job (red X) |\n| `wait-for-base-b` timeout | `expected_count` doesn't match matrix size | Verify job spec counts match `matrix:` array length |\n| `pr-test-finish` fails but all jobs green | A job was `cancelled` (counts as failure in finish) | Check concurrency cancellation, and whether a `wait-for-base-*` job hit the 6h hosted-runner cap |\n| Tests pass locally but fail in CI | Partition assignment, runner GPU type, or `est_time` inaccuracy | Check which partition the test lands in; verify runner label |\n| Flaky test retried and passed | Retriable failure (accuracy/perf) | Check `[CI Retry]` markers in job logs |\n| Flaky test NOT retried | Matched non-retriable pattern | Check if error matches `NON_RETRIABLE_PATTERNS` in `ci_utils.py` |\n\n---\n\n## Slash Commands\n\n| Command | Effect |\n|---------|--------|\n| `/tag-run-ci-label` | Adds `run-ci` label to PR |\n| `/tag-run-ci-label extra` | Adds both `run-ci` and `run-ci-extra` labels |\n| `/rerun-failed-ci` | Reruns failed jobs in the latest workflow run |\n| `/tag-and-rerun-ci` | Adds `run-ci` label + reruns failed |\n| `/tag-and-rerun-ci extra` | Adds both `run-ci` and `run-ci-extra` labels + reruns failed |\n| `/run-full-ci` | Short form of `/tag-and-rerun-ci extra` (baseline + extra). Alias: `/rerun-full-ci` |\n| `/run-extra-ci` | Adds both labels + reruns **only** `PR Test Extra`, leaving baseline runs alone. Alias: `/rerun-extra-ci` |\n| `/rerun-test [ ...]` | Reruns specific test file(s) via `rerun-test.yml`. A file arg containing a glob metacharacter (`*`, `?`, `[...]`) expands against `test/registered/` and the multimodal test dir to every matching `test_*.py` (e.g. `/rerun-test test_*backend*.py` — wrap in backticks so GitHub doesn't italicize the `*`); matches are deduped, grouped by dispatch shape, and can't carry a `::test` selector. No match → single ⛔ reply, nothing dispatched. |\n| `/rerun-group [ ...]` | Expands registered test groups, then reuses `/rerun-test` |\n| `/run-cpu-ci` | Dispatches `run-cpu-ci.yml` on main: only the `base-a-test-cpu` suite, at the PR head SHA, no caches saved. Same gate as `/rerun-test`. Alias: `/rerun-cpu-ci` |\n\nHandled by `scripts/ci/utils/slash_command_handler.py` → `.github/workflows/slash-command-handler.yml`.\n\n### Label-gated workflow dispatch (pr-test, pr-test-extra)\n\n`pr-test.yml` and `pr-test-extra.yml` both listen for `pull_request.labeled` (in addition to `opened`/`synchronize`/`reopened`). The `check-changes.if` gate has two clauses:\n\n1. **For `labeled` events**: the just-added label must be one of the gating labels (`run-ci` for pr-test, `run-ci` or `run-ci-extra` for pr-test-extra) — otherwise every unrelated label addition would dispatch a full CI run.\n2. **All events**: the PR must currently carry the required labels.\n\nThis is what lets `/tag-run-ci-label` (and the `extra` variant) trigger a fresh CI run without an extra push.\n\n**Caveat — skipped runs cannot be un-skipped by `run.rerun()`:** GitHub's rerun API reuses the original event payload, so rerunning a `pull_request`-event run that was skipped because of missing labels will skip again (label set in the frozen payload doesn't update). The only way to recover a label-skipped run is to add the missing label, which fires a fresh `labeled` event with the current label set. `handle_rerun_failed_ci` in the slash handler is for rerunning failed/non-label-skipped runs; it cannot revive label-skipped ones.\n","frontmatter":{"name":"ci-workflow-guide","description":"Guide to SGLang CI workflow orchestration — stage ordering, fail-fast, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages."},"isInternal":false,"tokens":6103,"sizeBytes":25002},{"name":"noise-sources.md","path":".agents/skills/clean-startup-log/references/noise-sources.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/references/noise-sources.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/clean-startup-log/references/noise-sources.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/references/noise-sources.md)","tokens":300,"sizeBytes":238},{"name":"SKILL.md","path":".agents/skills/clean-startup-log/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/SKILL.md","title":"clean-startup-log","category":"anthropic-skill","format":"markdown","content":"---\nname: clean-startup-log\ndescription: Audit SGLang startup logs, save evidence, and propose cleanup for user review. With no arguments, run Qwen3-8B at TP1 and TP2 plus gpt-oss-20b at TP1.\ndisable-model-invocation: true\n---\n\n# Audit SGLang Startup Logs\n\nThe default outcome is saved logs and a findings report. Apply runtime changes\nonly after the user selects them. A request to edit this skill does not itself\nlaunch servers.\n\n## Default runs\n\nA bare `$clean-startup-log` invocation runs these cases sequentially without\nasking for commands. Explicit commands, models, or TP sizes replace this matrix.\n\n| Case / log filename | Command |\n|---|---|\n| `qwen3-8b-tp1.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B` |\n| `qwen3-8b-tp2.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B --tp 2` |\n| `gpt-oss-20b-tp1.log` | `uv run sglang serve --model-path openai/gpt-oss-20b` |\n\nThese cover dense, tensor-parallel, and MoE/hybrid sliding-window attention\nstartup. Reuse complete captures from the current audit when code and environment\nhave not changed.\n\n## Capture logs\n\n1. Check the checkout, free GPUs, and ports once. Use the requested command when\n resources are free; otherwise select free GPUs with `CUDA_VISIBLE_DEVICES` and\n an unused `--port`, recording the adjustments. Leave existing servers alone.\n2. Create a unique directory with `mktemp -d /tmp/sglang-startup-audit-XXXXXX`.\n Save raw stdout and stderr together in a separate log for each case, for\n example with `set -o pipefail` and `COMMAND 2>&1 | tee LOG_PATH`. Record commands,\n GPU IDs, ports, commit, relevant overrides, and readiness status.\n3. Wait for `The server is fired up and ready to roll!`, then stop that server and\n its workers before the next case. First runs can spend many minutes downloading\n weights or compiling FlashInfer kernels; check download/compiler activity\n before treating a quiet log as a hang. Preserve partial logs for failed or\n stalled starts and report the last stage. Continue independent cases when possible.\n4. Preserve the user's logging configuration, including `NCCL_DEBUG`. If NCCL\n verbosity needs explaining, inspect relevant shell settings, `NCCL_CONF_FILE`,\n and `/etc/nccl.conf`. Do not override intentional diagnostics or recommend\n `NCCL_DEBUG=WARN` solely because the output is long. Avoid full environment dumps.\n\n## Investigate efficiently\n\n- Scan for deprecations, duplicate handler output, unrelated import failures,\n unformatted prints, and unexpected warnings. Read representative excerpts and\n counts instead of repeatedly dumping `server_args`, progress redraws, or NCCL\n diagnostics. Normalize carriage returns for analysis only; preserve raw logs.\n- Trace each candidate to its actual emitter with focused `rg` searches. Inspect\n its log level: SGLang's formatter may omit severity. Group shared signatures\n across cases and distinguish handler duplication from separate GPU/process calls.\n- Repetition, WARNING severity, or a different third-party format alone does not\n establish a cleanup need. Consider whether the message explains configuration,\n progress, resource use, or an operational limitation.\n- Consult [noise-source hints](references/noise-sources.md) only for a matching\n signature or an unresolved emitter. Verify current code rather than trusting\n historical line numbers, fix status, or assumptions about unrelated models.\n\n## Accepted output\n\nPreserve these reviewed messages unless the user requests a different policy:\n\n- NCCL diagnostics enabled by the user's environment or host configuration.\n- NUMA permission warnings, including one check per GPU in TP runs.\n- GPT-OSS MXFP4 backend-selection warnings and default page-size selection warnings.\n- `Init Unified Radix Cache. Components: ... Tree Core: ...`, tree-cache summaries,\n SWA allocation details, and per-rank memory/timing records.\n- Useful progress bars, warmup HTTP access logs, uv synchronization messages,\n isolated NCCL/Gloo startup lines, and one timestamped HF authentication warning.\n\nThese can appear in a clean startup log. Do not repeatedly propose the declined\nNUMA deduplication, backend/page-size level changes, or NCCL verbosity override.\nKeep real operational warnings visible: for example, a Harmony vocabulary failure\ncan disable `/v1/responses` even when server readiness and `/generate` succeed.\n\n## Report before changing code\n\nReturn a compact run table with readiness status and clickable raw-log links.\nFor each actual cleanup candidate, give an exact representative message, affected\ncases/counts, source file/function, and specific proposed behavior. Distinguish\nconfirmed findings from suspicions and operational failures from logging noise.\n\nIf there are no actionable cleanup findings, say the logs are clean and no\nfurther cleanup is needed. Otherwise, ask which numbered changes to adopt and\nwait for the user's selections before editing runtime code or preparing patches.\nHonor existing approvals and declined items without asking again.\n\n## Apply selected changes\n\n- Batch compatible approved edits, then verify affected cases once. Repeat a\n startup only for a new change, failure, or unresolved concern; do not relaunch\n after every one-line edit. Save verification logs separately from baselines.\n- Preserve useful warnings and application handlers. HF can warn during early\n CLI model detection before `configure_logger()`, and spawned processes have\n independent logger state. Keep `configure_hf_hub_logger()` in both\n `suppress_noisy_warnings()` and `configure_logger()`; make repeated setup safe.\n- Keep the legacy compiled-kernel cache migration notice at DEBUG. Avoid broad\n library-level suppression or fd redirection for a narrow logging problem.\n- Run relevant formatting and focused existing checks. Add tests only when they\n verify meaningful behavior, not a log-level spelling. Report changes and\n verification; create branches, commits, and PRs when requested.\n","frontmatter":{"name":"clean-startup-log","description":"Audit SGLang startup logs, save evidence, and propose cleanup for user review. With no arguments, run Qwen3-8B at TP1 and TP2 plus gpt-oss-20b at TP1.","disable-model-invocation":true,"disableModelInvocation":true},"isInternal":false,"tokens":1299,"sizeBytes":5978},{"name":"SKILL.md","path":".agents/skills/compute-mamba-ratio/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/compute-mamba-ratio/SKILL.md","title":"compute-mamba-ratio","category":"anthropic-skill","format":"markdown","content":"---\nname: compute-mamba-ratio\ndescription: Compute the optimal --mamba-full-memory-ratio (or --max-mamba-cache-size pin) for a hybrid attention + linear-attention (Mamba / GDN / KDA) model's two serving memory pools, from the workload and serving config. Use when a user asks what ratio to set, why concurrency is clamped, or how to size the state vs KV pools for a hybrid model.\n---\n\n# Optimal hybrid dual-pool ratio (`--mamba-full-memory-ratio`)\n\nA hybrid model (attention layers + linear-attention layers — the recurrent-state family: Mamba/SSM, GDN, KDA, etc.) splits serving memory into two independently-budgeted pools, fixed once at startup:\n\n- **state pool** (the linear-attention recurrent state) → caps **concurrency** (hard: whole slots, worst-case reserved, fail-loud)\n- **full-KV pool** (attention KV) → caps **context × concurrency** (soft: paged, over-committable via retraction)\n\n`--mamba-full-memory-ratio r` splits the post-weight budget: `mamba_budget = rest · r/(1+r)`, i.e. `mamba_budget : kv_budget = r`. This skill picks the `r` (or the pin-`--max-mamba-cache-size` alternative) at which **neither pool bottlenecks first** for the user's workload.\n\n## The formula\n\n```\nr* = (S + D) · token_equiv · dcp_size / L\ntoken_equiv = state_bytes_per_slot / kv_bytes_per_token\n```\n\n- `L` = average context length per request (input + output tokens)\n- `token_equiv` = full-KV token-equivalent of one state slot\n- `S` = state slots per running request (cache-strategy dependent, table below)\n- `D` = `--speculative-num-draft-tokens` (0 if NOSPEC); each running req carries `D` extra intermediate states\n- `dcp_size` = `--dcp-size` (1 without DCP). DCP shards the per-rank KV by `dcp_size`, so KV gets ~`dcp_size×` cheaper per request → the balance shifts that much toward the state pool.\n\n`r` is dimensionless (just the split). To also predict the actual concurrency you need `rest` (below).\n\n## Inputs to collect from the user\n\n1. **`L`** — average context (input + output) in tokens.\n2. **The two per-GPU byte constants** — one of:\n - **(a) measured (preferred, exact)** — from one boot log at *any* ratio:\n - `Mamba Cache is allocated. ... ssm_state size X GB` with `max_mamba_cache_size: N` → `state_bytes_per_slot = X / N`\n - `KV Cache is allocated. #tokens: M, KV size: Y GB` → `kv_bytes_per_token = Y / M`\n - **(b) derived** — model arch (linear-layer count + state dims `d_state/d_conv/heads/head_dim`; attention type + KV dims: MLA latent dim, or GQA `kv_heads·head_dim·layers`) **× the dtypes** below.\n3. **`S`** — from `--mamba-radix-cache-strategy`, the overlap scheduler, and `SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK` (table below).\n4. **`D`** — `--speculative-num-draft-tokens` (0 if NOSPEC; **also 0 when ReplaySSM spec-verify is enabled** — see caveats).\n5. **`dcp_size`** — `--dcp-size` (1 if no DCP). If DCP **and** spec with a replicated draft KV, apply the draft caveat (below).\n6. **KV dtype** (bf16 / fp8) and **ssm dtype** (fp32 / bf16) — they set the two byte constants (see the `token_equiv` 2×2).\n7. *(only to also predict the clamp, not just `r`)* **`rest`** = per-GPU memory − weights at the chosen `--mem-fraction-static`. Read `avail mem` after `Load weight end`, or `Memory pool end. avail mem` + pool sizes, from the boot log.\n\n### `S` (state slots per running request), set by `--mamba-radix-cache-strategy`\n\nAll strategies keep prefix caching on; they differ in the track buffer that snapshots chunk-boundary state under the overlap scheduler. `S = base + ping-pong`, where base is 3 (live state + radix retention/COW headroom), and ping-pong = 2 (overlap on, non-lazy) / 1 (lazy, **or** overlap off) / 0 (no track buffer at all). This mirrors `kv_cache_configurator._calculate_mamba_ratio`; read it there if a release moves the constants.\n\n`SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK=1` takes the base to 2, since the decode-time skip frees one resident slot per running request. `no_buffer` is the exception and stays at an effective 3: it adds that slot straight back, because its binding limit is the prefill-to-decode peak, which a decode-time change does not shrink.\n\n| strategy | S | S with decode-lock skip | notes |\n|---|---|---|---|\n| `--disable-radix-cache` | 1 | 1 | prefix cache off entirely; most concurrency, no reuse |\n| `no_buffer` | 3 | 3 | no track buffer; requires overlap **off** (+ `page_size 1`) → lower decode throughput |\n| `extra_buffer` (default), overlap on | 5 | 4 | track buffer reserved per running request |\n| `extra_buffer`, overlap off | 4 | 3 | same buffer, but ping-pong costs 1 slot instead of 2 |\n| `extra_buffer_lazy` (recommended) | 4 | 3 | track buffer allocated lazily at the boundary → 1 fewer slot/req; requires overlap **on** |\n\nThe skip is a real concurrency lever, not a rounding detail: at the default strategy it takes `S` from 5 to 4, so the same state budget admits 25% more requests and `r*` drops by the same factor.\n\nOverlap counts as off when `--disable-overlap-schedule` is passed **or** `--pp-size > 1` (pipeline parallelism turns it off for you). So a PP deployment left on the default strategy sits on the `S = 4` row, not `S = 5` — using 5 there overstates the state cost and starves KV.\n\n### `token_equiv` — compute it from the two byte constants\n\n`token_equiv = state_bytes_per_slot / kv_bytes_per_token`. Both bytes are **per-model, per-GPU** — read them from a boot log (or derive from the arch + dtypes) and divide. That's the whole thing; the dtype effects below are just for a quick mental estimate.\n\nThe recurrent state has two parts: `state_bytes = SSM_bytes + conv_bytes`. The **SSM** part scales with `--mamba-ssm-dtype`; the **conv** part is always bf16 (`SGLANG_MAMBA_CONV_DTYPE`, fixed). How the dtype knobs move `token_equiv`:\n\n- **fp8 KV → ×2** (universal, exact): fp8 is 1 byte vs bf16's 2, so `kv_bytes_per_token` halves.\n- **fp32 ssm vs bf16 ssm → ×≈2**: only the SSM tensor doubles, conv stays fixed, so the exact factor is `(2·SSM + conv)/(SSM + conv)` — just under 2. **conv is usually small relative to SSM, so estimate ≈2** (e.g. one measured model: conv ≈7% of state → factor ≈1.9). Only bother with the exact split if a model's conv is unusually large.\n- **fp32 ssm + fp8 KV → ×≈4**.\n\nFor an actual number, don't apply factors — read `state_bytes_per_slot` and `kv_bytes_per_token` for **your** model+dtype from the boot log and divide.\n\nThe `GB` the boot log prints for both pools is **GiB** (`1024³`). Convert both the same way and `token_equiv` is unaffected either way, since it is a ratio — but get it right before comparing an absolute byte figure against your own log.\n\nWorked example (one measured model, TP8, fp32 ssm + fp8 KV), read off one boot log:\n\n```\nMamba Cache is allocated. max_mamba_cache_size: 257, conv_state size: 0.46GB, ssm_state size: 13.04GB\nKV Cache is allocated. dtype: torch.float8_e4m3fn, #tokens: 1167552, KV size: 15.03 GB\n```\n\n`state_bytes_per_slot = (0.46 + 13.04) GiB / 257 ≈ 56.4 MB`, `kv_bytes_per_token = 15.03 GiB / 1167552 ≈ 13.8 KB` → `token_equiv ≈ **4080**`.\n\n## Compute\n\n```python\ndef optimal_ratio(L, state_bytes_per_slot, kv_bytes_per_token, S, D=0, dcp_size=1):\n token_equiv = state_bytes_per_slot / kv_bytes_per_token\n r = (S + D) * token_equiv * dcp_size / L\n return r # value of --mamba-full-memory-ratio (>1 is legal)\n\ndef predict_clamp(rest_bytes, r, state_bytes_per_slot, S, D=0):\n mamba_budget = rest_bytes * r / (1 + r)\n slots = mamba_budget / state_bytes_per_slot\n # spec: each running req reserves (S+D) worth; non-spec just S\n return int(slots // S) # NOSPEC; with spec the budget joint-solves for (S+D)·per_req per req\n```\n\nThen state the result three ways: the **`r` value**, the **predicted clamp** (if `rest` given), and **which pool binds** (`min(mamba_clamp, KV_cap)`, where `KV_cap = kv_tokens · dcp_size / L`).\n\n## Procedure\n\n1. **Free lever first**: if the boot log shows large idle in `Memory pool end. avail mem` (e.g. 30–40 GB at mem-frac 0.85), raise `--mem-fraction-static` (→0.92) before touching the split — it grows `rest` for both pools at no cost. Validate graph-capture headroom once.\n2. Collect the inputs. Prefer a real boot log for the two byte constants.\n3. Compute `r*`. If `r*` would drive the state pool below one request's worth (`mamba_budget < S · per_req`, happens at very long `L`), **switch to pinning `--max-mamba-cache-size = target_concurrency · S`** and let the rest go to KV — a sub-0.15 `r` is fragile.\n4. Report `r`, predicted clamp, binding pool, and any dtype accuracy gate that applies.\n\n## Worked examples (TP8, B300, validated against measured clamps)\n\nConstants (from boot logs): `state_bytes_per_slot ≈ 56.4 MB` (fp32 ssm), `kv_bytes_per_token ≈ 13.8 KB` (fp8) → `token_equiv ≈ 4080`. Cache strategy `extra_buffer_lazy` → `S = 4`. Workload `L = 9216` (8192 in + 1024 out). NOSPEC → `D = 0`.\n\n- **TP (dcp_size=1)**: `r = 4 · 4080 · 1 / 9216 ≈ 1.8`. Measured: clamp 88, KV cap 87 → balanced. ✅\n- **DCP8 (dcp_size=8)**: `r = 4 · 4080 · 8 / 9216 ≈ 14`. Measured at r=14: clamp 125, KV cap 129 → balanced. ✅ (At the naive r=1.8 the DCP KV pool is ~8× over-provisioned — cap 683 vs clamp 86 — wasting budget that should go to the state pool.)\n\n`predict_clamp` on the same model at the **default** `extra_buffer`, overlap on and the decode-lock skip off (`S = 5`), three boot logs:\n\n| config | `max_mamba_cache_size` | `mmcs // S` | measured `max_running_requests` |\n|---|---|---|---|\n| TP8, no DCP | 257 | 51 | 51 |\n| DCP8, `r = 5.97` | 451 | 90 | 90 |\n| DCP8, `r = 9` | 474 | 94 | 94 |\n\nReference `r` for **this example model** (`token_equiv ≈ 4080`, i.e. fp32 ssm + fp8 KV; S=5 default; multiply by `(S+D)/S` for spec, by `dcp_size` for DCP) — recompute with your own `token_equiv` for a different model:\n\n| L | 2K | 4K | 8K | 32K | 64K | 128K |\n|---|---|---|---|---|---|---|\n| r | 10.0 | 5.0 | 2.5 | 0.62 | 0.31 | 0.16 |\n\n## Caveats\n\n- **DCP + spec**: a replicated (non-DCP-sharded) draft KV does **not** shard by `dcp_size`; at long `L` it dominates per-token cost, so the clean `×dcp_size` overstates DCP's advantage — fall back to a per-token direct-solve (KV term `/dcp` + an un-sharded draft-KV term) when spec is on. (NOSPEC → clean `×dcp_size` holds.)\n- **Spec + ReplaySSM → use `D=0`, not the draft-token count.** When ReplaySSM spec-verify is enabled, the `D` intermediate SSM states move off the per-request slot budget onto a **fixed ring** (a one-time deduction from `rest`, not a per-req term). So the mamba-slot cost per running request drops back to `S` (clamp = `mmcs / S`, not `/(S+D)`), and the balance ratio returns to the NOSPEC value (`r* ≈ S·token_equiv·dcp/L`). Measured example (TP8, D=8, L≈9K): applying `D=8` computes `r*≈2.6` but the true optimum is `r≈1.0` — `D=8` lands KV-bound at ~45% below the achievable peak concurrency. Plain spec (no replayssm) keeps `D` = the draft-token count.\n- **Under DCP the balance point sits beyond any realistic `L`** (the state pool binds first for almost everything), so the practical recommendation is to **pin `--max-mamba-cache-size = target_concurrency · S` directly** rather than dial a large `r`.\n- **`--max-mamba-cache-size` overrides `r`**; bytes beyond the `r` budget come out of the KV pool one-for-one.\n- **Precision changes are behind accuracy gates**: `--kv-cache-dtype fp8_e4m3` (doubles `token_equiv` → doubles `r`) and `--mamba-ssm-dtype bfloat16` (~halves `token_equiv`; also silently switches the linear-attention decode backend on SM100+ — pin `--linear-attn-decode-backend triton`) shift outputs; validate accuracy for the workload before production.\n- **Asymmetry**: the state pool is worst-case-reserved and fail-loud; KV degrades gracefully (retraction). When `L` is spiky/uncertain, bias `r` **up** rather than starve the state pool.\n","frontmatter":{"name":"compute-mamba-ratio","description":"Compute the optimal --mamba-full-memory-ratio (or --max-mamba-cache-size pin) for a hybrid attention + linear-attention (Mamba / GDN / KDA) model's two serving memory pools, from the workload and serving config. Use when a user asks what ratio to set, why concurrency is clamped, or how to size the state vs KV pools for a hybrid model."},"isInternal":false,"tokens":3553,"sizeBytes":11941},{"name":"authoring-reference.md","path":".agents/skills/cookbook-add-model/references/authoring-reference.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/authoring-reference.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/authoring-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/authoring-reference.md)","tokens":300,"sizeBytes":252},{"name":"diffusion-authoring.md","path":".agents/skills/cookbook-add-model/references/diffusion-authoring.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/diffusion-authoring.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/diffusion-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/diffusion-authoring.md)","tokens":300,"sizeBytes":252},{"name":"engine-axis.md","path":".agents/skills/cookbook-add-model/references/engine-axis.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/engine-axis.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/engine-axis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/engine-axis.md)","tokens":300,"sizeBytes":236},{"name":"mintlify-authoring.md","path":".agents/skills/cookbook-add-model/references/mintlify-authoring.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/mintlify-authoring.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/mintlify-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/mintlify-authoring.md)","tokens":300,"sizeBytes":250},{"name":"vendor-logo.md","path":".agents/skills/cookbook-add-model/references/vendor-logo.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/vendor-logo.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/vendor-logo.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/vendor-logo.md)","tokens":300,"sizeBytes":236},{"name":"SKILL.md","path":".agents/skills/cookbook-add-model/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/SKILL.md","title":"cookbook-add-model","category":"anthropic-skill","format":"markdown","content":"---\nname: cookbook-add-model\ndescription: Add a new model to the SGLang Cookbook (docs/, Mintlify), config-driven format — instantiate the model-agnostic template into a per-model config (+ benchmarks) JSX under src/snippets/configs/, an MDX page, the docs.json nav entry, NEW-tag hygiene, and the homepage vendor card. Interactive, multi-phase. Run with /cookbook-add-model.\ndisable-model-invocation: true\n---\n\n# Add a model to the SGLang Cookbook\n\n> Migrating an **existing legacy-template page** (one that imports a monolithic\n> `…/autoregressive/-deployment.jsx` generator)? Use the\n> `cookbook-migrate-model` skill instead — same target format, but the legacy\n> page (not the user) is the source of truth.\n\nThe cookbook is **config-driven**: two shared engines contain NO model-specific code —\n`docs/src/snippets/_deployment.jsx` (the 5-dim deploy matrix) and\n`_playground.jsx` (the diff-based override Playground). Adding a model = adding **data**:\na per-model `config` (+ optional `benchmarks`) consumed by both engines, plus an MDX page\nthat imports them. No engine edits.\n\n**Instantiate the model-agnostic template** (NOT a clone of any live cookbook — the\ntemplate is decoupled and covers all hardware + all axes):\n- `templates/config.jsx.tmpl` → `docs/src/snippets/configs//.jsx`\n- `templates/benchmarks.jsx.tmpl` → `…/-benchmarks.jsx` (skip if no numbers)\n- `templates/page.mdx.tmpl` → `docs/cookbook///.mdx`\n\nThe template uses explicit `__TOKEN__` placeholders; you fill them, prune what the model\nlacks, and replace the EXAMPLE cells with verified recipes. DeepSeek-V4 is a populated\n*instance* you can consult, but is not the template.\n\n**Deep references (read on demand, don't inline):**\n- [references/authoring-reference.md](references/authoring-reference.md) — field-by-field config / cells / playground / MDX contract.\n- [references/mintlify-authoring.md](references/mintlify-authoring.md) — MDX rules (forbidden syntax, JSX tables, labeled fences) + invocation-example patterns. Read before writing §1–§3 prose.\n- [references/engine-axis.md](references/engine-axis.md) — adding a new Playground feature axis (rare engine work).\n- [references/vendor-logo.md](references/vendor-logo.md) — new-vendor card logo: ask the user for the brand logo, then generate the icon-only 940×525 RGBA PNG (spec + Pillow recipe + `git add -f`).\n- [references/diffusion-authoring.md](references/diffusion-authoring.md) — required opening, tag, command-picker, and feature-overlay contract for diffusion model pages. Read before editing any `docs/cookbook/diffusion//.mdx` page.\n\n## Architecture at a glance\n\n```\n┌─────────────────────────────────────────────────────────────────┐\n│ cookbook///.mdx │\n│ import { Deployment } from \"/src/snippets/_deployment.jsx\"; │\n│ import { Playground } from \"/src/snippets/_playground.jsx\"; │\n│ import { config } from \"/src/snippets/configs/.../X.jsx\"; │\n│ │\n└─────────────────────────────────────────────────────────────────┘\n │ (config passed as React prop)\n ▼\n┌─────────────────────────────────────────────────────────────────┐\n│ src/snippets/configs//.jsx │\n│ export const config = { │\n│ supportedHardware, variants, quantizations, strategies, ... │\n│ cells: [ { match:{hw,variant,quant,strategy,nodes}, │\n│ env:[...], flags:[...] }, ... ], // 5-dim matrix │\n│ playgroundFeatures: { attention, moe, parsers, ... }, │\n│ }; │\n└─────────────────────────────────────────────────────────────────┘\n │ (consumed by BOTH engines — no model code in engines)\n ▼\n┌──────────────────────────────────┬──────────────────────────────┐\n│ _deployment.jsx │ _playground.jsx │\n│ Renders the verified matrix; │ Renders override chips + │\n│ one cell → its env/flags. │ diff against the cell. │\n└──────────────────────────────────┴──────────────────────────────┘\n```\n\nThe two widgets stay in sync via: the **URL hash** (deploy mirrors its selection;\nplayground reads it), the **`sglang-deploy-sel` custom event** (deploy dispatches on\nevery change; playground listens — `replaceState` doesn't fire `hashchange`), and the\nshared **`sglang-deploy-env` localStorage key** (HOST/PORT placeholders).\n\n> The main template is **autoregressive**. Diffusion pages use\n> `templates/diffusion-page.mdx.tmpl` plus\n> [references/diffusion-authoring.md](references/diffusion-authoring.md); do not force the\n> autoregressive deployment matrix on them. Omni pages follow their own category structure.\n> All categories still obey the Mintlify / NEW-tag / docs.json / category-card / validation\n> rules below.\n\n---\n\n**Interactive, multi-step workflow. Collect inputs incrementally — don't ask for\neverything upfront.** The real work is the verified `cells[]` recipes + measured\nbenchmarks (Phases 2 + 4); everything else is filling the template.\n\n## Phase 1 — Collect inputs\n\n1. **Model card** — HuggingFace repo/URL. **Fetch the page** and extract description,\n param count, architecture, context length, **license**. (Fetching guards factual bugs\n like an off-by-a-few-B param count.) If the model isn't public, ask the user.\n2. **Variants / quantizations** — keep separate: variants are size/mode (e.g. Flash/Pro,\n Instruct/Thinking); quantizations come from the HF card / linked repos (BF16/FP8/FP4/…).\n Default to BF16 when a full-precision repo exists.\n3. **Tested hardware + parallelism** — which platforms are actually tested, and TP/EP/DP\n for each. List only tested hw (unlisted greys out).\n4. **Verified launch recipes** — the full `sglang serve` flags per\n (hw × variant × quant × strategy × nodes) combo → these become `cells[]`. Rewrite any\n `python -m sglang.launch_server` to `sglang serve` form.\n5. **sglang version / image tag** — ask which sglang build the recipes + benchmarks ran on\n (a release like `0.5.x`, or `main`/nightly). **Never guess or hallucinate it.** This one\n tag fills `dockerImages` and the benchmarks' `sglang_version`; when the user is unsure,\n default the image to `lmsysorg/sglang:dev` (nightly) rather than inventing a release.\n6. **Pre-flight**: `gh pr list --repo sgl-project/sglang --search \"\"` (dup check).\n\n**Hardware reference** (the shared `HARDWARE_CATALOG` in `_deployment.jsx`). A GPU **not** in\nthis table (RTX PRO 6000, GH200, future chips) goes in the model's own `config.hardware`\n(`{id,label,vram,vendor}`) — the engine merges it in; don't edit the engine catalog:\n\n| Platform | Vendor | VRAM | Docker image |\n|---|---|---|---|\n| H100 | NVIDIA | 80GB | `lmsysorg/sglang:` |\n| H200 | NVIDIA | 141GB | `lmsysorg/sglang:` |\n| B200 | NVIDIA | 192GB | `lmsysorg/sglang:` |\n| B300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130` when required) |\n| GB200 | NVIDIA | 192GB | `lmsysorg/sglang:` (or `-cu130`) |\n| GB300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130`) |\n| DGX Spark | NVIDIA | 128GB (unified) | `lmsysorg/sglang:` — needs a **CUDA 13** build |\n| MI300X | AMD | 192GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI325X | AMD | 256GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI350X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n| MI355X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n\n- **Image tag (``)**: don't guess — ask the user for the tag the recipes ran on, or\n default to `dev` (nightly). The same tag goes in `dockerImages` and benchmarks'\n `sglang_version`; the engine falls back to `lmsysorg/sglang:dev` for any unmapped hw.\n- **TP sizing** (sanity-check recipes): `weight_GB / gpu_mem`, round up to a power of 2,\n ~20–30% headroom. BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5. MoE → **total** weight, not\n active params. FP4 is Blackwell-only (B200/B300/GB200/GB300/DGX Spark). GB200/GB300\n single-node hosts are typically **4 GPUs** (TP=4 ceiling); a DGX Spark is **1 GPU**, so\n its only multi-GPU topology is TP=2 across 2 nodes.\n- **Platform flags**: Blackwell may need `--attention-backend trtllm_mha`; AMD typically\n needs `--attention-backend triton` + env `SGLANG_USE_AITER=1` /\n `SGLANG_ROCM_FUSED_DECODE_MLA=0` (check AITER TP constraints, e.g. `heads_per_gpu % 16 == 0`).\n- **EP** (MoE): 8-GPU NVIDIA `--tp 8 --ep 8`; AMD `EP = TP`; small NVIDIA (TP≤4) omit\n `--ep` unless benchmarked. (The template's AMD example cell shows these.)\n\n## Phase 2 — Instantiate the template\n\nFor a diffusion model, instantiate `templates/diffusion-page.mdx.tmpl` and keep the\nshared `DiffusionModelTags` component, plus `templates/diffusion-config.jsx.tmpl` for the\nopt-in scoped command builder. Put the compact install command and builder in §1 Quick\nstart. The first two paragraphs in §2 Model capabilities are not generic filler: they must\nstate the model's capability range, strongest differentiator, when to choose it, and at\nleast one real deployment or capability boundary. Put orthogonal runtime features in\n`scope: \"serve\"` or `scope: \"request\"`, not in the base recipe; use the schema from\n`references/diffusion-authoring.md`.\n\n1. **Copy** the three template files to their target paths (above). Note the two\n vendor-folder conventions: under `configs/` the folder is the **HuggingFace org**\n (`deepseek-ai`); under `cookbook/` it's the **display vendor** (`DeepSeek`).\n2. **Replace every `__TOKEN__`**: `__MODEL_DISPLAY__`, `__MODEL_SLUG__`, `__HF_ORG__`,\n `__HF_REPO__`, `__REASONING_PARSER__`, `__TOOLCALL_PARSER__`, `__ONE_LINER__`. Verify\n none remain: `grep -rn '__[A-Z_]*__' `.\n3. **Prune** to what the model supports (delete, don't stub) — using\n [references/authoring-reference.md](references/authoring-reference.md):\n - `supportedHardware` + the EXAMPLE cells: keep your tested families; **delete the\n `mi*` ids + AMD example cell if no AMD recipe**, etc. A GPU not in the shared catalog\n (e.g. RTX PRO 6000) → declare it in `config.hardware` and add its id here.\n - `playgroundFeatures` axes: remove the `megamoe` backend option + the\n `megamoeQuant` block from the `moe` axis (non-Blackwell-MoE), delete `hisparse`\n (non-DSA), `pdDisagg`/`router` (no PD), the `parsers` axis (no parsers), etc.\n - `quantizations` / `variants`: drop what the model doesn't ship; collapse `variants`\n to single `default` if there's no variant axis (then drop the `variant` half of\n `modelNames`/`defaultAccuracy` keys).\n4. **Fill `cells[]`** with the verified recipes from Phase 1 (replace every EXAMPLE cell;\n set `verified: true` only on tested combos), and `modelNames` with real HF slugs,\n `dockerImages` for your hw (use the Phase-1 tag, or default `lmsysorg/sglang:dev` — never\n a guessed release; key by `hw`, or a more specific key — `hw|variant|quant`, `variant|quant`,\n `hw|quant|strategy`, `hw|quant`, resolved in that order before `hw` — when one\n variant/quant/strategy needs its own image), `multiNodeHints` only for fabric-specific hw\n (e.g. gb200).\n\n5. **Diffusion pages: add the ComfyUI section.** Every diffusion cookbook page ends with\n a `## . Run in ComfyUI` section so a reader never has to guess whether the model is\n reachable from ComfyUI. It is one component; the per-model facts live in the component,\n not the page:\n\n ```mdx\n ## . Run in ComfyUI\n\n import { ComfyUISupport } from '/src/snippets/diffusion/comfyui-support.jsx';\n\n \" />\n ```\n\n Pick `model` from the table in\n `docs/src/snippets/diffusion/comfyui-support.jsx`. Use the model's own key when it has\n an entry (its executor or dedicated node differs); otherwise use the generic `image` or\n `video`. A model gets its own key only when the plugin actually treats it specially —\n an entry in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a\n dedicated node. Adding a key without the matching plugin support makes the page lie.\n\n### Site-wiring (do all three)\n\n- **`docs/docs.json`** — add the page under Cookbook → `` → ``, at\n the **top** of that vendor's `pages` (root-relative, no `.mdx`:\n `cookbook///`). New vendor group → insert in the section's\n local ordering.\n- **NEW-tag hygiene** — the new page keeps `tag: NEW` (from the template). Scan the\n vendor dir for existing NEW and strip it from siblings; verify ≤1:\n `grep -rn 'tag: NEW' docs/cookbook///` → at most one result. (Scan\n files; don't assume the first `docs.json` entry holds NEW.)\n- **Homepage card** — `docs/cookbook//intro.mdx`: if the org already has a\n ``, update only its `href` (keep `img`). If the org is **new**, add a ``\n (title = nav-group name; keep card order aligned with `docs.json`) **and create its logo**:\n ask the user for the brand logo, then generate the conforming **icon-only 940×525 RGBA\n transparent** PNG → `docs/cards/logos/.png` per\n [references/vendor-logo.md](references/vendor-logo.md) (track with `git add -f` — `*.png`\n is gitignored repo-wide). Never invent or copy a logo.\n\n## Phase 3 — Validate\n\n```bash\ncd docs\nmint validate # frontmatter, missing nav entries, MDX/JSX errors\nmint broken-links\nmint dev # visual smoke test at http://localhost:3000/cookbook///\n```\n\nSpot-check: cells render sensible commands; URL-hash nav persists across reload; the\nPlayground inherits the Deploy selection live; each axis toggle produces the expected\ndiff; Docker mode wraps in `docker run` with the right image; multi-node cells emit the\nhints + `--nnodes N`; cURL resolves the model name; the NEW badge shows on the new page\nand is gone from same-vendor siblings; the homepage card points to the new model.\n\n## Phase 4 — Interactive testing\n\nThe user deploys each cell, runs the benches, and pastes results; you fill the data:\n- mark each tested `cells[]` entry `verified: true` (absent = yellow/unverified badge);\n- fill the `-benchmarks.jsx` entries (one per cell `match`) with measured\n speed/accuracy + the `sglang_version` the user reports (don't invent one — the template's\n `0.0.0` is a deliberate TODO); set model-level `defaultAccuracy` per variant. Leave a\n cell's entry as a bare `match` stub if it has no numbers yet (the card shows \"pending\").\n\n## Phase 5 — Prose & config tips\n\n**Read [references/mintlify-authoring.md](references/mintlify-authoring.md) first** (it\ncarries the parser-output-shape / thinking-mode / Output-Example / no-hardcoded-sampling\nrules + the Mintlify forbidden-syntax list). Then rewrite the MDX prose from the HF card +\nuser notes: §1 Model Introduction (description, links, params, license, variants table),\n§2 Configuration Tips (hw-specific tuning, caveats), §3 Advanced Usage (Reasoning /\nTool-Calling / HiCache — keep only what applies; match the reasoning example to the\nparser's output shape; each runnable block gets an `**Output Example:**`).\n\nFor diffusion pages, follow the category-specific Quick start and capability contract in\n`references/diffusion-authoring.md`. Run `node docs/scripts/check_cookbook_configs.mjs` to\nverify the tag widget and introduction structure before rendering the page.\n\n## Phase 6 — Review\n\n```\n/cookbook-review-pr \n```\n\n## Git workflow\n\nAlways branch — never commit to main directly.\n\n```bash\ngit checkout -b add--cookbook\ngit add docs/src/snippets/configs//.jsx \\\n docs/src/snippets/configs//-benchmarks.jsx \\\n docs/cookbook///.mdx \\\n docs/docs.json docs/cookbook//intro.mdx\ngit commit -m \"Add cookbook\"\ngit push -u origin add--cookbook\ngh pr create --title \"[Doc] cookbook: Add cookbook\" --body \"...\"\n```\n","frontmatter":{"name":"cookbook-add-model","description":"Add a new model to the SGLang Cookbook (docs/, Mintlify), config-driven format — instantiate the model-agnostic template into a per-model config (+ benchmarks) JSX under src/snippets/configs/, an MDX page, the docs.json nav entry, NEW-tag hygiene, and the homepage vendor card. Interactive, multi-phase. Run with /cookbook-add-model.","disable-model-invocation":true,"disableModelInvocation":true},"isInternal":false,"tokens":4373,"sizeBytes":17120},{"name":"dimension-mapping.md","path":".agents/skills/cookbook-migrate-model/references/dimension-mapping.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/references/dimension-mapping.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/cookbook-migrate-model/references/dimension-mapping.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/references/dimension-mapping.md)","tokens":300,"sizeBytes":256},{"name":"SKILL.md","path":".agents/skills/cookbook-migrate-model/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/SKILL.md","title":"cookbook-migrate-model","category":"anthropic-skill","format":"markdown","content":"---\nname: cookbook-migrate-model\ndescription: Migrate a legacy-template SGLang cookbook page (monolithic per-model generator under docs/src/snippets/autoregressive/) onto the config-driven template (shared _deployment.jsx / _playground.jsx engines + per-model config). Use when asked to migrate, convert, or port an existing cookbook page — NOT for brand-new models (use cookbook-add-model for those). Run with /cookbook-migrate-model .\n---\n\n# Cookbook Migrate Model\n\nConvert one legacy cookbook page to the config-driven format, faithfully. The\nlegacy page — its generator widget and its measured benchmark blocks — is the\n**single source of truth**. You are transcribing it into the new data model,\nnot improving it.\n\nReuses the `cookbook-add-model` skill's assets (read them on demand):\n- `../cookbook-add-model/templates/config.jsx.tmpl`, `page.mdx.tmpl`, `benchmarks.jsx.tmpl`\n- `../cookbook-add-model/references/authoring-reference.md` (config/cells/playground contract)\n- `../cookbook-add-model/references/mintlify-authoring.md` (MDX rules)\n\nMigration-specific references in this skill:\n- [references/dimension-mapping.md](references/dimension-mapping.md) — legacy-control → new-dimension mapping rules, command rewrite table, per-family strategy sets, and the Qwen3.5 pilot as a worked example (PR #27848).\n\nThe round's per-model inventory (scope, batch order, quirks, measured-data\nsurvey) is tracked by the migration maintainer outside the repo — expect it in\nyour dispatch prompt, or ask for it.\n\n## Hard rules (non-negotiable)\n\n1. **Never modernize.** Env vars, flags, TP values, docker tags, version strings\n are copied verbatim from the legacy page — even when today's defaults differ\n (e.g. a flag whose behavior has since become the default; keep it anyway). The recipe\n that was verified is the recipe as written. Allowed normalizations are ONLY\n the five alias rewrites in dimension-mapping.md §2 (`launch_server`→`sglang\n serve`, `--model`→`--model-path`, `--tp-size`→`--tp`, abbreviated\n `--speculative-algo`→`--speculative-algorithm`,\n `--expert-parallel-size`→`--ep`). **Accuracy-degrading flags**\n (`--kv-cache-dtype fp8_e4m3`, W4A4-style runtime quant) follow a\n deterministic rule — enforced in migration, no asking: offered as a\n legacy **selectable option** → never select it (cells mirror the\n accuracy-safe side), and the option **survives as a Playground axis** —\n an existing one where it fits, else add one via a separate prior engine\n PR (rule 4); a legacy choice never degrades to a tips mention. Baked\n into the recipe's **default/unconditional command** → keep it verbatim\n (the recipe was measured with it, and fp8 KV halves KV memory —\n stripping could OOM it). See dimension-mapping.md §2 caveats.\n2. **Never invent versions or numbers.** Benchmark numbers only from the\n legacy page's measured blocks, and a result migrates ONLY when its\n `sglang_version` is a **reproducible anchor** — the bar is reproducibility,\n not \"must be a release\":\n - ✅ release tag/version (`v0.5.9` / `0.5.9`), commit hash, OR — for\n **Day-0 support** (the enabling PR isn't merged and no release is cut\n yet) — a specific **PR (`PR #27944`) or commit** you can `gh pr checkout`\n / `git checkout `. Commit is most precise; a PR pin is fine for day-0.\n - ❌ a moving ref — `\"main branch\"`, `\"main (2026-06-11)\"`, open-ended\n `\"0.5.8+\"` — is NOT reproducible: **drop the WHOLE result (speed AND\n accuracy)**, not just speed. Keep `benchmarkCommands` so ⚡Reproduce still\n guides re-measurement against a pinned build.\n **Never inherit cross-model numbers** — measurements the legacy page\n attributes to a *different model* (e.g. a K2.6 page carrying K2.5-measured\n speed) are dropped regardless of version. When kept, `sglang_version` is the\n legacy page's string verbatim. Docker tags only the ones the legacy page\n pinned (unmapped hw falls back to `:dev`).\n3. **Verified policy (strictest tier).** `verified: true` ONLY when (a) the\n legacy page has concrete measured data for that exact 5-dim combo AND\n (b) the cell's flags equal the deployment command used for that measurement\n — modulo `{{HOST_IP}}`/`{{PORT}}`, the five alias rewrites, and **parser\n flags**: `--reasoning-parser`/`--tool-call-parser` are stripped from every\n cell (Playground-only feature; when the measured run had them on, say so in\n the benchmarks file header). When the measured command diverges from the\n generator default, **the verified cell follows the measured command**; the\n generator default stays as the sibling strategy/cell or a tips note. Everything else is unverified (yellow) —\n including combos that look memory-infeasible; keep them verbatim and list\n them in the PR body for the re-verification track.\n4. **Engines are read-only.** `_deployment.jsx` / `_playground.jsx` must not\n change in a migration PR. Model-specific features are config DATA consumed\n by generic axis handlers (MegaMoE precedent), so they need NO engine change.\n A **titled single-select that strips a flag family** — KV Cache DType\n (`--kv-cache-dtype`), mamba (`--mamba-radix-cache-strategy`), … — is already\n covered by the merged generic **`flagSelects`** axis: declare it in the\n config (a list of `{ id, title, stripPrefixes, options }`; see the Qwen3.5\n mamba example), **no engine PR**. Only a genuinely new control *shape* that\n `flagSelects` can't express would need a one-time generic primitive (never a\n model-named handler) on a separate prior PR (engine-axis.md).\n5. **`github.cookbookModel` must be set** (`/`, e.g.\n `qwen/qwen3.5`) and the block never pruned — without it Submit ↗ mislabels\n as deepseek-v4. The issue template itself needs NO edits (free-form input).\n6. **Strategy tiers are signal-driven.** A cell goes under `low-latency` /\n `high-throughput` ONLY on a signal present in the legacy source (an\n explicit performance toggle, a named recipe, or prose stating the\n operating point); **no signal → `balanced`**. Never derive a slant from\n your own hardware intuition — re-tiering on measured evidence is the\n hardware owner's follow-up PR, not part of a migration\n (dimension-mapping.md §4).\n\n## Workflow (one model = one PR)\n\n### 1. Inventory the legacy assets\n- Read the legacy generator (`docs/src/snippets/autoregressive/-deployment.jsx`)\n end-to-end: every option dimension (radio vs checkbox vs dynamic), every\n gate/SUPPORT matrix, the full emitted command per reachable combo (env\n prefixes, `# Error` pseudo-commands included).\n- Read the legacy MDX: §2 install docker tags → `dockerImages` (pinned, not\n upgraded); §3.2 tips → new §2; §4 invocation examples → new §3 (keep real\n Output Examples verbatim); §5 benchmark blocks → transcribe each measured\n block: deploy command used, bench command (dataset/isl/osl/num-prompts/\n concurrency), P50 (median) TTFT/TPOT, output tok/s, hardware, version string.\n- Inbound-anchor sweep: `grep -rn \"\" docs/ --include='*.mdx'` —\n find links/`#fragments` into this page (`mint broken-links` does NOT check\n fragments). Fix referrers or add `` shims in the same PR.\n- Check the maintainer-provided inventory notes for this model's known quirks —\n but treat them (and the §4 family table) as a **survey snapshot**: re-verify\n every dimension against the live legacy files before mapping. Pages keep\n receiving updates (precedent: Kimi-K2.6's live generator has a speculative\n toggle the 2026-06-10 survey notes lack).\n\n### 2. Design the 5-dim mapping\nApply [references/dimension-mapping.md](references/dimension-mapping.md). Key\ndecision — which legacy toggle becomes the `strategies` dimension: a toggle\nthat **changes other parts of the command** (TP, mem) must (the Playground\ncan't do coupled changes), and so does a toggle the legacy page itself labels\nwith operating-point words — e.g. a `dpattention` radio whose options are\nsubtitled \"Low Latency\" / \"High Throughput\" (GLM-5.1 / Kimi-K2.6 pattern) —\neven when its flags are uncoupled (`--attn-dp-size N` is a pure\nflag add). Any other toggle that only adds/removes its own flags becomes a\nPlayground axis with the flags baked into cells when the legacy default was\nON — EXCEPT parsers:\n`--reasoning-parser`/`--tool-call-parser` are NEVER baked into cells, they are\nPlayground-only (DSv4 convention). **Every legacy control survives as an\ninteractive control** — a dimension or a Playground axis, never a tips-only\nmention — and a model-specific control is **config data, not engine code**\n(MegaMoE W4A4 is all DSv4 config on the existing `moe` axis). It's pure\nconfig whenever it fits an existing axis's data schema. A **titled\nsingle-select that strips a flag family** (Nemotron3's \"KV Cache DType\",\nmamba `--mamba-radix-cache-strategy`, …) fits the merged generic **`flagSelects`**\naxis — so it too is config-only (declare a `flagSelects` list). Only a control\nwhose *shape* `flagSelects` still can't express would need a ONE-TIME generic\nprimitive (never a model-named handler) on a separate PRIOR engine PR, keeping\nthe migration PR data-only (hard rule 4, engine-axis.md). The strategy count follows the page's\noperating points: **one recipe → a single `balanced`; two → `low-latency` +\n`high-throughput`; three → the full trio (the ideal)**. The tiers apply per\n(hw × variant × quant) combination — a single-recipe combination on a\nmulti-strategy page parks under its semantically honest tier (no\nlatency/throughput slant → `balanced`; the page's list is the union). When the\nlegacy toggle is MTP / speculative decoding, the direction is a deterministic\ndefault — apply without asking: **MTP on → `low-latency`, MTP off →\n`high-throughput`** (reversed only with maintainer confirmation). Tier\nplacement is signal-driven (hard rule 6). Never invent a recipe just to fill\nstrategy chips (see dimension-mapping.md §4). Record the outcome as a\n**strategy mapping table** for the PR body — one row per group of\ncombinations sharing the same legacy signal (e.g. \"all GPU combos: MTP\ntoggle → low-latency / high-throughput\"; \"xeon: (none) → balanced\"), with a\none-line rationale each; don't enumerate 60 identical rows. The table is what\nhardware owners sign off on at review.\n\n### 3. Generate the config (codegen, then audit)\n- For >~30 cells, port the legacy `generateCommand()` into a throwaway Node\n script that enumerates combos and emits the `cells:[...]` literal (output\n must stay a pure literal — Mintlify forbids runtime spreads/calls). Apply the\n verified-cell override in the script. See the pilot scripts embedded in the\n worked example of dimension-mapping.md §5.\n- **Independent equivalence audit (required):** extract the ORIGINAL generator\n from git (`git show main:` — NOT `HEAD:`, the migration branch deletes\n the file, see dimension-mapping.md §5 item 7), stub React hooks, run it for\n every combo, and diff token-by-token against the new cells. Expected deltas\n only: the appended `--host {{HOST_IP}}`/`--port {{PORT}}`, the\n engine-injected multi-node trio, the §2 alias rewrites (the entrypoint\n rewrite doesn't appear in cell tokens — cells hold flags only; the audit\n script normalizes it on the legacy side), and the intentional verified-cell\n override. Paste the PASS count + the audit script in the PR body\n (collapsed `
`).\n- Hand-author the non-cells fields per authoring-reference.md. Structural\n self-checks: every cell resolves a `modelNames` key; no `--nnodes/--node-rank/\n --dist-init-addr/--host/--port` literals; every `{{KEY}}` declared; every\n `supportedHardware` id has ≥1 cell.\n\n### 4. Benchmarks file\nOne entry per measured block only (cells without entries already render\n\"pending\" — bare `{match}` stubs are unnecessary). `tokens_per_sec_per_gpu` =\ntotal (in+out) tok/s/GPU = `output tok/s ÷ (tp × nnodes) ×\n(isl+osl)/osl` — stored directly (the card shows it as-is). TTFT/TPOT\ntake the P50 (median) rows; set `config.latencyPercentile` (default `\"P50\"`; use\n`\"Mean\"` only for legacy Mean-recorded data — temporary, being migrated to P50; an\nentry-level `latencyPercentile` overrides the page value per cell).\nPut the workload's\n`num_prompts` into `workload`. **`config.accuracyLabels` is required whenever\nthe benchmarks carry accuracy data** — the engine ships no default eval set\n(#27842), so missing labels means the accuracy rows silently don't render;\nextra context (sample counts, suites that don't fit) goes in the entry's\n`notes`. Zero-measured-data pages: skip the file and the `benchmarks` prop\nentirely, but keep `benchmarkCommands` so ⚡Reproduce still guides users.\n\n### 5. Rewrite the MDX\nFrom `page.mdx.tmpl`: keep the original `title` (nav identity), write a fresh\nSEO `description` (top-level — delete any legacy `metatags.description`), **no\n`tag: NEW`** (a migration is not a launch), **no `mode:`**. Install accordion\ncarries the legacy install content + pinned images. Keep the template's\nDSv4-style strategy bullets — serving semantics first (single-user chat /\ntypical multi-user / batch throughput), trimmed to the strategies the page\nships, plus at most a one-line note on what each strategy changes on this\nmodel; do NOT rewrite them as toggle-/migration-centric explanations.\nLegacy §5 benchmark prose is deleted (numbers → benchmark card, commands →\n⚡Reproduce); legacy prose deploy commands are deleted (doc↔config parity —\nfold their unique flags into §2 tips). Invocation examples + real outputs carry\nover verbatim, but **wrapped in Accordions** — §3 commands and outputs are\ncollapsible (required, DeepSeek-V4 pattern): code in an\n``, output in a following\n``; legacy pages kept them inline.\n\n### 6. Delete the legacy generator\nRemove `docs/src/snippets/autoregressive/-deployment.jsx` and its\nimport. `grep -rn \"-deployment\" docs/` must return nothing (config\nprovenance comments must not name the deleted path). Site wiring needs **no\nchanges**: docs.json path/title unchanged, vendor card + logo already exist.\n\n### 7. Validate\n- `grep -rn '__[A-Z_]*__'` on the new files (no template tokens).\n- `cd docs && mint validate && mint broken-links` (pre-existing breaks on\n main are not yours — say so in the PR).\n- `mint dev` browser smoke: initial selection = the verified cell (first in\n `cells[]`) with green badge; multi-node cells show the injected trio +\n header; AMD cells show env prefixes; Docker mode wraps with the pinned image\n and passes cell env as `--env`; condition-hidden combos grey out; benchmark\n card values; NO parser flags in any Deploy command; Playground parser\n toggles ADD the parser flags (green additions) while spec toggles strike\n the baked spec flags (red); Submit ↗ prefills this model. **Probe pitfall:** drive at most\n ONE programmatic click per evaluation and wait for React to settle —\n multiple clicks in one synchronous script batch and read stale DOM.\n- Token-level audit from step 3 passes.\n\n### 8. PR + review\nOne PR per model. PR body: migration framing, verified policy applied, the\nstrategy mapping table (step 2), the audit PASS count + script, any\ninherited-infeasible combos flagged for re-verification. Then run `/cookbook-review-pr ` and fix findings.\nFYI: docs previews only build for in-repo (`sgl-project/sglang`) branches —\na fork-headed PR is perfectly fine but renders no preview; a maintainer can\nre-push the branch in-repo if a preview is wanted for review.\n\n### 9. Keep this skill current\nAny new convention, engine behavior, or pitfall you discover while migrating\n(naming decisions, audit-script gotchas, review-rule conflicts, …) MUST be fed\nback into this skill — same PR if it's skill-file-only, or an immediate\nfollow-up commit on the skill's branch/PR. The next agent runs on what's\nwritten here, not on your session's context.\n","frontmatter":{"name":"cookbook-migrate-model","description":"Migrate a legacy-template SGLang cookbook page (monolithic per-model generator under docs/src/snippets/autoregressive/) onto the config-driven template (shared _deployment.jsx / _playground.jsx engines + per-model config). Use when asked to migrate, convert, or port an existing cookbook page — NOT for brand-new models (use cookbook-add-model for those). Run with /cookbook-migrate-model ."},"isInternal":false,"tokens":4070,"sizeBytes":16129},{"name":"SKILL.md","path":".agents/skills/cookbook-review-pr/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-review-pr/SKILL.md","title":"cookbook-review-pr","category":"anthropic-skill","format":"markdown","content":"---\nname: cookbook-review-pr\ndescription: Review a pull request against the SGLang Cookbook (docs/, Mintlify) contribution checklist — the config-driven format (per-model config + benchmarks JSX consumed by the shared _deployment.jsx / _playground.jsx engines). Run with /cookbook-review-pr .\n---\n\n# Cookbook Review PR\n\nFetch the diff, run the checklist, report what you find. The cookbook is **config-driven**:\nshared engines (`_deployment.jsx`, `_playground.jsx`) with NO model-specific code; each\nmodel is a data `config` (+ optional `benchmarks`) under `src/snippets/configs//`\nplus an MDX page. This checklist targets that layout. Field-schema detail lives in\n`.claude/skills/cookbook-add-model/references/authoring-reference.md` — defer to it rather\nthan restating.\n\n## Usage\n\n```\n/cookbook-review-pr \n```\n\n## Steps\n\n1. `gh pr view --repo sgl-project/sglang --json title,body,files,author,baseRefName,headRefName,commits,reviews`\n2. `gh pr diff --repo sgl-project/sglang`\n3. `gh pr list --repo sgl-project/sglang --state open --search \"\"` (duplicate check)\n4. Run every checklist item against the diff.\n5. Output per-file verdicts + overall recommendation.\n\n## Checklist\n\n### 1. File hygiene\n- A cookbook PR should only touch: `docs/src/snippets/configs//*.jsx`\n (config + benchmarks), `docs/cookbook/**/*.mdx`, `docs/docs.json`,\n `docs/cookbook//intro.mdx` (vendor card), `docs/cards/logos/.png`\n (new vendor only). Flag stray files (`settings.local.json`, lockfiles, IDE configs).\n- Pages must be `.mdx`, not `.md`. Files end with a trailing newline. Check commit history\n for unrelated commits accidentally included.\n- **Engines untouched**: `_deployment.jsx` / `_playground.jsx` should NOT change in a\n model-add PR (adding a model is data-only). Engine edits = a separate axis/feature PR\n (see `cookbook-add-model/references/engine-axis.md`); review them against that checklist.\n\n### 2. Config quality (the per-model `config`)\n- Single `export const config = { ... }` literal — **no** function calls, spreads,\n fragment refs, or IIFE (Mintlify re-evals at hydration → `ReferenceError`).\n- No `!(x in y)` anywhere (Mintlify AST walker crashes) — use `obj.key === undefined`.\n- `supportedHardware` ⊆ `HARDWARE_CATALOG` (in `_deployment.jsx`) ∪ `config.hardware`. A\n model-specific GPU the shared catalog lacks must be declared in `config.hardware`\n (`{id,label,vram,vendor}`), **not** added to the engine catalog — unless the platform is\n genuinely shared, which is an engine change the PR body has to argue (and the\n `cookbook-add-model` hardware table moves with it). A new catalog entry is inert for\n other pages: each filters the catalog by its own `supportedHardware`.\n- `placeholders` declares every `{{KEY}}` used in `curl` or any cell.\n- `modelNames` covers every cell (by `hw|variant|quant` triple or `variant|quant` pair).\n- `strategies` count matches the page's operating points — 1 recipe → a single `balanced`;\n 2 → `low-latency` + `high-throughput`; 3 → the full trio. Tiers apply per\n (hw × variant × quant) combination: a single-recipe combination must park under its\n semantically honest tier (clear slant → that tier, e.g. a workstation card under\n `low-latency`; no slant → `balanced`, e.g. a CPU platform) — **flag a no-slant recipe\n parked under low-latency/high-throughput**. Mixed unions\n like [low-latency, balanced, high-throughput] with per-selection greying are fine. Also\n flag model-specific ids (e.g. `mtp`), and flag an INVERTED speculative mapping — the\n deterministic default is MTP/spec-decoding ON → `low-latency`, OFF → `high-throughput`\n (at saturation the draft+verify overhead outweighs the speedup); the reverse needs an\n explicit maintainer-confirmed justification in the PR. The MDX strategy bullets describe serving semantics\n in the DSv4 style (single-user chat / typical multi-user / batch jobs), not internal\n toggles.\n- `dockerImages` covers the hw ids that have cells (else users hit the `:dev` fallback); a\n `hw|quant` key (resolved before the plain `hw`) is valid when one quant on a shared GPU needs\n a different image (e.g. an FP4 dev build) — don't flag those.\n- `multiNodeHints` present ONLY for hw whose fabric needs manual NIC env (e.g. `gb200`\n NVL72) — NOT every `multi-N` hw (standard-IB DeepEP / Marlin multi-node don't need it).\n Hints render above BOTH run modes, so flag one that reads as docker-only (\"add these\n docker run flags\") — `docker run` flags belong in the hardware entry's\n `multiNodeDockerFlags`, which the engine emits into the Docker command itself.\n- `github.cookbookModel` is set to the model's HF id (`/`). The issue\n template's `model` field is a free-form input prefilled from this value; if the config\n omits the `github` block, the engine falls back to `deepseek-ai/deepseek-v4` and the\n page's submissions get mislabeled.\n- `playgroundFeatures` is opt-OUT: the **general axes ship on every cookbook by default**\n (`attention` TP/CP/DP-Attn, `moe` backend+EP for MoE models, `parsers`, `speculative`,\n `pdDisagg`, `hicache`) — flag a missing general axis unless the model genuinely cannot\n use it. Model-specific axes only where applicable (MegaMoE backend + `megamoeQuant`\n only on Blackwell MoE, gated by `requiresHw`; `hisparse` only DSA-style). Knobs that are\n meaningless for a subset of variants/hw are `disable`d with a reason, not silently live\n (e.g. MoE knobs greyed on dense variants). No empty/stub axes.\n- **No leftover `__TOKEN__`** — the config was stamped from the template and every\n placeholder is filled (`grep -rn '__[A-Z_]*__'` on the new config/benchmarks/MDX returns\n nothing).\n- **All-hardware considered**: every `supportedHardware` id (from the catalog or `config.hardware`) has ≥1 cell OR is a deliberate\n greyed \"coming soon\"; AMD was pruned or kept on purpose (not a leftover template family).\n\n### 3. Cells / 5-dim matrix\n- Every cell `match` has EXACTLY the 5 keys (`hw`, `variant`, `quant`, `strategy`, `nodes`).\n- `env` / `flags` are flat literals (only `{{PLACEHOLDER}}` subst) — no shared\n `commonFlags` reference (Mintlify won't inline it).\n- NO `--nnodes` / `--node-rank` / `--dist-init-addr` literals in multi-node cells\n (the renderer injects them from `match.nodes`).\n- NO literal `--host` / `--port` — use `{{HOST_IP}}` / `{{PORT}}`.\n- NO `--reasoning-parser` / `--tool-call-parser` in any cell — parsers are a\n Playground-only feature added on top of the base command (DSv4 convention);\n flag any cell that bakes them in.\n- Accuracy-degrading flags in cells — runtime quant below the checkpoint\n (e.g. MegaMoE W4A4 — DSv4 gates it behind the Playground's `megamoeQuant`)\n and lossy `--kv-cache-dtype` (e.g. `fp8_e4m3` over a higher-precision-KV\n checkpoint): **flag for explicit maintainer confirmation**. Output quality\n should be exactly what the quant chip declares, so absent a recorded\n sign-off in the PR (e.g. carried verbatim from a measured legacy recipe's\n default command), request the flag move to Playground/tips.\n- Flag order: `--model-path` first (an optional `--trust-remote-code` may precede it —\n the DSv4 cells do), then parallelism, then MoE, then tuning, `--host`/`--port`\n last (the playground's insert anchors assume this).\n- TP/memory sanity: `model_weight_GB / (tp × gpu_mem)` fits with ~20–30% headroom\n (BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5; MoE uses **total** weight, not active params).\n\n### 4. Benchmarks\n- Each `benchmarks[]` entry's `match` tuple corresponds to a real cell.\n- `accuracyLabels` is present whenever the benchmarks carry accuracy data — the engine\n ships NO default eval set; without it the accuracy rows silently don't render.\n `defaultAccuracy` / per-cell `accuracy` / `benchmarkCommands.accuracy` keys all\n ∈ `config.accuracyLabels`.\n- A benchmark's quantization must match a variant actually listed — `(BF16)` on a model\n that only released FP8/FP4 is a factual bug.\n- `benchmarkCommands.speed` is `python3 -m sglang.bench_serving` (the workload), separate\n from the `sglang serve` deploy command, and should carry `--flush-cache`: bench_serving's\n `random` prompts are deterministic, so a warm rerun hits the radix cache and inflates\n throughput — speed numbers are measured cache-cold.\n- `sglang_version` is a real build the author ran (a release, or `dev`/nightly) — not a\n guessed/placeholder value (no leftover `0.0.0`).\n- **Latency percentile**: `config.latencyPercentile` (default `\"P50\"`, or `\"Mean\"`) matches the\n percentile the TTFT/TPOT values actually are — the card renders `TTFT ()`. A benchmarks\n entry may carry its own `latencyPercentile` to override the page value per cell\n (entry → config → `\"P50\"`): on a P50 page, kept legacy Mean cells must set it — a\n `sglang_version` tag alone doesn't convey the percentile. (`\"Mean\"` is temporary — legacy\n data is being re-measured to P50.)\n- **Throughput convention**: `tokens_per_sec_per_gpu` is stored as **total (in+out)/GPU**\n = `output tok/s/GPU × (isl+osl)/osl`, shown by the card as-is. Flag output-only values.\n- **Consistent accuracy harness across entries**: every value under one `accuracyLabels`\n column must be produced by the SAME harness — flag a page that, say, measures one\n platform's GSM8K with a historical raw completion harness and another's with\n `sgl-eval run gsm8k --num-examples 1319` and shows both as one \"GSM8K %\"\n (the scores aren't comparable). Either standardize on one harness (matching\n `benchmarkCommands.accuracy`) or require an explicit per-entry note. Common when folding\n a second contributor's measurements (e.g. an AMD/ROCm PR) into the page.\n\n### 5. Doc ↔ config parity (the #1 finding)\n- Any `sglang serve` command shown in MDX prose (config tips, benchmark section) must\n equal what the engine emits from the corresponding cell — same flags, same order. Drift\n here is the most common review miss.\n\n### 5b. ComfyUI section (diffusion pages)\n- A diffusion page ends with `## . Run in ComfyUI` rendering ``. A\n reader must not have to guess whether the model is reachable from ComfyUI.\n- The `model` prop is a key that exists in `docs/src/snippets/diffusion/comfyui-support.jsx`.\n A model-specific key is only correct when the plugin really treats it specially — an entry\n in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a dedicated\n node. Otherwise the generic `image` / `video` key is the honest one; a model-specific key\n without matching plugin support makes the page claim support that does not exist.\n- Prose describing ComfyUI support inline instead of using the component is a finding: the\n facts drift from the plugin.\n\n### 6. Commands / port\n- Launch uses `sglang serve` — flag any `python -m sglang.launch_server` /\n `python3 -m sglang.launch_server` (deprecated). The engine already emits `sglang serve`;\n guard against prose/cells reintroducing the old launcher.\n- Port `30000` everywhere (launch, curl, client `base_url`, bench) — flag `8000`.\n Launch port must match client/curl port on the same page.\n\n### 7. Frontmatter\n- Every new MDX page has `title:` and a **top-level** `description:` (a real one-line value\n prop, not copied from another vendor) — NOT `metatags.description` (non-canonical; the\n top-level field is what renders as the subtitle and SEO meta — see mintlify-authoring).\n- **No `mode: wide` on a model page** — it hides the right-hand \"On this page\" ToC that every\n other model page has. Leave `mode` unset (the Deploy/Playground panels self-cap at 900px, so\n the default column holds them fine). `mode: wide` belongs only on category `intro.mdx` grids.\n- `tag: NEW` only for genuine new launches; when one is added, stale `tag: NEW` on older\n pages should be dropped in the same PR (`grep -RlE \"^tag: NEW\" docs/cookbook/`).\n- MDX imports BOTH `Deployment` and `Playground` from `/src/snippets/...` (absolute).\n- Deploy heading slugs to `deployment` (or `deploy`), Playground to `playground` — so\n \"↑ Switch base\" and \"Open the Playground →\" scroll. No numbered headings for these two.\n\n### 8. Navigation & homepage\n- New page → `docs/docs.json` updated: under the right vendor group inside\n `navigation` → Cookbook → Autoregressive Models, root-relative, **no `.mdx`**:\n `cookbook///`.\n- Homepage `` in `docs/cookbook//intro.mdx` points to the vendor's\n flagship; new vendors get a new `` + a logo at `docs/cards/logos/.png` —\n **940×525 RGBA transparent, icon-only (no wordmark)**, lowercase filename, tracked via\n `git add -f` (`*.png` is gitignored repo-wide). Card order matches the `docs.json` nav order.\n- Don't change `docs/cookbook/intro.mdx` for individual model adds (top-level only).\n\n### 9. Links & factual\n- HuggingFace URLs resolve to a real model. License section matches the actual HF license\n (don't copy from another model). Docker images from `lmsysorg/sglang`; no `sgl-project-dev`.\n The image **tag** is a real build (a release the author ran, or `:dev`/nightly) — not a\n guessed version.\n- Internal links root-relative, no extension (`/cookbook/.../`); flag `.md`/`.mdx`\n or `../`-relative links. `docs.sglang.io` is canonical.\n- No Google-Drive image links (don't render). Shell placeholders are `export VAR=`,\n not `${VAR}` (a bash no-op).\n- **Parser ids must exist in the code registries** on the PR's target branch: every\n `--reasoning-parser X` / `--tool-call-parser Y` named in prose or in\n `playgroundFeatures.parsers` flags is a registered key in\n `python/sglang/srt/parser/reasoning_parser.py` (DetectorMap) /\n `python/sglang/srt/function_call/function_call_parser.py` (ToolCallParserEnum) —\n prose naming a near-miss id (e.g. the reasoning id where the tool id differs) is a\n factual bug. `--…-parser auto` is acceptable ONLY if the template-detection rules\n (`python/sglang/srt/parser/template_detection.py`) actually resolve THIS model's\n chat template to the right parser — no rule match means auto silently disables the\n parser; when in doubt require explicit ids (the DSv4 page pins explicit ids).\n\n### 9b. MDX authoring (Mintlify) — detail in `cookbook-add-model/references/mintlify-authoring.md`\n- **Forbidden syntax**: no Docusaurus admonitions (`:::`), `@site`/`@theme`, GitHub alert\n blocks (`> [!NOTE]`), markdown **pipe tables** (use JSX `
`), inline `
`,\n or unknown components. ``/`` only on category `intro.mdx`, not model pages.\n- Code fences are **labeled** (e.g. `python Example` / `bash Command` / `text Output` after\n the opening fence); a fenced block nested inside another uses four backticks outside.\n- §3 commands and outputs are **collapsible** (DeepSeek-V4 pattern): every runnable\n example wrapped in an ``, its real output in a following\n `` (`Pending update...` only with user\n acknowledgement). Flag bare/inline example blocks and `**Output Example:**` headings.\n- Reasoning-parser example matches the parser's **output shape**: separate-field\n (`reasoning_content` + `content`) vs inline `` tags parsed out of `content`.\n- No hardcoded sampling params (`temperature` / `top_p`) in sample code (SGLang uses\n `generation_config.json` defaults); listing them in §1 informationally is fine.\n\n### 10. Quantization rules\n- **NVFP4** checkpoints are Blackwell-only (B200/B300/GB300) — never AMD. An AMD FP4 cell\n is legitimate ONLY when the vendor published an **MXFP4** checkpoint for it (e.g.\n `amd/Qwen3.5-397B-A17B-MXFP4` on MI355X) — verify the HF repo resolves; otherwise the\n AMD FP4 chip must be absent/`disabled`.\n- BF16 / FP8 work on NVIDIA and AMD. `--kv-cache-dtype fp8_e4m3` in a cell is an\n accuracy-degrading flag — see §3 (needs explicit maintainer sign-off; default\n home is Playground/tips).\n\n### 11. Scope\n- Changes match the PR title. Flag global changes hiding behind a platform-specific title\n (e.g. an \"H200 FP8\" PR that adds a flag to ALL cells). Unmentioned side-fixes belong in\n the PR body.\n\n### 12. Duplicate PRs\n- Another open PR for the same model? Flag it; compare completeness; note merge-conflict\n risk on `docs.json` + the vendor card; flag a superseded older PR by the same author.\n\n### 13. Build / validate\n```bash\ncd docs\nmint validate\nmint broken-links\n```\nOptional: `mint dev` for a visual smoke test.\n\n### 14. Reviewer feedback\n- `gh api repos/sgl-project/sglang/pulls//comments` — have prior reviewer requests been\n addressed? Unresolved requested-changes should be flagged.\n\n### 15. Grammar & spelling\n- Check added/changed prose for typos and grammar (e.g. \"recommend\" vs \"recommended\").\n Flag each with the exact wrong text + correction.\n\n## Output\n\nPer file:\n- ✅ PASS\n- ⚠️ ISSUE: \\\n- 🔴 BLOCK: \\\n\nOverall: **APPROVE** / **REQUEST CHANGES** / **BLOCKED**\n","frontmatter":{"name":"cookbook-review-pr","description":"Review a pull request against the SGLang Cookbook (docs/, Mintlify) contribution checklist — the config-driven format (per-model config + benchmarks JSX consumed by the shared _deployment.jsx / _playground.jsx engines). Run with /cookbook-review-pr ."},"isInternal":false,"tokens":4561,"sizeBytes":17126},{"name":"SKILL.md","path":".agents/skills/debug-cuda-crash/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/debug-cuda-crash/SKILL.md","title":"debug-cuda-crash","category":"anthropic-skill","format":"markdown","content":"---\nname: debug-cuda-crash\ndescription: Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging\n---\n\n# Tutorial: Debugging CUDA Crashes with Kernel API Logging\n\nThis tutorial shows you how to debug CUDA crashes and errors in SGLang using the `@debug_kernel_api` logging decorator.\n\n## Goal\n\nWhen your code crashes with CUDA errors such as illegal memory access, device-side assert, out-of-bounds, or NaN/Inf, use kernel API logging to:\n- Capture input tensors BEFORE the crash occurs\n- Understand what data caused the problem\n- Track tensor shapes, dtypes, and values through the call boundary that triggered the crash\n- Detect numerical issues such as NaN, Inf, or obviously wrong shapes\n\n## Why Use Kernel API Logging?\n\n**Problem**: CUDA errors often crash the program before normal debugging output is flushed.\n\n**Solution**: SGLang's `@debug_kernel_api` decorator logs inputs before execution, so you can still see what caused the crash even after the program aborts.\n\n## What Is Covered?\n\nThe current logging coverage focuses on the highest-value kernel boundaries in SGLang:\n- Custom ops registered through `register_custom_op(...)`\n- External custom ops registered through `register_custom_op_from_extern(...)`\n- LLM attention, linear, quantization, and multi-platform wrapper entry points\n- Diffusion attention impl, linear, rotary, and custom-op wrapper entry points\n- Selected direct `torch.ops.sglang.*` hotspots and model-specific bypasses\n\nThis means the logging is useful for both LLM and diffusion kernel debugging, but it does not automatically cover every pure PyTorch call in the repository.\n\n## Step 1: Enable Kernel API Logging\n\n### Basic Logging (Function Names Only)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=1\nexport SGLANG_KERNEL_API_LOGDEST=stdout\n\npython my_script.py\n```\n\nOutput:\n```\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: RMSNorm.forward\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.custom_op.fused_inplace_qknorm\n```\n\nThis is a real level-1 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Detailed Logging (Inputs with Metadata)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nOutput in `debug.log`:\n```\n================================================================================\n[2026-03-19 00:47:30] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[0]=QKVParallelLinear(\n repr=QKVParallelLinear(in_features=1024, output_features=4096, bias=False, tp_size=1, gather_output=False)\n )\n arg[1]=Tensor(\n shape=(1, 1024)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n arg[2]=None\nOutput:\n return=Tensor(\n shape=(1, 4096)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n```\n\nThis is a real level-3 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Full Logging (With Tensor Statistics)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nAdditional output:\n```\n================================================================================\n[2026-03-19 01:00:42] SGLang Kernel API Call: diffusion.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[1]=Tensor(\n shape=(1, 77, 768)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-27.250000\n max=28.500000\n mean=0.011723\n nan_count=0\n inf_count=0\n )\nOutput:\n return=Tensor(\n shape=(1, 77, 2304)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-8.937500\n max=9.375000\n mean=0.009460\n nan_count=0\n inf_count=0\n )\n```\n\nThis is a real level-5 excerpt captured from `black-forest-labs/FLUX.1-dev`.\n\n### Crash-Safe Dumps (Inputs Saved Before Execution)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\n\npython my_script.py\n```\n\nAt level 10, SGLang saves the inputs before execution. If the kernel crashes, the dump directory still contains the inputs and exception metadata.\n\nIf CUDA graph capture is active, tensor dumps are skipped automatically to avoid capture-time CUDA errors. In that case, you still get the kernel API call log, but not `inputs.pt` / `outputs.pt`.\n\nLevel-10 dumps are best understood as crash-safe call snapshots. They always preserve the observed call boundary. They do not guarantee one-click replay for every method, because some methods depend on module state that is not serialized into the dump.\n\nReal level-10 dump layout from `Qwen/Qwen3-0.6B`:\n\n```text\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/inputs.pt\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/metadata.json\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/outputs.pt\n```\n\nReal `metadata.json` excerpt:\n\n```json\n{\n \"function_name\": \"RotaryEmbedding.forward\",\n \"timestamp\": \"20260319_004821_182\",\n \"process_id\": 919286,\n \"execution_status\": \"completed\",\n \"input_tensor_keys\": [\"arg_0\", \"arg_1\", \"arg_2\"],\n \"output_tensor_keys\": [\"result_0\", \"result_1\"]\n}\n```\n\n## Step 2: Reproduce an LLM CUDA Crash\n\nCreate a temporary reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_llm_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.srt.utils.custom_op import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(indices, table):\\\\n\"\n \" return torch.empty((*indices.shape, table.shape[-1]), device=table.device, dtype=table.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_llm_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_llm_cuda_crash(indices, table):\\\\n\"\n \" out = F.embedding(indices, table)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"table = torch.randn(4, 8, device='cuda', dtype=torch.float16)\\\\n\"\n \"indices = torch.tensor([0, 7], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_llm_cuda_crash(indices, table)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level1.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nWhat to expect:\n- The script exits with a CUDA `device-side assert`\n- The log still contains the last API boundary before the crash\n\nTry the same example at level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level3.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow the log shows tensor metadata before the crash.\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_llm_level10_dumps \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow you should see:\n- A log entry for `sglang.custom_op.mock_llm_cuda_crash`\n- A dump directory with `inputs.pt`\n- `metadata.json` showing `execution_status: \"exception\"`\n- No `outputs.pt`, because the kernel crashed before producing output\n\nFor real-model success-path level-10 dumps, it is often easier to temporarily disable CUDA graph and piecewise CUDA graph for the debug run.\n\n## Step 3: Reproduce a Diffusion CUDA Crash\n\nCreate a temporary diffusion-side reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_diffusion_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.multimodal_gen.runtime.layers.utils import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(positions, cache):\\\\n\"\n \" return torch.empty((*positions.shape, cache.shape[-1]), device=cache.device, dtype=cache.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_diffusion_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_diffusion_cuda_crash(positions, cache):\\\\n\"\n \" out = F.embedding(positions, cache)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"cache = torch.randn(4, 64, device='cuda', dtype=torch.float16)\\\\n\"\n \"positions = torch.tensor([0, 9], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_diffusion_cuda_crash(positions, cache)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level1.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level3.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_diffusion_level10_dumps \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nIf your local environment has unrelated FlashInfer import issues, resolve them in the shell before running the example. The example itself does not set any `FLASHINFER_*` environment variable.\n\n## Step 4: Multi-Process Debugging\n\nWhen running with multiple GPUs or worker processes, use `%i` in the log path:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\n\ntorchrun --nproc_per_node=4 my_script.py\n```\n\nThis creates separate logs such as:\n- `debug_rank_12345.log`\n- `debug_rank_12346.log`\n- `debug_rank_12347.log`\n- `debug_rank_12348.log`\n\nReal multi-process example from a 2-GPU `Qwen/Qwen2.5-0.5B-Instruct` run:\n\n```text\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950201.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950349.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950350.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950351.log\n```\n\nYou should usually do the same for level-10 dump directories:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps_%i\n```\n\nThis avoids multiple ranks writing into the same dump directory tree.\n\n## Step 5: Filter Level-10 Dumps\n\nIf level 10 is too noisy, restrict dumps to specific APIs:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\nexport SGLANG_KERNEL_API_DUMP_INCLUDE='sglang.custom_op.*'\nexport SGLANG_KERNEL_API_DUMP_EXCLUDE='*.fake_impl'\n```\n\n`SGLANG_KERNEL_API_DUMP_INCLUDE` and `SGLANG_KERNEL_API_DUMP_EXCLUDE` use shell-style wildcard matching.\n\n## Step 6: Common CUDA Errors and What to Check\n\n### Illegal Memory Access or Device-Side Assert\n\n**Typical errors**:\n```\nRuntimeError: CUDA error: an illegal memory access was encountered\ntorch.AcceleratorError: CUDA error: device-side assert triggered\n```\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck in the logs:\n- ✅ Tensor shapes\n- ✅ Tensor dtypes\n- ✅ CUDA vs CPU device placement\n- ✅ Tensor stride / contiguity\n- ✅ Whether the failing call has inputs logged but no outputs logged\n\nTypical shape-mismatch pattern:\n\n```text\nSGLang Kernel API Call: ...\narg[0]=Tensor(shape=(..., 128), ...) # ✅ expected dimension\narg[1]=Tensor(shape=(..., 64), ...) # ❌ mismatch\n```\n\nThis often points to head-dim, hidden-dim, or cache-layout mismatch rather than a random CUDA failure.\n\n### NaN or Inf\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nCheck:\n- `min`\n- `max`\n- `mean`\n- `nan_count`\n- `inf_count`\n\nTypical bad pattern:\n\n```text\nTensor(\n ...\n min=-1234567.000000 # ❌ suspiciously large\n max=9876543.000000 # ❌ suspiciously large\n mean=nan # ❌ bad\n nan_count=128 # ❌ found NaNs\n inf_count=0 # ✅ no Infs here\n)\n```\n\nThis usually means the bad values were already present before the crashing kernel.\n\n### Out of Memory\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck:\n- Unexpectedly large tensor shapes\n- Batch size\n- Sequence length\n- Frame count or image resolution in diffusion workloads\n\nAlso check whether a supposedly per-token or per-frame tensor accidentally became full-sequence or full-image sized.\n\nTypical bad pattern:\n\n```text\nTensor(\n shape=(1024, 8192, 128, 128) # ❌ way too large\n ...\n)\n```\n\n### Example: Spot a Shape Bug from the Log\n\nSuppose the failing API log looks like this:\n\n```text\n[2026-03-19 00:47:30] SGLang Kernel API Call: RotaryEmbedding.forward\nPositional input arguments:\n arg[0]=Tensor(shape=(1, 8), dtype=torch.int64, ...)\n arg[1]=Tensor(shape=(1, 8, 8, 256), dtype=torch.bfloat16, ...) # ✅ query\n arg[2]=Tensor(shape=(1, 8, 4, 64), dtype=torch.bfloat16, ...) # ❌ key head_dim mismatch\n```\n\nWhat this tells you:\n- ✅ positions look reasonable\n- ✅ query looks plausible\n- ❌ key last dimension is inconsistent with the expected rotary/head dimension\n\nThat usually means the bug is in projection layout, head packing, or cache format rather than in the rotary kernel itself.\n\n## Step 7: Combine with compute-sanitizer\n\nFor harder bugs, combine kernel API logging with CUDA memory checking:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncompute-sanitizer --tool memcheck python3 /tmp/sglang_llm_crash.py\n```\n\nUse `debug.log` to see the exact inputs that reached the crashing API boundary.\n\nTypical `compute-sanitizer` output:\n\n```text\n========= COMPUTE-SANITIZER\n========= Invalid __global__ write of size 4 bytes\n========= at 0x1234 in SomeKernel\n========= by thread (256,0,0) in block (10,0,0)\n========= Address 0x... is out of bounds\n```\n\nUse the sanitizer output to identify the failing kernel and use `debug.log` to identify the exact tensors that reached the API boundary right before it.\n\nIf you need more synchronous host-side error reporting, you can try `CUDA_LAUNCH_BLOCKING=1` as a separate follow-up experiment. It is not part of the default workflow because it changes execution timing and can hide concurrency-related behavior.\n\n## Step 8: Combine with cuda-gdb\n\nFor crashes that need a stack trace instead of only memory diagnostics:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncuda-gdb --args python3 /tmp/sglang_llm_crash.py\n```\n\nInside `cuda-gdb`:\n\n```text\n(cuda-gdb) run\n(cuda-gdb) where\n```\n\nThen correlate the backtrace with `debug.log`.\n\n## Step 9: Kernel-Level Debugging with printf()\n\nWhen you own the CUDA kernel, `printf()` is still useful for narrowing down bad indices, bad launch geometry, or broken state propagation.\n\nBasic pattern:\n\n```cpp\n__global__ void MyKernel(const float* input, float* output, int n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n\n if (threadIdx.x == 0 && blockIdx.x == 0) {\n printf(\"n=%d input0=%f\\n\", n, input[0]);\n }\n\n if (idx < n) {\n output[idx] = input[idx] * 2.0f;\n }\n}\n```\n\nAfter launch, force the output to flush:\n\n```python\nmy_kernel(...)\ntorch.cuda.synchronize()\n```\n\nFor warp-specialized kernels, do not blindly print only on `threadIdx.x == 0`. Pick one representative thread per warp or per specialization group instead.\n\n### Warp-Specialized Kernels: Choosing the Right Print Thread\n\nProblem:\n- `threadIdx.x == 0` only prints from the first warp in the block\n- for warp-specialized kernels, that often misses the warp or group that is actually wrong\n\nBetter pattern:\n\n```cpp\n__global__ void WarpSpecializedKernel(...) {\n // Example: first lane of each warp\n if ((threadIdx.x % 32) == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n }\n}\n```\n\nOr, if the kernel is organized in larger specialization groups, print once per group instead of once per block.\n\nCommon mistake:\n\n```cpp\n// Only warp 0 prints\nif (threadIdx.x == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n}\n```\n\n### Quick Reference\n\n| Kernel Type | Print Condition | Notes |\n|----------|----------|-------------|\n| Simple kernel | `threadIdx.x == 0` | One thread per block is usually enough |\n| Warp-specialized kernel | one representative lane per warp | e.g. `threadIdx.x % 32 == 0` |\n| Group-specialized kernel | one representative lane per group | choose based on the kernel's scheduling layout |\n\n### Other Kernel Debugging Tools\n\n```cpp\nassert(value >= 0.0f && \"value must be non-negative\");\nstatic_assert(BLOCK_SIZE % 32 == 0, \"BLOCK_SIZE must be warp aligned\");\n```\n\n## Environment Variables Reference\n\n| Variable | Values | Description |\n|----------|--------|-------------|\n| `SGLANG_KERNEL_API_LOGLEVEL` | `0` | No logging (default) |\n| | `1` | Function names only |\n| | `3` | Inputs and outputs with metadata |\n| | `5` | Level 3 plus tensor statistics |\n| | `10` | Level 5 plus crash-safe tensor dumps |\n| `SGLANG_KERNEL_API_LOGDEST` | `stdout` | Log to stdout |\n| | `stderr` | Log to stderr |\n| | `` | Log to file |\n| | `log_%i.txt` | `%i` expands to process ID |\n| `SGLANG_KERNEL_API_DUMP_DIR` | `` | Directory for level-10 dumps |\n| `SGLANG_KERNEL_API_DUMP_INCLUDE` | wildcard list | Only dump matching API names |\n| `SGLANG_KERNEL_API_DUMP_EXCLUDE` | wildcard list | Skip matching API names |\n\n## Best Practices\n\n### 1. Start with Level 3\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nLevel 3 is usually enough to catch wrong shapes, wrong dtypes, and wrong devices.\n\n### 2. Use Level 5 for Numerical Issues\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nUse it when you suspect NaN or Inf values.\n\n### 3. Use Level 10 for Crash Reproduction\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\n```\n\nThis is the most useful mode when the process crashes before you can inspect live tensors.\n\nIf you need successful input/output dumps from a real model run, temporarily disable CUDA graph for that debug session.\n\nWhen level 10 is too noisy, pair it with `SGLANG_KERNEL_API_DUMP_INCLUDE` / `SGLANG_KERNEL_API_DUMP_EXCLUDE` instead of dumping every covered API.\n\n### 4. Log to File for Crashes\n\n```bash\nexport SGLANG_KERNEL_API_LOGDEST=crash.log\n```\n\nFile logs are safer than stdout when the process aborts.\n\n### 5. Disable Logging in Production\n\n```bash\nunset SGLANG_KERNEL_API_LOGLEVEL\n```\n\nWhen disabled, the decorator returns the original callable and adds no runtime logging overhead.\n\n## Troubleshooting\n\n### No Logs Appear\n\nCheck:\n1. `echo $SGLANG_KERNEL_API_LOGLEVEL`\n2. `echo $SGLANG_KERNEL_API_LOGDEST`\n3. Whether the failing path goes through a covered API boundary\n\n### Too Much Output\n\nReduce the level:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\n### Statistics Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nstatistics=[skipped: CUDA graph capture in progress]\n```\n\nThat is expected. Level-5 statistics are intentionally skipped during CUDA graph capture to avoid synchronization side effects.\n\n### Tensor Dumps Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nTensor dump skipped: CUDA graph capture in progress\n```\n\nThat is also expected. Level-10 dumps require copying tensors to CPU, which is not allowed during CUDA graph capture.\n","frontmatter":{"name":"debug-cuda-crash","description":"Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging"},"isInternal":false,"tokens":5257,"sizeBytes":19737},{"name":"SKILL.md","path":".agents/skills/debug-distributed-hang/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/debug-distributed-hang/SKILL.md","title":"debug-distributed-hang","category":"anthropic-skill","format":"markdown","content":"---\nname: debug-distributed-hang\ndescription: Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations.\n---\n\n# Debugging Distributed Hangs in SGLang\n\n## Overview\n\nHangs in distributed inference happen when ranks diverge in state, causing collective operations (AllGather, AllReduce, Broadcast, Barrier) to deadlock. Common causes:\n\n- **Size mismatch**: ranks pass different tensor sizes to a collective\n- **Branch divergence**: one rank enters a collective, another skips it\n- **Cascading state drift**: a small non-determinism (e.g., floating-point) propagates into different batch structures\n- **Resource exhaustion**: one rank OOMs or crashes, others wait forever\n\n## Prerequisites\n\n- **py-spy**: `pip install py-spy` or system package. Requires root or `CAP_SYS_PTRACE` to attach to running processes.\n- **cuda-gdb**: Ships with the CUDA toolkit. Ensure it's on your `PATH`.\n\n## Step 1: Confirm and Locate the Hang\n\n### 1a. Watchdog / py-spy\n\nSGLang's watchdog automatically dumps py-spy traces on timeout. Look for:\n\n```\nScheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)\n```\n\nThe py-spy dump shows the stack trace of each thread. The hanging thread is typically blocked in a CUDA synchronize or NCCL collective:\n\n```\nThread (active): \"MainThread\"\n cuStreamSynchronize (libcuda.so)\n ...\n forward_extend (model_runner.py)\n```\n\nSGLang has two watchdog modes (see `python/sglang/srt/utils/watchdog.py`):\n- **Hard watchdog** (`soft=False`, default): dumps py-spy traces then sends `SIGQUIT` to kill the parent process.\n- **Soft watchdog** (`soft=True`): only logs the timeout without killing the process, giving you more time to manually attach debuggers or collect coredumps.\n\nIf the watchdog doesn't trigger, manually dump:\n\n```bash\npy-spy dump --pid \n```\n\n### 1b. NCCL Debug Logging\n\n```bash\nexport NCCL_DEBUG=INFO\nexport NCCL_DEBUG_SUBSYS=COLL\n```\n\nLook for the last collective logged before the hang. Mismatched sizes show up as one rank waiting and another never entering.\n\n### 1c. CUDA Coredump\n\nWhen a process hangs, you can trigger a GPU coredump on demand to see which kernel is stuck. Set these env vars before launching:\n\n```bash\nexport CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1\nexport CUDA_COREDUMP_PIPE=\"/tmp/cuda_pipe_%h_%p\"\nexport CUDA_COREDUMP_FILE=\"/tmp/cuda_coredump_%h_%p\"\nexport CUDA_COREDUMP_SHOW_PROGRESS=1\nexport CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'\n```\n\nWhile the process is hanging, find the pipe via `/proc//fd/` and write to it to trigger the dump:\n\n```bash\nls /proc//fd/ -la 2>/dev/null | grep cuda_pipe\ndd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe__\n```\n\nAlternatively, if you don't need to keep the process alive, `kill -SIGABRT ` also triggers a CUDA coredump (but terminates the process).\n\nThen open with `cuda-gdb --batch -ex \"target cudacore \"`. On load, it immediately shows which kernel is stuck. For example:\n\n```\nOpening GPU coredump: \n[Current focus set to CUDA kernel 0, grid 622721, cluster (4,0,0), block (16,0,0), thread (64,0,0), device 0, sm 0, warp 0, lane 0]\n#0 0x00007f8029b2b040 in ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)<<<(24,1,1),(512,1,1)>>> ()\n```\n\nThis told us the hang was in an NCCL AllGather — not a compute kernel. Combined with the py-spy stack pointing to `LogitsProcessor.forward` → `tensor_model_parallel_all_gather`, we knew it was an AllGather size mismatch between TP ranks.\n\n\n### 1d. Identify the Collective\n\nFrom the stack traces and logs, identify:\n- Which collective hangs (AllGather, AllReduce, Broadcast)\n- Which code path invokes it (e.g., `LogitsProcessor`, `tensor_model_parallel_all_gather`)\n- Whether it's a size mismatch or a missing participant\n\n## Step 2: Per-Rank Logging\n\nThe key technique: each rank writes its own log file so you can diff them.\n\n### Setup Pattern\n\n```python\nimport os\n\n_debug_files = {}\n\ndef get_debug_file(rank):\n key = f\"rank{rank}\"\n if key not in _debug_files:\n _debug_files[key] = open(f\"/tmp/debug_rank{rank}.log\", \"w\")\n return _debug_files[key]\n```\n\nGate logging behind an env var to avoid overhead in production. `SGLANG_DEBUG_HANG` is not a built-in SGLang env var — you need to add this check yourself in the code you're instrumenting:\n\n```python\nif os.environ.get(\"SGLANG_DEBUG_HANG\"):\n f = get_debug_file(rank)\n f.write(f\"EVENT_NAME key1={val1} key2={val2}\\n\")\n f.flush()\n```\n\n### What to Log\n\nLog structured events at key state-mutation points:\n\n```python\nf.write(f\"SCHED_BATCH step={step} num_reqs={n} extend_lens={lens}\\n\")\nf.write(f\"VERIFY predict_hash={hash} accept_len={alen}\\n\")\nf.write(f\"CACHE_INSERT rid={rid} num_tokens={n}\\n\")\n```\n\nUse consistent event names (uppercase prefix) for easy grep/diff.\n\n### Hash Large Tensors\n\nFor tensor values, compute a hash instead of dumping raw data:\n\n```python\nimport hashlib\nh = hashlib.md5(tensor.cpu().numpy().tobytes()).hexdigest()[:8]\nf.write(f\"LOGITS logits_hash={h}\\n\")\n```\n\nFor token ID lists, `str(list).encode()` works:\n\n```python\nh = hashlib.md5(str(tensor.tolist()).encode()).hexdigest()[:8]\n```\n\n### Avoid Implicit Synchronization\n\n`tensor.cpu()`, `tensor.tolist()`, and `tensor.numpy()` all trigger CUDA synchronization. This can:\n- Change timing and mask or move the hang\n- Deadlock if the log point is between two collectives that must run back-to-back\n\nPrefer logging values that are already on CPU (e.g., Python ints, list lengths, request IDs). When you must hash a GPU tensor, do it at a point where the GPU is already idle (e.g., between scheduler steps, not inside a model forward pass).\n\n## Step 3: Diff to Find the Diverge Point\n\n### Basic Diff\n\n```bash\n# Extract specific event type\ngrep \"^VERIFY\" /tmp/debug_rank0.log > /tmp/v_r0.txt\ngrep \"^VERIFY\" /tmp/debug_rank1.log > /tmp/v_r1.txt\ndiff /tmp/v_r0.txt /tmp/v_r1.txt | head -20\n```\n\n### Count Events\n\n```bash\ngrep -c \"^VERIFY\" /tmp/debug_rank*.log\n```\n\nIf counts differ, one rank executed more iterations — that's already a diverge signal.\n\n### Find First Diverge\n\nThe first diff line tells you the exact step where ranks diverge. All lines before it are identical — the root cause is at or before this step.\n\n## Step 4: Binary-Search the Root Cause\n\nOnce you find the diverging event, trace backwards:\n\n### 4a. Identify Inputs\n\nFor the diverging operation, list all its inputs. Add hash logging for each:\n\n```python\nf.write(\n f\"OP_INPUTS input_a_hash={h_a} input_b_hash={h_b} \"\n f\"input_c_hash={h_c} input_d_hash={h_d}\\n\"\n)\n```\n\n### 4b. Diff Inputs Across Ranks\n\nCompare the hashes. Some inputs will match, some won't. The non-matching input is where divergence entered.\n\n### 4c. Recurse\n\nFor the non-matching input, trace where it was produced and repeat: hash its inputs, diff across ranks, find the divergent one. Continue until you reach the root cause.\n\n## Step 5: Common Root Causes and Fixes\n\n### Floating-Point Non-Determinism\n\n**Symptom**: All \"logical\" inputs are identical (same logits after all-gather), but derived floating-point values (softmax, probabilities) differ across GPUs.\n\n**Example**: EAGLE speculative decoding — `F.softmax` → `top_k_renorm_prob` → `top_p_renorm_prob` produces slightly different `target_probs` on each GPU. The sampling kernel then picks different tokens. These flow into `output_ids` → radix cache → different prefix match depths → different `extend_seq_lens` → AllGather size mismatch → hang.\n\n### Random Number Divergence\n\n**Symptom**: Operations using `torch.rand` produce different values on each rank.\n\n**Fix**: Generate on rank 0 and broadcast, or use a shared seed.\n\n### Conditional Code Paths\n\n**Symptom**: A condition (e.g., memory check, queue length) evaluates differently on different ranks, causing one rank to enter a collective while another skips it.\n\n**Fix**: Synchronize the condition value before branching, or restructure to ensure all ranks take the same path.\n\n### Pipeline Parallel (PP) Send/Recv Mismatch\n\n**Symptom**: In PP setups, one stage issues a `send` that the next stage never `recv`s (or vice versa), causing both to block indefinitely. Unlike TP hangs (collective mismatches), PP hangs typically involve point-to-point operations.\n\n**Fix**: Ensure all stages agree on the number of microbatches and the sequence of send/recv calls for each microbatch.\n\n## Step 6: Verify the Fix\n\nRun the failing test multiple times to confirm the fix is stable. Intermittent hangs require many runs. A test that hung ~30% of the time needs at least 10 clean passes to be confident.\n\n## Quick Reference\n\n| Technique | When to Use |\n|-----------|-------------|\n| py-spy dump | First step — see where each rank is stuck |\n| `NCCL_DEBUG=INFO` | Identify which collective and sizes |\n| CUDA coredump + `cuda-gdb` | See which GPU kernel is blocked |\n| Per-rank log files | Compare rank states over time |\n| Hash of tensors | Efficiently compare large tensors across ranks |\n| `diff` on extracted events | Find the exact step of divergence |\n| `broadcast(result, src=0)` | Fix floating-point or sampling non-determinism |\n","frontmatter":{"name":"debug-distributed-hang","description":"Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations."},"isInternal":false,"tokens":2413,"sizeBytes":9504},{"name":"SKILL.md","path":".agents/skills/env-var-conventions/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/env-var-conventions/SKILL.md","title":"env-var-conventions","category":"anthropic-skill","format":"markdown","content":"# env-var-conventions\n\nPath: `.agents/skills/env-var-conventions/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/env-var-conventions/SKILL.md)","tokens":300,"sizeBytes":207},{"name":"SKILL.md","path":".agents/skills/generate-profile/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/generate-profile/SKILL.md","title":"generate-profile","category":"anthropic-skill","format":"markdown","content":"# generate-profile\n\nPath: `.agents/skills/generate-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/generate-profile/SKILL.md)","tokens":300,"sizeBytes":198},{"name":"SKILL.md","path":".agents/skills/kernel-organization/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kernel-organization/SKILL.md","title":"kernel-organization","category":"anthropic-skill","format":"markdown","content":"# kernel-organization\n\nPath: `.agents/skills/kernel-organization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kernel-organization/SKILL.md)","tokens":300,"sizeBytes":207},{"name":"SKILL.md","path":".agents/skills/kl-consistency-test/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kl-consistency-test/SKILL.md","title":"kl-consistency-test","category":"anthropic-skill","format":"markdown","content":"# kl-consistency-test\n\nPath: `.agents/skills/kl-consistency-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kl-consistency-test/SKILL.md)","tokens":300,"sizeBytes":207},{"name":"SKILL.md","path":".agents/skills/large-class-style/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/large-class-style/SKILL.md","title":"large-class-style","category":"anthropic-skill","format":"markdown","content":"# large-class-style\n\nPath: `.agents/skills/large-class-style/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/large-class-style/SKILL.md)","tokens":300,"sizeBytes":201},{"name":"fuse-overlap-catalog.md","path":".agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md)","tokens":300,"sizeBytes":272},{"name":"heuristics.md","path":".agents/skills/llm-torch-profiler-analysis/references/heuristics.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/heuristics.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/heuristics.md)","tokens":300,"sizeBytes":252},{"name":"overlap-catalog.md","path":".agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/overlap-catalog.md)","tokens":300,"sizeBytes":262},{"name":"source-map.md","path":".agents/skills/llm-torch-profiler-analysis/references/source-map.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/source-map.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/source-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/source-map.md)","tokens":300,"sizeBytes":252},{"name":"vllm-torch-compile-fusions.md","path":".agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/vllm-torch-compile-fusions.md)","tokens":300,"sizeBytes":284},{"name":"SKILL.md","path":".agents/skills/llm-torch-profiler-analysis/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/SKILL.md","title":"llm-torch-profiler-analysis","category":"anthropic-skill","format":"markdown","content":"# llm-torch-profiler-analysis\n\nPath: `.agents/skills/llm-torch-profiler-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/SKILL.md)","tokens":300,"sizeBytes":231},{"name":"guide-construct-proof.md","path":".agents/skills/mechanical-refactor-verify/guide-construct-proof.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-construct-proof.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-construct-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-construct-proof.md)","tokens":300,"sizeBytes":266},{"name":"guide-modify-skill.md","path":".agents/skills/mechanical-refactor-verify/guide-modify-skill.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-modify-skill.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-modify-skill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-modify-skill.md)","tokens":300,"sizeBytes":260},{"name":"guide-split.md","path":".agents/skills/mechanical-refactor-verify/guide-split.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-split.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-split.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-split.md)","tokens":300,"sizeBytes":246},{"name":"guide-verify-proof.md","path":".agents/skills/mechanical-refactor-verify/guide-verify-proof.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-verify-proof.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/guide-verify-proof.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/guide-verify-proof.md)","tokens":300,"sizeBytes":260},{"name":"SKILL.md","path":".agents/skills/mechanical-refactor-verify/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/SKILL.md","title":"mechanical-refactor-verify","category":"anthropic-skill","format":"markdown","content":"# mechanical-refactor-verify\n\nPath: `.agents/skills/mechanical-refactor-verify/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/SKILL.md)","tokens":300,"sizeBytes":228},{"name":"spec-reproduction-cli.md","path":".agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-cli.md)","tokens":300,"sizeBytes":266},{"name":"spec-reproduction-utils.md","path":".agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md","title":"Mechanical-refactor-verify Skill","category":"anthropic-skill","format":"markdown","content":"# Mechanical-refactor-verify Skill\n\nPath: `.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/mechanical-refactor-verify/spec-reproduction-utils.md)","tokens":300,"sizeBytes":270},{"name":"SKILL.md","path":".agents/skills/processor-model-parity/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/processor-model-parity/SKILL.md","title":"processor-model-parity","category":"anthropic-skill","format":"markdown","content":"# processor-model-parity\n\nPath: `.agents/skills/processor-model-parity/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/processor-model-parity/SKILL.md)","tokens":300,"sizeBytes":216},{"name":"SKILL.md","path":".agents/skills/scripted-runtime-notes/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/scripted-runtime-notes/SKILL.md","title":"scripted-runtime-notes","category":"anthropic-skill","format":"markdown","content":"# scripted-runtime-notes\n\nPath: `.agents/skills/scripted-runtime-notes/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/scripted-runtime-notes/SKILL.md)","tokens":300,"sizeBytes":216},{"name":"SKILL.md","path":".agents/skills/sglang-bisect-ci-regression/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-bisect-ci-regression/SKILL.md","title":"sglang-bisect-ci-regression","category":"anthropic-skill","format":"markdown","content":"# sglang-bisect-ci-regression\n\nPath: `.agents/skills/sglang-bisect-ci-regression/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-bisect-ci-regression/SKILL.md)","tokens":300,"sizeBytes":231},{"name":"SKILL.md","path":".agents/skills/sglang-cherrypick/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-cherrypick/SKILL.md","title":"sglang-cherrypick","category":"anthropic-skill","format":"markdown","content":"# sglang-cherrypick\n\nPath: `.agents/skills/sglang-cherrypick/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-cherrypick/SKILL.md)","tokens":300,"sizeBytes":201},{"name":"case-studies.md","path":".agents/skills/sglang-prod-incident-triage/references/case-studies.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/case-studies.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/case-studies.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/case-studies.md)","tokens":300,"sizeBytes":256},{"name":"decision-tree.md","path":".agents/skills/sglang-prod-incident-triage/references/decision-tree.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/decision-tree.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/decision-tree.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/decision-tree.md)","tokens":300,"sizeBytes":258},{"name":"endpoints-and-signals.md","path":".agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/endpoints-and-signals.md)","tokens":300,"sizeBytes":274},{"name":"replay-trace-profile.md","path":".agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/references/replay-trace-profile.md)","tokens":300,"sizeBytes":272},{"name":"SKILL.md","path":".agents/skills/sglang-prod-incident-triage/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/SKILL.md","title":"sglang-prod-incident-triage","category":"anthropic-skill","format":"markdown","content":"# sglang-prod-incident-triage\n\nPath: `.agents/skills/sglang-prod-incident-triage/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-prod-incident-triage/SKILL.md)","tokens":300,"sizeBytes":231},{"name":"SKILL.md","path":".agents/skills/sglang-runtime-context/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-runtime-context/SKILL.md","title":"sglang-runtime-context","category":"anthropic-skill","format":"markdown","content":"# sglang-runtime-context\n\nPath: `.agents/skills/sglang-runtime-context/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/sglang-runtime-context/SKILL.md)","tokens":300,"sizeBytes":216},{"name":"SKILL.md","path":".agents/skills/speculative-naming/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/speculative-naming/SKILL.md","title":"speculative-naming","category":"anthropic-skill","format":"markdown","content":"# speculative-naming\n\nPath: `.agents/skills/speculative-naming/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/speculative-naming/SKILL.md)","tokens":300,"sizeBytes":204},{"name":"SKILL.md","path":".agents/skills/write-sglang-test/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/write-sglang-test/SKILL.md","title":"write-sglang-test","category":"anthropic-skill","format":"markdown","content":"# write-sglang-test\n\nPath: `.agents/skills/write-sglang-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/write-sglang-test/SKILL.md)","tokens":300,"sizeBytes":201},{"name":"comment-style.md","path":".claude/rules/comment-style.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/comment-style.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/comment-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/comment-style.md)","tokens":300,"sizeBytes":173},{"name":"cookbook-first.md","path":".claude/rules/cookbook-first.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/cookbook-first.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/cookbook-first.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/cookbook-first.md)","tokens":300,"sizeBytes":175},{"name":"forward-batch-init-new-purity.md","path":".claude/rules/forward-batch-init-new-purity.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/forward-batch-init-new-purity.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/forward-batch-init-new-purity.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/forward-batch-init-new-purity.md)","tokens":300,"sizeBytes":205},{"name":"general-code-style.md","path":".claude/rules/general-code-style.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/general-code-style.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/general-code-style.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/general-code-style.md)","tokens":300,"sizeBytes":183},{"name":"kernel-benchmark.md","path":".claude/rules/kernel-benchmark.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/kernel-benchmark.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/kernel-benchmark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/kernel-benchmark.md)","tokens":300,"sizeBytes":179},{"name":"modify-component-must-read.md","path":".claude/rules/modify-component-must-read.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/modify-component-must-read.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/modify-component-must-read.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/modify-component-must-read.md)","tokens":300,"sizeBytes":199},{"name":"no-dataclasses.md","path":".claude/rules/no-dataclasses.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-dataclasses.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/no-dataclasses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-dataclasses.md)","tokens":300,"sizeBytes":175},{"name":"no-getattr-defensive.md","path":".claude/rules/no-getattr-defensive.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-getattr-defensive.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/no-getattr-defensive.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/no-getattr-defensive.md)","tokens":300,"sizeBytes":187},{"name":"prefer-sgl-eval.md","path":".claude/rules/prefer-sgl-eval.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/prefer-sgl-eval.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/prefer-sgl-eval.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/prefer-sgl-eval.md)","tokens":300,"sizeBytes":177},{"name":"schedule-batch-out-of-place-mutation.md","path":".claude/rules/schedule-batch-out-of-place-mutation.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/schedule-batch-out-of-place-mutation.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/schedule-batch-out-of-place-mutation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/schedule-batch-out-of-place-mutation.md)","tokens":300,"sizeBytes":219},{"name":"unit-test-admission.md","path":".claude/rules/unit-test-admission.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/unit-test-admission.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `.claude/rules/unit-test-admission.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.claude/rules/unit-test-admission.md)","tokens":300,"sizeBytes":185},{"name":"testing-and-accuracy.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/references/testing-and-accuracy.md)","tokens":300,"sizeBytes":328},{"name":"SKILL.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md","title":"sglang-diffusion-add-model","category":"anthropic-skill","format":"markdown","content":"# sglang-diffusion-add-model\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-add-model/SKILL.md)","tokens":300,"sizeBytes":286},{"name":"benchmark-and-profile.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md","title":"Sglang-diffusion-benchmark-profile Skill","category":"anthropic-skill","format":"markdown","content":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/benchmark-and-profile.md)","tokens":300,"sizeBytes":348},{"name":"existing-fast-paths.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md","title":"Sglang-diffusion-benchmark-profile Skill","category":"anthropic-skill","format":"markdown","content":"# Sglang-diffusion-benchmark-profile Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/existing-fast-paths.md)","tokens":300,"sizeBytes":344},{"name":"SKILL.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md","title":"sglang-diffusion-benchmark-profile","category":"anthropic-skill","format":"markdown","content":"# sglang-diffusion-benchmark-profile\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-benchmark-profile/SKILL.md)","tokens":300,"sizeBytes":310},{"name":"block_adapter.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/references/block_adapter.md)","tokens":300,"sizeBytes":314},{"name":"SKILL.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md","title":"sglang-diffusion-cache-dit","category":"anthropic-skill","format":"markdown","content":"# sglang-diffusion-cache-dit\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-cache-dit/SKILL.md)","tokens":300,"sizeBytes":286},{"name":"SKILL.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md","title":"sglang-diffusion-modelopt-quant","category":"anthropic-skill","format":"markdown","content":"# sglang-diffusion-modelopt-quant\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-modelopt-quant/SKILL.md)","tokens":300,"sizeBytes":301},{"name":"SKILL.md","path":"python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md","title":"sglang-diffusion-performance","category":"anthropic-skill","format":"markdown","content":"# sglang-diffusion-performance\n\nPath: `python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/python/sglang/multimodal_gen/.agents/skills/sglang-diffusion-performance/SKILL.md)","tokens":300,"sizeBytes":292},{"name":"replay_table.json","path":"tools/sglang-simulator/examples/assets/replay_table.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/replay_table.json","title":"Assets Skill","category":"anthropic-skill","format":"json","content":"# Assets Skill\n\nPath: `tools/sglang-simulator/examples/assets/replay_table.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/replay_table.json)","tokens":300,"sizeBytes":226},{"name":"tokenizer_config.json","path":"tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json","title":"Tokenizer Skill","category":"anthropic-skill","format":"json","content":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer_config.json)","tokens":300,"sizeBytes":257},{"name":"tokenizer.json","path":"tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json","title":"Tokenizer Skill","category":"anthropic-skill","format":"json","content":"# Tokenizer Skill\n\nPath: `tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/assets/tokenizer/tokenizer.json)","tokens":300,"sizeBytes":243},{"name":"aic_silicon.json","path":"tools/sglang-simulator/examples/sim_configs/aic_silicon.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_silicon.json","title":"Sim_configs Skill","category":"anthropic-skill","format":"json","content":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_silicon.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_silicon.json)","tokens":300,"sizeBytes":239},{"name":"aic_sol.json","path":"tools/sglang-simulator/examples/sim_configs/aic_sol.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_sol.json","title":"Sim_configs Skill","category":"anthropic-skill","format":"json","content":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/aic_sol.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/aic_sol.json)","tokens":300,"sizeBytes":231},{"name":"ml.json","path":"tools/sglang-simulator/examples/sim_configs/ml.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/ml.json","title":"Sim_configs Skill","category":"anthropic-skill","format":"json","content":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/ml.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/ml.json)","tokens":300,"sizeBytes":221},{"name":"replay.json","path":"tools/sglang-simulator/examples/sim_configs/replay.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/replay.json","title":"Sim_configs Skill","category":"anthropic-skill","format":"json","content":"# Sim_configs Skill\n\nPath: `tools/sglang-simulator/examples/sim_configs/replay.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/sim_configs/replay.json)","tokens":300,"sizeBytes":229},{"name":"sharegpt-example.json","path":"tools/sglang-simulator/examples/workloads/sharegpt-example.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/workloads/sharegpt-example.json","title":"Workloads Skill","category":"anthropic-skill","format":"json","content":"# Workloads Skill\n\nPath: `tools/sglang-simulator/examples/workloads/sharegpt-example.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/workloads/sharegpt-example.json)","tokens":300,"sizeBytes":243},{"name":"config.json","path":"tools/sglang-simulator/test/assets/qwen3-8b/config.json","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/test/assets/qwen3-8b/config.json","title":"Qwen3-8b Skill","category":"anthropic-skill","format":"json","content":"# Qwen3-8b Skill\n\nPath: `tools/sglang-simulator/test/assets/qwen3-8b/config.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/test/assets/qwen3-8b/config.json)","tokens":300,"sizeBytes":226},{"name":"README.md","path":"examples/profiler/nsys_profile_tools/README.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/examples/profiler/nsys_profile_tools/README.md","title":"nsys_profile_tools Documentation","category":"plugin-manifest","format":"markdown","content":"# nsys_profile_tools Documentation\n\nPath: `examples/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/examples/profiler/nsys_profile_tools/README.md)","tokens":300,"sizeBytes":226},{"name":"README.md","path":"tools/sglang-simulator/examples/README.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/README.md","title":"examples Documentation","category":"plugin-manifest","format":"markdown","content":"# examples Documentation\n\nPath: `tools/sglang-simulator/examples/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/examples/README.md)","tokens":300,"sizeBytes":206},{"name":"README.md","path":"tools/sglang-simulator/README.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/README.md","title":"sglang-simulator Documentation","category":"plugin-manifest","format":"markdown","content":"# sglang-simulator Documentation\n\nPath: `tools/sglang-simulator/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/README.md)","tokens":300,"sizeBytes":196},{"name":"README.md","path":"tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md","title":"collect_baseline Documentation","category":"plugin-manifest","format":"markdown","content":"# collect_baseline Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/collect_baseline/README.md)","tokens":300,"sizeBytes":284},{"name":"README.md","path":"tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md","rawUrl":"https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md","title":"train_ml_time_predictor Documentation","category":"plugin-manifest","format":"markdown","content":"# train_ml_time_predictor Documentation\n\nPath: `tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/tools/sglang-simulator/src/sglang_simulator/tools/train_ml_time_predictor/README.md)","tokens":300,"sizeBytes":305}],"systemPromptSnippet":"\n\n# docs - AI Agent Protocol & Instructions\n\nPath: `docs/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/docs/AGENTS.md)\n\n\n# sgl-router - System Architecture & Module Boundaries\n\nPath: `sgl-router/ARCHITECTURE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/sgl-router/ARCHITECTURE.md)\n\n\n---\nname: add-jit-kernel\ndescription: Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups\n---\n\n# Tutorial: Adding a New JIT Kernel to SGLang\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as a JIT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float, passed at runtime)\n- Output: `x * factor` (element-wise), allocated internally\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n\n## When to use JIT vs AOT (`sgl-kernel`)\n\n- **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.\n- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `python/sglang/kernels/aot/` and participate in the wheel build / torch op registration flow.\n- **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.\n\n---\n\n## Conventions\n\nThese hold for every step below.\n\n- **`namespace sglang` is where JIT code lives.** Open it after the include block and close it at the end of the file, with the device kernels, traits and host wrapper inside. The shared `host::` / `device::` helpers are nested in it too, so they resolve unqualified. `load_jit` emits the `TVM_FFI_DLL_EXPORT_TYPED_FUNC` wrapper inside `namespace sglang` as well, so the `kernel_name` you pass from Python needs no `sglang::` prefix.\n- **Tuning policy belongs in Python; C++ gets the answer, not the decision.** For a hyperparameter — split factor, threads per item, block size, vector width, an algorithm variant — prefer (not required) making it a **template parameter** and letting the `@cache_once` module factory choose the value, over a runtime `if`/`switch` in the launcher that picks among pre-instantiated kernels. The kernel then compiles exactly one specialisation and asserts its own preconditions with `static_assert`, while the heuristic that produced the value sits in Python where it is readable, adjustable, and inspectable without recompiling CUDA. `store_cache`'s `num_threads` is the worked example: a `get_kernel(num_split)` ladder guarded by byte-alignment `if constexpr`s became one template argument plus a heuristic in `_jit_kvcache_module`. The exception is a knob that genuinely varies per call with the runtime shape — that has to stay a kernel argument.\n- **Check where the check is cheapest: `static_assert` > C++ host check > cached Python > per-call Python.** Anything fixed at compile time is a `static_assert`. Anything about the tensors is a `TensorMatcher` / `CHECK_HOST` in the C++ launcher, free next to a kernel launch. A check Python cannot delegate goes inside the `@cache_once` module factory, where it runs once per specialisation. What remains in the per-call entry point costs interpreter time on *every* forward, so it should be nothing but picking the module and allocating `out`.\n- **Fixed-width integer types.** Prefer `int32_t` / `int64_t` / `uint32_t` / `size_t` over `int`, `long`, or `long long`, so an index has the same width on both sides of the FFI boundary. Bare `int` is fine only where the width plainly cannot matter — an unrolled loop counter over a `constexpr` bound, a template `int` parameter. Shapes arrive as `int64_t` (`SymbolicSize::unwrap()`); narrowing to `uint32_t` for in-kernel indexing is a deliberate act, so write the `static_cast` explicitly and only where the range is known.\n- **Doxygen comments in C++.** Document exported entities with `///` or `/** ... */` blocks using `\\brief`, `\\param`, `\\tparam`, `\\return`, the way `include/sgl_kernel/` does. `python -m sglang.kernels.jit` writes `CommentFormat: Doxygen` into `.clangd` when clangd is 21 or newer, so these render on hover in the editor. Plain `//` remains fine for implementation notes inside a function body.\n- **ASCII only in C++ and CUDA sources.** Write `--`, `->`, `<=` instead of `—`, `→`, `≤`, including in comments. `grep -nP '[^\\x00-\\x7F]' ` before committing.\n- **`namespace details` is private.** Anything inside one is an implementation detail of its own header, free to change without notice. Do not name `details::` from another module, and never `using namespace details`. If you find yourself wanting something in there, that is the signal to promote it to a documented name instead.\n- **`const T* __restrict__` for read-only pointers.** This is what `csrc/` does throughout, and it lets the compiler emit non-coherent (`LDG`) loads.\n- **Watch the register budget.** For memory-bound kernels, keep to roughly 64 registers per thread so occupancy does not become the limit; prefer recomputing a value over letting it spill. Run `SGLANG_JIT_LOG_RESOURCE_USAGE=1 SGLANG_JIT_FORCE_RECOMPILE=1` to log per-kernel registers, spills and shared memory at INFO. Both halves are load-bearing: a cache hit has no compiler output to report, and passing `-Xptxas -v` by hand does nothing on its own because the build captures compiler output and replays it only on failure.\n- **Alignment is a contract you enforce, not a property you hope for.** A vectorized copy derives its access width from the row *size* (`load_bytes` picks 16B for a 1024B row), but the address it lands on is `base + index * row_stride` — and a stride is not constrained by the size. `TensorMatcher` admits any stride unless you say otherwise, so a padded cache row silently produces `misaligned address`. Every kernel that vectorizes is highly recommended to end its validation with `.ensure_alignment(w)`, where `w` is the width the kernel actually uses.\n\n---\n\n## Common Abstractions in `python/sglang/kernels/jit/include/sgl_kernel/`\n\n**Always prefer these abstractions over raw CUDA primitives.** They provide safety, readability, and consistency with the rest of the codebase. The only reason to drop to raw primitives is performance the abstraction cannot reach — a trade you make deliberately, and justify in a comment.\n\n### `utils.h` — Host-side utilities\n\n```cpp\n#include \n```\n\n- **`CHECK_HOST(cond) << \"msg \" << value`** — **Preferred** runtime check: stream-style, throws `PanicError` with file/line info on failure. Zero overhead on the true path — the message expressions are only evaluated when the check fails.\n- **`host::RuntimeCheck(cond, args...)`** — Function-style alternative to `CHECK_HOST`. Note its message args are always evaluated (even when the check passes), so prefer `CHECK_HOST` — especially on hot paths.\n- **`host::Panic(args...)`** — Unconditionally throw a `PanicError` with a descriptive message.\n- **`host::div_ceil(a, b)`** — Integer ceiling division `(a + b - 1) / b`.\n- **`host::irange(n)`** / **`host::irange(start, end)`** — Range views for cleaner loops.\n- **`host::pointer::offset(ptr, offsets...)`** — Byte-safe pointer arithmetic on `void*`. Use this instead of raw casts.\n\n### `utils.cuh` — Device-side utilities + `LaunchKernel`\n\n```cpp\n#include \n```\n\n- **Type aliases**: `fp16_t`, `bf16_t`, `fp32_t`, `fp8_e4m3_t`, `fp8_e5m2_t` and their packed variants `fp16x2_t`, `bf16x2_t`, `fp32x2_t`, etc.\n- **`SGL_DEVICE`** — Expands to `__forceinline__ __device__`. Use on all device functions.\n- **`SGL_DEVICE_HOST`** — `__forceinline__ __device__ __host__`. Use it when a `constexpr` helper must be callable from both the kernel and its launcher, so the two cannot drift.\n- **`device::kWarpThreads`** — Constant `32` (the codebase's *logical* warp width). `kWarpSize` is an alias. Note this is a group size, not the hardware wave: on gfx950 `warpSize` is 64, and `warp::kFullWidth` reflects that.\n- **`device::get_lane_id()`** — this thread's index within its logical group. **Prefer it over `threadIdx.x % kNumThreads` whenever the value feeds an address.** On CUDA the lane register is one read that folds into the address computation, whereas the modulo is a derived value the compiler re-materialises at *every* address scale — 8 instructions and 2 registers on a two-tile warp copy, measured on sm_100a. It picks the cheaper form per platform, so callers never need to care. Only equals the true in-warp lane when `blockDim.x` is a multiple of `kNumThreads`.\n- **`device::load_as(ptr, offset)`** / **`device::store_as(ptr, val, offset)`** — Type-safe loads/stores from `void*`.\n- **`device::pointer::offset(ptr, offsets...)`** — Pointer arithmetic on device.\n- **`host::LaunchKernel(grid, block, device_or_stream [, smem])`** — RAII kernel launcher that:\n - Resolves the CUDA stream from a `DLDevice` via TVM-FFI automatically.\n - Checks the CUDA error with file/line info after launch via `operator()(kernel, args...)`.\n - Supports `.enable_pdl(bool)` for PDL (Programmatic Dependent Launch, SM90+).\n- **`device::PDLWaitPrimary()`** / **`device::PDLTriggerSecondary()`** — The two halves of PDL, on sm_90+ (no-ops on older archs and ROCm). Their guarantees are **not** symmetric:\n - `PDLTriggerSecondary` (`griddepcontrol.launch_dependents`) only lets the next kernel in the stream *start* early. It carries no memory ordering and publishes nothing — matching that, the header's asm has no `\"memory\"` clobber.\n - `PDLWaitPrimary` (`griddepcontrol.wait`) is the ordering point: it waits until the preceding kernel has fully finished and its writes are visible.\n\n So every read of data the preceding kernel produced must come after `PDLWaitPrimary()`. What overlaps with the primary's tail is whatever you put *before* the wait — loading parameters, computing indices, touching buffers the primary never wrote — so a kernel that waits on its first line gains nothing. Neither call is a barrier: threads may reach or skip them independently. See \"Programmatic Dependent Launch and Synchronization\" in the CUDA C++ Programming Guide.\n- **`CHECK_CUDA(expr) << \"context\"`** — Stream-style CUDA error check; evaluates `expr` once and throws `PanicError` with `cudaGetErrorString` + file/line info if it is not `cudaSuccess`. Extra streamed context is optional.\n- **`host::RuntimeDeviceCheck(cudaError_t)`** — Function-style alternative to `CHECK_CUDA`. It takes no context message, so prefer `CHECK_CUDA`, which builds its error object only on the failure path.\n\n### `tensor.h` — Tensor validation (`TensorMatcher`, Symbolic types)\n\n```cpp\n#include \n```\n\nThis is the **primary validation API** for all kernel launchers. Use it to validate every `tvm::ffi::TensorView` argument.\n\n- **`host::SymbolicSize{\"name\"}`** — A named symbolic dimension. Call `.set_value(n)` to pin it, `.unwrap()` to extract after verification.\n- **`host::SymbolicDType`** — Symbolic dtype. Use `.set_options()` to restrict allowed types.\n- **`host::SymbolicDevice`** — Symbolic device. Use `.set_options()` to restrict to CUDA.\n- **`host::TensorMatcher({dims...})`** — Fluent builder for tensor validation:\n - `.with_dtype()` — require a specific C++ type (e.g. `fp16_t`)\n - `.with_dtype()` — allow a set of types\n - `.with_device(device_sym)` — require CUDA and bind the checked device to a `SymbolicDevice`\n - `.with_strides({strides...})` — validate strides (omit to require contiguous)\n - `.ensure_alignment(bytes)` — require the data pointer **and every stride except the innermost** to be a multiple of `bytes` (a power of two). Size-1 dimensions are skipped, since their strides are arbitrary. This is how a vectorized kernel states its precondition; see the alignment convention above.\n - `.verify(tensor_view)` — execute the check; throws `PanicError` with full context on failure; **chainable** (`verify(a).verify(b)` to check multiple tensors with the same shape)\n- **`host::is_type(dtype)`** — whether a `DLDataType` denotes the C++ type `T` (e.g. `fp16_t`).\n\n**Typical pattern:**\n```cpp\nauto N = SymbolicSize{\"num_elements\"};\nauto device = SymbolicDevice{};\ndevice.set_options();\nTensorMatcher({N}) //\n .with_dtype()\n .with_device(device)\n .ensure_alignment(16) // e.g. for 128-bit vectorized load/store\n .verify(dst)\n .verify(src); // same shape, dtype, device as dst\nconst int64_t n = N.unwrap();\nconst DLDevice dev = device.unwrap();\nconst int64_t last_dim = 128;\nTensorMatcher({N, last_dim}) // a fixed dimension can be a plain integer\n .with_dtype()\n .with_device(device)\n .verify(tensor_2d);\n```\n\n### `ffi.h` — Tensor allocation and blob wrapping (`host::ffi::`)\n\n```cpp\n#include \n```\n\nThe counterpart to `tensor.h`: that one validates what came in, this one produces new `tvm::ffi::Tensor` values. Allocation goes through the environment allocator (`TVMFFIEnvTensorAlloc`), so buffers come from PyTorch's caching allocator rather than a raw `cudaMalloc`.\n\n- **`host::alloc_workspace_tensor(nbytes, device)`** (declared in `utils.cuh`) — **the way to get scratch memory**: a 1-D `uint8` tensor of `nbytes`, or an empty tensor when `nbytes == 0`. Hold the returned `Tensor` in a local across every launch that touches it — it frees on destruction.\n- **`host::ffi::empty(shape, dtype, device)`** — Uninitialized tensor; `shape` accepts a braced list, so `ffi::empty({rows, sizeof(Plan)}, dtype, device)` works for a typed scratch array.\n- **`host::ffi::empty_like(tensor_view)`** — Same shape, dtype, and device as an existing tensor.\n- **`host::ffi::from_blob(data, shape, dtype, device[, deleter, stride, byte_offset])`** / **`from_blob_like(data, tensor_view, ...)`** — View memory you already own as a `Tensor`, no copy. The default deleter does nothing, so ownership stays with the caller; pass one only when the `Tensor` should own the block. Strides default to contiguous.\n\n### `type.cuh` — `DTypeTrait`, `packed_t`, and reduction traits\n\n```cpp\n#include \n```\n\n- **`DTypeTrait`** — Static trait struct, specialized for integral types, `fp32_t`, `fp16_t`, `bf16_t`, `fp8_e4m3_t`, and their packed x2/x4 variants. Provides:\n - `DTypeTrait::from(value)` — convert from another type via the right CUDA intrinsic (e.g. `fp32_t` → `fp16_t`)\n - `DTypeTrait::abs/max/min` — type-dispatched math (fp32, fp16/bf16 scalar and x2, integrals)\n - `DTypeTrait::sqrt/rsqrt/exp/sin/cos(x)` — `fp32_t` only\n - Metadata: `packed_t` / `unpacked_t` / `kVecSize` (packed layout), `kFloatMax` (dtype max as float, e.g. 448.0f for fp8-e4m3), `kZeroBits`\n- **`packed_t`** — Two-element packed alias: `packed_t` = `fp16x2_t`, `packed_t` = `bf16x2_t`, `packed_t` = `fp32x2_t`. Use for vectorized loads/stores.\n- **`device::cast(value)`** — Type-safe cast using `DTypeTrait`, e.g. `cast(v)`.\n- **`device::unpack(value)`** — View a packed value as an `unpacked_t[kVecSize]` array reference (e.g. `fp32x2_t` → `fp32_t[2]`); element writes propagate back to the packed value.\n- **`device::ReductionOp` (`SUM`/`MAX`/`MIN`) and `device::ReductionTrait::reduce(x, y)`** — One binary reduction step, dispatched through `DTypeTrait` (packed types reduce elementwise). This is the engine behind `warp::reduce`; use it directly when writing custom reductions.\n\n### `vec.cuh` — Vectorized memory access (`AlignedVector`)\n\n```cpp\n#include \n```\n\n- **`device::AlignedVector`** — Aligned storage for N elements of type T. N must be a power of two, `sizeof(T)*N <= 32`. Enables vectorized loads/stores for bandwidth efficiency. In terms of API/codegen constraints, the upper bound is 256-bit; in practice, 128-bit is the portable default, while 256-bit vectorization is typically only viable on `SM100+` and should be gated by an architecture check when needed.\n - `.load(ptr, offset)` — vectorized load from `ptr[offset]`\n - `.store(ptr, offset)` — vectorized store to `ptr[offset]`\n - `.fill(value)` — fill all N elements with `value`\n - `operator[](i)` — element access\n- **`device::LoadStoreBytes`** — named widths to pass around instead of bare integers: `MAX_GMEM` (arch-dependent, 32 on Blackwell), `MAX_SMEM` (16), `MAX_PORTABLE` (16, safe on CUDA and HIP), `MIN_COALALESCED` (4), and `RAW_1B` .. `RAW_32B`.\n\n### `tile.cuh` — `tile::Memory` (strided memory access pattern)\n\n```cpp\n#include \n```\n\n- `tile::Memory` is fundamentally a **1D cooperative accessor** over a contiguous region.\n- **`device::tile::Memory::cta(blockDim.x)`** — Creates a tile accessor where each thread handles `tid = threadIdx.x` with stride `tsize` (for `cta(blockDim.x)`, this is `blockDim.x`). Common for loops over a 1D array.\n- **`device::tile::Memory::warp()`** — the 32-lane flavour, and what `warp::load_bytes` is built on. It takes its lane index from `get_lane_id()` rather than `threadIdx.x % 32`, for the addressing reason described under `utils.cuh`. Use `warp(int n)` for a narrower sub-group; that overload keeps the modulo, since a hardware lane index is the wrong grouping below 32.\n- **`device::tile::Memory::thread()`** — single-thread accessor, no cooperation.\n- **`.load(ptr, offset)`** — loads `ptr[tid + offset * tsize]`\n- **`.store(ptr, val, offset)`** — stores to `ptr[tid + offset * tsize]`\n- **`.in_bound(n, offset)`** — boundary check\n\nFor a **2D tile**, either flatten `(row, col)` into a linear tile index first, or compute the address manually with `ptr[row * stride + col]` using your thread/block coordinates.\n\n### `bits.h` — Compile-time bit helpers (`host::`)\n\n```cpp\n#include \n```\n\n`constexpr` wrappers over ``, usable in `static_assert` and in template arguments: `host::is_pow2(x)`, `log2_floor(x)`, `log2_ceil(x)` (both `-1` for `x == 0`), `round_up_pow2(x)`, `round_down_pow2(x)`. They live in `host::` but are constant-expression-only, so device-side `static_assert` can call them.\n\n### `math.cuh` — Device math (`device::math::`)\n\n```cpp\n#include \n```\n\n- `device::math::max/min(a, b)` — type-dispatched binary math via `DTypeTrait`\n- `device::math::abs/sqrt/rsqrt/exp/sin/cos(x)` — type-dispatched unary math via `DTypeTrait`\n\n### `warp.cuh` — Warp-level primitives\n\n```cpp\n#include \n```\n\n**Vectorized row copy — reach for this before hand-rolling.** A warp cooperatively moving one contiguous row is the single most common shape in this tree, and it used to be re-implemented per kernel.\n\n- **`warp::load_bytes(src)`** / **`warp::store_bytes(dst, val)`** — the whole warp moves `kBytes` from/to one contiguous row. `kBytes` is arbitrary, down to 1; the helper picks the widest vector that divides it and handles the ragged tail. The returned value is opaque and encodes the width, so a load and its matching store **must use the same ``**.\n- **`warp::LoadStorePattern`** — the `kPattern` values, and the sign carries the meaning. A **negative** one (`WARP_UNIFORM_16B` and friends: `_GMEM`, `_SMEM`, `_4B`/`_8B`/`_32B`) says *the whole warp splits this row*, so the width is derived from the per-lane share (`gcd(kBytes / 32, |kPattern|)`), falling back to `gcd(kBytes, 4)` when the row is too narrow to give every lane 4 bytes. A **positive** one (a bare `LoadStoreBytes` value) makes no warp-splitting assumption and just caps the per-thread vector at `gcd(kBytes, kPattern)`. `WARP_UNIFORM_16B` is the usual choice for a row copy.\n- **`LoadStorePattern::get_vec_bytes()`** — the width actually chosen. `SGL_DEVICE_HOST`, so **the launcher must call this to feed `.ensure_alignment(...)`**. Pass it the same value the kernel passes `load_bytes` — for a split row that is the *per-warp share*, not the whole row; the full row can resolve to a different width and would under-constrain the strides.\n\n**Reductions.**\n\n- `warp::reduce(value, active_mask)` — `Op` is a `device::ReductionOp` (`SUM`/`MAX`/`MIN`). `kStart` and `kFinish` bound a range of lane-index bits: `` reduces within contiguous groups of `N` lanes, `` reduces *across* groups at the same offset. The operation is **symmetric** — `` and `` reduce the same lane set.\n- `warp::reduce_sum/reduce_max/reduce_min(value)` — wrappers. Work for any type with a `ReductionTrait`: floats, integers, packed x2. This is more widely used in practice than generic `reduce`.\n- `warp::inclusive_reduce(val, lane_id, mask)` and `inclusive_sum` / `inclusive_max` / `inclusive_min` — segmented inclusive scan: every lane keeps its own running total, unlike `reduce`. Forward when `kStart < kFinish`, backward when `kStart > kFinish`. `lane_id` defaults to `get_lane_id()`, which is correct by construction — only pass it if you already have the **segment-relative** index (`threadIdx.x % kWidth`); a warp-relative one silently corrupts every segment past the first when `kWidth < 32`.\n- `warp::broadcast(value, src_lane)` — `src_lane` is **segment-relative**: each `kWidth` segment reads its own lane, not one warp-wide source.\n- `warp::get_lane_id`, `warp::elect_one_lane()` (one elected lane, for gating a single-thread TMA issue; CUDA-only), `warp::kFullMask` and `warp::kFullWidth` (32 on CUDA, 64 on HIP — note this differs from `kWarpThreads`).\n\n### `cta.cuh` — CTA-level primitives\n\n```cpp\n#include \n```\n\n- `device::cta::reduce_max(value, smem, min_value)` — CTA-wide max using shared memory + warp reduction. Caller is responsible for a `__syncthreads()` after if the result in `smem[0]` is needed.\n\n### `atomic.cuh` — Atomic operations\n\n```cpp\n#include \n```\n\n- `device::atomic::max(float* addr, float value)` — float atomic max (handles negative values correctly via bit tricks).\n- **`device::atomic::Event`** — a cross-CTA arrive/wait counter in one 32-bit word. Producers call `arrive()`; consumers call `wait(num_producers)` (exactly one consumer) or `wait_multi(num_producers, num_consumers, n)` (several). The last released consumer resets the word, so one `Event` is reusable across launches with no host re-zero. Preconditions, all of them load-bearing:\n - the word must be **zeroed by the host** before first use, and must live in **global** memory — a shared or local `Event` is an illegal address;\n - `arrive()` is a release and the waits are acquires, so a producer's writes before `arrive()` are visible after the wait;\n - **generations must be explicitly ordered** — the word has no phase bit, so overlapping generation N+1 with a consumer still in generation N is undefined behaviour (a lagging consumer never exits);\n - `wait()` and `wait_multi()` lay the word out incompatibly; mixing them on one `Event` is undefined behaviour.\n\n### `runtime.cuh` — Occupancy and device info\n\n```cpp\n#include \n```\n\n- `host::runtime::get_blocks_per_sm(kernel, block_dim)` — max active blocks per SM (occupancy)\n- `host::runtime::get_sm_count(device_id [, use_cache])` — number of SMs on the device. Memoized per device ordinal, so it is cheap enough to call on a launch path; pass `use_cache=false` to force a driver query.\n- `host::runtime::get_cc_major` / `get_cc_minor` / `get_sm_version` — compute capability, same caching.\n\n**Do not query the architecture at runtime.** The JIT compiles for the exact local GPU, so the arch is a *compile-time* fact: `SGL_CUDA_ARCH` is injected by `load_jit`, and `SGL_ARCH_HOPPER_OR_GREATER` / `SGL_ARCH_BLACKWELL_OR_GREATER` / `device::kMaxVecBytes` are derived from it. Reach for `get_cc_*` only for something the arch genuinely does not determine. SM *count* is the opposite case — it varies within an arch (B200 vs B300 vs a MIG slice), so it has to stay a runtime query.\n\n**Persistent kernel pattern** (cap blocks to SM count × occupancy):\n```cpp\nstatic const uint32_t max_occ = runtime::get_blocks_per_sm(kernel, kBlockSize);\nstatic const uint32_t num_sm = runtime::get_sm_count(device.unwrap().device_id);\nconst auto num_blocks = std::min(num_sm * max_occ, div_ceil(n, kBlockSize));\nLaunchKernel(num_blocks, kBlockSize, device.unwrap())(kernel, params);\n```\n\n---\n\n## Step 0 (optional): Generate a `.clangd` config for better IDE support\n\n```bash\npython -m sglang.kernels.jit -h # for verbose help info about clangd configuration\npython -m sglang.kernels.jit\npython -m sglang.kernels.jit --dep cutlass flashinfer # with cutlass/flashinfer dependency\n```\n\n---\n\n## Step 1: Implement the CUDA kernel in `kernels/jit/csrc/`\n\nCreate `python/sglang/kernels/jit/csrc/elementwise/scale.cuh`.\n\nThe implementation fully uses the project abstractions described above:\n\n```cpp\n// NOTE: Comments for headers are not common in practice.\n// It is only shown here for tutorial purposes to highlight the key abstractions.\n#include // For TensorMatcher, SymbolicSize, SymbolicDevice\n#include // For DTypeTrait, fp16_t, bf16_t, fp32_t\n#include // For CHECK_HOST, div_ceil\n#include // For LaunchKernel, SGL_DEVICE\n#include // For AlignedVector\n\n#include \n#include \n\nnamespace sglang {\n\n/**\n * \\brief Element-wise scale using vectorized 128-bit loads/stores.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kVecN Elements per vector load (e.g. 8 for fp16)\n * \\tparam kUsePDL Whether to emit the PDL wait/trigger pair\n * \\param dst Output buffer, `n_total` elements\n * \\param src Input buffer, `n_total` elements\n * \\param factor Runtime scale factor\n * \\param n_total Number of elements to scale\n */\ntemplate \n__global__ void scale_kernel(T* __restrict__ dst,\n const T* __restrict__ src,\n float factor,\n uint32_t n_total) {\n using vec_t = device::AlignedVector;\n const uint32_t n_vecs = n_total / kVecN;\n\n // If using PDL, wait for primary kernel before any global memory load.\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLWaitPrimary();\n\n // --- vectorised body ---\n const uint32_t vec_stride = blockDim.x * gridDim.x;\n for (uint32_t vi = blockIdx.x * blockDim.x + threadIdx.x;\n vi < n_vecs;\n vi += vec_stride) {\n vec_t v;\n v.load(src, vi);\n#pragma unroll\n for (int i = 0; i < kVecN; ++i) {\n v[i] = static_cast(static_cast(v[i]) * factor);\n }\n v.store(dst, vi);\n }\n\n // --- scalar tail ---\n const uint32_t base = n_vecs * kVecN;\n const uint32_t scalar_stride = blockDim.x * gridDim.x;\n for (uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;\n base + i < n_total;\n i += scalar_stride) {\n dst[base + i] = static_cast(static_cast(src[base + i]) * factor);\n }\n\n // If using PDL, signal for the secondary kernel to start after all threads have finished\n // This is NOT a synchronization point, which means some threads can early exit before this.\n device::PDLTriggerSecondary();\n}\n\n/**\n * \\brief Validate the tensors, select the vector width, launch `scale_kernel`.\n *\n * \\tparam T Element type: fp16_t | bf16_t | fp32_t\n * \\tparam kUsePDL Whether to launch with PDL enabled\n * \\param dst Output tensor; same shape / dtype / device as `src`\n * \\param src Input tensor on CUDA\n * \\param factor Runtime scale factor\n */\ntemplate \nvoid scale(tvm::ffi::TensorView dst, tvm::ffi::TensorView src, float factor) {\n using namespace host;\n\n // 1. Validate input tensors with TensorMatcher\n SymbolicSize N = {\"num_elements\"};\n SymbolicDevice device_;\n device_.set_options();\n\n TensorMatcher({N}) //\n .with_dtype()\n .with_device(device_)\n .verify(dst)\n .verify(src); // same shape / dtype / device as dst\n\n const uint32_t n = static_cast(N.unwrap());\n const DLDevice device = device_.unwrap();\n\n CHECK_HOST(n > 0) << \"scale: num_elements must be > 0, got \" << n;\n\n // 2. Choose vector width for 128-bit loads (16 bytes)\n // fp16/bf16: 8 elements x 2 bytes = 16 bytes\n // fp32: 4 elements x 4 bytes = 16 bytes\n // We encourage using `device::kMaxVecBytes`, which will change according to\n // the target architecture and can enable 256-bit vectorization on SM100+ if desired.\n // But 128-bit is more commonly adapted for better compatibility,\n // so it's still ok to hardcode 16 here just for simplicity.\n constexpr int kVecN = 16 / sizeof(T);\n const uint32_t n_work_items = div_ceil(n, static_cast(kVecN));\n\n // 3. Launch\n constexpr uint32_t kBlockSize = 256;\n const uint32_t grid = div_ceil(n_work_items, kBlockSize);\n\n // PDL feature is 100% optional. Without `enable_pdl`, the code should still be correct.\n // Try to enable it if profiling shows that it can benefit the performance of this kernel.\n LaunchKernel(grid, kBlockSize, device).enable_pdl(kUsePDL)(\n scale_kernel,\n static_cast(dst.data_ptr()),\n static_cast(src.data_ptr()),\n factor,\n n);\n}\n\n} // namespace sglang\n```\n\n**Key points:**\n\n- Include headers from `sgl_kernel/` — **not** raw CUDA headers for anything already covered\n- Use `TensorMatcher` for all tensor validation; never manually check shape/dtype/device\n- Use `AlignedVector` for vectorised 128-bit loads/stores — significant bandwidth win\n- Use `LaunchKernel` — it resolves the stream and checks errors automatically\n- Use `CHECK_HOST(cond) << ...` for runtime assertions with useful error messages (zero overhead when the check passes)\n- Prefer passing runtime scalars like `factor` directly unless compile-time specialisation is genuinely required\n- `fp16_t` / `bf16_t` / `fp32_t` are the project's type aliases (from `utils.cuh`)\n- `device::cast` or `DTypeTrait::from(val)` for cross-type conversions\n- `device::math::` functions for device math instead of bare `__` intrinsics if possible.\n- Consider PDL — it can help when the kernel has prologue work to overlap. Place `PDLWaitPrimary()` right before the first read of upstream data, not at the top of the kernel\n\n---\n\n## Step 2: Add the Python wrapper in `kernels/ops/`\n\nThe wrapper lives next to its functional group under `python/sglang/kernels/ops/`, not beside the CUDA source — `kernels/jit/` holds only the JIT infrastructure (`csrc/`, `include/`, `utils/`, `benchmark/`). Create `python/sglang/kernels/ops/elementwise/scale.py`:\n\n```python\nfrom __future__ import annotations\n\nfrom typing import TYPE_CHECKING\n\nimport torch\n\nfrom sglang.kernels.jit.utils import (\n cache_once,\n is_arch_support_pdl,\n load_jit,\n make_cpp_args,\n)\n\nif TYPE_CHECKING:\n from tvm_ffi.module import Module\n\n\n@cache_once\ndef _jit_scale_module(dtype: torch.dtype) -> Module:\n \"\"\"Compile and cache the JIT scale module for a given dtype.\"\"\"\n # Checks on the compile key live here, not in `scale`: `cache_once` keys on\n # `dtype`, so this runs once per specialisation instead of once per call.\n if dtype not in (torch.float16, torch.bfloat16, torch.float32):\n raise RuntimeError(\n f\"Unsupported dtype {dtype}. Supported: float16, bfloat16, float32\"\n )\n args = make_cpp_args(dtype, is_arch_support_pdl())\n return load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n )\n\n\ndef scale(src: torch.Tensor, factor: float, out: torch.Tensor | None = None) -> torch.Tensor:\n \"\"\"\n Element-wise scale: dst = src * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n src : CUDA tensor (FP16 / BF16 / FP32)\n factor : scale factor\n out : optional pre-allocated output tensor (same shape/dtype as src)\n\n Returns\n -------\n Scaled tensor (dst = src * factor).\n \"\"\"\n # DO NOT add proactive validation here: every check costs interpreter time\n # on a per-forward path. Tensor invariants belong in the C++ launcher, and\n # anything about the compile key belongs in `_jit_scale_module`.\n if out is None:\n out = torch.empty_like(src)\n\n module = _jit_scale_module(src.dtype)\n module.scale(out, src, factor)\n return out\n```\n\n**Key points:**\n\n- Use `cache_once` — **not** `functools.lru_cache` (incompatible with `torch.compile`)\n- `load_jit` first arg(s) form the unique build marker; same marker = same cached binary\n- Only include compile-time specialisation knobs in the build marker; runtime values like `factor` should stay runtime unless the kernel truly needs templating\n- `cuda_wrappers`: `(export_name, kernel_symbol)` — `export_name` is called from Python\n- `make_cpp_args(dtype, ...)` converts `torch.dtype` to C++ type alias:\n- `is_arch_support_pdl()` checks if the current architecture supports PDL, which is typically passed as a template argument to the kernel.\n- Keep the entry point thin (see Conventions). What Python must still check goes in the `@cache_once` module factory, not in the entry point: `cache_once` keys on its arguments, so a check there costs one evaluation per specialisation instead of one per call — that is where the supported-dtype guard lives. Tensor invariants belong in the C++ launcher; if something here never reaches a `.verify(...)`, close that gap on the C++ side rather than in Python\n\n| `torch.dtype` | C++ type |\n|--------------------|------------|\n| `torch.float16` | `fp16_t` |\n| `torch.bfloat16` | `bf16_t` |\n| `torch.float32` | `fp32_t` |\n\n---\n\n## Register the public entry point (required)\n\nAdd metadata in `python/sglang/kernels/ops/elementwise/__init__.py` so the\noperator appears in the registry without importing its implementation:\n\n```python\nfrom sglang.kernels.registry import register_kernel\nfrom sglang.kernels.spec import CapabilityRequirement, KernelBackend, KernelSpec\n\nregister_kernel(\n KernelSpec(\n op=\"elementwise.scale\",\n backend=KernelBackend.JIT,\n target=\"sglang.kernels.ops.elementwise.scale:scale\",\n capabilities=frozenset({CapabilityRequirement.CUDA}),\n )\n)\n```\n\nRecord the devices the implementation actually supports; the CUDA-only scale\nexample above does not imply that every JIT kernel is CUDA-only. Runtime and\ncorrectness tests import `scale` from `sglang.kernels.ops.elementwise.scale`.\nDo not import the implementation eagerly in the group initializer. When an\nexisting `BaseFusedOp` owns the operation, add the implementation there instead\nof registering a second conflicting op/backend pair.\n\n## Step 3 (optional): Tune JIT build flags\n\nIf your kernel uses some math functions like `expf` or `sinf`, consider enabling `--use_fast_math` for better performance (with a potential precision tradeoff):\n\n```python\nreturn load_jit(\n \"scale\",\n *args,\n cuda_files=[\"elementwise/scale.cuh\"],\n cuda_wrappers=[(\"scale\", f\"scale<{args}>\")],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n)\n```\n\nIf your kernel requires SM90+, raise a clear Python error before calling `load_jit`. Arch gating is one of the checks that has to live in Python — it decides whether to compile at all, so the C++ launcher never gets to run:\n\n```python\nif torch.cuda.get_device_capability()[0] < 9:\n raise RuntimeError(\"This kernel requires SM90 (Hopper) or later\")\n```\n\n---\n\n## Step 4: Write tests (required)\n\nJIT kernel correctness tests and benchmarks live under `test/registered/kernels/ops//` and `test/registered/kernels/benchmark//`, mirroring the wrapper's group under `python/sglang/kernels/ops/` (NOT inside the `sglang` package -- a `register_*_ci(...)` call anywhere under `python/sglang/` is rejected by the `check-no-registered-tests-in-package` pre-commit hook). Only their test-only helpers (e.g. `benchmark/marker.py`) stay alongside the kernel source under `python/sglang/kernels/jit/` and are imported by absolute path. **CI does not run `pytest` in those directories directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` and `bench_*.py` under `test/registered/`, collects `register_*_ci(...)` calls by **statically parsing each file's AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check.\n\n- **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`.\n- **Nightly kernel suite**: register with `stage=\"nightly\"` plus the `runner_config` of the machine it needs (e.g. `1-gpu-large`), giving the `nightly-test-1-gpu-large` suite. `.github/workflows/nightly-test-nvidia.yml` sets `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` for the whole nightly run, so the expanded parameter grids apply automatically (see `python/sglang/kernels/jit/utils/common.py` → `should_run_full_tests` / `get_ci_test_range`). There is no separate kernel-only nightly job: every nightly test on one machine type shares that machine's suite.\n\nRegistration pattern (module level, **literal** `est_time`, `stage`, and `runner_config` values — required for AST parsing):\n\n```python\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n# Optional B200/SM100 registration for tests that cover Blackwell-specific code paths\n# register_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"4-gpu-b200\")\n# Optional second registration: same file also runs nightly, same form,\n# stage is just \"nightly\" there (and no `nightly=True`)\n# register_cuda_ci(est_time=120, stage=\"nightly\", runner_config=\"1-gpu-large\")\n```\n\nCI generates the suite name as `{stage}-test-{runner_config}`, so `stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\"` becomes the `base-b-kernel-unit-test-1-gpu-large` suite you pass to `run_suite.py` below — don't put the `-test-` infix in `register_cuda_ci`. Nightly uses the same shape with `stage=\"nightly\"`; the single-string `suite=` form is left only for `stress` and non-CUDA pools.\n\nKeep `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.\n\nUse `register_cuda_ci(..., disabled=\"reason\")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only).\n\n**Run like CI** (from repo root):\n\n```bash\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-1-gpu-large)\n# For B200/SM100-specific coverage:\n(cd test && python3 run_suite.py --hw cuda --suite base-b-kernel-unit-test-4-gpu-b200)\n```\n\nFor fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`.\n\nCreate `test/registered/kernels/ops/elementwise/test_scale.py`:\n\n```python\nimport pytest\nimport torch\nfrom sglang.kernels.ops.elementwise.scale import scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=30, stage=\"base-b-kernel-unit\", runner_config=\"1-gpu-large\")\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [1, 127, 128, 1024, 4097]) # cover tail remainder\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0, 3.0])\ndef test_scale_correctness(dtype, size, factor):\n src = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = scale(src, factor)\n expected = src * factor\n\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\ndef test_scale_out_param(dtype):\n src = torch.randn(1024, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(src)\n result = scale(src, 2.0, out=out)\n assert result is out\n torch.testing.assert_close(out, src * 2.0, rtol=1e-2, atol=1e-2)\n\n\ndef test_scale_cpu_error():\n src = torch.randn(128, dtype=torch.float16) # CPU tensor\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n scale(src, 2.0)\n\n\ndef test_scale_unsupported_dtype():\n src = torch.randint(0, 10, (128,), dtype=torch.int32, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"dtype\"):\n scale(src, 2.0)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-v\", \"-s\"]))\n```\n\n---\n\n## Step 5: Add a benchmark (required)\n\nBenchmarks are `bench_*.py` files under `test/registered/kernels/benchmark//`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`base-b-kernel-benchmark-test-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large`).\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nBenchmarks use the project's own `marker` framework (in `python/sglang/kernels/jit/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`):\n\n- **`@marker.benchmark(line_arg, line_vals, *, unit=\"us\")`** — the **innermost** decorator (bottom of the stack, directly above `def benchmark`). Declares the column axis: each value in `line_vals` becomes a result column, and `line_arg` is the parameter name passed into the benchmark function. `unit` is one of `\"us\" | \"ms\" | \"s\"`.\n- **`@marker.parametrize(names, vals, ci_vals=None)`** — stackable decorator that adds a row axis (pytest-style). Each `@parametrize` adds one (or more, correlated) parameter the benchmark is swept over (Cartesian product across all `parametrize` decorators). `names` may be a single name (`\"size\"`) or a comma-separated correlated tuple axis (`\"h,d\"`, with `vals` then a list of tuples like `[(1, 64), (2, 128)]`). Pass the optional third `ci_vals` for a smaller sweep that is auto-selected under `is_in_ci()` — this is the built-in CI-shrinking mechanism, so you usually don't need `get_benchmark_range` for swept axes.\n- **`marker.do_bench(fn, *, input_args=(), input_kwargs={}, ...)`** — runs `fn` under CUDA graph (default) or a naive loop, returns a `BenchResult`. Key knobs:\n - `memory_args`: defaults to `\"all\"` (footprint derived from all input args/kwargs). Pass an explicit tuple of tensors (e.g. `(k, v, indices)`) to count only the inputs the kernel actually touches.\n - `memory_output`: defaults to `\"out\"` — re-runs `fn` once to capture its **returned** tensor and counts it. For in-place kernels (which return `None`), pass the written tensors explicitly (e.g. `memory_output=(k, v)`); the re-run is then skipped. Set to `None` to count no output.\n - Together `memory_args` + `memory_output` give the GB/s column; with both defaults a function `out = f(src)` already reports `bytes(src) + bytes(out)`.\n - `graph_clone_args` / `graph_clone_kwargs`: which inputs to clone per CUDA-graph iteration to defeat L2 cache reuse. Defaults to `\"all\"`; if you narrow it, keep every read, in-place, and output tensor.\n - `use_cuda_graph=False` for kernels that can't be captured.\n - `metrics=(0.5, \"avg\")` controls reported quantiles (the first metric becomes the table latency column).\n - `disable_log_bandwidth` (defaults from `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1`) skips the bandwidth column entirely.\n- **`utils.create_random(*shape)` / `utils.create_empty(*shape)`** — shorthand for `torch.randn` / `torch.empty` with `DEFAULT_DTYPE` (`bfloat16`) and `DEFAULT_DEVICE` (`\"cuda\"`). Override via the `dtype=` / `device=` kwargs.\n- **`utils.get_benchmark_range(full_range, ci_range)`** — returns the smaller `ci_range` under CI (`is_in_ci()`), the `full_range` locally. Still available for the `benchmark(...)` column axis (which has no `ci_vals`); for `parametrize` row axes prefer the built-in `ci_vals` argument.\n\nCreate `test/registered/kernels/benchmark/elementwise/bench_scale.py`:\n\n```python\nimport torch\n\nfrom sglang.kernels.jit.benchmark import marker\nfrom sglang.kernels.jit.benchmark.utils import create_random\nfrom sglang.kernels.ops.elementwise.scale import scale as jit_scale\nfrom sglang.test.ci.ci_register import register_cuda_ci\n\nregister_cuda_ci(est_time=6, stage=\"base-b-kernel-benchmark\", runner_config=\"1-gpu-large\")\n\n\n@torch.compile()\ndef torch_impl_scale(src: torch.Tensor, factor: float) -> torch.Tensor:\n return src * factor\n\n\nFN_MAP = {\n \"jit\": jit_scale,\n \"torch\": torch_impl_scale,\n}\n\n\n# `parametrize(name, full_vals, ci_vals)`: the 3rd arg is the smaller sweep\n# auto-selected under CI; the full range runs locally.\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096, 65536]) # 1K .. 512K\n@marker.benchmark(\"impl\", [\"jit\", \"torch\"])\ndef benchmark(size: int, impl: str):\n src = create_random(size)\n factor = 2.0\n return marker.do_bench(\n FN_MAP[impl],\n input_args=(src, factor),\n # `src` is read -> clone it per iter to avoid L2 reuse; factor is a scalar.\n graph_clone_args=(0,),\n # Defaults already report bandwidth: memory_args=\"all\" counts src,\n # memory_output=\"out\" counts the returned tensor -> bytes(src)+bytes(out).\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n**Key points:**\n\n- The `line_arg` name passed to `benchmark` (`\"impl\"` here) must match a parameter on `benchmark(...)`; same for every `parametrize` name (`\"size\"`).\n- Stack `@parametrize` once per swept axis. The required `@marker.benchmark` is the **innermost** decorator (bottom of the stack, directly above the function) — `@parametrize` rows go above it.\n- Prefer `create_random` / `create_empty` from `utils.py` over open-coding `torch.randn(..., dtype=..., device=...)`.\n- The GB/s column appears by default (`memory_args=\"all\"` + `memory_output=\"out\"`). For memory-bound kernels it's the most informative number; scope `memory_args` / `memory_output` to the tensors actually touched if the defaults over- or under-count. For compute-bound kernels where bandwidth is misleading, set `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH=1` (or `disable_log_bandwidth=True`).\n- For in-place kernels (which return `None`), pass the written tensors via `memory_output=(...)` since the `\"out\"` default would capture nothing.\n- Tune `graph_clone_args` / `graph_clone_kwargs` to all the arguments that might be read by the kernel. We can only skip cloning for write-only args. For in-place modified args, we still need to clone them to get accurate timing (reusing the same buffer keeps it L2-hot and skews results).\n- Call `benchmark.run()` (no `print_data=` kwarg — the marker framework prints directly).\n\nRun locally:\n\n```bash\npython test/registered/kernels/benchmark/elementwise/bench_scale.py\n```\n\nRun the benchmark suite the way CI does:\n\n```bash\ncd test && python3 run_suite.py --hw cuda --suite base-b-kernel-benchmark-test-1-gpu-large\n```\n\n---\n\n## Troubleshooting\n\n- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time`, `stage`, and `runner_config`; starred args and non-literal values break AST collection\n- **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/kernels/jit/csrc/`; reduce template argument combinations\n- **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...`\n- **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. `graph_clone_args` defaults to `\"all\"`; if you narrow it, it must still cover every *read* tensor — reusing a single buffer keeps it L2-hot and skews results. Keep *write* tensors in it too: they are what sets the rotation count, and a shared output buffer stays L2-hot the same way.\n- **Missing GB/s column**: the column is on by default; check that `SGLANG_JIT_BENCHMARK_DISABLE_LOG_BANDWIDTH` is not `1` and `disable_log_bandwidth` is not `True`. For in-place kernels (return `None`) the `memory_output=\"out\"` default counts nothing — pass the written tensors via `memory_output=(...)`\n\n---\n\n## References\n\n- `docs/docs/developer_guide/development_jit_kernel_guide.mdx`\n- `test/run_suite.py` — suite names, discovery of `test/registered/`, execution entrypoint for CI\n- `python/sglang/test/ci/ci_register.py` — `register_cuda_ci` and AST registration rules\n- `python/sglang/kernels/jit/utils/compile/` — `load_jit` (`loader.py`), `make_cpp_args` (`cpp_args.py`)\n- `python/sglang/kernels/jit/utils/common.py` — `cache_once`, `should_run_full_tests`, `get_ci_test_range`\n- `python/sglang/kernels/jit/include/sgl_kernel/tensor.h` — `TensorMatcher`, `SymbolicSize/DType/Device`, `is_type`\n- `python/sglang/kernels/jit/include/sgl_kernel/ffi.h` — `ffi::empty`, `ffi::empty_like`, `ffi::from_blob`\n- `python/sglang/kernels/jit/include/sgl_kernel/utils.cuh` — type aliases, `LaunchKernel`, `SGL_DEVICE`\n- `python/sglang/kernels/jit/include/sgl_kernel/vec.cuh` — `AlignedVector`\n- `python/sglang/kerne\n\n\n---\nname: add-sgl-kernel\ndescription: Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)\n---\n\n# Tutorial: Adding a New Kernel to `sgl-kernel` (AOT / Heavyweight)\n\nApply [kernel-organization](../kernel-organization/SKILL.md) for the public\noperator namespace, logical grouping, lazy registry metadata, and test placement.\nThe implementation tutorial below does not replace that API contract.\n\nThis tutorial walks through adding a simple element-wise scale operation as an AOT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.\n\n## Goal\n\nAdd a new operation that scales each element of a tensor by a scalar factor:\n\n- Input: tensor `x` (CUDA) and scalar `factor` (float)\n- Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)\n- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**\n - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `python/sglang/kernels/aot/include/utils.h`)\n\n## Two rules of thumb (must follow)\n\n1. **Prefer `python/sglang/kernels/jit` first** when the kernel does **not** depend on CUTLASS or another large C++ project. This is the default path for lightweight kernels that benefit from rapid iteration.\n2. **Prefer `sgl-kernel`** when the kernel **does** depend on CUTLASS or another large C++ project, or when it should be part of the AOT wheel / torch op registration flow.\n3. **Exception**: if the dependency is `flashinfer`, or CUTLASS that is already provided through `flashinfer`, the kernel can still be implemented as `jit_kernel`.\n\nIn addition, every new kernel must ship with:\n\n- **Tests** (pytest)\n- **A benchmark script** (`marker.do_bench` from `sglang.kernels.jit.benchmark` — it works for any callable, not only JIT kernels)\n\n---\n\n## Repository integration map\n\nYou will typically touch these files/areas:\n\n- Implementation: `python/sglang/kernels/aot/csrc/elementwise/scale.cu` (pick the right subdirectory)\n- Public declarations: `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- Torch extension registration: `python/sglang/kernels/aot/csrc/common_extension.cc`\n- Build: `python/sglang/kernels/aot/CMakeLists.txt` (`set(SOURCES ...)`)\n- Python API: `python/sglang/kernels/aot/python/sgl_kernel/` and `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n- Tests: `python/sglang/kernels/aot/tests/test_scale.py`\n- Benchmarks: `python/sglang/kernels/aot/benchmark/bench_scale.py`\n\n---\n\n## Step 1: Implement the kernel in `csrc/`\n\nPick the right subdirectory:\n\n- `csrc/elementwise/` — for element-wise ops (our example)\n- `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories\n\nCreate `python/sglang/kernels/aot/csrc/elementwise/scale.cu`:\n\n```cpp\n#include \n#include \n#include \n\n#include \"utils.h\" // DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16\n\n// scale_kernel: out[i] = input[i] * factor\n// Supports float, half (__half), __nv_bfloat16 via template T\ntemplate \n__global__ void scale_kernel(T* __restrict__ out,\n const T* __restrict__ input,\n float factor,\n int64_t n) {\n int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x;\n if (idx < n) {\n out[idx] = static_cast(static_cast(input[idx]) * factor);\n }\n}\n\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor) {\n TORCH_CHECK(input.is_cuda(), \"input must be a CUDA tensor\");\n TORCH_CHECK(input.is_contiguous(), \"input must be contiguous\");\n TORCH_CHECK(out.is_cuda(), \"out must be a CUDA tensor\");\n TORCH_CHECK(out.is_contiguous(), \"out must be contiguous\");\n TORCH_CHECK(out.sizes() == input.sizes(), \"out and input must have the same shape\");\n TORCH_CHECK(out.scalar_type() == input.scalar_type(),\n \"out and input must have the same dtype\");\n\n const int64_t n = input.numel();\n const int threads = 256;\n const int blocks = (n + threads - 1) / threads;\n\n const cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n const at::cuda::OptionalCUDAGuard device_guard(device_of(input));\n\n // Dispatches over float, float16, bfloat16\n DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {\n scale_kernel<<>>(\n static_cast(out.data_ptr()),\n static_cast(input.data_ptr()),\n static_cast(factor),\n n);\n cudaError_t status = cudaGetLastError();\n TORCH_CHECK(status == cudaSuccess,\n \"scale_kernel launch failed: \", cudaGetErrorString(status));\n return true;\n });\n}\n```\n\n**Key points:**\n\n- Use `at::Tensor` (PyTorch tensors), `TORCH_CHECK` for validation, `at::cuda::getCurrentCUDAStream()` for stream\n- Keep Python wrappers thin; do shape/dtype/device validation in C++ right around the launch path\n- `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` covers `float`, `half` (FP16), `__nv_bfloat16` (BF16)\n- Add device error checking after every kernel launch\n- If a kernel only works on certain architectures, enforce that with `TORCH_CHECK` and skip logic in tests\n\n---\n\n## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`\n\nEdit `python/sglang/kernels/aot/include/sgl_kernel_ops.h`, add to the elementwise section:\n\n```cpp\nvoid scale(at::Tensor& out, const at::Tensor& input, double factor);\n```\n\n---\n\n## Step 3: Register the op in `csrc/common_extension.cc`\n\nEdit `python/sglang/kernels/aot/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:\n\n```cpp\n// From csrc/elementwise\nm.def(\"scale(Tensor! out, Tensor input, float factor) -> ()\");\nm.impl(\"scale\", torch::kCUDA, &scale);\n```\n\n**Key points:**\n\n- `Tensor!` means in-place / mutable output argument\n- The schema is important for `torch.compile` and for consistent call signatures\n- Keep the torch schema in PyTorch scalar types (`float` here), but note that the C++ launcher signature still needs `double` for scalar arguments accepted by `torch::Library`\n\n---\n\n## Step 4: Add the new source file to `CMakeLists.txt`\n\nEdit `python/sglang/kernels/aot/CMakeLists.txt`, add to `set(SOURCES ...)`:\n\n```cmake\ncsrc/elementwise/scale.cu\n```\n\n**Key points:**\n\n- Keep the list **alphabetically sorted** (the file explicitly requires this)\n- If the kernel has arch constraints, reflect that in tests/benchmarks via skip logic\n\n---\n\n## Step 5: Expose a Python API under `python/sglang/kernels/aot/python/sgl_kernel/`\n\nPrefer following the existing module organization first. For elementwise kernels, the usual pattern is:\n\n- implement the Python wrapper in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`\n- then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`\n\nFor example, in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`, add:\n\n```python\nimport torch\n\ndef scale(\n input: torch.Tensor,\n factor: float,\n out: torch.Tensor | None = None,\n) -> torch.Tensor:\n \"\"\"\n Element-wise scale: out = input * factor.\n\n Supported dtypes: torch.float16, torch.bfloat16, torch.float32.\n\n Parameters\n ----------\n input : CUDA input tensor\n factor : scale factor (float)\n out : optional pre-allocated CUDA output tensor (same shape/dtype as input)\n \"\"\"\n if out is None:\n out = torch.empty_like(input)\n torch.ops.sgl_kernel.scale.default(out, input, factor)\n return out\n```\n\nThen re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.\n\n---\n\n## SGLang integration (required for runtime use)\n\nAfter exposing the AOT wheel symbol, add a lazy wrapper and a `KernelSpec` under\n`python/sglang/kernels/ops//`. Set `backend=KernelBackend.AOT`, record\nactual device/architecture support, and keep `sgl_kernel` imports inside the\nimplementation path. SGLang runtime and integration tests import that wrapper.\nFor this example, use `sglang.kernels.ops.elementwise.scale`.\n\nThe wheel-level tests below validate its standalone API/build. They do not\nreplace CI-registered SGLang correctness tests in\n`test/registered/kernels/ops/elementwise/test_scale.py` and benchmarks in\n`test/registered/kernels/benchmark/elementwise/bench_scale.py`. Follow\n`write-sglang-test` for their registration and CI budget; do not register tests\nunder the shipped `python/sglang/` package.\n\n## Step 6: Write tests (required)\n\nCreate `python/sglang/kernels/aot/tests/test_scale.py`:\n```python\nimport pytest\n\nimport torch\nimport sgl_kernel\n\n@pytest.mark.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32])\n@pytest.mark.parametrize(\"size\", [128, 1024, 4096, 65536])\n@pytest.mark.parametrize(\"factor\", [0.5, 1.0, 2.0])\ndef test_scale_correctness(dtype, size, factor):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n\n result = sgl_kernel.scale(input, factor, out=out)\n assert result is out\n\n expected = input * factor\n rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)\n torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)\n\n\ndef test_scale_shape_mismatch():\n input = torch.randn(128, dtype=torch.float16, device=\"cuda\")\n out = torch.empty(256, dtype=torch.float16, device=\"cuda\")\n with pytest.raises(RuntimeError, match=\"same shape\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\ndef test_scale_cpu_input():\n input = torch.randn(128, dtype=torch.float16) # CPU\n out = torch.empty_like(input)\n with pytest.raises(RuntimeError, match=\"CUDA\"):\n sgl_kernel.scale(input, 2.0, out=out)\n\n\nif __name__ == \"__main__\":\n import sys\n sys.exit(pytest.main([__file__, \"-q\"]))\n```\n\n---\n\n## Step 7: Add a benchmark (required)\n\nEvery benchmark must account for L2 cache reuse — see [`rules/kernel-benchmark.md`](../../rules/kernel-benchmark.md).\n\nCreate `python/sglang/kernels/aot/benchmark/bench_scale.py`:\n\n```python\nimport torch\n\nimport sgl_kernel\nfrom sglang.kernels.jit.benchmark import marker\n\n\ndef sglang_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n sgl_kernel.scale(input, factor, out=out)\n\n\ndef torch_scale(input: torch.Tensor, factor: float, out: torch.Tensor) -> None:\n torch.mul(input, factor, out=out)\n\n\nFN_MAP = {\"sglang\": sglang_scale, \"torch\": torch_scale}\n\n\n@marker.parametrize(\"dtype\", [torch.float16, torch.bfloat16, torch.float32], [torch.float16])\n@marker.parametrize(\"size\", [2**n for n in range(10, 20)], [4096]) # 1K .. 512K\n@marker.benchmark(\"provider\", [\"sglang\", \"torch\"])\ndef benchmark(dtype: torch.dtype, size: int, provider: str):\n input = torch.randn(size, dtype=dtype, device=\"cuda\")\n out = torch.empty_like(input)\n return marker.do_bench(\n FN_MAP[provider],\n # Pass every tensor through input_args (not a closure) so marker rotates\n # them across CUDA-graph calls to defeat L2 reuse.\n input_args=(input, 2.0, out),\n # Bandwidth = bytes(input) + bytes(out), both already in input_args.\n memory_output=None,\n )\n\n\nif __name__ == \"__main__\":\n benchmark.run()\n```\n\n---\n\n## Step 8: Build\n\nBuild:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j16\n```\n\nIf you need to limit host resource usage:\n\n```bash\ncd python/sglang/kernels/aot\nmake build -j1 MAX_JOBS=2 CMAKE_ARGS=\"-DSGL_KERNEL_COMPILE_THREADS=1\"\n```\n\n---\n\n## Step 9: Validate\n\nAfter building successfully, run the test and benchmark:\n\n```bash\npytest python/sglang/kernels/aot/tests/test_scale.py -q\npython python/sglang/kernels/aot/benchmark/bench_scale.py\n```\n\nPR CI also runs `pr-test-sgl-kernel.yml`, including the B200 job\n`sgl-kernel-b200-test` when kernel changes are detected. Use that job as the\nBlackwell coverage signal for AOT `sgl-kernel` changes.\n\n---\n\n## Troubleshooting\n\n- **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`\n- **Memory errors**: `compute-sanitizer --tool memcheck python ...`\n- **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`\n- **Binary bloat**: use `python/sglang/kernels/aot/analyze_whl_kernel_sizes.py`\n- **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time\n\n---\n\n## References\n\n- `python/sglang/kernels/aot/README.md`\n- `python/sglang/kernels/aot/include/sgl_kernel_ops.h`\n- `python/sglang/kernels/aot/csrc/common_extension.cc`\n- `python/sglang/kernels/aot/CMakeLists.txt`\n- `python/sglang/kernels/aot/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends\n- `python/sglang/kernels/aot/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern\n\n## Summary of Files Created/Modified\n\n```\npython/sglang/kernels/aot/csrc/elementwise/scale.cu # NEW: CUDA kernel + launcher\npython/sglang/kernels/aot/include/sgl_kernel_ops.h # MODIFIED: C++ declaration\npython/sglang/kernels/aot/csrc/common_extension.cc # MODIFIED: schema + dispatch registration\npython/sglang/kernels/aot/CMakeLists.txt # MODIFIED: add source file (alphabetical)\npython/sglang/kernels/aot/python/sgl_kernel/elementwise.py # MODIFIED: Python wrapper\npython/sglang/kernels/aot/python/sgl_kernel/__init__.py # MODIFIED: re-export Python API\npython/sglang/kernels/aot/tests/test_scale.py # NEW: tests\npython/sglang/kernels/aot/benchmark/bench_scale.py # NEW: benchmark\n```\n\n\n# Agents Skill\n\nPath: `.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/babysit-pr-to-pass-ci/agents/openai.yaml)\n\n\n---\nname: babysit-pr-to-pass-ci\ndescription: Start and persistently pursue a goal to babysit an SGLang pull request until selected GitHub Actions workflows pass on the latest PR head. Use when asked to monitor, babysit, retry, or fix PR CI for lint.yml, pr-test.yml, pr-test-extra.yml, AMD, or other named workflows; classify failures as PR-related versus flaky or infrastructural, auto-fix and push only small clean fixes, rerun failed jobs only up to 10 times, and ignore unselected workflows.\n---\n\n# Babysit a PR Until Selected CI Passes\n\nDrive only the selected workflows to green on the PR's latest head commit. Keep working across turns by starting a durable goal; do not stop merely because a run was dispatched, a retry began, or a fix was pushed.\n\n## Invocation and arguments\n\nAccept:\n\n```text\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] [WORKFLOW ...]\n$babysit-pr-to-pass-ci [PR_NUMBER|PR_URL] --only WORKFLOW [WORKFLOW ...]\n```\n\n- If the PR is omitted, resolve the PR associated with the current branch.\n- If workflows are omitted, select `lint.yml` and `pr-test.yml`.\n- Without `--only`, add named workflows to the two defaults.\n- With `--only`, monitor exactly the named workflows.\n- Accept a workflow basename such as `pr-test-extra.yml`, a repository path such as `.github/workflows/pr-test-extra.yml`, or an unambiguous workflow display name.\n- Treat `extra` as `pr-test-extra.yml`.\n- Accept natural-language equivalents, such as “PR 12345 plus extra and AMD.” Resolve “AMD” to an exact workflow file from the user's wording or repository context; do not guess when multiple AMD workflows fit.\n\nExamples:\n\n```text\n$babysit-pr-to-pass-ci\n$babysit-pr-to-pass-ci 12345\n$babysit-pr-to-pass-ci https://github.com/sgl-project/sglang/pull/12345 pr-test-extra.yml\n$babysit-pr-to-pass-ci 12345 --only pr-test-amd.yml\n```\n\n## Start or continue the durable goal\n\nTreat explicit invocation of this skill as authorization to create a goal and perform the scoped CI actions below.\n\n1. Resolve the PR and exact workflow set first.\n2. Inspect the current goal.\n3. If no unfinished goal exists, create one immediately using the product's goal mechanism. Include the resolved PR URL, selected workflow files, and this stopping condition: every selected workflow has a successful run for the PR's latest head SHA.\n4. Include these constraints in the goal objective:\n - Ignore every unselected workflow.\n - Diagnose each selected-workflow failure as PR-related, unrelated/infrastructure, or still uncertain.\n - Implement, validate, commit, and push a PR-related fix only when it is clean, non-tricky, and at most 100 changed lines.\n - Stop for user review before editing when the proposed fix is tricky, dirty, or over 100 changed lines.\n - Retry unrelated failures at most 10 times per workflow and head SHA, using failed-jobs-only reruns.\n - Never rerun a full workflow merely to recover failed tests.\n5. If the active goal already represents this same PR and workflow set, continue it instead of creating another goal.\n6. If a different unfinished goal is active, do not replace it silently. Report the conflict and ask the user to pause or clear it.\n\nDo not mark the goal complete until the completion check at the end of this skill succeeds.\n\n## Preflight\n\n1. Verify GitHub authentication and resolve the repository without changing remote configuration.\n2. Resolve the PR and record at least:\n - PR number and URL\n - open/draft state\n - base branch and base SHA\n - head branch and head SHA\n - head repository owner/name, including whether it is a fork\n - labels\n - whether maintainers may modify the branch\n3. Require an open PR. If it is draft and a selected workflow is blocked by the draft gate, stop and ask the user to mark it ready; do not change draft state automatically.\n4. Normalize every selected workflow to an existing file under `.github/workflows/`. Fail fast on ambiguous or nonexistent names.\n5. Read each selected workflow's triggers and gates. Do not assume all optional workflows behave like `pr-test.yml`.\n6. Inspect `git status`. Preserve all user changes. If a fix becomes necessary and the current checkout is dirty, on a different branch, or otherwise unsafe, use an isolated worktree for the PR head.\n\nKnown SGLang gates:\n\n- `pr-test.yml` requires the `run-ci` label for relevant PR changes. If its existing gate job failed because the label was absent, add `run-ci` when authorized and rerun failed jobs in that existing run; adding the label alone does not necessarily create a new `pr-test.yml` run.\n- `pr-test-extra.yml` requires both `run-ci` and `run-ci-extra`. It listens for relevant label events, so adding a missing label may create a fresh run. Prefer that fresh current-head run; rerun failed jobs in an existing current run only when needed.\n- AMD-extra workflows may use the same `run-ci` plus `run-ci-extra` gate. Verify the actual selected file.\n- Treat draft, cooldown, maintenance, and missing-label failures as gate/infrastructure conditions, not code regressions. Fix only the safe gate condition in scope; wait for cooldown or maintenance windows rather than burning retries immediately.\n\nAdding the known CI opt-in labels is within scope. Do not add unrelated labels, mark the PR ready, alter reviewers, merge, or mutate other PR metadata.\n\n## Track only current-head runs\n\nUse both PR checks and workflow-run data so workflow names, run IDs, attempts, jobs, and logs can be correlated. Prefer workflow file IDs over display names when filtering runs.\n\nFor each selected workflow:\n\n1. Select the newest run associated with the current PR head SHA.\n2. Ignore runs for older head SHAs, including green runs made stale by a new push.\n3. Treat `queued`, `requested`, `waiting`, `pending`, and `in_progress` as active; keep monitoring.\n4. Treat a workflow as passing only when its current-head run concludes `success`.\n5. If no current-head run exists, wait briefly for event delivery, then inspect the workflow's path filters, event triggers, and label gates.\n6. Trigger only through a mechanism explicitly supported by that workflow. Never invent `workflow_dispatch` inputs or dispatch a workflow against the default branch when the intent is to test the PR head.\n7. If a selected workflow cannot safely be triggered for the PR head, report the exact blocker and request direction.\n\nAfter any user push or skill-authored push, re-read the PR head SHA, discard stale run selections, and start tracking the new head. Reset retry counters for the new head SHA.\n\n## Monitor persistently\n\n- Maintain a compact ledger containing workflow file, head SHA, run ID, attempt, status, conclusion, skill-initiated retry count, and latest failure signature.\n- Poll at reasonable intervals, normally 30–60 seconds while runs are active. Give concise progress updates during long waits.\n- Re-check the PR head SHA on each monitoring cycle and before every rerun or push.\n- Ignore failures, cancellations, and pending checks from unselected workflows. Do not diagnose, retry, fix, or wait for them.\n- Do not declare success while a selected workflow is absent, stale, pending, skipped, cancelled, neutral, or failed.\n\n## Diagnose a selected-workflow failure\n\n1. Inspect the failed jobs and failed-step logs. In `pr-test.yml`, distinguish the root failure from cascade failures caused by `check-pr-test-health`, wait jobs, or final aggregation.\n2. Record the exact failing test/check, error text, affected runner or hardware, and relevant log URL.\n3. Compare the failure against the PR's changed files and diff.\n4. Classify it:\n\n **PR-related** when evidence ties the failure to changed code, tests, dependencies, configuration, generated files, or behavior introduced by the PR. A reproducible failure on the PR that does not occur on the base is strong evidence.\n\n **Unrelated/infrastructure** when evidence points to a transient network or download error, runner loss, GPU/driver instability, service outage, pre-existing main-branch failure, known flaky threshold, unrelated test area, or another environmental cause.\n\n **Uncertain** when evidence is insufficient. Do not label a failure flaky merely because it might be intermittent. Narrowly reproduce it, compare recent main/scheduled CI, inspect relevant code, or otherwise gather enough evidence to choose a category.\n\n5. When several root failures exist, classify each one. Address PR-related failures before spending retries on unrelated cascades.\n\n## Handle a PR-related failure\n\nEstimate the proposed fix before editing.\n\nA fix is eligible for automatic execution only when all are true:\n\n- The fix is well understood, localized, and clean.\n- The fix is not tricky, hacky, speculative, or a broad workaround.\n- The fix changes at most 100 lines, counting additions plus deletions in the fix patch, not the pre-existing PR diff.\n- The fix does not require a broad API or architecture redesign, risky dependency migration, large generated-file rewrite, or unrelated cleanup.\n- A focused validation is available.\n\nIf any condition fails, stop before editing. Report the diagnosis, proposed approach, expected files, estimated size, risks, and validation plan, then wait for user review.\n\nFor an eligible fix:\n\n1. Work from the latest PR head in a clean checkout or isolated worktree.\n2. Make only the necessary change and add or update focused tests when appropriate.\n3. Run the smallest meaningful local validation. For lint failures, reproduce the specific hook/check before considering broader lint validation.\n4. Review the resulting patch and count changed lines. If it becomes over 100 lines or turns tricky/dirty, do not commit or push it; stop for review and clearly identify the uncommitted skill-authored changes.\n5. Commit only skill-authored files. Never use `git add .`, rewrite user commits, amend without request, or include pre-existing local changes.\n6. Push a new commit non-forcibly to the exact PR head repository and branch. Never force-push. If permission or fork routing is unclear, stop and report the push command/target that needs authorization.\n7. Resolve the new PR head SHA and return to current-head monitoring. A push is progress, not completion.\n\n## Handle an unrelated or infrastructure failure\n\nRetry only the failed jobs in the current selected-workflow run:\n\n```bash\ngh run rerun --failed --repo \n```\n\nRules:\n\n- Never use a full-run rerun as a substitute for `--failed`.\n- Never rerun a stale run after the PR head SHA changes or a newer current-head run exists.\n- Count only successfully dispatched skill-initiated reruns.\n- Allow at most 10 such reruns per selected workflow and head SHA.\n- Wait for each new attempt to finish, inspect its result, and update the failure signature before deciding on another retry.\n- If the failure changes or becomes plausibly PR-related, return to diagnosis instead of blindly consuming all retries.\n- If 10 retries are exhausted, stop. Report the attempts, recurring signatures, job/run links, and evidence that the failure is unrelated or still uncertain. Do not exceed the cap without a new explicit user instruction.\n\nIf GitHub cannot rerun only the failed portion of the selected run, stop and ask the user rather than rerunning the full workflow.\n\n## Completion check\n\nBefore completing the goal:\n\n1. Fetch the PR head SHA again.\n2. For every selected workflow, verify that its newest applicable run is for that exact SHA and concludes `success`.\n3. Verify no selected workflow is still pending and no success was invalidated by a newer push.\n4. Ignore all unselected workflows even if they are failing.\n5. Mark the goal complete only after all selected workflows satisfy these checks.\n\nReport the final PR URL, head SHA, selected workflows, passing run links, any commits pushed, and retry counts. Keep the report scoped to the selected workflows.\n\n\n# Ci-test-audit Skill\n\nPath: `.agents/skills/ci-test-audit/action-items.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/ci-test-audit/action-items.md)\n\n\n---\nname: ci-test-audit\ndescription: Audit the existing test tree and CI configuration for improvements, using a catalog of patterns previously applied in this repo. Use when asked to audit tests or CI, shrink CI time or cost, find redundant or misplaced tests, clean up a test group, or review whether a CI change follows established practice.\n---\n\n# CI / Test Audit\n\nAudits what is already in `test/`, `.github/` and `scripts/ci/`. For writing a new test see\n[write-sglang-test](../write-sglang-test/SKILL.md); for how the pipeline dispatches and\ngates work see [ci-workflow-guide](../ci-workflow-guide/SKILL.md).\n\n## How to use it\n\nRead [action-items.md](action-items.md) first. It is a catalog of patterns that have\nbeen applied to this repo, each with a way to spot it and example PRs. The catalog is the substance of this skill; everything below is only scaffolding.\n\nThe user says what to audit -- a test group, a stage, a workflow file, or a complaint\nlike \"this suite is slow\" -- and how thoroughly. Read that scope and decide which\npatterns apply.\n\nAn audit is only useful if it is specific. \"This suite could be trimmed\" is not a\nfinding; \"`TestFooLargePage` is a strict subset of `TestFooRetractLargePage`, drop it\"\nis. Every finding needs the file, the pattern it matches, and what to do.\n\n## Reporting\n\nGroup findings by confidence in the claim, not by pattern id. Lead with the ones where\nthe evidence is in the file you just read; keep the speculative ones separate and say\nwhat you would need to check. Skip anything you are not reasonably sure about -- a long\nlist of maybes costs more to triage than it saves.\n\nPropose, do not apply. Deleting a test, moving a registration to another stage, and\nchanging a threshold are all decisions for the user. Land them as separate changes,\nsince a trim and a threshold change fail for different reasons.\n\n## Keeping the catalog current\n\nExample PRs age. When a pattern shows up in newer work, or a new pattern recurs, update\n`action-items.md`: the pattern text should stay general, the examples are\nreplaceable evidence.\n\n\n---\nname: ci-workflow-guide\ndescription: Guide to SGLang CI workflow orchestration — stage ordering, fail-fast, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages.\n---\n\n# SGLang CI Workflow Orchestration Guide\n\nThis skill covers the CI **infrastructure** layer — how tests are dispatched, gated, and aborted on failure across stages. For test authoring (templates, fixtures, registration, model selection), see the [write-sglang-test skill](../write-sglang-test/SKILL.md).\n\n---\n\n## Naming Conventions\n\n- **Suite**: `base-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `base-b-test-1-gpu-small`)\n- **Test group**: Directory-level registered test group under `test/registered/` (e.g., `hicache` maps to `test/registered/hicache/test_*.py`)\n- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)\n\n---\n\n## Key Files\n\n| File | Role |\n|------|------|\n| `.github/workflows/pr-test.yml` | Main workflow — all stages, jobs, conditions, matrix definitions |\n| `.github/workflows/pr-test-extra.yml` | Extra workflow — gated by BOTH `run-ci` and `run-ci-extra` labels |\n| `.github/workflows/pr-gate.yml` | PR gating: draft check, `run-ci` label, per-user rate limiting |\n| `.github/actions/check-pr-test-health/action.yml` | Cross-job fail-fast: queries API for any failed job |\n| `.github/actions/wait-for-jobs/action.yml` | Stage gating: polls API until stage jobs complete |\n| `.github/actions/check-maintenance/action.yml` | Maintenance mode check |\n| `.github/scripts/ci-labels.cjs` | Resolves the four CI control labels into dispatch axes |\n| `test/run_suite.py` | Suite runner: collects, filters, partitions, executes tests |\n| `python/sglang/test/ci/ci_register.py` | Test registration (AST-parsed markers), LPT auto-partition |\n| `python/sglang/test/ci/ci_utils.py` | `run_unittest_files()`: execution, retry, continue-on-error |\n| `scripts/ci/utils/slash_command_handler.py` | Handles slash commands from PR comments |\n\n---\n\n## Architecture Overview\n\n```\n ┌──────────────┐\n │ build kernel │\n └──────┬───────┘\n │\n ├─ check-changes ──── detects which packages changed\n │ (main_package, sgl_kernel, jit_kernel, multimodal_gen)\n │\n ├─ call-gate ──────── pr-gate.yml (draft? label? rate limit?)\n │\n ├─────────────────────────────────────────────────────┐\n │ │\n ▼ │\n ┌─────────────────────────────────────┐ │\n │ Base A (~3 min) │ │\n │ pre-flight check │ │\n │ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-1-gpu-small │ │ │\n │ │ (small GPUs) │ │ │\n │ └─────────────────────────────┘ │ │\n │ ┌─────────────────────────────┐ │ │\n │ │ base-a-test-cpu │ │ │\n │ │ (CPU) │ │ │\n │ └─────────────────────────────┘ │ │\n └──────┬──────────────────────────────┘ │\n │ │\n ▼ ▼\n ┌─────────────────────────────────────┐ ┌──────────────────────────┐\n │ Base B (~30 min) │ │ kernel test │\n │ base tests │ └──────────────────────────┘\n │ │ ┌──────────────────────────┐\n │ ┌─────────────────────────────┐ │ │ multimodal gen test │\n │ │ base-b-test-1-gpu-small │ │ └──────────────────────────┘\n │ │ (small GPUs, e.g. 5090) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-1-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-b-test-2-gpu-large │ │\n │ │ (large GPUs, e.g. H100) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ Base C (~30 min) │\n │ advanced tests │\n │ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-h100 │ │\n │ │ (H100 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-8-gpu-h200 │ │\n │ │ (8 x H200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ base-c-test-4-gpu-b200 │ │\n │ │ (4 x B200 GPUs) │ │\n │ └─────────────────────────────┘ │\n │ ┌─────────────────────────────┐ │\n │ │ Other advanced tests │ │\n │ │ (DeepEP, PD Disagg, GB300) │ │\n │ └─────────────────────────────┘ │\n └──────┬──────────────────────────────┘\n │\n ▼\n ┌─────────────────────────────────────┐\n │ pr-test-finish │\n │ aggregates all results, fails if │\n │ any job failed/cancelled │\n └─────────────────────────────────────┘\n```\n\n**Every stage test job** includes a `check-pr-test-health` step after checkout — if any job in the run has already failed, the job fails fast (red X) with a root cause annotation.\n\n**Scheduled runs** skip `wait-for-base-*` jobs, running all stages in parallel. Fail-fast is also disabled.\n\n---\n\n## Fail-Fast Layers\n\n4 layers of fail-fast, from fine to coarse:\n\n| Layer | Mechanism | Granularity | Disabled on schedule? |\n|-------|-----------|-------------|----------------------|\n| **1. Test method → file** | `unittest -f` (failfast) | One test method fails → entire test file stops immediately | Yes |\n| **2. File → suite** | `run_unittest_files()` default | One test file fails → entire suite stops (`--continue-on-error` off) | Yes |\n| **3. Job → job (same stage)** | `check-pr-test-health` action | One job fails → other waiting jobs in same stage fail-fast (red X) | Yes |\n| **4. Stage → stage (cross-stage)** | `wait-for-base-*` + `needs` | Base A fails → base B/C jobs skip entirely (never get a runner) | Yes (wait jobs skipped) |\n\n- **Layer 1**: `-f` flag appended to all `python3 -m pytest` / `unittest` invocations in `ci_utils.py`\n- **Layer 2**: `--continue-on-error` flag in `run_suite.py` — off for PRs, on for scheduled runs\n- **Layer 3**: `check-pr-test-health` auto-detects `schedule` event and skips; filters out cascade failures to show only root cause jobs\n- **Layer 4**: `wait-for-base-*` jobs are conditioned on `github.event_name == 'pull_request'` — skipped for scheduled runs\n\n---\n\n## Execution Modes\n\n| Aspect | PR (`pull_request`) | Scheduled (`cron`, every 6h) | Manual dispatch (`workflow_dispatch`) |\n|--------|---------------------|------------------------------|--------------------------------------|\n| **Stage ordering** | Sequential: A → B → C via `wait-for-base-*` | Parallel (all at once) | Single target stage only |\n| **Cross-job fail-fast** | Yes (`check-pr-test-health`) | Yes | Yes |\n| **continue-on-error** | No (stop at first failure within suite) | Yes (run all tests) | No |\n| **Retry** | Enabled | Enabled | Enabled |\n| **max_parallel** | 3 (default), 14 if `max-concurrency` label | 14 | 3 (default), 14 if `max-concurrency` |\n| **PR gate** | Yes (draft, label, rate limit) | Skipped | Skipped |\n| **Concurrency** | `cancel-in-progress: true` per PR | Queue (no cancel) | Isolated per stage+SHA |\n\nFour labels relax these limits for one PR: `bypass-fail-fast`, `parallel-stages`,\n`max-concurrency`, and `highest-priority` (all three). `.github/scripts/ci-labels.cjs`\nresolves them; the [contribution guide](https://docs.sglang.io/developer_guide/contribution_guide.html#ci-control-labels)\ndescribes what each one does.\n\n---\n\n## Stage Gating (`wait-for-jobs` action)\n\n`wait-for-base-a` and `wait-for-base-b` are lightweight `ubuntu-latest` jobs that poll the GitHub Actions API.\n\n**How it works:**\n1. Calls `listJobsForWorkflowRun` to list all jobs in the current run\n2. Matches jobs by exact name or prefix (for matrix jobs, e.g., `base-b-test-1-gpu-small (3)`)\n3. If any matched job has `conclusion === 'failure'` → fail immediately (fail-fast)\n4. If all matched jobs are completed and count matches `expected_count` → success\n5. Otherwise → sleep `poll-interval-seconds` (default: 60s) and retry\n6. Timeout after `max-wait-minutes` of wall-clock time (240 min for base-a, 350 min for base-b). Keep budgets under the 6h hard cap on GitHub-hosted runners: a wait killed by the cap ends `cancelled`, not `timeout`, and `timeout-minutes` cannot raise it. The base-b wait sets `fail-on-timeout: 'false'`, so a timeout lets base-c start; `pr-test-finish` fails the run on any base-b failure regardless\n\n**Job specs example** (base-b):\n```json\n[\n {\"prefix\": \"base-b-test-1-gpu-small\", \"expected_count\": 8},\n {\"prefix\": \"base-b-test-1-gpu-large\", \"expected_count\": 14},\n {\"prefix\": \"base-b-test-2-gpu-large\", \"expected_count\": 4},\n {\"prefix\": \"base-b-test-4-gpu-b200\", \"expected_count\": 1}\n]\n```\n\n> **Critical**: `expected_count` must match the matrix size. If you add/remove matrix entries, update the wait job's spec accordingly.\n\n**PR only**: Condition `github.event_name == 'pull_request' && !inputs.target_stage` — scheduled runs and manual dispatches skip these entirely, allowing parallel execution.\n\n---\n\n## Cross-Job Fail-Fast (`check-pr-test-health` action)\n\nComposite action called after checkout in every stage test job (21 jobs total across `pr-test.yml`, `pr-test-multimodal-gen.yml`, `pr-test-sgl-kernel.yml`, `pr-test-jit-kernel.yml`).\n\n**How it works:**\n1. Queries `listJobsForWorkflowRun` for the current workflow run\n2. Filters for **root cause failures only** — jobs with `conclusion === 'failure'` whose failing step is NOT `check-pr-test-health` (excludes cascade failures)\n3. If root cause failures found → calls `core.setFailed()` with the list of root cause job names\n4. If none → does nothing (step succeeds)\n\n**Cascade filtering**: When job A fails fast due to the health check, it also has `conclusion: failure`. Without filtering, job B would list both the original failure AND job A's fail-fast. The filter checks each failed job's `steps` array — if the failing step name contains `check-pr-test-health` or `Check PR test health`, it's excluded from the root cause list.\n\n**Usage pattern:**\n```yaml\nsteps:\n - name: Checkout code\n uses: actions/checkout@v4\n ...\n\n - uses: ./.github/actions/check-pr-test-health\n id: pr-test-health\n\n - name: Install dependencies # skipped automatically if health check failed\n ... # (default if: success() is false)\n\n - name: Run test # also skipped\n ...\n```\n\n**Visual effect**: Job shows **red X** (failure) with error annotation showing root cause job names. Subsequent steps are naturally skipped (default `if: success()` is false after a failed step). No per-step `if` guards needed.\n\n**No stage filtering**: Checks ALL jobs in the run, not just the current stage. Any failure anywhere triggers fail-fast.\n\n**Error message example:**\n```\nFail-fast: skipping — root cause job(s): base-b-test-1-gpu-small (0), base-b-test-1-gpu-small (1)\n```\n\n---\n\n## Within-Suite Failure Handling\n\nControlled by `run_unittest_files()` in `python/sglang/test/ci/ci_utils.py`.\n\n### Flags\n\n| Flag | PR default | Scheduled default | Effect |\n|------|------------|-------------------|--------|\n| `--continue-on-error` | Off | On | Off: stop at first failure. On: run all files, report all failures at end |\n| `--enable-retry` | On | On | Retry retriable failures (accuracy/perf assertions) |\n| `--max-attempts` | 2 | 2 | Max attempts per file including initial run |\n\n### Retry Classification\n\nWhen a test fails and retry is enabled, the output is classified:\n\n**Non-retriable** (checked first — real code errors):\n`SyntaxError`, `ImportError`, `ModuleNotFoundError`, `NameError`, `TypeError`, `AttributeError`, `RuntimeError`, `CUDA out of memory`, `OOM`, `Segmentation fault`, `core dumped`, `ConnectionRefusedError`, `FileNotFoundError`\n\n**Retriable** (accuracy/performance):\n`AssertionError` with comparison patterns (`not greater than`, `not less than`, `not equal to`), `accuracy`, `score`, `latency`, `throughput`, `timeout`\n\n**Default**: Unknown `AssertionError` → retriable. Other unknown failures → not retriable.\n\n### How `continue_on_error` is set\n\nIn `pr-test.yml`'s `check-changes` job:\n- `schedule` runs or `run_all_tests` flag → `continue_on_error = 'true'`\n- PR runs → `continue_on_error = 'false'`\n\nEach test job propagates via:\n```yaml\nenv:\n CONTINUE_ON_ERROR_FLAG: ${{ needs.check-changes.outputs.continue_on_error == 'true' && '--continue-on-error' || '' }}\nrun: |\n python3 run_suite.py --hw cuda --suite $CONTINUE_ON_ERROR_FLAG\n```\n\n---\n\n## Test Partitioning\n\nLarge suites are split across matrix jobs using the **LPT (Longest Processing Time) heuristic** in `ci_register.py:auto_partition()`:\n\n1. Sort tests by `est_time` descending, filename as tie-breaker (deterministic)\n2. Greedily assign each test to the partition with smallest cumulative time\n3. Result: roughly equal total time per partition\n\n**Partition table** (CUDA per-commit suites):\n\n| Suite | Partitions | Runner | max_parallel |\n|-------|-----------|--------|-------------|\n| `base-a-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — |\n| `base-a-test-cpu` | 4 | `ubuntu-latest` | — |\n| `base-b-test-1-gpu-small` | 8 | `1-gpu-5090` | 8 |\n| `base-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) |\n| `base-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — |\n| `base-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-b-kernel-unit-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |\n| `base-b-kernel-unit-test-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |\n| `base-b-kernel-benchmark-test-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |\n| `base-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — |\n| `base-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — |\n| `base-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — |\n| `base-c-test-deepep-4-gpu-h100` | 1 (no matrix) | `4-gpu-h100` | — |\n| `base-c-test-4-gpu-b200` | 3 | `4-gpu-b200` | — |\n| `base-c-test-4-gpu-b200-small` | 3 | `4-gpu-b200-low-disk` | — |\n| `base-c-test-8-gpu-b200` | registered only | `8-gpu-b200` | — |\n| `base-c-test-4-gpu-gb200` | registered only | `4-gpu-gb200` | — |\n\n> **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`.\n>\n> **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`.\n\n**Workflow usage:**\n```yaml\nstrategy:\n matrix:\n partition: [0, 1, 2, 3, 4, 5, 6, 7]\nsteps:\n - run: python3 run_suite.py --hw cuda --suite base-b-test-1-gpu-small \\\n --auto-partition-id ${{ matrix.partition }} --auto-partition-size 8\n```\n\n---\n\n## check-changes Job\n\nDetermines which test suites to run based on file changes.\n\n### Detection Methods\n\n| Trigger | Method | Details |\n|---------|--------|---------|\n| `pull_request` | `dorny/paths-filter` | Detects changes via GitHub diff |\n| `workflow_dispatch` (with `pr_head_sha`) | GitHub API | `repos/{repo}/compare/main...{sha}` |\n| `schedule` / `run_all_tests` | Force all true | Runs everything |\n\n### Output Flags\n\n| Output | Triggers |\n|--------|----------|\n| `main_package` | Base A/B/C test suites |\n| `sgl_kernel` | Kernel wheel builds + kernel test suites; also switches B200 jobs to kernel-build runner labels outside `target_stage` mode |\n| `jit_kernel` | JIT kernel test workflow |\n| `multimodal_gen` | Multimodal-gen test workflow |\n\n> **Note**: In `target_stage` mode, `sgl_kernel` is only active when `include_wheel_build=true`. Without that opt-in, kernel-change reruns fail validation instead of running a target stage without freshly built wheels. Outside `target_stage`, `sgl_kernel=true` switches B200 jobs from `4-gpu-b200` / `4-gpu-b200-low-disk` to `4-gpu-b200-kernel` / `4-gpu-b200-kernel-low-disk`.\n\n---\n\n## Concurrency Control\n\n```\ngroup: pr-test-{event_name}-{branch}-{pr_sha}-{stage}\n```\n\n| Segment | Source | Purpose |\n|---------|--------|---------|\n| `event_name` | `github.event_name` | Prevents scheduled runs colliding with fork PRs named `main` |\n| `branch` | `github.head_ref \\|\\| github.ref_name` | Per-branch isolation |\n| `pr_sha` | `inputs.pr_head_sha \\|\\| 'current'` | Isolates manual dispatches from main runs |\n| `stage` | `inputs.target_stage \\|\\| 'all'` | Allows parallel stage dispatches |\n\n`cancel-in-progress: true` for `pull_request` events (new push cancels old run), `false` for `workflow_call`.\n\n---\n\n## How To: Add a New Stage Job\n\n1. Define the job in `pr-test.yml` with `needs: [check-changes, call-gate, wait-for-base-X, ...]`\n2. Copy the `if:` condition pattern from an existing same-stage job (handles `target_stage`, `schedule`, `main_package`)\n3. Add `checkout` step\n4. Add `check-pr-test-health` step (after checkout) — if any prior job failed, `core.setFailed()` fires and all subsequent steps auto-skip via default `if: success()`\n5. Add `check-maintenance` step\n6. Add `download-artifact` step if `sgl_kernel` changed\n7. Add `install dependencies` step\n8. Add `run test` step with `$CONTINUE_ON_ERROR_FLAG`\n9. Add `upload-cuda-coredumps` step with `if: always()`\n10. Register the suite name in `PER_COMMIT_SUITES` in `test/run_suite.py`\n11. If using matrix, add `--auto-partition-id` and `--auto-partition-size` to the run command\n12. **Update `wait-for-base-X`** job spec with the new job name and `expected_count` (if matrix)\n13. **Add the job to `pr-test-finish.needs`** list\n\n---\n\n## How To: Debug CI Failures\n\n| Symptom | Likely cause | What to check |\n|---------|-------------|---------------|\n| All stage-B/C jobs green but steps skipped | Earlier job failed, `check-pr-test-health` triggered | Find the actual failed job (red X) |\n| `wait-for-base-b` timeout | `expected_count` doesn't match matrix size | Verify job spec counts match `matrix:` array length |\n| `pr-test-finish` fails but all jobs green | A job was `cancelled` (counts as failure in finish) | Check concurrency cancellation, and whether a `wait-for-base-*` job hit the 6h hosted-runner cap |\n| Tests pass locally but fail in CI | Partition assignment, runner GPU type, or `est_time` inaccuracy | Check which partition the test lands in; verify runner label |\n| Flaky test retried and passed | Retriable failure (accuracy/perf) | Check `[CI Retry]` markers in job logs |\n| Flaky test NOT retried | Matched non-retriable pattern | Check if error matches `NON_RETRIABLE_PATTERNS` in `ci_utils.py` |\n\n---\n\n## Slash Commands\n\n| Command | Effect |\n|---------|--------|\n| `/tag-run-ci-label` | Adds `run-ci` label to PR |\n| `/tag-run-ci-label extra` | Adds both `run-ci` and `run-ci-extra` labels |\n| `/rerun-failed-ci` | Reruns failed jobs in the latest workflow run |\n| `/tag-and-rerun-ci` | Adds `run-ci` label + reruns failed |\n| `/tag-and-rerun-ci extra` | Adds both `run-ci` and `run-ci-extra` labels + reruns failed |\n| `/run-full-ci` | Short form of `/tag-and-rerun-ci extra` (baseline + extra). Alias: `/rerun-full-ci` |\n| `/run-extra-ci` | Adds both labels + reruns **only** `PR Test Extra`, leaving baseline runs alone. Alias: `/rerun-extra-ci` |\n| `/rerun-test [ ...]` | Reruns specific test file(s) via `rerun-test.yml`. A file arg containing a glob metacharacter (`*`, `?`, `[...]`) expands against `test/registered/` and the multimodal test dir to every matching `test_*.py` (e.g. `/rerun-test test_*backend*.py` — wrap in backticks so GitHub doesn't italicize the `*`); matches are deduped, grouped by dispatch shape, and can't carry a `::test` selector. No match → single ⛔ reply, nothing dispatched. |\n| `/rerun-group [ ...]` | Expands registered test groups, then reuses `/rerun-test` |\n| `/run-cpu-ci` | Dispatches `run-cpu-ci.yml` on main: only the `base-a-test-cpu` suite, at the PR head SHA, no caches saved. Same gate as `/rerun-test`. Alias: `/rerun-cpu-ci` |\n\nHandled by `scripts/ci/utils/slash_command_handler.py` → `.github/workflows/slash-command-handler.yml`.\n\n### Label-gated workflow dispatch (pr-test, pr-test-extra)\n\n`pr-test.yml` and `pr-test-extra.yml` both listen for `pull_request.labeled` (in addition to `opened`/`synchronize`/`reopened`). The `check-changes.if` gate has two clauses:\n\n1. **For `labeled` events**: the just-added label must be one of the gating labels (`run-ci` for pr-test, `run-ci` or `run-ci-extra` for pr-test-extra) — otherwise every unrelated label addition would dispatch a full CI run.\n2. **All events**: the PR must currently carry the required labels.\n\nThis is what lets `/tag-run-ci-label` (and the `extra` variant) trigger a fresh CI run without an extra push.\n\n**Caveat — skipped runs cannot be un-skipped by `run.rerun()`:** GitHub's rerun API reuses the original event payload, so rerunning a `pull_request`-event run that was skipped because of missing labels will skip again (label set in the frozen payload doesn't update). The only way to recover a label-skipped run is to add the missing label, which fires a fresh `labeled` event with the current label set. `handle_rerun_failed_ci` in the slash handler is for rerunning failed/non-label-skipped runs; it cannot revive label-skipped ones.\n\n\n# References Skill\n\nPath: `.agents/skills/clean-startup-log/references/noise-sources.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/clean-startup-log/references/noise-sources.md)\n\n\n---\nname: clean-startup-log\ndescription: Audit SGLang startup logs, save evidence, and propose cleanup for user review. With no arguments, run Qwen3-8B at TP1 and TP2 plus gpt-oss-20b at TP1.\ndisable-model-invocation: true\n---\n\n# Audit SGLang Startup Logs\n\nThe default outcome is saved logs and a findings report. Apply runtime changes\nonly after the user selects them. A request to edit this skill does not itself\nlaunch servers.\n\n## Default runs\n\nA bare `$clean-startup-log` invocation runs these cases sequentially without\nasking for commands. Explicit commands, models, or TP sizes replace this matrix.\n\n| Case / log filename | Command |\n|---|---|\n| `qwen3-8b-tp1.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B` |\n| `qwen3-8b-tp2.log` | `uv run sglang serve --model-path Qwen/Qwen3-8B --tp 2` |\n| `gpt-oss-20b-tp1.log` | `uv run sglang serve --model-path openai/gpt-oss-20b` |\n\nThese cover dense, tensor-parallel, and MoE/hybrid sliding-window attention\nstartup. Reuse complete captures from the current audit when code and environment\nhave not changed.\n\n## Capture logs\n\n1. Check the checkout, free GPUs, and ports once. Use the requested command when\n resources are free; otherwise select free GPUs with `CUDA_VISIBLE_DEVICES` and\n an unused `--port`, recording the adjustments. Leave existing servers alone.\n2. Create a unique directory with `mktemp -d /tmp/sglang-startup-audit-XXXXXX`.\n Save raw stdout and stderr together in a separate log for each case, for\n example with `set -o pipefail` and `COMMAND 2>&1 | tee LOG_PATH`. Record commands,\n GPU IDs, ports, commit, relevant overrides, and readiness status.\n3. Wait for `The server is fired up and ready to roll!`, then stop that server and\n its workers before the next case. First runs can spend many minutes downloading\n weights or compiling FlashInfer kernels; check download/compiler activity\n before treating a quiet log as a hang. Preserve partial logs for failed or\n stalled starts and report the last stage. Continue independent cases when possible.\n4. Preserve the user's logging configuration, including `NCCL_DEBUG`. If NCCL\n verbosity needs explaining, inspect relevant shell settings, `NCCL_CONF_FILE`,\n and `/etc/nccl.conf`. Do not override intentional diagnostics or recommend\n `NCCL_DEBUG=WARN` solely because the output is long. Avoid full environment dumps.\n\n## Investigate efficiently\n\n- Scan for deprecations, duplicate handler output, unrelated import failures,\n unformatted prints, and unexpected warnings. Read representative excerpts and\n counts instead of repeatedly dumping `server_args`, progress redraws, or NCCL\n diagnostics. Normalize carriage returns for analysis only; preserve raw logs.\n- Trace each candidate to its actual emitter with focused `rg` searches. Inspect\n its log level: SGLang's formatter may omit severity. Group shared signatures\n across cases and distinguish handler duplication from separate GPU/process calls.\n- Repetition, WARNING severity, or a different third-party format alone does not\n establish a cleanup need. Consider whether the message explains configuration,\n progress, resource use, or an operational limitation.\n- Consult [noise-source hints](references/noise-sources.md) only for a matching\n signature or an unresolved emitter. Verify current code rather than trusting\n historical line numbers, fix status, or assumptions about unrelated models.\n\n## Accepted output\n\nPreserve these reviewed messages unless the user requests a different policy:\n\n- NCCL diagnostics enabled by the user's environment or host configuration.\n- NUMA permission warnings, including one check per GPU in TP runs.\n- GPT-OSS MXFP4 backend-selection warnings and default page-size selection warnings.\n- `Init Unified Radix Cache. Components: ... Tree Core: ...`, tree-cache summaries,\n SWA allocation details, and per-rank memory/timing records.\n- Useful progress bars, warmup HTTP access logs, uv synchronization messages,\n isolated NCCL/Gloo startup lines, and one timestamped HF authentication warning.\n\nThese can appear in a clean startup log. Do not repeatedly propose the declined\nNUMA deduplication, backend/page-size level changes, or NCCL verbosity override.\nKeep real operational warnings visible: for example, a Harmony vocabulary failure\ncan disable `/v1/responses` even when server readiness and `/generate` succeed.\n\n## Report before changing code\n\nReturn a compact run table with readiness status and clickable raw-log links.\nFor each actual cleanup candidate, give an exact representative message, affected\ncases/counts, source file/function, and specific proposed behavior. Distinguish\nconfirmed findings from suspicions and operational failures from logging noise.\n\nIf there are no actionable cleanup findings, say the logs are clean and no\nfurther cleanup is needed. Otherwise, ask which numbered changes to adopt and\nwait for the user's selections before editing runtime code or preparing patches.\nHonor existing approvals and declined items without asking again.\n\n## Apply selected changes\n\n- Batch compatible approved edits, then verify affected cases once. Repeat a\n startup only for a new change, failure, or unresolved concern; do not relaunch\n after every one-line edit. Save verification logs separately from baselines.\n- Preserve useful warnings and application handlers. HF can warn during early\n CLI model detection before `configure_logger()`, and spawned processes have\n independent logger state. Keep `configure_hf_hub_logger()` in both\n `suppress_noisy_warnings()` and `configure_logger()`; make repeated setup safe.\n- Keep the legacy compiled-kernel cache migration notice at DEBUG. Avoid broad\n library-level suppression or fd redirection for a narrow logging problem.\n- Run relevant formatting and focused existing checks. Add tests only when they\n verify meaningful behavior, not a log-level spelling. Report changes and\n verification; create branches, commits, and PRs when requested.\n\n\n---\nname: compute-mamba-ratio\ndescription: Compute the optimal --mamba-full-memory-ratio (or --max-mamba-cache-size pin) for a hybrid attention + linear-attention (Mamba / GDN / KDA) model's two serving memory pools, from the workload and serving config. Use when a user asks what ratio to set, why concurrency is clamped, or how to size the state vs KV pools for a hybrid model.\n---\n\n# Optimal hybrid dual-pool ratio (`--mamba-full-memory-ratio`)\n\nA hybrid model (attention layers + linear-attention layers — the recurrent-state family: Mamba/SSM, GDN, KDA, etc.) splits serving memory into two independently-budgeted pools, fixed once at startup:\n\n- **state pool** (the linear-attention recurrent state) → caps **concurrency** (hard: whole slots, worst-case reserved, fail-loud)\n- **full-KV pool** (attention KV) → caps **context × concurrency** (soft: paged, over-committable via retraction)\n\n`--mamba-full-memory-ratio r` splits the post-weight budget: `mamba_budget = rest · r/(1+r)`, i.e. `mamba_budget : kv_budget = r`. This skill picks the `r` (or the pin-`--max-mamba-cache-size` alternative) at which **neither pool bottlenecks first** for the user's workload.\n\n## The formula\n\n```\nr* = (S + D) · token_equiv · dcp_size / L\ntoken_equiv = state_bytes_per_slot / kv_bytes_per_token\n```\n\n- `L` = average context length per request (input + output tokens)\n- `token_equiv` = full-KV token-equivalent of one state slot\n- `S` = state slots per running request (cache-strategy dependent, table below)\n- `D` = `--speculative-num-draft-tokens` (0 if NOSPEC); each running req carries `D` extra intermediate states\n- `dcp_size` = `--dcp-size` (1 without DCP). DCP shards the per-rank KV by `dcp_size`, so KV gets ~`dcp_size×` cheaper per request → the balance shifts that much toward the state pool.\n\n`r` is dimensionless (just the split). To also predict the actual concurrency you need `rest` (below).\n\n## Inputs to collect from the user\n\n1. **`L`** — average context (input + output) in tokens.\n2. **The two per-GPU byte constants** — one of:\n - **(a) measured (preferred, exact)** — from one boot log at *any* ratio:\n - `Mamba Cache is allocated. ... ssm_state size X GB` with `max_mamba_cache_size: N` → `state_bytes_per_slot = X / N`\n - `KV Cache is allocated. #tokens: M, KV size: Y GB` → `kv_bytes_per_token = Y / M`\n - **(b) derived** — model arch (linear-layer count + state dims `d_state/d_conv/heads/head_dim`; attention type + KV dims: MLA latent dim, or GQA `kv_heads·head_dim·layers`) **× the dtypes** below.\n3. **`S`** — from `--mamba-radix-cache-strategy`, the overlap scheduler, and `SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK` (table below).\n4. **`D`** — `--speculative-num-draft-tokens` (0 if NOSPEC; **also 0 when ReplaySSM spec-verify is enabled** — see caveats).\n5. **`dcp_size`** — `--dcp-size` (1 if no DCP). If DCP **and** spec with a replicated draft KV, apply the draft caveat (below).\n6. **KV dtype** (bf16 / fp8) and **ssm dtype** (fp32 / bf16) — they set the two byte constants (see the `token_equiv` 2×2).\n7. *(only to also predict the clamp, not just `r`)* **`rest`** = per-GPU memory − weights at the chosen `--mem-fraction-static`. Read `avail mem` after `Load weight end`, or `Memory pool end. avail mem` + pool sizes, from the boot log.\n\n### `S` (state slots per running request), set by `--mamba-radix-cache-strategy`\n\nAll strategies keep prefix caching on; they differ in the track buffer that snapshots chunk-boundary state under the overlap scheduler. `S = base + ping-pong`, where base is 3 (live state + radix retention/COW headroom), and ping-pong = 2 (overlap on, non-lazy) / 1 (lazy, **or** overlap off) / 0 (no track buffer at all). This mirrors `kv_cache_configurator._calculate_mamba_ratio`; read it there if a release moves the constants.\n\n`SGLANG_OPT_MAMBA_SKIP_DECODE_LOCK=1` takes the base to 2, since the decode-time skip frees one resident slot per running request. `no_buffer` is the exception and stays at an effective 3: it adds that slot straight back, because its binding limit is the prefill-to-decode peak, which a decode-time change does not shrink.\n\n| strategy | S | S with decode-lock skip | notes |\n|---|---|---|---|\n| `--disable-radix-cache` | 1 | 1 | prefix cache off entirely; most concurrency, no reuse |\n| `no_buffer` | 3 | 3 | no track buffer; requires overlap **off** (+ `page_size 1`) → lower decode throughput |\n| `extra_buffer` (default), overlap on | 5 | 4 | track buffer reserved per running request |\n| `extra_buffer`, overlap off | 4 | 3 | same buffer, but ping-pong costs 1 slot instead of 2 |\n| `extra_buffer_lazy` (recommended) | 4 | 3 | track buffer allocated lazily at the boundary → 1 fewer slot/req; requires overlap **on** |\n\nThe skip is a real concurrency lever, not a rounding detail: at the default strategy it takes `S` from 5 to 4, so the same state budget admits 25% more requests and `r*` drops by the same factor.\n\nOverlap counts as off when `--disable-overlap-schedule` is passed **or** `--pp-size > 1` (pipeline parallelism turns it off for you). So a PP deployment left on the default strategy sits on the `S = 4` row, not `S = 5` — using 5 there overstates the state cost and starves KV.\n\n### `token_equiv` — compute it from the two byte constants\n\n`token_equiv = state_bytes_per_slot / kv_bytes_per_token`. Both bytes are **per-model, per-GPU** — read them from a boot log (or derive from the arch + dtypes) and divide. That's the whole thing; the dtype effects below are just for a quick mental estimate.\n\nThe recurrent state has two parts: `state_bytes = SSM_bytes + conv_bytes`. The **SSM** part scales with `--mamba-ssm-dtype`; the **conv** part is always bf16 (`SGLANG_MAMBA_CONV_DTYPE`, fixed). How the dtype knobs move `token_equiv`:\n\n- **fp8 KV → ×2** (universal, exact): fp8 is 1 byte vs bf16's 2, so `kv_bytes_per_token` halves.\n- **fp32 ssm vs bf16 ssm → ×≈2**: only the SSM tensor doubles, conv stays fixed, so the exact factor is `(2·SSM + conv)/(SSM + conv)` — just under 2. **conv is usually small relative to SSM, so estimate ≈2** (e.g. one measured model: conv ≈7% of state → factor ≈1.9). Only bother with the exact split if a model's conv is unusually large.\n- **fp32 ssm + fp8 KV → ×≈4**.\n\nFor an actual number, don't apply factors — read `state_bytes_per_slot` and `kv_bytes_per_token` for **your** model+dtype from the boot log and divide.\n\nThe `GB` the boot log prints for both pools is **GiB** (`1024³`). Convert both the same way and `token_equiv` is unaffected either way, since it is a ratio — but get it right before comparing an absolute byte figure against your own log.\n\nWorked example (one measured model, TP8, fp32 ssm + fp8 KV), read off one boot log:\n\n```\nMamba Cache is allocated. max_mamba_cache_size: 257, conv_state size: 0.46GB, ssm_state size: 13.04GB\nKV Cache is allocated. dtype: torch.float8_e4m3fn, #tokens: 1167552, KV size: 15.03 GB\n```\n\n`state_bytes_per_slot = (0.46 + 13.04) GiB / 257 ≈ 56.4 MB`, `kv_bytes_per_token = 15.03 GiB / 1167552 ≈ 13.8 KB` → `token_equiv ≈ **4080**`.\n\n## Compute\n\n```python\ndef optimal_ratio(L, state_bytes_per_slot, kv_bytes_per_token, S, D=0, dcp_size=1):\n token_equiv = state_bytes_per_slot / kv_bytes_per_token\n r = (S + D) * token_equiv * dcp_size / L\n return r # value of --mamba-full-memory-ratio (>1 is legal)\n\ndef predict_clamp(rest_bytes, r, state_bytes_per_slot, S, D=0):\n mamba_budget = rest_bytes * r / (1 + r)\n slots = mamba_budget / state_bytes_per_slot\n # spec: each running req reserves (S+D) worth; non-spec just S\n return int(slots // S) # NOSPEC; with spec the budget joint-solves for (S+D)·per_req per req\n```\n\nThen state the result three ways: the **`r` value**, the **predicted clamp** (if `rest` given), and **which pool binds** (`min(mamba_clamp, KV_cap)`, where `KV_cap = kv_tokens · dcp_size / L`).\n\n## Procedure\n\n1. **Free lever first**: if the boot log shows large idle in `Memory pool end. avail mem` (e.g. 30–40 GB at mem-frac 0.85), raise `--mem-fraction-static` (→0.92) before touching the split — it grows `rest` for both pools at no cost. Validate graph-capture headroom once.\n2. Collect the inputs. Prefer a real boot log for the two byte constants.\n3. Compute `r*`. If `r*` would drive the state pool below one request's worth (`mamba_budget < S · per_req`, happens at very long `L`), **switch to pinning `--max-mamba-cache-size = target_concurrency · S`** and let the rest go to KV — a sub-0.15 `r` is fragile.\n4. Report `r`, predicted clamp, binding pool, and any dtype accuracy gate that applies.\n\n## Worked examples (TP8, B300, validated against measured clamps)\n\nConstants (from boot logs): `state_bytes_per_slot ≈ 56.4 MB` (fp32 ssm), `kv_bytes_per_token ≈ 13.8 KB` (fp8) → `token_equiv ≈ 4080`. Cache strategy `extra_buffer_lazy` → `S = 4`. Workload `L = 9216` (8192 in + 1024 out). NOSPEC → `D = 0`.\n\n- **TP (dcp_size=1)**: `r = 4 · 4080 · 1 / 9216 ≈ 1.8`. Measured: clamp 88, KV cap 87 → balanced. ✅\n- **DCP8 (dcp_size=8)**: `r = 4 · 4080 · 8 / 9216 ≈ 14`. Measured at r=14: clamp 125, KV cap 129 → balanced. ✅ (At the naive r=1.8 the DCP KV pool is ~8× over-provisioned — cap 683 vs clamp 86 — wasting budget that should go to the state pool.)\n\n`predict_clamp` on the same model at the **default** `extra_buffer`, overlap on and the decode-lock skip off (`S = 5`), three boot logs:\n\n| config | `max_mamba_cache_size` | `mmcs // S` | measured `max_running_requests` |\n|---|---|---|---|\n| TP8, no DCP | 257 | 51 | 51 |\n| DCP8, `r = 5.97` | 451 | 90 | 90 |\n| DCP8, `r = 9` | 474 | 94 | 94 |\n\nReference `r` for **this example model** (`token_equiv ≈ 4080`, i.e. fp32 ssm + fp8 KV; S=5 default; multiply by `(S+D)/S` for spec, by `dcp_size` for DCP) — recompute with your own `token_equiv` for a different model:\n\n| L | 2K | 4K | 8K | 32K | 64K | 128K |\n|---|---|---|---|---|---|---|\n| r | 10.0 | 5.0 | 2.5 | 0.62 | 0.31 | 0.16 |\n\n## Caveats\n\n- **DCP + spec**: a replicated (non-DCP-sharded) draft KV does **not** shard by `dcp_size`; at long `L` it dominates per-token cost, so the clean `×dcp_size` overstates DCP's advantage — fall back to a per-token direct-solve (KV term `/dcp` + an un-sharded draft-KV term) when spec is on. (NOSPEC → clean `×dcp_size` holds.)\n- **Spec + ReplaySSM → use `D=0`, not the draft-token count.** When ReplaySSM spec-verify is enabled, the `D` intermediate SSM states move off the per-request slot budget onto a **fixed ring** (a one-time deduction from `rest`, not a per-req term). So the mamba-slot cost per running request drops back to `S` (clamp = `mmcs / S`, not `/(S+D)`), and the balance ratio returns to the NOSPEC value (`r* ≈ S·token_equiv·dcp/L`). Measured example (TP8, D=8, L≈9K): applying `D=8` computes `r*≈2.6` but the true optimum is `r≈1.0` — `D=8` lands KV-bound at ~45% below the achievable peak concurrency. Plain spec (no replayssm) keeps `D` = the draft-token count.\n- **Under DCP the balance point sits beyond any realistic `L`** (the state pool binds first for almost everything), so the practical recommendation is to **pin `--max-mamba-cache-size = target_concurrency · S` directly** rather than dial a large `r`.\n- **`--max-mamba-cache-size` overrides `r`**; bytes beyond the `r` budget come out of the KV pool one-for-one.\n- **Precision changes are behind accuracy gates**: `--kv-cache-dtype fp8_e4m3` (doubles `token_equiv` → doubles `r`) and `--mamba-ssm-dtype bfloat16` (~halves `token_equiv`; also silently switches the linear-attention decode backend on SM100+ — pin `--linear-attn-decode-backend triton`) shift outputs; validate accuracy for the workload before production.\n- **Asymmetry**: the state pool is worst-case-reserved and fail-loud; KV degrades gracefully (retraction). When `L` is spiky/uncertain, bias `r` **up** rather than starve the state pool.\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/authoring-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/authoring-reference.md)\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/diffusion-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/diffusion-authoring.md)\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/engine-axis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/engine-axis.md)\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/mintlify-authoring.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/mintlify-authoring.md)\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-add-model/references/vendor-logo.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-add-model/references/vendor-logo.md)\n\n\n---\nname: cookbook-add-model\ndescription: Add a new model to the SGLang Cookbook (docs/, Mintlify), config-driven format — instantiate the model-agnostic template into a per-model config (+ benchmarks) JSX under src/snippets/configs/, an MDX page, the docs.json nav entry, NEW-tag hygiene, and the homepage vendor card. Interactive, multi-phase. Run with /cookbook-add-model.\ndisable-model-invocation: true\n---\n\n# Add a model to the SGLang Cookbook\n\n> Migrating an **existing legacy-template page** (one that imports a monolithic\n> `…/autoregressive/-deployment.jsx` generator)? Use the\n> `cookbook-migrate-model` skill instead — same target format, but the legacy\n> page (not the user) is the source of truth.\n\nThe cookbook is **config-driven**: two shared engines contain NO model-specific code —\n`docs/src/snippets/_deployment.jsx` (the 5-dim deploy matrix) and\n`_playground.jsx` (the diff-based override Playground). Adding a model = adding **data**:\na per-model `config` (+ optional `benchmarks`) consumed by both engines, plus an MDX page\nthat imports them. No engine edits.\n\n**Instantiate the model-agnostic template** (NOT a clone of any live cookbook — the\ntemplate is decoupled and covers all hardware + all axes):\n- `templates/config.jsx.tmpl` → `docs/src/snippets/configs//.jsx`\n- `templates/benchmarks.jsx.tmpl` → `…/-benchmarks.jsx` (skip if no numbers)\n- `templates/page.mdx.tmpl` → `docs/cookbook///.mdx`\n\nThe template uses explicit `__TOKEN__` placeholders; you fill them, prune what the model\nlacks, and replace the EXAMPLE cells with verified recipes. DeepSeek-V4 is a populated\n*instance* you can consult, but is not the template.\n\n**Deep references (read on demand, don't inline):**\n- [references/authoring-reference.md](references/authoring-reference.md) — field-by-field config / cells / playground / MDX contract.\n- [references/mintlify-authoring.md](references/mintlify-authoring.md) — MDX rules (forbidden syntax, JSX tables, labeled fences) + invocation-example patterns. Read before writing §1–§3 prose.\n- [references/engine-axis.md](references/engine-axis.md) — adding a new Playground feature axis (rare engine work).\n- [references/vendor-logo.md](references/vendor-logo.md) — new-vendor card logo: ask the user for the brand logo, then generate the icon-only 940×525 RGBA PNG (spec + Pillow recipe + `git add -f`).\n- [references/diffusion-authoring.md](references/diffusion-authoring.md) — required opening, tag, command-picker, and feature-overlay contract for diffusion model pages. Read before editing any `docs/cookbook/diffusion//.mdx` page.\n\n## Architecture at a glance\n\n```\n┌─────────────────────────────────────────────────────────────────┐\n│ cookbook///.mdx │\n│ import { Deployment } from \"/src/snippets/_deployment.jsx\"; │\n│ import { Playground } from \"/src/snippets/_playground.jsx\"; │\n│ import { config } from \"/src/snippets/configs/.../X.jsx\"; │\n│ │\n└─────────────────────────────────────────────────────────────────┘\n │ (config passed as React prop)\n ▼\n┌─────────────────────────────────────────────────────────────────┐\n│ src/snippets/configs//.jsx │\n│ export const config = { │\n│ supportedHardware, variants, quantizations, strategies, ... │\n│ cells: [ { match:{hw,variant,quant,strategy,nodes}, │\n│ env:[...], flags:[...] }, ... ], // 5-dim matrix │\n│ playgroundFeatures: { attention, moe, parsers, ... }, │\n│ }; │\n└─────────────────────────────────────────────────────────────────┘\n │ (consumed by BOTH engines — no model code in engines)\n ▼\n┌──────────────────────────────────┬──────────────────────────────┐\n│ _deployment.jsx │ _playground.jsx │\n│ Renders the verified matrix; │ Renders override chips + │\n│ one cell → its env/flags. │ diff against the cell. │\n└──────────────────────────────────┴──────────────────────────────┘\n```\n\nThe two widgets stay in sync via: the **URL hash** (deploy mirrors its selection;\nplayground reads it), the **`sglang-deploy-sel` custom event** (deploy dispatches on\nevery change; playground listens — `replaceState` doesn't fire `hashchange`), and the\nshared **`sglang-deploy-env` localStorage key** (HOST/PORT placeholders).\n\n> The main template is **autoregressive**. Diffusion pages use\n> `templates/diffusion-page.mdx.tmpl` plus\n> [references/diffusion-authoring.md](references/diffusion-authoring.md); do not force the\n> autoregressive deployment matrix on them. Omni pages follow their own category structure.\n> All categories still obey the Mintlify / NEW-tag / docs.json / category-card / validation\n> rules below.\n\n---\n\n**Interactive, multi-step workflow. Collect inputs incrementally — don't ask for\neverything upfront.** The real work is the verified `cells[]` recipes + measured\nbenchmarks (Phases 2 + 4); everything else is filling the template.\n\n## Phase 1 — Collect inputs\n\n1. **Model card** — HuggingFace repo/URL. **Fetch the page** and extract description,\n param count, architecture, context length, **license**. (Fetching guards factual bugs\n like an off-by-a-few-B param count.) If the model isn't public, ask the user.\n2. **Variants / quantizations** — keep separate: variants are size/mode (e.g. Flash/Pro,\n Instruct/Thinking); quantizations come from the HF card / linked repos (BF16/FP8/FP4/…).\n Default to BF16 when a full-precision repo exists.\n3. **Tested hardware + parallelism** — which platforms are actually tested, and TP/EP/DP\n for each. List only tested hw (unlisted greys out).\n4. **Verified launch recipes** — the full `sglang serve` flags per\n (hw × variant × quant × strategy × nodes) combo → these become `cells[]`. Rewrite any\n `python -m sglang.launch_server` to `sglang serve` form.\n5. **sglang version / image tag** — ask which sglang build the recipes + benchmarks ran on\n (a release like `0.5.x`, or `main`/nightly). **Never guess or hallucinate it.** This one\n tag fills `dockerImages` and the benchmarks' `sglang_version`; when the user is unsure,\n default the image to `lmsysorg/sglang:dev` (nightly) rather than inventing a release.\n6. **Pre-flight**: `gh pr list --repo sgl-project/sglang --search \"\"` (dup check).\n\n**Hardware reference** (the shared `HARDWARE_CATALOG` in `_deployment.jsx`). A GPU **not** in\nthis table (RTX PRO 6000, GH200, future chips) goes in the model's own `config.hardware`\n(`{id,label,vram,vendor}`) — the engine merges it in; don't edit the engine catalog:\n\n| Platform | Vendor | VRAM | Docker image |\n|---|---|---|---|\n| H100 | NVIDIA | 80GB | `lmsysorg/sglang:` |\n| H200 | NVIDIA | 141GB | `lmsysorg/sglang:` |\n| B200 | NVIDIA | 192GB | `lmsysorg/sglang:` |\n| B300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130` when required) |\n| GB200 | NVIDIA | 192GB | `lmsysorg/sglang:` (or `-cu130`) |\n| GB300 | NVIDIA | 288GB | `lmsysorg/sglang:` (or `-cu130`) |\n| DGX Spark | NVIDIA | 128GB (unified) | `lmsysorg/sglang:` — needs a **CUDA 13** build |\n| MI300X | AMD | 192GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI325X | AMD | 256GB | `lmsysorg/sglang:-rocm720-mi30x` |\n| MI350X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n| MI355X | AMD | 288GB | `lmsysorg/sglang:-rocm720-mi35x` |\n\n- **Image tag (``)**: don't guess — ask the user for the tag the recipes ran on, or\n default to `dev` (nightly). The same tag goes in `dockerImages` and benchmarks'\n `sglang_version`; the engine falls back to `lmsysorg/sglang:dev` for any unmapped hw.\n- **TP sizing** (sanity-check recipes): `weight_GB / gpu_mem`, round up to a power of 2,\n ~20–30% headroom. BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5. MoE → **total** weight, not\n active params. FP4 is Blackwell-only (B200/B300/GB200/GB300/DGX Spark). GB200/GB300\n single-node hosts are typically **4 GPUs** (TP=4 ceiling); a DGX Spark is **1 GPU**, so\n its only multi-GPU topology is TP=2 across 2 nodes.\n- **Platform flags**: Blackwell may need `--attention-backend trtllm_mha`; AMD typically\n needs `--attention-backend triton` + env `SGLANG_USE_AITER=1` /\n `SGLANG_ROCM_FUSED_DECODE_MLA=0` (check AITER TP constraints, e.g. `heads_per_gpu % 16 == 0`).\n- **EP** (MoE): 8-GPU NVIDIA `--tp 8 --ep 8`; AMD `EP = TP`; small NVIDIA (TP≤4) omit\n `--ep` unless benchmarked. (The template's AMD example cell shows these.)\n\n## Phase 2 — Instantiate the template\n\nFor a diffusion model, instantiate `templates/diffusion-page.mdx.tmpl` and keep the\nshared `DiffusionModelTags` component, plus `templates/diffusion-config.jsx.tmpl` for the\nopt-in scoped command builder. Put the compact install command and builder in §1 Quick\nstart. The first two paragraphs in §2 Model capabilities are not generic filler: they must\nstate the model's capability range, strongest differentiator, when to choose it, and at\nleast one real deployment or capability boundary. Put orthogonal runtime features in\n`scope: \"serve\"` or `scope: \"request\"`, not in the base recipe; use the schema from\n`references/diffusion-authoring.md`.\n\n1. **Copy** the three template files to their target paths (above). Note the two\n vendor-folder conventions: under `configs/` the folder is the **HuggingFace org**\n (`deepseek-ai`); under `cookbook/` it's the **display vendor** (`DeepSeek`).\n2. **Replace every `__TOKEN__`**: `__MODEL_DISPLAY__`, `__MODEL_SLUG__`, `__HF_ORG__`,\n `__HF_REPO__`, `__REASONING_PARSER__`, `__TOOLCALL_PARSER__`, `__ONE_LINER__`. Verify\n none remain: `grep -rn '__[A-Z_]*__' `.\n3. **Prune** to what the model supports (delete, don't stub) — using\n [references/authoring-reference.md](references/authoring-reference.md):\n - `supportedHardware` + the EXAMPLE cells: keep your tested families; **delete the\n `mi*` ids + AMD example cell if no AMD recipe**, etc. A GPU not in the shared catalog\n (e.g. RTX PRO 6000) → declare it in `config.hardware` and add its id here.\n - `playgroundFeatures` axes: remove the `megamoe` backend option + the\n `megamoeQuant` block from the `moe` axis (non-Blackwell-MoE), delete `hisparse`\n (non-DSA), `pdDisagg`/`router` (no PD), the `parsers` axis (no parsers), etc.\n - `quantizations` / `variants`: drop what the model doesn't ship; collapse `variants`\n to single `default` if there's no variant axis (then drop the `variant` half of\n `modelNames`/`defaultAccuracy` keys).\n4. **Fill `cells[]`** with the verified recipes from Phase 1 (replace every EXAMPLE cell;\n set `verified: true` only on tested combos), and `modelNames` with real HF slugs,\n `dockerImages` for your hw (use the Phase-1 tag, or default `lmsysorg/sglang:dev` — never\n a guessed release; key by `hw`, or a more specific key — `hw|variant|quant`, `variant|quant`,\n `hw|quant|strategy`, `hw|quant`, resolved in that order before `hw` — when one\n variant/quant/strategy needs its own image), `multiNodeHints` only for fabric-specific hw\n (e.g. gb200).\n\n5. **Diffusion pages: add the ComfyUI section.** Every diffusion cookbook page ends with\n a `## . Run in ComfyUI` section so a reader never has to guess whether the model is\n reachable from ComfyUI. It is one component; the per-model facts live in the component,\n not the page:\n\n ```mdx\n ## . Run in ComfyUI\n\n import { ComfyUISupport } from '/src/snippets/diffusion/comfyui-support.jsx';\n\n \" />\n ```\n\n Pick `model` from the table in\n `docs/src/snippets/diffusion/comfyui-support.jsx`. Use the model's own key when it has\n an entry (its executor or dedicated node differs); otherwise use the generic `image` or\n `video`. A model gets its own key only when the plugin actually treats it specially —\n an entry in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a\n dedicated node. Adding a key without the matching plugin support makes the page lie.\n\n### Site-wiring (do all three)\n\n- **`docs/docs.json`** — add the page under Cookbook → `` → ``, at\n the **top** of that vendor's `pages` (root-relative, no `.mdx`:\n `cookbook///`). New vendor group → insert in the section's\n local ordering.\n- **NEW-tag hygiene** — the new page keeps `tag: NEW` (from the template). Scan the\n vendor dir for existing NEW and strip it from siblings; verify ≤1:\n `grep -rn 'tag: NEW' docs/cookbook///` → at most one result. (Scan\n files; don't assume the first `docs.json` entry holds NEW.)\n- **Homepage card** — `docs/cookbook//intro.mdx`: if the org already has a\n ``, update only its `href` (keep `img`). If the org is **new**, add a ``\n (title = nav-group name; keep card order aligned with `docs.json`) **and create its logo**:\n ask the user for the brand logo, then generate the conforming **icon-only 940×525 RGBA\n transparent** PNG → `docs/cards/logos/.png` per\n [references/vendor-logo.md](references/vendor-logo.md) (track with `git add -f` — `*.png`\n is gitignored repo-wide). Never invent or copy a logo.\n\n## Phase 3 — Validate\n\n```bash\ncd docs\nmint validate # frontmatter, missing nav entries, MDX/JSX errors\nmint broken-links\nmint dev # visual smoke test at http://localhost:3000/cookbook///\n```\n\nSpot-check: cells render sensible commands; URL-hash nav persists across reload; the\nPlayground inherits the Deploy selection live; each axis toggle produces the expected\ndiff; Docker mode wraps in `docker run` with the right image; multi-node cells emit the\nhints + `--nnodes N`; cURL resolves the model name; the NEW badge shows on the new page\nand is gone from same-vendor siblings; the homepage card points to the new model.\n\n## Phase 4 — Interactive testing\n\nThe user deploys each cell, runs the benches, and pastes results; you fill the data:\n- mark each tested `cells[]` entry `verified: true` (absent = yellow/unverified badge);\n- fill the `-benchmarks.jsx` entries (one per cell `match`) with measured\n speed/accuracy + the `sglang_version` the user reports (don't invent one — the template's\n `0.0.0` is a deliberate TODO); set model-level `defaultAccuracy` per variant. Leave a\n cell's entry as a bare `match` stub if it has no numbers yet (the card shows \"pending\").\n\n## Phase 5 — Prose & config tips\n\n**Read [references/mintlify-authoring.md](references/mintlify-authoring.md) first** (it\ncarries the parser-output-shape / thinking-mode / Output-Example / no-hardcoded-sampling\nrules + the Mintlify forbidden-syntax list). Then rewrite the MDX prose from the HF card +\nuser notes: §1 Model Introduction (description, links, params, license, variants table),\n§2 Configuration Tips (hw-specific tuning, caveats), §3 Advanced Usage (Reasoning /\nTool-Calling / HiCache — keep only what applies; match the reasoning example to the\nparser's output shape; each runnable block gets an `**Output Example:**`).\n\nFor diffusion pages, follow the category-specific Quick start and capability contract in\n`references/diffusion-authoring.md`. Run `node docs/scripts/check_cookbook_configs.mjs` to\nverify the tag widget and introduction structure before rendering the page.\n\n## Phase 6 — Review\n\n```\n/cookbook-review-pr \n```\n\n## Git workflow\n\nAlways branch — never commit to main directly.\n\n```bash\ngit checkout -b add--cookbook\ngit add docs/src/snippets/configs//.jsx \\\n docs/src/snippets/configs//-benchmarks.jsx \\\n docs/cookbook///.mdx \\\n docs/docs.json docs/cookbook//intro.mdx\ngit commit -m \"Add cookbook\"\ngit push -u origin add--cookbook\ngh pr create --title \"[Doc] cookbook: Add cookbook\" --body \"...\"\n```\n\n\n# References Skill\n\nPath: `.agents/skills/cookbook-migrate-model/references/dimension-mapping.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/cookbook-migrate-model/references/dimension-mapping.md)\n\n\n---\nname: cookbook-migrate-model\ndescription: Migrate a legacy-template SGLang cookbook page (monolithic per-model generator under docs/src/snippets/autoregressive/) onto the config-driven template (shared _deployment.jsx / _playground.jsx engines + per-model config). Use when asked to migrate, convert, or port an existing cookbook page — NOT for brand-new models (use cookbook-add-model for those). Run with /cookbook-migrate-model .\n---\n\n# Cookbook Migrate Model\n\nConvert one legacy cookbook page to the config-driven format, faithfully. The\nlegacy page — its generator widget and its measured benchmark blocks — is the\n**single source of truth**. You are transcribing it into the new data model,\nnot improving it.\n\nReuses the `cookbook-add-model` skill's assets (read them on demand):\n- `../cookbook-add-model/templates/config.jsx.tmpl`, `page.mdx.tmpl`, `benchmarks.jsx.tmpl`\n- `../cookbook-add-model/references/authoring-reference.md` (config/cells/playground contract)\n- `../cookbook-add-model/references/mintlify-authoring.md` (MDX rules)\n\nMigration-specific references in this skill:\n- [references/dimension-mapping.md](references/dimension-mapping.md) — legacy-control → new-dimension mapping rules, command rewrite table, per-family strategy sets, and the Qwen3.5 pilot as a worked example (PR #27848).\n\nThe round's per-model inventory (scope, batch order, quirks, measured-data\nsurvey) is tracked by the migration maintainer outside the repo — expect it in\nyour dispatch prompt, or ask for it.\n\n## Hard rules (non-negotiable)\n\n1. **Never modernize.** Env vars, flags, TP values, docker tags, version strings\n are copied verbatim from the legacy page — even when today's defaults differ\n (e.g. a flag whose behavior has since become the default; keep it anyway). The recipe\n that was verified is the recipe as written. Allowed normalizations are ONLY\n the five alias rewrites in dimension-mapping.md §2 (`launch_server`→`sglang\n serve`, `--model`→`--model-path`, `--tp-size`→`--tp`, abbreviated\n `--speculative-algo`→`--speculative-algorithm`,\n `--expert-parallel-size`→`--ep`). **Accuracy-degrading flags**\n (`--kv-cache-dtype fp8_e4m3`, W4A4-style runtime quant) follow a\n deterministic rule — enforced in migration, no asking: offered as a\n legacy **selectable option** → never select it (cells mirror the\n accuracy-safe side), and the option **survives as a Playground axis** —\n an existing one where it fits, else add one via a separate prior engine\n PR (rule 4); a legacy choice never degrades to a tips mention. Baked\n into the recipe's **default/unconditional command** → keep it verbatim\n (the recipe was measured with it, and fp8 KV halves KV memory —\n stripping could OOM it). See dimension-mapping.md §2 caveats.\n2. **Never invent versions or numbers.** Benchmark numbers only from the\n legacy page's measured blocks, and a result migrates ONLY when its\n `sglang_version` is a **reproducible anchor** — the bar is reproducibility,\n not \"must be a release\":\n - ✅ release tag/version (`v0.5.9` / `0.5.9`), commit hash, OR — for\n **Day-0 support** (the enabling PR isn't merged and no release is cut\n yet) — a specific **PR (`PR #27944`) or commit** you can `gh pr checkout`\n / `git checkout `. Commit is most precise; a PR pin is fine for day-0.\n - ❌ a moving ref — `\"main branch\"`, `\"main (2026-06-11)\"`, open-ended\n `\"0.5.8+\"` — is NOT reproducible: **drop the WHOLE result (speed AND\n accuracy)**, not just speed. Keep `benchmarkCommands` so ⚡Reproduce still\n guides re-measurement against a pinned build.\n **Never inherit cross-model numbers** — measurements the legacy page\n attributes to a *different model* (e.g. a K2.6 page carrying K2.5-measured\n speed) are dropped regardless of version. When kept, `sglang_version` is the\n legacy page's string verbatim. Docker tags only the ones the legacy page\n pinned (unmapped hw falls back to `:dev`).\n3. **Verified policy (strictest tier).** `verified: true` ONLY when (a) the\n legacy page has concrete measured data for that exact 5-dim combo AND\n (b) the cell's flags equal the deployment command used for that measurement\n — modulo `{{HOST_IP}}`/`{{PORT}}`, the five alias rewrites, and **parser\n flags**: `--reasoning-parser`/`--tool-call-parser` are stripped from every\n cell (Playground-only feature; when the measured run had them on, say so in\n the benchmarks file header). When the measured command diverges from the\n generator default, **the verified cell follows the measured command**; the\n generator default stays as the sibling strategy/cell or a tips note. Everything else is unverified (yellow) —\n including combos that look memory-infeasible; keep them verbatim and list\n them in the PR body for the re-verification track.\n4. **Engines are read-only.** `_deployment.jsx` / `_playground.jsx` must not\n change in a migration PR. Model-specific features are config DATA consumed\n by generic axis handlers (MegaMoE precedent), so they need NO engine change.\n A **titled single-select that strips a flag family** — KV Cache DType\n (`--kv-cache-dtype`), mamba (`--mamba-radix-cache-strategy`), … — is already\n covered by the merged generic **`flagSelects`** axis: declare it in the\n config (a list of `{ id, title, stripPrefixes, options }`; see the Qwen3.5\n mamba example), **no engine PR**. Only a genuinely new control *shape* that\n `flagSelects` can't express would need a one-time generic primitive (never a\n model-named handler) on a separate prior PR (engine-axis.md).\n5. **`github.cookbookModel` must be set** (`/`, e.g.\n `qwen/qwen3.5`) and the block never pruned — without it Submit ↗ mislabels\n as deepseek-v4. The issue template itself needs NO edits (free-form input).\n6. **Strategy tiers are signal-driven.** A cell goes under `low-latency` /\n `high-throughput` ONLY on a signal present in the legacy source (an\n explicit performance toggle, a named recipe, or prose stating the\n operating point); **no signal → `balanced`**. Never derive a slant from\n your own hardware intuition — re-tiering on measured evidence is the\n hardware owner's follow-up PR, not part of a migration\n (dimension-mapping.md §4).\n\n## Workflow (one model = one PR)\n\n### 1. Inventory the legacy assets\n- Read the legacy generator (`docs/src/snippets/autoregressive/-deployment.jsx`)\n end-to-end: every option dimension (radio vs checkbox vs dynamic), every\n gate/SUPPORT matrix, the full emitted command per reachable combo (env\n prefixes, `# Error` pseudo-commands included).\n- Read the legacy MDX: §2 install docker tags → `dockerImages` (pinned, not\n upgraded); §3.2 tips → new §2; §4 invocation examples → new §3 (keep real\n Output Examples verbatim); §5 benchmark blocks → transcribe each measured\n block: deploy command used, bench command (dataset/isl/osl/num-prompts/\n concurrency), P50 (median) TTFT/TPOT, output tok/s, hardware, version string.\n- Inbound-anchor sweep: `grep -rn \"\" docs/ --include='*.mdx'` —\n find links/`#fragments` into this page (`mint broken-links` does NOT check\n fragments). Fix referrers or add `` shims in the same PR.\n- Check the maintainer-provided inventory notes for this model's known quirks —\n but treat them (and the §4 family table) as a **survey snapshot**: re-verify\n every dimension against the live legacy files before mapping. Pages keep\n receiving updates (precedent: Kimi-K2.6's live generator has a speculative\n toggle the 2026-06-10 survey notes lack).\n\n### 2. Design the 5-dim mapping\nApply [references/dimension-mapping.md](references/dimension-mapping.md). Key\ndecision — which legacy toggle becomes the `strategies` dimension: a toggle\nthat **changes other parts of the command** (TP, mem) must (the Playground\ncan't do coupled changes), and so does a toggle the legacy page itself labels\nwith operating-point words — e.g. a `dpattention` radio whose options are\nsubtitled \"Low Latency\" / \"High Throughput\" (GLM-5.1 / Kimi-K2.6 pattern) —\neven when its flags are uncoupled (`--attn-dp-size N` is a pure\nflag add). Any other toggle that only adds/removes its own flags becomes a\nPlayground axis with the flags baked into cells when the legacy default was\nON — EXCEPT parsers:\n`--reasoning-parser`/`--tool-call-parser` are NEVER baked into cells, they are\nPlayground-only (DSv4 convention). **Every legacy control survives as an\ninteractive control** — a dimension or a Playground axis, never a tips-only\nmention — and a model-specific control is **config data, not engine code**\n(MegaMoE W4A4 is all DSv4 config on the existing `moe` axis). It's pure\nconfig whenever it fits an existing axis's data schema. A **titled\nsingle-select that strips a flag family** (Nemotron3's \"KV Cache DType\",\nmamba `--mamba-radix-cache-strategy`, …) fits the merged generic **`flagSelects`**\naxis — so it too is config-only (declare a `flagSelects` list). Only a control\nwhose *shape* `flagSelects` still can't express would need a ONE-TIME generic\nprimitive (never a model-named handler) on a separate PRIOR engine PR, keeping\nthe migration PR data-only (hard rule 4, engine-axis.md). The strategy count follows the page's\noperating points: **one recipe → a single `balanced`; two → `low-latency` +\n`high-throughput`; three → the full trio (the ideal)**. The tiers apply per\n(hw × variant × quant) combination — a single-recipe combination on a\nmulti-strategy page parks under its semantically honest tier (no\nlatency/throughput slant → `balanced`; the page's list is the union). When the\nlegacy toggle is MTP / speculative decoding, the direction is a deterministic\ndefault — apply without asking: **MTP on → `low-latency`, MTP off →\n`high-throughput`** (reversed only with maintainer confirmation). Tier\nplacement is signal-driven (hard rule 6). Never invent a recipe just to fill\nstrategy chips (see dimension-mapping.md §4). Record the outcome as a\n**strategy mapping table** for the PR body — one row per group of\ncombinations sharing the same legacy signal (e.g. \"all GPU combos: MTP\ntoggle → low-latency / high-throughput\"; \"xeon: (none) → balanced\"), with a\none-line rationale each; don't enumerate 60 identical rows. The table is what\nhardware owners sign off on at review.\n\n### 3. Generate the config (codegen, then audit)\n- For >~30 cells, port the legacy `generateCommand()` into a throwaway Node\n script that enumerates combos and emits the `cells:[...]` literal (output\n must stay a pure literal — Mintlify forbids runtime spreads/calls). Apply the\n verified-cell override in the script. See the pilot scripts embedded in the\n worked example of dimension-mapping.md §5.\n- **Independent equivalence audit (required):** extract the ORIGINAL generator\n from git (`git show main:` — NOT `HEAD:`, the migration branch deletes\n the file, see dimension-mapping.md §5 item 7), stub React hooks, run it for\n every combo, and diff token-by-token against the new cells. Expected deltas\n only: the appended `--host {{HOST_IP}}`/`--port {{PORT}}`, the\n engine-injected multi-node trio, the §2 alias rewrites (the entrypoint\n rewrite doesn't appear in cell tokens — cells hold flags only; the audit\n script normalizes it on the legacy side), and the intentional verified-cell\n override. Paste the PASS count + the audit script in the PR body\n (collapsed `
`).\n- Hand-author the non-cells fields per authoring-reference.md. Structural\n self-checks: every cell resolves a `modelNames` key; no `--nnodes/--node-rank/\n --dist-init-addr/--host/--port` literals; every `{{KEY}}` declared; every\n `supportedHardware` id has ≥1 cell.\n\n### 4. Benchmarks file\nOne entry per measured block only (cells without entries already render\n\"pending\" — bare `{match}` stubs are unnecessary). `tokens_per_sec_per_gpu` =\ntotal (in+out) tok/s/GPU = `output tok/s ÷ (tp × nnodes) ×\n(isl+osl)/osl` — stored directly (the card shows it as-is). TTFT/TPOT\ntake the P50 (median) rows; set `config.latencyPercentile` (default `\"P50\"`; use\n`\"Mean\"` only for legacy Mean-recorded data — temporary, being migrated to P50; an\nentry-level `latencyPercentile` overrides the page value per cell).\nPut the workload's\n`num_prompts` into `workload`. **`config.accuracyLabels` is required whenever\nthe benchmarks carry accuracy data** — the engine ships no default eval set\n(#27842), so missing labels means the accuracy rows silently don't render;\nextra context (sample counts, suites that don't fit) goes in the entry's\n`notes`. Zero-measured-data pages: skip the file and the `benchmarks` prop\nentirely, but keep `benchmarkCommands` so ⚡Reproduce still guides users.\n\n### 5. Rewrite the MDX\nFrom `page.mdx.tmpl`: keep the original `title` (nav identity), write a fresh\nSEO `description` (top-level — delete any legacy `metatags.description`), **no\n`tag: NEW`** (a migration is not a launch), **no `mode:`**. Install accordion\ncarries the legacy install content + pinned images. Keep the template's\nDSv4-style strategy bullets — serving semantics first (single-user chat /\ntypical multi-user / batch throughput), trimmed to the strategies the page\nships, plus at most a one-line note on what each strategy changes on this\nmodel; do NOT rewrite them as toggle-/migration-centric explanations.\nLegacy §5 benchmark prose is deleted (numbers → benchmark card, commands →\n⚡Reproduce); legacy prose deploy commands are deleted (doc↔config parity —\nfold their unique flags into §2 tips). Invocation examples + real outputs carry\nover verbatim, but **wrapped in Accordions** — §3 commands and outputs are\ncollapsible (required, DeepSeek-V4 pattern): code in an\n``, output in a following\n``; legacy pages kept them inline.\n\n### 6. Delete the legacy generator\nRemove `docs/src/snippets/autoregressive/-deployment.jsx` and its\nimport. `grep -rn \"-deployment\" docs/` must return nothing (config\nprovenance comments must not name the deleted path). Site wiring needs **no\nchanges**: docs.json path/title unchanged, vendor card + logo already exist.\n\n### 7. Validate\n- `grep -rn '__[A-Z_]*__'` on the new files (no template tokens).\n- `cd docs && mint validate && mint broken-links` (pre-existing breaks on\n main are not yours — say so in the PR).\n- `mint dev` browser smoke: initial selection = the verified cell (first in\n `cells[]`) with green badge; multi-node cells show the injected trio +\n header; AMD cells show env prefixes; Docker mode wraps with the pinned image\n and passes cell env as `--env`; condition-hidden combos grey out; benchmark\n card values; NO parser flags in any Deploy command; Playground parser\n toggles ADD the parser flags (green additions) while spec toggles strike\n the baked spec flags (red); Submit ↗ prefills this model. **Probe pitfall:** drive at most\n ONE programmatic click per evaluation and wait for React to settle —\n multiple clicks in one synchronous script batch and read stale DOM.\n- Token-level audit from step 3 passes.\n\n### 8. PR + review\nOne PR per model. PR body: migration framing, verified policy applied, the\nstrategy mapping table (step 2), the audit PASS count + script, any\ninherited-infeasible combos flagged for re-verification. Then run `/cookbook-review-pr ` and fix findings.\nFYI: docs previews only build for in-repo (`sgl-project/sglang`) branches —\na fork-headed PR is perfectly fine but renders no preview; a maintainer can\nre-push the branch in-repo if a preview is wanted for review.\n\n### 9. Keep this skill current\nAny new convention, engine behavior, or pitfall you discover while migrating\n(naming decisions, audit-script gotchas, review-rule conflicts, …) MUST be fed\nback into this skill — same PR if it's skill-file-only, or an immediate\nfollow-up commit on the skill's branch/PR. The next agent runs on what's\nwritten here, not on your session's context.\n\n\n---\nname: cookbook-review-pr\ndescription: Review a pull request against the SGLang Cookbook (docs/, Mintlify) contribution checklist — the config-driven format (per-model config + benchmarks JSX consumed by the shared _deployment.jsx / _playground.jsx engines). Run with /cookbook-review-pr .\n---\n\n# Cookbook Review PR\n\nFetch the diff, run the checklist, report what you find. The cookbook is **config-driven**:\nshared engines (`_deployment.jsx`, `_playground.jsx`) with NO model-specific code; each\nmodel is a data `config` (+ optional `benchmarks`) under `src/snippets/configs//`\nplus an MDX page. This checklist targets that layout. Field-schema detail lives in\n`.claude/skills/cookbook-add-model/references/authoring-reference.md` — defer to it rather\nthan restating.\n\n## Usage\n\n```\n/cookbook-review-pr \n```\n\n## Steps\n\n1. `gh pr view --repo sgl-project/sglang --json title,body,files,author,baseRefName,headRefName,commits,reviews`\n2. `gh pr diff --repo sgl-project/sglang`\n3. `gh pr list --repo sgl-project/sglang --state open --search \"\"` (duplicate check)\n4. Run every checklist item against the diff.\n5. Output per-file verdicts + overall recommendation.\n\n## Checklist\n\n### 1. File hygiene\n- A cookbook PR should only touch: `docs/src/snippets/configs//*.jsx`\n (config + benchmarks), `docs/cookbook/**/*.mdx`, `docs/docs.json`,\n `docs/cookbook//intro.mdx` (vendor card), `docs/cards/logos/.png`\n (new vendor only). Flag stray files (`settings.local.json`, lockfiles, IDE configs).\n- Pages must be `.mdx`, not `.md`. Files end with a trailing newline. Check commit history\n for unrelated commits accidentally included.\n- **Engines untouched**: `_deployment.jsx` / `_playground.jsx` should NOT change in a\n model-add PR (adding a model is data-only). Engine edits = a separate axis/feature PR\n (see `cookbook-add-model/references/engine-axis.md`); review them against that checklist.\n\n### 2. Config quality (the per-model `config`)\n- Single `export const config = { ... }` literal — **no** function calls, spreads,\n fragment refs, or IIFE (Mintlify re-evals at hydration → `ReferenceError`).\n- No `!(x in y)` anywhere (Mintlify AST walker crashes) — use `obj.key === undefined`.\n- `supportedHardware` ⊆ `HARDWARE_CATALOG` (in `_deployment.jsx`) ∪ `config.hardware`. A\n model-specific GPU the shared catalog lacks must be declared in `config.hardware`\n (`{id,label,vram,vendor}`), **not** added to the engine catalog — unless the platform is\n genuinely shared, which is an engine change the PR body has to argue (and the\n `cookbook-add-model` hardware table moves with it). A new catalog entry is inert for\n other pages: each filters the catalog by its own `supportedHardware`.\n- `placeholders` declares every `{{KEY}}` used in `curl` or any cell.\n- `modelNames` covers every cell (by `hw|variant|quant` triple or `variant|quant` pair).\n- `strategies` count matches the page's operating points — 1 recipe → a single `balanced`;\n 2 → `low-latency` + `high-throughput`; 3 → the full trio. Tiers apply per\n (hw × variant × quant) combination: a single-recipe combination must park under its\n semantically honest tier (clear slant → that tier, e.g. a workstation card under\n `low-latency`; no slant → `balanced`, e.g. a CPU platform) — **flag a no-slant recipe\n parked under low-latency/high-throughput**. Mixed unions\n like [low-latency, balanced, high-throughput] with per-selection greying are fine. Also\n flag model-specific ids (e.g. `mtp`), and flag an INVERTED speculative mapping — the\n deterministic default is MTP/spec-decoding ON → `low-latency`, OFF → `high-throughput`\n (at saturation the draft+verify overhead outweighs the speedup); the reverse needs an\n explicit maintainer-confirmed justification in the PR. The MDX strategy bullets describe serving semantics\n in the DSv4 style (single-user chat / typical multi-user / batch jobs), not internal\n toggles.\n- `dockerImages` covers the hw ids that have cells (else users hit the `:dev` fallback); a\n `hw|quant` key (resolved before the plain `hw`) is valid when one quant on a shared GPU needs\n a different image (e.g. an FP4 dev build) — don't flag those.\n- `multiNodeHints` present ONLY for hw whose fabric needs manual NIC env (e.g. `gb200`\n NVL72) — NOT every `multi-N` hw (standard-IB DeepEP / Marlin multi-node don't need it).\n Hints render above BOTH run modes, so flag one that reads as docker-only (\"add these\n docker run flags\") — `docker run` flags belong in the hardware entry's\n `multiNodeDockerFlags`, which the engine emits into the Docker command itself.\n- `github.cookbookModel` is set to the model's HF id (`/`). The issue\n template's `model` field is a free-form input prefilled from this value; if the config\n omits the `github` block, the engine falls back to `deepseek-ai/deepseek-v4` and the\n page's submissions get mislabeled.\n- `playgroundFeatures` is opt-OUT: the **general axes ship on every cookbook by default**\n (`attention` TP/CP/DP-Attn, `moe` backend+EP for MoE models, `parsers`, `speculative`,\n `pdDisagg`, `hicache`) — flag a missing general axis unless the model genuinely cannot\n use it. Model-specific axes only where applicable (MegaMoE backend + `megamoeQuant`\n only on Blackwell MoE, gated by `requiresHw`; `hisparse` only DSA-style). Knobs that are\n meaningless for a subset of variants/hw are `disable`d with a reason, not silently live\n (e.g. MoE knobs greyed on dense variants). No empty/stub axes.\n- **No leftover `__TOKEN__`** — the config was stamped from the template and every\n placeholder is filled (`grep -rn '__[A-Z_]*__'` on the new config/benchmarks/MDX returns\n nothing).\n- **All-hardware considered**: every `supportedHardware` id (from the catalog or `config.hardware`) has ≥1 cell OR is a deliberate\n greyed \"coming soon\"; AMD was pruned or kept on purpose (not a leftover template family).\n\n### 3. Cells / 5-dim matrix\n- Every cell `match` has EXACTLY the 5 keys (`hw`, `variant`, `quant`, `strategy`, `nodes`).\n- `env` / `flags` are flat literals (only `{{PLACEHOLDER}}` subst) — no shared\n `commonFlags` reference (Mintlify won't inline it).\n- NO `--nnodes` / `--node-rank` / `--dist-init-addr` literals in multi-node cells\n (the renderer injects them from `match.nodes`).\n- NO literal `--host` / `--port` — use `{{HOST_IP}}` / `{{PORT}}`.\n- NO `--reasoning-parser` / `--tool-call-parser` in any cell — parsers are a\n Playground-only feature added on top of the base command (DSv4 convention);\n flag any cell that bakes them in.\n- Accuracy-degrading flags in cells — runtime quant below the checkpoint\n (e.g. MegaMoE W4A4 — DSv4 gates it behind the Playground's `megamoeQuant`)\n and lossy `--kv-cache-dtype` (e.g. `fp8_e4m3` over a higher-precision-KV\n checkpoint): **flag for explicit maintainer confirmation**. Output quality\n should be exactly what the quant chip declares, so absent a recorded\n sign-off in the PR (e.g. carried verbatim from a measured legacy recipe's\n default command), request the flag move to Playground/tips.\n- Flag order: `--model-path` first (an optional `--trust-remote-code` may precede it —\n the DSv4 cells do), then parallelism, then MoE, then tuning, `--host`/`--port`\n last (the playground's insert anchors assume this).\n- TP/memory sanity: `model_weight_GB / (tp × gpu_mem)` fits with ~20–30% headroom\n (BF16 ≈ params×2 GB, FP8 ≈ ×1, FP4 ≈ ×0.5; MoE uses **total** weight, not active params).\n\n### 4. Benchmarks\n- Each `benchmarks[]` entry's `match` tuple corresponds to a real cell.\n- `accuracyLabels` is present whenever the benchmarks carry accuracy data — the engine\n ships NO default eval set; without it the accuracy rows silently don't render.\n `defaultAccuracy` / per-cell `accuracy` / `benchmarkCommands.accuracy` keys all\n ∈ `config.accuracyLabels`.\n- A benchmark's quantization must match a variant actually listed — `(BF16)` on a model\n that only released FP8/FP4 is a factual bug.\n- `benchmarkCommands.speed` is `python3 -m sglang.bench_serving` (the workload), separate\n from the `sglang serve` deploy command, and should carry `--flush-cache`: bench_serving's\n `random` prompts are deterministic, so a warm rerun hits the radix cache and inflates\n throughput — speed numbers are measured cache-cold.\n- `sglang_version` is a real build the author ran (a release, or `dev`/nightly) — not a\n guessed/placeholder value (no leftover `0.0.0`).\n- **Latency percentile**: `config.latencyPercentile` (default `\"P50\"`, or `\"Mean\"`) matches the\n percentile the TTFT/TPOT values actually are — the card renders `TTFT ()`. A benchmarks\n entry may carry its own `latencyPercentile` to override the page value per cell\n (entry → config → `\"P50\"`): on a P50 page, kept legacy Mean cells must set it — a\n `sglang_version` tag alone doesn't convey the percentile. (`\"Mean\"` is temporary — legacy\n data is being re-measured to P50.)\n- **Throughput convention**: `tokens_per_sec_per_gpu` is stored as **total (in+out)/GPU**\n = `output tok/s/GPU × (isl+osl)/osl`, shown by the card as-is. Flag output-only values.\n- **Consistent accuracy harness across entries**: every value under one `accuracyLabels`\n column must be produced by the SAME harness — flag a page that, say, measures one\n platform's GSM8K with a historical raw completion harness and another's with\n `sgl-eval run gsm8k --num-examples 1319` and shows both as one \"GSM8K %\"\n (the scores aren't comparable). Either standardize on one harness (matching\n `benchmarkCommands.accuracy`) or require an explicit per-entry note. Common when folding\n a second contributor's measurements (e.g. an AMD/ROCm PR) into the page.\n\n### 5. Doc ↔ config parity (the #1 finding)\n- Any `sglang serve` command shown in MDX prose (config tips, benchmark section) must\n equal what the engine emits from the corresponding cell — same flags, same order. Drift\n here is the most common review miss.\n\n### 5b. ComfyUI section (diffusion pages)\n- A diffusion page ends with `## . Run in ComfyUI` rendering ``. A\n reader must not have to guess whether the model is reachable from ComfyUI.\n- The `model` prop is a key that exists in `docs/src/snippets/diffusion/comfyui-support.jsx`.\n A model-specific key is only correct when the plugin really treats it specially — an entry\n in `executor_class_dict`\n (`python/sglang/multimodal_gen/apps/ComfyUI_SGLDiffusion/core/generator.py`) or a dedicated\n node. Otherwise the generic `image` / `video` key is the honest one; a model-specific key\n without matching plugin support makes the page claim support that does not exist.\n- Prose describing ComfyUI support inline instead of using the component is a finding: the\n facts drift from the plugin.\n\n### 6. Commands / port\n- Launch uses `sglang serve` — flag any `python -m sglang.launch_server` /\n `python3 -m sglang.launch_server` (deprecated). The engine already emits `sglang serve`;\n guard against prose/cells reintroducing the old launcher.\n- Port `30000` everywhere (launch, curl, client `base_url`, bench) — flag `8000`.\n Launch port must match client/curl port on the same page.\n\n### 7. Frontmatter\n- Every new MDX page has `title:` and a **top-level** `description:` (a real one-line value\n prop, not copied from another vendor) — NOT `metatags.description` (non-canonical; the\n top-level field is what renders as the subtitle and SEO meta — see mintlify-authoring).\n- **No `mode: wide` on a model page** — it hides the right-hand \"On this page\" ToC that every\n other model page has. Leave `mode` unset (the Deploy/Playground panels self-cap at 900px, so\n the default column holds them fine). `mode: wide` belongs only on category `intro.mdx` grids.\n- `tag: NEW` only for genuine new launches; when one is added, stale `tag: NEW` on older\n pages should be dropped in the same PR (`grep -RlE \"^tag: NEW\" docs/cookbook/`).\n- MDX imports BOTH `Deployment` and `Playground` from `/src/snippets/...` (absolute).\n- Deploy heading slugs to `deployment` (or `deploy`), Playground to `playground` — so\n \"↑ Switch base\" and \"Open the Playground →\" scroll. No numbered headings for these two.\n\n### 8. Navigation & homepage\n- New page → `docs/docs.json` updated: under the right vendor group inside\n `navigation` → Cookbook → Autoregressive Models, root-relative, **no `.mdx`**:\n `cookbook///`.\n- Homepage `` in `docs/cookbook//intro.mdx` points to the vendor's\n flagship; new vendors get a new `` + a logo at `docs/cards/logos/.png` —\n **940×525 RGBA transparent, icon-only (no wordmark)**, lowercase filename, tracked via\n `git add -f` (`*.png` is gitignored repo-wide). Card order matches the `docs.json` nav order.\n- Don't change `docs/cookbook/intro.mdx` for individual model adds (top-level only).\n\n### 9. Links & factual\n- HuggingFace URLs resolve to a real model. License section matches the actual HF license\n (don't copy from another model). Docker images from `lmsysorg/sglang`; no `sgl-project-dev`.\n The image **tag** is a real build (a release the author ran, or `:dev`/nightly) — not a\n guessed version.\n- Internal links root-relative, no extension (`/cookbook/.../`); flag `.md`/`.mdx`\n or `../`-relative links. `docs.sglang.io` is canonical.\n- No Google-Drive image links (don't render). Shell placeholders are `export VAR=`,\n not `${VAR}` (a bash no-op).\n- **Parser ids must exist in the code registries** on the PR's target branch: every\n `--reasoning-parser X` / `--tool-call-parser Y` named in prose or in\n `playgroundFeatures.parsers` flags is a registered key in\n `python/sglang/srt/parser/reasoning_parser.py` (DetectorMap) /\n `python/sglang/srt/function_call/function_call_parser.py` (ToolCallParserEnum) —\n prose naming a near-miss id (e.g. the reasoning id where the tool id differs) is a\n factual bug. `--…-parser auto` is acceptable ONLY if the template-detection rules\n (`python/sglang/srt/parser/template_detection.py`) actually resolve THIS model's\n chat template to the right parser — no rule match means auto silently disables the\n parser; when in doubt require explicit ids (the DSv4 page pins explicit ids).\n\n### 9b. MDX authoring (Mintlify) — detail in `cookbook-add-model/references/mintlify-authoring.md`\n- **Forbidden syntax**: no Docusaurus admonitions (`:::`), `@site`/`@theme`, GitHub alert\n blocks (`> [!NOTE]`), markdown **pipe tables** (use JSX `
`), inline `
`,\n or unknown components. ``/`` only on category `intro.mdx`, not model pages.\n- Code fences are **labeled** (e.g. `python Example` / `bash Command` / `text Output` after\n the opening fence); a fenced block nested inside another uses four backticks outside.\n- §3 commands and outputs are **collapsible** (DeepSeek-V4 pattern): every runnable\n example wrapped in an ``, its real output in a following\n `` (`Pending update...` only with user\n acknowledgement). Flag bare/inline example blocks and `**Output Example:**` headings.\n- Reasoning-parser example matches the parser's **output shape**: separate-field\n (`reasoning_content` + `content`) vs inline `` tags parsed out of `content`.\n- No hardcoded sampling params (`temperature` / `top_p`) in sample code (SGLang uses\n `generation_config.json` defaults); listing them in §1 informationally is fine.\n\n### 10. Quantization rules\n- **NVFP4** checkpoints are Blackwell-only (B200/B300/GB300) — never AMD. An AMD FP4 cell\n is legitimate ONLY when the vendor published an **MXFP4** checkpoint for it (e.g.\n `amd/Qwen3.5-397B-A17B-MXFP4` on MI355X) — verify the HF repo resolves; otherwise the\n AMD FP4 chip must be absent/`disabled`.\n- BF16 / FP8 work on NVIDIA and AMD. `--kv-cache-dtype fp8_e4m3` in a cell is an\n accuracy-degrading flag — see §3 (needs explicit maintainer sign-off; default\n home is Playground/tips).\n\n### 11. Scope\n- Changes match the PR title. Flag global changes hiding behind a platform-specific title\n (e.g. an \"H200 FP8\" PR that adds a flag to ALL cells). Unmentioned side-fixes belong in\n the PR body.\n\n### 12. Duplicate PRs\n- Another open PR for the same model? Flag it; compare completeness; note merge-conflict\n risk on `docs.json` + the vendor card; flag a superseded older PR by the same author.\n\n### 13. Build / validate\n```bash\ncd docs\nmint validate\nmint broken-links\n```\nOptional: `mint dev` for a visual smoke test.\n\n### 14. Reviewer feedback\n- `gh api repos/sgl-project/sglang/pulls//comments` — have prior reviewer requests been\n addressed? Unresolved requested-changes should be flagged.\n\n### 15. Grammar & spelling\n- Check added/changed prose for typos and grammar (e.g. \"recommend\" vs \"recommended\").\n Flag each with the exact wrong text + correction.\n\n## Output\n\nPer file:\n- ✅ PASS\n- ⚠️ ISSUE: \\\n- 🔴 BLOCK: \\\n\nOverall: **APPROVE** / **REQUEST CHANGES** / **BLOCKED**\n\n\n---\nname: debug-cuda-crash\ndescription: Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging\n---\n\n# Tutorial: Debugging CUDA Crashes with Kernel API Logging\n\nThis tutorial shows you how to debug CUDA crashes and errors in SGLang using the `@debug_kernel_api` logging decorator.\n\n## Goal\n\nWhen your code crashes with CUDA errors such as illegal memory access, device-side assert, out-of-bounds, or NaN/Inf, use kernel API logging to:\n- Capture input tensors BEFORE the crash occurs\n- Understand what data caused the problem\n- Track tensor shapes, dtypes, and values through the call boundary that triggered the crash\n- Detect numerical issues such as NaN, Inf, or obviously wrong shapes\n\n## Why Use Kernel API Logging?\n\n**Problem**: CUDA errors often crash the program before normal debugging output is flushed.\n\n**Solution**: SGLang's `@debug_kernel_api` decorator logs inputs before execution, so you can still see what caused the crash even after the program aborts.\n\n## What Is Covered?\n\nThe current logging coverage focuses on the highest-value kernel boundaries in SGLang:\n- Custom ops registered through `register_custom_op(...)`\n- External custom ops registered through `register_custom_op_from_extern(...)`\n- LLM attention, linear, quantization, and multi-platform wrapper entry points\n- Diffusion attention impl, linear, rotary, and custom-op wrapper entry points\n- Selected direct `torch.ops.sglang.*` hotspots and model-specific bypasses\n\nThis means the logging is useful for both LLM and diffusion kernel debugging, but it does not automatically cover every pure PyTorch call in the repository.\n\n## Step 1: Enable Kernel API Logging\n\n### Basic Logging (Function Names Only)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=1\nexport SGLANG_KERNEL_API_LOGDEST=stdout\n\npython my_script.py\n```\n\nOutput:\n```\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: RMSNorm.forward\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\n================================================================================\n[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.custom_op.fused_inplace_qknorm\n```\n\nThis is a real level-1 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Detailed Logging (Inputs with Metadata)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nOutput in `debug.log`:\n```\n================================================================================\n[2026-03-19 00:47:30] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[0]=QKVParallelLinear(\n repr=QKVParallelLinear(in_features=1024, output_features=4096, bias=False, tp_size=1, gather_output=False)\n )\n arg[1]=Tensor(\n shape=(1, 1024)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n arg[2]=None\nOutput:\n return=Tensor(\n shape=(1, 4096)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n )\n```\n\nThis is a real level-3 excerpt captured from `Qwen/Qwen3-0.6B`.\n\n### Full Logging (With Tensor Statistics)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\npython my_script.py\n```\n\nAdditional output:\n```\n================================================================================\n[2026-03-19 01:00:42] SGLang Kernel API Call: diffusion.quant_method.UnquantizedLinearMethod.apply\nPositional input arguments:\n arg[1]=Tensor(\n shape=(1, 77, 768)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-27.250000\n max=28.500000\n mean=0.011723\n nan_count=0\n inf_count=0\n )\nOutput:\n return=Tensor(\n shape=(1, 77, 2304)\n dtype=torch.bfloat16\n device=cuda:0\n requires_grad=False\n is_contiguous=True\n min=-8.937500\n max=9.375000\n mean=0.009460\n nan_count=0\n inf_count=0\n )\n```\n\nThis is a real level-5 excerpt captured from `black-forest-labs/FLUX.1-dev`.\n\n### Crash-Safe Dumps (Inputs Saved Before Execution)\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\n\npython my_script.py\n```\n\nAt level 10, SGLang saves the inputs before execution. If the kernel crashes, the dump directory still contains the inputs and exception metadata.\n\nIf CUDA graph capture is active, tensor dumps are skipped automatically to avoid capture-time CUDA errors. In that case, you still get the kernel API call log, but not `inputs.pt` / `outputs.pt`.\n\nLevel-10 dumps are best understood as crash-safe call snapshots. They always preserve the observed call boundary. They do not guarantee one-click replay for every method, because some methods depend on module state that is not serialized into the dump.\n\nReal level-10 dump layout from `Qwen/Qwen3-0.6B`:\n\n```text\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/inputs.pt\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/metadata.json\n/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/outputs.pt\n```\n\nReal `metadata.json` excerpt:\n\n```json\n{\n \"function_name\": \"RotaryEmbedding.forward\",\n \"timestamp\": \"20260319_004821_182\",\n \"process_id\": 919286,\n \"execution_status\": \"completed\",\n \"input_tensor_keys\": [\"arg_0\", \"arg_1\", \"arg_2\"],\n \"output_tensor_keys\": [\"result_0\", \"result_1\"]\n}\n```\n\n## Step 2: Reproduce an LLM CUDA Crash\n\nCreate a temporary reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_llm_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.srt.utils.custom_op import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(indices, table):\\\\n\"\n \" return torch.empty((*indices.shape, table.shape[-1]), device=table.device, dtype=table.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_llm_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_llm_cuda_crash(indices, table):\\\\n\"\n \" out = F.embedding(indices, table)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"table = torch.randn(4, 8, device='cuda', dtype=torch.float16)\\\\n\"\n \"indices = torch.tensor([0, 7], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_llm_cuda_crash(indices, table)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level1.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nWhat to expect:\n- The script exits with a CUDA `device-side assert`\n- The log still contains the last API boundary before the crash\n\nTry the same example at level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level3.log \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow the log shows tensor metadata before the crash.\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_llm_level10_dumps \\\npython3 /tmp/sglang_llm_crash.py\n```\n\nNow you should see:\n- A log entry for `sglang.custom_op.mock_llm_cuda_crash`\n- A dump directory with `inputs.pt`\n- `metadata.json` showing `execution_status: \"exception\"`\n- No `outputs.pt`, because the kernel crashed before producing output\n\nFor real-model success-path level-10 dumps, it is often easier to temporarily disable CUDA graph and piecewise CUDA graph for the debug run.\n\n## Step 3: Reproduce a Diffusion CUDA Crash\n\nCreate a temporary diffusion-side reproducer:\n\n```bash\npython3 - <<'PY'\nfrom pathlib import Path\nPath(\"/tmp/sglang_diffusion_crash.py\").write_text(\n \"import torch\\\\n\"\n \"import torch.nn.functional as F\\\\n\"\n \"from sglang.multimodal_gen.runtime.layers.utils import register_custom_op\\\\n\\\\n\"\n \"def _fake_embedding(positions, cache):\\\\n\"\n \" return torch.empty((*positions.shape, cache.shape[-1]), device=cache.device, dtype=cache.dtype)\\\\n\\\\n\"\n \"@register_custom_op(op_name='mock_diffusion_cuda_crash', fake_impl=_fake_embedding)\\\\n\"\n \"def mock_diffusion_cuda_crash(positions, cache):\\\\n\"\n \" out = F.embedding(positions, cache)\\\\n\"\n \" torch.cuda.synchronize()\\\\n\"\n \" return out\\\\n\\\\n\"\n \"cache = torch.randn(4, 64, device='cuda', dtype=torch.float16)\\\\n\"\n \"positions = torch.tensor([0, 9], device='cuda', dtype=torch.long)\\\\n\"\n \"mock_diffusion_cuda_crash(positions, cache)\\\\n\"\n)\nPY\n\nSGLANG_KERNEL_API_LOGLEVEL=1 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level1.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 3:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=3 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level3.log \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nTry level 10:\n\n```bash\nSGLANG_KERNEL_API_LOGLEVEL=10 \\\nSGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level10.log \\\nSGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_diffusion_level10_dumps \\\npython3 /tmp/sglang_diffusion_crash.py\n```\n\nIf your local environment has unrelated FlashInfer import issues, resolve them in the shell before running the example. The example itself does not set any `FLASHINFER_*` environment variable.\n\n## Step 4: Multi-Process Debugging\n\nWhen running with multiple GPUs or worker processes, use `%i` in the log path:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\n\ntorchrun --nproc_per_node=4 my_script.py\n```\n\nThis creates separate logs such as:\n- `debug_rank_12345.log`\n- `debug_rank_12346.log`\n- `debug_rank_12347.log`\n- `debug_rank_12348.log`\n\nReal multi-process example from a 2-GPU `Qwen/Qwen2.5-0.5B-Instruct` run:\n\n```text\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950201.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950349.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950350.log\n/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950351.log\n```\n\nYou should usually do the same for level-10 dump directories:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps_%i\n```\n\nThis avoids multiple ranks writing into the same dump directory tree.\n\n## Step 5: Filter Level-10 Dumps\n\nIf level 10 is too noisy, restrict dumps to specific APIs:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\nexport SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps\nexport SGLANG_KERNEL_API_DUMP_INCLUDE='sglang.custom_op.*'\nexport SGLANG_KERNEL_API_DUMP_EXCLUDE='*.fake_impl'\n```\n\n`SGLANG_KERNEL_API_DUMP_INCLUDE` and `SGLANG_KERNEL_API_DUMP_EXCLUDE` use shell-style wildcard matching.\n\n## Step 6: Common CUDA Errors and What to Check\n\n### Illegal Memory Access or Device-Side Assert\n\n**Typical errors**:\n```\nRuntimeError: CUDA error: an illegal memory access was encountered\ntorch.AcceleratorError: CUDA error: device-side assert triggered\n```\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck in the logs:\n- ✅ Tensor shapes\n- ✅ Tensor dtypes\n- ✅ CUDA vs CPU device placement\n- ✅ Tensor stride / contiguity\n- ✅ Whether the failing call has inputs logged but no outputs logged\n\nTypical shape-mismatch pattern:\n\n```text\nSGLang Kernel API Call: ...\narg[0]=Tensor(shape=(..., 128), ...) # ✅ expected dimension\narg[1]=Tensor(shape=(..., 64), ...) # ❌ mismatch\n```\n\nThis often points to head-dim, hidden-dim, or cache-layout mismatch rather than a random CUDA failure.\n\n### NaN or Inf\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nCheck:\n- `min`\n- `max`\n- `mean`\n- `nan_count`\n- `inf_count`\n\nTypical bad pattern:\n\n```text\nTensor(\n ...\n min=-1234567.000000 # ❌ suspiciously large\n max=9876543.000000 # ❌ suspiciously large\n mean=nan # ❌ bad\n nan_count=128 # ❌ found NaNs\n inf_count=0 # ✅ no Infs here\n)\n```\n\nThis usually means the bad values were already present before the crashing kernel.\n\n### Out of Memory\n\nUse:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nCheck:\n- Unexpectedly large tensor shapes\n- Batch size\n- Sequence length\n- Frame count or image resolution in diffusion workloads\n\nAlso check whether a supposedly per-token or per-frame tensor accidentally became full-sequence or full-image sized.\n\nTypical bad pattern:\n\n```text\nTensor(\n shape=(1024, 8192, 128, 128) # ❌ way too large\n ...\n)\n```\n\n### Example: Spot a Shape Bug from the Log\n\nSuppose the failing API log looks like this:\n\n```text\n[2026-03-19 00:47:30] SGLang Kernel API Call: RotaryEmbedding.forward\nPositional input arguments:\n arg[0]=Tensor(shape=(1, 8), dtype=torch.int64, ...)\n arg[1]=Tensor(shape=(1, 8, 8, 256), dtype=torch.bfloat16, ...) # ✅ query\n arg[2]=Tensor(shape=(1, 8, 4, 64), dtype=torch.bfloat16, ...) # ❌ key head_dim mismatch\n```\n\nWhat this tells you:\n- ✅ positions look reasonable\n- ✅ query looks plausible\n- ❌ key last dimension is inconsistent with the expected rotary/head dimension\n\nThat usually means the bug is in projection layout, head packing, or cache format rather than in the rotary kernel itself.\n\n## Step 7: Combine with compute-sanitizer\n\nFor harder bugs, combine kernel API logging with CUDA memory checking:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncompute-sanitizer --tool memcheck python3 /tmp/sglang_llm_crash.py\n```\n\nUse `debug.log` to see the exact inputs that reached the crashing API boundary.\n\nTypical `compute-sanitizer` output:\n\n```text\n========= COMPUTE-SANITIZER\n========= Invalid __global__ write of size 4 bytes\n========= at 0x1234 in SomeKernel\n========= by thread (256,0,0) in block (10,0,0)\n========= Address 0x... is out of bounds\n```\n\nUse the sanitizer output to identify the failing kernel and use `debug.log` to identify the exact tensors that reached the API boundary right before it.\n\nIf you need more synchronous host-side error reporting, you can try `CUDA_LAUNCH_BLOCKING=1` as a separate follow-up experiment. It is not part of the default workflow because it changes execution timing and can hide concurrency-related behavior.\n\n## Step 8: Combine with cuda-gdb\n\nFor crashes that need a stack trace instead of only memory diagnostics:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\nexport SGLANG_KERNEL_API_LOGDEST=debug.log\n\ncuda-gdb --args python3 /tmp/sglang_llm_crash.py\n```\n\nInside `cuda-gdb`:\n\n```text\n(cuda-gdb) run\n(cuda-gdb) where\n```\n\nThen correlate the backtrace with `debug.log`.\n\n## Step 9: Kernel-Level Debugging with printf()\n\nWhen you own the CUDA kernel, `printf()` is still useful for narrowing down bad indices, bad launch geometry, or broken state propagation.\n\nBasic pattern:\n\n```cpp\n__global__ void MyKernel(const float* input, float* output, int n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n\n if (threadIdx.x == 0 && blockIdx.x == 0) {\n printf(\"n=%d input0=%f\\n\", n, input[0]);\n }\n\n if (idx < n) {\n output[idx] = input[idx] * 2.0f;\n }\n}\n```\n\nAfter launch, force the output to flush:\n\n```python\nmy_kernel(...)\ntorch.cuda.synchronize()\n```\n\nFor warp-specialized kernels, do not blindly print only on `threadIdx.x == 0`. Pick one representative thread per warp or per specialization group instead.\n\n### Warp-Specialized Kernels: Choosing the Right Print Thread\n\nProblem:\n- `threadIdx.x == 0` only prints from the first warp in the block\n- for warp-specialized kernels, that often misses the warp or group that is actually wrong\n\nBetter pattern:\n\n```cpp\n__global__ void WarpSpecializedKernel(...) {\n // Example: first lane of each warp\n if ((threadIdx.x % 32) == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n }\n}\n```\n\nOr, if the kernel is organized in larger specialization groups, print once per group instead of once per block.\n\nCommon mistake:\n\n```cpp\n// Only warp 0 prints\nif (threadIdx.x == 0) {\n printf(\"warp=%d\\n\", threadIdx.x / 32);\n}\n```\n\n### Quick Reference\n\n| Kernel Type | Print Condition | Notes |\n|----------|----------|-------------|\n| Simple kernel | `threadIdx.x == 0` | One thread per block is usually enough |\n| Warp-specialized kernel | one representative lane per warp | e.g. `threadIdx.x % 32 == 0` |\n| Group-specialized kernel | one representative lane per group | choose based on the kernel's scheduling layout |\n\n### Other Kernel Debugging Tools\n\n```cpp\nassert(value >= 0.0f && \"value must be non-negative\");\nstatic_assert(BLOCK_SIZE % 32 == 0, \"BLOCK_SIZE must be warp aligned\");\n```\n\n## Environment Variables Reference\n\n| Variable | Values | Description |\n|----------|--------|-------------|\n| `SGLANG_KERNEL_API_LOGLEVEL` | `0` | No logging (default) |\n| | `1` | Function names only |\n| | `3` | Inputs and outputs with metadata |\n| | `5` | Level 3 plus tensor statistics |\n| | `10` | Level 5 plus crash-safe tensor dumps |\n| `SGLANG_KERNEL_API_LOGDEST` | `stdout` | Log to stdout |\n| | `stderr` | Log to stderr |\n| | `` | Log to file |\n| | `log_%i.txt` | `%i` expands to process ID |\n| `SGLANG_KERNEL_API_DUMP_DIR` | `` | Directory for level-10 dumps |\n| `SGLANG_KERNEL_API_DUMP_INCLUDE` | wildcard list | Only dump matching API names |\n| `SGLANG_KERNEL_API_DUMP_EXCLUDE` | wildcard list | Skip matching API names |\n\n## Best Practices\n\n### 1. Start with Level 3\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\nLevel 3 is usually enough to catch wrong shapes, wrong dtypes, and wrong devices.\n\n### 2. Use Level 5 for Numerical Issues\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=5\n```\n\nUse it when you suspect NaN or Inf values.\n\n### 3. Use Level 10 for Crash Reproduction\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=10\n```\n\nThis is the most useful mode when the process crashes before you can inspect live tensors.\n\nIf you need successful input/output dumps from a real model run, temporarily disable CUDA graph for that debug session.\n\nWhen level 10 is too noisy, pair it with `SGLANG_KERNEL_API_DUMP_INCLUDE` / `SGLANG_KERNEL_API_DUMP_EXCLUDE` instead of dumping every covered API.\n\n### 4. Log to File for Crashes\n\n```bash\nexport SGLANG_KERNEL_API_LOGDEST=crash.log\n```\n\nFile logs are safer than stdout when the process aborts.\n\n### 5. Disable Logging in Production\n\n```bash\nunset SGLANG_KERNEL_API_LOGLEVEL\n```\n\nWhen disabled, the decorator returns the original callable and adds no runtime logging overhead.\n\n## Troubleshooting\n\n### No Logs Appear\n\nCheck:\n1. `echo $SGLANG_KERNEL_API_LOGLEVEL`\n2. `echo $SGLANG_KERNEL_API_LOGDEST`\n3. Whether the failing path goes through a covered API boundary\n\n### Too Much Output\n\nReduce the level:\n\n```bash\nexport SGLANG_KERNEL_API_LOGLEVEL=3\n```\n\n### Statistics Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nstatistics=[skipped: CUDA graph capture in progress]\n```\n\nThat is expected. Level-5 statistics are intentionally skipped during CUDA graph capture to avoid synchronization side effects.\n\n### Tensor Dumps Are Skipped During CUDA Graph Capture\n\nIf you see:\n```text\nTensor dump skipped: CUDA graph capture in progress\n```\n\nThat is also expected. Level-10 dumps require copying tensors to CPU, which is not allowed during CUDA graph capture.\n\n\n---\nname: debug-distributed-hang\ndescription: Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations.\n---\n\n# Debugging Distributed Hangs in SGLang\n\n## Overview\n\nHangs in distributed inference happen when ranks diverge in state, causing collective operations (AllGather, AllReduce, Broadcast, Barrier) to deadlock. Common causes:\n\n- **Size mismatch**: ranks pass different tensor sizes to a collective\n- **Branch divergence**: one rank enters a collective, another skips it\n- **Cascading state drift**: a small non-determinism (e.g., floating-point) propagates into different batch structures\n- **Resource exhaustion**: one rank OOMs or crashes, others wait forever\n\n## Prerequisites\n\n- **py-spy**: `pip install py-spy` or system package. Requires root or `CAP_SYS_PTRACE` to attach to running processes.\n- **cuda-gdb**: Ships with the CUDA toolkit. Ensure it's on your `PATH`.\n\n## Step 1: Confirm and Locate the Hang\n\n### 1a. Watchdog / py-spy\n\nSGLang's watchdog automatically dumps py-spy traces on timeout. Look for:\n\n```\nScheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)\n```\n\nThe py-spy dump shows the stack trace of each thread. The hanging thread is typically blocked in a CUDA synchronize or NCCL collective:\n\n```\nThread (active): \"MainThread\"\n cuStreamSynchronize (libcuda.so)\n ...\n forward_extend (model_runner.py)\n```\n\nSGLang has two watchdog modes (see `python/sglang/srt/utils/watchdog.py`):\n- **Hard watchdog** (`soft=False`, default): dumps py-spy traces then sends `SIGQUIT` to kill the parent process.\n- **Soft watchdog** (`soft=True`): only logs the timeout without killing the process, giving you more time to manually attach debuggers or collect coredumps.\n\nIf the watchdog doesn't trigger, manually dump:\n\n```bash\npy-spy dump --pid \n```\n\n### 1b. NCCL Debug Logging\n\n```bash\nexport NCCL_DEBUG=INFO\nexport NCCL_DEBUG_SUBSYS=COLL\n```\n\nLook for the last collective logged before the hang. Mismatched sizes show up as one rank waiting and another never entering.\n\n### 1c. CUDA Coredump\n\nWhen a process hangs, you can trigger a GPU coredump on demand to see which kernel is stuck. Set these env vars before launching:\n\n```bash\nexport CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1\nexport CUDA_COREDUMP_PIPE=\"/tmp/cuda_pipe_%h_%p\"\nexport CUDA_COREDUMP_FILE=\"/tmp/cuda_coredump_%h_%p\"\nexport CUDA_COREDUMP_SHOW_PROGRESS=1\nexport CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'\n```\n\nWhile the process is hanging, find the pipe via `/proc//fd/` and write to it to trigger the dump:\n\n```bash\nls /proc//fd/ -la 2>/dev/null | grep cuda_pipe\ndd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe__\n```\n\nAlternatively, if you don't need to keep the process alive, `kill -SIGABRT ` also triggers a CUDA coredump (but terminates the process).\n\nThen open with `cuda-gdb --batch -ex \"target cudacore \"`. On load, it immediately shows which kernel is stuck. For example:\n\n```\nOpening GPU coredump: \n[Current focus set to CUDA kernel 0, grid 622721, cluster (4,0,0), block (16,0,0), thread (64,0,0), device 0, sm 0, warp 0, lane 0]\n#0 0x00007f8029b2b040 in ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)<<<(24,1,1),(512,1,1)>>> ()\n```\n\nThis told us the hang was in an NCCL AllGather — not a compute kernel. Combined with the py-spy stack pointing to `LogitsProcessor.forward` → `tensor_model_parallel_all_gather`, we knew it was an AllGather size mismatch between TP ranks.\n\n\n### 1d. Identify the Collective\n\nFrom the stack traces and logs, identify:\n- Which collective hangs (AllGather, AllReduce, Broadcast)\n- Which code path invokes it (e.g., `LogitsProcessor`, `tensor_model_parallel_all_gather`)\n- Whether it's a size mismatch or a missing participant\n\n## Step 2: Per-Rank Logging\n\nThe key technique: each rank writes its own log file so you can diff them.\n\n### Setup Pattern\n\n```python\nimport os\n\n_debug_files = {}\n\ndef get_debug_file(rank):\n key = f\"rank{rank}\"\n if key not in _debug_files:\n _debug_files[key] = open(f\"/tmp/debug_rank{rank}.log\", \"w\")\n return _debug_files[key]\n```\n\nGate logging behind an env var to avoid overhead in production. `SGLANG_DEBUG_HANG` is not a built-in SGLang env var — you need to add this check yourself in the code you're instrumenting:\n\n```python\nif os.environ.get(\"SGLANG_DEBUG_HANG\"):\n f = get_debug_file(rank)\n f.write(f\"EVENT_NAME key1={val1} key2={val2}\\n\")\n f.flush()\n```\n\n### What to Log\n\nLog structured events at key state-mutation points:\n\n```python\nf.write(f\"SCHED_BATCH step={step} num_reqs={n} extend_lens={lens}\\n\")\nf.write(f\"VERIFY predict_hash={hash} accept_len={alen}\\n\")\nf.write(f\"CACHE_INSERT rid={rid} num_tokens={n}\\n\")\n```\n\nUse consistent event names (uppercase prefix) for easy grep/diff.\n\n### Hash Large Tensors\n\nFor tensor values, compute a hash instead of dumping raw data:\n\n```python\nimport hashlib\nh = hashlib.md5(tensor.cpu().numpy().tobytes()).hexdigest()[:8]\nf.write(f\"LOGITS logits_hash={h}\\n\")\n```\n\nFor token ID lists, `str(list).encode()` works:\n\n```python\nh = hashlib.md5(str(tensor.tolist()).encode()).hexdigest()[:8]\n```\n\n### Avoid Implicit Synchronization\n\n`tensor.cpu()`, `tensor.tolist()`, and `tensor.numpy()` all trigger CUDA synchronization. This can:\n- Change timing and mask or move the hang\n- Deadlock if the log point is between two collectives that must run back-to-back\n\nPrefer logging values that are already on CPU (e.g., Python ints, list lengths, request IDs). When you must hash a GPU tensor, do it at a point where the GPU is already idle (e.g., between scheduler steps, not inside a model forward pass).\n\n## Step 3: Diff to Find the Diverge Point\n\n### Basic Diff\n\n```bash\n# Extract specific event type\ngrep \"^VERIFY\" /tmp/debug_rank0.log > /tmp/v_r0.txt\ngrep \"^VERIFY\" /tmp/debug_rank1.log > /tmp/v_r1.txt\ndiff /tmp/v_r0.txt /tmp/v_r1.txt | head -20\n```\n\n### Count Events\n\n```bash\ngrep -c \"^VERIFY\" /tmp/debug_rank*.log\n```\n\nIf counts differ, one rank executed more iterations — that's already a diverge signal.\n\n### Find First Diverge\n\nThe first diff line tells you the exact step where ranks diverge. All lines before it are identical — the root cause is at or before this step.\n\n## Step 4: Binary-Search the Root Cause\n\nOnce you find the diverging event, trace backwards:\n\n### 4a. Identify Inputs\n\nFor the diverging operation, list all its inputs. Add hash logging for each:\n\n```python\nf.write(\n f\"OP_INPUTS input_a_hash={h_a} input_b_hash={h_b} \"\n f\"input_c_hash={h_c} input_d_hash={h_d}\\n\"\n)\n```\n\n### 4b. Diff Inputs Across Ranks\n\nCompare the hashes. Some inputs will match, some won't. The non-matching input is where divergence entered.\n\n### 4c. Recurse\n\nFor the non-matching input, trace where it was produced and repeat: hash its inputs, diff across ranks, find the divergent one. Continue until you reach the root cause.\n\n## Step 5: Common Root Causes and Fixes\n\n### Floating-Point Non-Determinism\n\n**Symptom**: All \"logical\" inputs are identical (same logits after all-gather), but derived floating-point values (softmax, probabilities) differ across GPUs.\n\n**Example**: EAGLE speculative decoding — `F.softmax` → `top_k_renorm_prob` → `top_p_renorm_prob` produces slightly different `target_probs` on each GPU. The sampling kernel then picks different tokens. These flow into `output_ids` → radix cache → different prefix match depths → different `extend_seq_lens` → AllGather size mismatch → hang.\n\n### Random Number Divergence\n\n**Symptom**: Operations using `torch.rand` produce different values on each rank.\n\n**Fix**: Generate on rank 0 and broadcast, or use a shared seed.\n\n### Conditional Code Paths\n\n**Symptom**: A condition (e.g., memory check, queue length) evaluates differently on different ranks, causing one rank to enter a collective while another skips it.\n\n**Fix**: Synchronize the condition value before branching, or restructure to ensure all ranks take the same path.\n\n### Pipeline Parallel (PP) Send/Recv Mismatch\n\n**Symptom**: In PP setups, one stage issues a `send` that the next stage never `recv`s (or vice versa), causing both to block indefinitely. Unlike TP hangs (collective mismatches), PP hangs typically involve point-to-point operations.\n\n**Fix**: Ensure all stages agree on the number of microbatches and the sequence of send/recv calls for each microbatch.\n\n## Step 6: Verify the Fix\n\nRun the failing test multiple times to confirm the fix is stable. Intermittent hangs require many runs. A test that hung ~30% of the time needs at least 10 clean passes to be confident.\n\n## Quick Reference\n\n| Technique | When to Use |\n|-----------|-------------|\n| py-spy dump | First step — see where each rank is stuck |\n| `NCCL_DEBUG=INFO` | Identify which collective and sizes |\n| CUDA coredump + `cuda-gdb` | See which GPU kernel is blocked |\n| Per-rank log files | Compare rank states over time |\n| Hash of tensors | Efficiently compare large tensors across ranks |\n| `diff` on extracted events | Find the exact step of divergence |\n| `broadcast(result, src=0)` | Fix floating-point or sampling non-determinism |\n\n\n# env-var-conventions\n\nPath: `.agents/skills/env-var-conventions/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/env-var-conventions/SKILL.md)\n\n\n# generate-profile\n\nPath: `.agents/skills/generate-profile/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/generate-profile/SKILL.md)\n\n\n# kernel-organization\n\nPath: `.agents/skills/kernel-organization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kernel-organization/SKILL.md)\n\n\n# kl-consistency-test\n\nPath: `.agents/skills/kl-consistency-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/kl-consistency-test/SKILL.md)\n\n\n# large-class-style\n\nPath: `.agents/skills/large-class-style/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/large-class-style/SKILL.md)\n\n\n# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md)\n\n\n# References Skill\n\nPath: `.agents/skills/llm-torch-profiler-analysis/references/heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/sgl-project/sglang/HEAD/.agents/skills/llm-torch-profiler-analysis/references/heuristics.md)\n\n\n"}