Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
55d8239
perf: precapture bounded decode cuda graphs
kuma-loong Aug 21, 2026
798dbbf
perf: capture larger cuda graph shapes first
kuma-loong Aug 22, 2026
e6ef95a
feat: restore batch-only decode graph providers
kuma-loong Aug 24, 2026
014d524
feat: add typed batch-only decode graph contract
kuma-loong Aug 25, 2026
6dac096
fix: use live decode graph runner during startup
kuma-loong Aug 25, 2026
67b05d1
fix: inspect captured graphs through live runner
kuma-loong Aug 25, 2026
f09ef4c
fix: read current decode graph probe summary
kuma-loong Aug 25, 2026
84dda47
test: close batch-only stage two coverage gaps
kuma-loong Aug 25, 2026
f997ede
feat: productionize batch-only decode providers
kuma-loong Aug 25, 2026
dcef2db
feat: productionize batch-only tilelang mla scores
kuma-loong Aug 25, 2026
01afe24
fix: support strided tilelang mla score views
kuma-loong Aug 25, 2026
491116b
fix: preserve sparse mla score contracts
kuma-loong Aug 25, 2026
cc3192d
test: cover sparse decode topology partitioning
kuma-loong Aug 25, 2026
cb22fed
test: cover tilelang mla replay through 64k
kuma-loong Aug 25, 2026
5f8d962
test: cover sparse graph path transitions
kuma-loong Aug 25, 2026
378d177
refactor: normalize batch-only decode providers
kuma-loong Aug 25, 2026
add354a
fix: pin sglang kernel abi
kuma-loong Aug 26, 2026
7c6e11a
fix: require native multiprocessor count for gemma4
kuma-loong Aug 26, 2026
86d98fe
docs: add batch-only operator review rules
kuma-loong Aug 26, 2026
e47c689
refactor: remove redundant decode score method set
kuma-loong Aug 26, 2026
6f7b847
feat: support deltakv batch-only decode graphs
kuma-loong Aug 26, 2026
4aada47
feat: default decode graphs to batch-only
kuma-loong Aug 26, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 13 additions & 1 deletion .agents/skills/review-operator-organization/SKILL.md
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
---
name: review-operator-organization
description: Review Sparse-vLLM operator architecture, provider selection, platform capability boundaries, kernel ownership, dependency compatibility, weight layouts, fallback semantics, and validation. Use for diffs touching src/sparsevllm/operators, src/sparsevllm/platforms, Triton or external kernels, model-to-operator call sites, quantized weight loading, CUDA Graph constraints, optional kernel dependencies, or backend removal and migration.
description: Review Sparse-vLLM operator architecture, provider selection, platform capability boundaries, kernel ownership, dependency compatibility, weight layouts, batch-only CUDA Graph adaptation, fallback semantics, and validation. Use for diffs touching src/sparsevllm/operators, src/sparsevllm/platforms, Triton or external kernels, model-to-operator call sites, quantized weight loading, CUDA Graph constraints, optional kernel dependencies, or backend removal and migration.
---

# Review Operator Organization
Expand Down Expand Up @@ -98,6 +98,18 @@ fallback path from model construction through execution.
- Ensure one provider's build or JIT failure does not disable unrelated
operators.

### Decode CUDA Graph

Batch-only is the only maintained decode CUDA Graph shape policy. For any
review touching captured decode, graph input preparation, provider graph state,
or sparse topology paths, read and enforce
[references/batch-only-decode-graph.md](references/batch-only-decode-graph.md).
That reference defines graph identity, static versus dynamic metadata, unified
input ownership, participant lifecycles, external wrappers, validation, and
finding severity. Eager may remain as a separate correctness path or for
operators that do not support graph capture; do not preserve a second bucketed
graph architecture.

### Kernel Portfolio

- Treat standard operations as upstream-first. Prefer a mature upstream public
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,207 @@
# Batch-Only Decode CUDA Graph Review

Read this reference when a review touches an operator reachable from captured
decode, decode graph input preparation, provider graph state, or sparse
short/long topology paths.

## Scope and Vocabulary

Batch-only is the only maintained decode CUDA Graph shape policy. Do not add or
preserve bucketed-only graph implementations, context-bucket routing, parallel
provider families, or configuration surfaces merely to keep a second graph
architecture alive. Eager may remain as an independent correctness or
unsupported-graph path; it must not leak context-dependent dispatch into
captured decode.

Use these terms consistently:

- **batch-only graph**: graph identity depends on batch capacity but not actual
per-step context lengths;
- **strict batch-only**: one forward graph per batch and sampling topology;
- **path-scoped batch-only**: one forward graph per batch, sampling topology,
and finite semantic topology path when kernel chains genuinely differ;
- **context capacity**: a capture-time storage and launch upper bound, not a
replay-time graph bucket;
- **static launch plan**: capture-time tile, warp, stage, split envelope, grid
envelope, compiled variant, and workspace capacity;
- **replay-before metadata**: dynamic state prepared outside the captured graph
before replay, also called graph-out preparation;
- **graph-in preparation**: fixed device work captured before operator forward;
- **stable graph state**: typed inputs, provider state, workspaces, wrappers,
and keepalive owners whose addresses and capacities remain fixed.

Reading `context_lens` or exposing `plan()` does not by itself violate
batch-only. The violation is allowing actual context to change graph identity,
captured topology, static launch plan, workspace shape, tensor addresses, or
provider binding.

## Operator and Provider Adaptation

- Define graph identity from batch capacity, finite semantic topology path,
sampling topology, and capture-time tensor/layout contract. Actual
`context_lens` must not enter graph keys or cause runtime capture.
- Resolve model/hardware tuning tables and compile-time choices before capture.
A table selected for a fixed model architecture and hardware combination is
valid static configuration. Tile, warp, stage, compiled variant, grid
envelope, and workspace shape are not replay metadata.
- Flag replay-time host thresholds that switch kernel chains, launch variants,
split envelopes, or workspaces. Replace them with a fixed envelope plus
device-side effective scheduling, bind another batch-only provider, or reject
the unsupported contract during resolution/preparation.
- Dynamic lengths may drive device masking, effective split/range metadata, or
an explicit replay-before provider plan when those updates write only stable
graph state and leave the captured launch contract unchanged.
- Permit separate startup-captured short/long paths only when the semantic
kernel chain truly differs. Merge methods or length regimes with identical
topology. Seal the startup plan; transitions among declared paths must not
JIT, reselect a provider or variant, grow workspace, or recapture.
- Require `supports(spec, caps)` and preparation to validate dtype, shape,
layout, capacity, padding, workspace, and batch-only compatibility before
forward. Do not treat a few fixed-shape experiments as production support.
- When a standard upstream provider already exposes a graph-stable lifecycle,
adapt that lifecycle instead of cloning its kernel. Use a repository-owned
fixed-grid provider for missing Sparse-vLLM semantics, portable fallback, or
an exact measured override—not as an automatic replacement for upstream.
- Fail unsupported capacity or layout before cache mutation. Once bound, do not
switch provider, allocate a larger workspace, or fall back after execution
begins.

## Unified Inputs and Participant Lifecycle

The unified input mechanism standardizes public replay inputs and update order;
it does not combine every tensor into one allocation or expose provider and
sparse-algorithm internals to the graph runner.

### Common input contract

Keep shared replay inputs in typed, fixed-address runner-owned state. At minimum
distinguish token ids, positions, context lengths, request indices, KV
write-slot mappings, and valid-row state. Every registered slot declares:

- shape, dtype, and device;
- batch axis and capacity;
- padding policy;
- semantic/value source and copy policy;
- stable-address requirement.

Prefer explicit `DecodeGraphInputs`-style fields. Flag an indefinitely growing
`dict[str, Tensor]`, an untyped memory blob, or a positional runner API carrying
method- and provider-private tensors.

### Ownership

For every field distinguish storage owner, semantic owner, and per-step value
producer:

- graph runner: common decode input storage, padding, capture/replay, and graph
identity;
- cache manager: physical KV storage, page/slot metadata, and physical cache
views;
- `SparseController`: logical sparse selection, cross-layer observation, and
attention-score coordination;
- provider: static kernel plan, schedule buffers, private graph state,
workspace, external wrapper, and physical weight/layout;
- model/attention layer: stable operator semantics only.

Do not move provider workspace into the common registry or physical cache
metadata into `SparseController`. The runner coordinates lifecycle and copy
order without taking ownership of private algorithms or layouts.

### Participant lifecycle

Use a typed lifecycle equivalent to:

```text
init_graph_state(contract, topology_path)
prepare_out_graph(step, state)
prepare_in_graph(state)
graph_keepalive_tensors(state)
```

- initialization allocates stable private buffers/workspaces, resolves the
static plan, initializes wrappers/JIT once, and records capacity;
- graph-out preparation updates dynamic host metadata or executes a documented
provider plan, writing only stable state;
- graph-in preparation contains fixed device work captured before forward;
- keepalive ownership prevents captured tensors, workspaces, wrappers, or
outputs from being released or replaced.

Coordinate provider preparation once before each model replay, outside
per-layer attention forward. Model and attention code consume prepared state and
must not contain sparse-method branches, provider names, external-wrapper
access, or graph lifecycle calls.

### Padding

Pad real batches to their capture bucket with an explicit active-row contract.
Padding rows use safe token, position, slot, page, and score metadata. Prove
they cannot access or mutate a live request's KV cache, sparse score, or
controller state. Do not rely on an incidental sentinel that a kernel still
dereferences before masking.

## External Graph Wrappers

For FlashInfer paged decode and comparable external providers with a public
CUDA Graph lifecycle:

- Use the upstream graph-enabled wrapper instead of an ordinary eager wrapper,
raw internal kernel, or repository reimplementation of its planner. Bind one
wrapper to each captured batch/topology state that needs distinct storage.
- Provider state owns fixed-capacity page indptr, page indices, last-page
lengths, integer/float workspaces, output owners, and the wrapper. The runner
invokes the participant lifecycle but never reads wrapper-private fields or
constructs provider-specific page metadata.
- Run context-dependent `plan()` or the documented fast-plan path during
replay-before preparation. Planning may change contents, not wrapper/workspace
identity, input/output addresses, launch contract, or captured `run()`
topology.
- Captured forward calls only the already-bound wrapper `run()`. Flag planning
in forward, wrapper recreation, real-length-driven `masked_select`, `cat`, or
allocation, workspace replacement, and runtime backend switching.
- Reuse persistent host/GPU staging. If the public API requires D2H or host
planning, keep the synchronization boundary explicit and report its p50/p95
cost separately; it must not alter captured addresses.
- If the minimum supported upstream version has no public wrapper contract that
satisfies these invariants, reject the provider for batch-only during binding.
Do not reach through private APIs or silently fall back after replay starts.
- Validate constructor, plan/fast-plan, and run with a real installation at the
declared minimum version. Mocks do not prove lifecycle compatibility.

## Required Review Evidence

For every claimed model/method/provider topology path require:

- one startup-captured graph per batch/topology/sampling state and no actual
context bucket in graph keys;
- repeated replay across representative, historical-threshold boundary,
ragged, padded, and maximum-capacity contexts;
- unchanged graph count, `recapture_count == 0`, and stable registered input,
workspace, output, and wrapper-owner addresses;
- no replay-time JIT, static plan/variant reselection, context-sized allocation,
workspace growth, provider switch, or per-layer host planning;
- independent numerical comparison for output and every required score, LSE,
cache mutation, or other side effect;
- padding and maximum-capacity memory-safety tests;
- real-model fixed and churn coverage, selected-provider/binding evidence, and
matched performance results;
- isolated timing for CPU metadata preparation, H2D/D2H, provider planning,
waits, and graph replay when replay-before work is nontrivial.

For semantic short/long paths, test the algorithm threshold below, at, and
above the boundary and transition between the already captured paths without
state loss or new capture. Do not preserve a historical kernel-tuning threshold
as a semantic topology path.

## Finding Severity Additions

- P1: a claimed batch-only path performs runtime recapture, changes captured
topology/variant from actual context, replaces captured addresses, switches
provider after binding, or lets padding access/mutate live request state.
- P2: a changed provider rejects batch-only cleanly but leaves a required
model/method without a production batch-only provider; replay preparation has
avoidable per-layer or allocation overhead; or new code extends only the
retired bucketed graph path without an explicit migration purpose.
- P3: terminology, binding-report, or ownership documentation is unclear while
behavior remains correct and observable.

Use the main skill's P0-P3 definitions for all other findings.
6 changes: 4 additions & 2 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -137,8 +137,10 @@ uv pip install flashinfer-cubin --index-url https://flashinfer.ai/whl

Use `cu129` instead of `cu130` for CUDA 12.9.

`einops`, `sglang-kernel`, and the training, benchmark, and test packages are all
part of the main installation; no workflow-specific extras are required.
`einops`, `sglang-kernel==0.4.5`, and the training, benchmark, and test packages
are all part of the main installation; no workflow-specific extras are required.
The SGL kernel package is pinned because its compiled operators must match the
validated PyTorch/CUDA ABI; other versions are rejected during provider setup.

Sparse-vLLM supports Qwen3.5/Qwen3.6/Qwen3.8 checkpoints in unquantized BF16
and block-scaled FP8 formats. These releases share the `qwen3_5` runtime
Expand Down
3 changes: 3 additions & 0 deletions README_zh.md
Original file line number Diff line number Diff line change
Expand Up @@ -119,6 +119,9 @@ uv pip install flashinfer-cubin --index-url https://flashinfer.ai/whl

CUDA 12.9 环境将 `cu130` 换成 `cu129`。

主依赖固定使用 `sglang-kernel==0.4.5`,因为其编译算子必须匹配已经验证的
PyTorch/CUDA ABI;其他版本会在 Provider 准备阶段明确失败,不会静默 fallback。

Sparse-vLLM 支持未量化 BF16 和 block-scaled FP8 格式的
Qwen3.5/Qwen3.6/Qwen3.8 checkpoint。三者共享 `qwen3_5` 运行时架构,以及
相同的精度、并行方式、稀疏方法和多模态支持。其 prefill causal Conv1D 和
Expand Down
38 changes: 38 additions & 0 deletions benchmark/efficiency/bench_probe.py
Original file line number Diff line number Diff line change
Expand Up @@ -163,6 +163,26 @@ def _percentile(values: list[float], quantile: float) -> float:
return ordered[lower] * (1.0 - weight) + ordered[upper] * weight


_DECODE_GRAPH_COUNTERS = (
"capture_count",
"replay_count",
"eager_static_count",
"force_eager_count",
"eviction_count",
"recapture_count",
)


def _decode_graph_counter_delta(
before: dict[str, Any],
after: dict[str, Any],
) -> dict[str, int]:
return {
name: int(after.get(name, 0)) - int(before.get(name, 0))
for name in _DECODE_GRAPH_COUNTERS
}


def _monitor_gpu_ids(explicit: str | None) -> list[int]:
value = explicit or os.environ.get("CUDA_VISIBLE_DEVICES", "")
if not value:
Expand Down Expand Up @@ -877,7 +897,12 @@ def run_sparsevllm_churn(
"[Sparse-vLLM Churn] Initializing "
f"method={args.sparse_method}, max_concurrency={concurrency}..."
)
engine_init_started = time.perf_counter()
llm = LLM(args.model_path, **engine_kwargs)
engine_init_s = time.perf_counter() - engine_init_started
startup_graph_summary = llm.debug_sparse_state_summaries()[0][
"decode_graph"
]
try:
request_count = concurrency * args.churn_request_multiplier
for p_len in args.prompt_lens:
Expand Down Expand Up @@ -917,6 +942,9 @@ def run_sparsevllm_churn(
try:
for iteration in range(args.num_iters):
profiler.reset()
graph_before = llm.debug_sparse_state_summaries()[0][
"decode_graph"
]
trace = _trace_for_iteration(
args,
model_specs,
Expand Down Expand Up @@ -979,6 +1007,9 @@ def run_sparsevllm_churn(
finished_times[seq_id] = now
generated_counts[seq_id] = len(token_ids)
elapsed_s = time.perf_counter() - started
graph_after = llm.debug_sparse_state_summaries()[0][
"decode_graph"
]

expected_seq_ids = set(seq_to_request)
for name, observed in (
Expand Down Expand Up @@ -1034,6 +1065,13 @@ def run_sparsevllm_churn(
"status": "success",
"elapsed_s": elapsed_s,
"step_count": step_count,
"engine_init_s": engine_init_s,
"startup_decode_cuda_graph": startup_graph_summary,
"decode_cuda_graph_before": graph_before,
"decode_cuda_graph_after": graph_after,
"decode_cuda_graph_counter_delta": (
_decode_graph_counter_delta(graph_before, graph_after)
),
"request_throughput_rps": request_count / elapsed_s,
"input_token_throughput_tps": total_input / elapsed_s,
"output_token_throughput_tps": total_output / elapsed_s,
Expand Down
26 changes: 13 additions & 13 deletions benchmark/efficiency/metrics_calculator.py
Original file line number Diff line number Diff line change
Expand Up @@ -154,22 +154,22 @@ def from_config_dict(cls, cfg: dict[str, Any], bytes_per_param: int = 2) -> Mode
"Model config must have positive num_attention_heads, got "
f"num_attention_heads={num_attention_heads}."
)
configured_head_dim = cfg.get("head_dim")
if configured_head_dim is None:
explicit_head_dim = cfg.get("head_dim")
if explicit_head_dim is None and cfg.get("qk_nope_head_dim") is not None:
explicit_head_dim = int(cfg["qk_nope_head_dim"]) + int(
cfg.get("qk_rope_head_dim", 0)
)
if explicit_head_dim is None:
if hidden_size % num_attention_heads != 0:
raise ValueError(
"Model config without an explicit head_dim requires "
"num_attention_heads to divide hidden_size, got "
f"hidden_size={hidden_size}, "
f"num_attention_heads={num_attention_heads}."
)
head_dim = hidden_size // num_attention_heads
else:
head_dim = int(configured_head_dim)
if head_dim <= 0:
raise ValueError(
f"Model config head_dim must be positive, got {head_dim}."
"Model config must define head_dim when hidden_size is not "
"divisible by num_attention_heads, got "
f"hidden_size={hidden_size}, num_attention_heads={num_attention_heads}."
)
explicit_head_dim = hidden_size // num_attention_heads
head_dim = int(explicit_head_dim)
if head_dim <= 0:
raise ValueError(f"Model config must have positive head_dim, got {head_dim}.")
vocab_size = int(cfg["vocab_size"])

# MoE parameters
Expand Down
6 changes: 6 additions & 0 deletions configs/debug/minimax_m2_tiny_random.json
Original file line number Diff line number Diff line change
@@ -0,0 +1,6 @@
{
"num_hidden_layers": 1,
"hidden_size": 3072,
"intermediate_size": 1536,
"max_position_embeddings": 4096
}
6 changes: 4 additions & 2 deletions docs/en/getting_started/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -34,8 +34,10 @@ MAX_JOBS=8 uv pip install flash-attn --no-build-isolation

Use `cu129` instead of `cu130` for CUDA 12.9.

`einops`, `sglang-kernel`, and the training, benchmark, and test packages are all
runtime dependencies, so workflow-specific extras are not required.
`einops`, `sglang-kernel==0.4.5`, and the training, benchmark, and test packages
are runtime dependencies, so workflow-specific extras are not required. The SGL
kernel package is pinned to the validated PyTorch/CUDA ABI; other versions fail
provider setup instead of falling back silently.

Sparse-vLLM supports Qwen3.5/Qwen3.6/Qwen3.8 checkpoints in unquantized BF16
and block-scaled FP8 formats. All three share the `qwen3_5` runtime architecture
Expand Down
Loading
Loading