Skip to content

Sparse decode attention: per-token reference lifecycle for async self-speculation - #157

Open
kunkunblueberry wants to merge 27 commits into
InferMatrixLab:mainfrom
kunkunblueberry:kunkun/rfc156-sparse-decode-async
Open

kunkunblueberry wants to merge 27 commits into
InferMatrixLab:mainfrom
kunkunblueberry:kunkun/rfc156-sparse-decode-async

Conversation

@kunkunblueberry

@kunkunblueberry kunkunblueberry commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

Status: draft — the per-key work is solved, the per-layer fixed cost is not

This branch is no longer blocked on anyone else's work. It carries the lifecycle contract, the numerical contract, and now both fused kernels. What it does not do is beat dense attention end to end, and the reason has changed twice during review, so the current one is worth stating precisely.

Per-layer, one attention call (32 rows, 24L/16H/128D, BF16, r=0.1, median of 10, dense = Triton):

ctx dense capture consume 4-loop dense 4-loop sparse ratio
512 2.55 5.51 3.02 10.20 14.11 1.38x
2048 2.82 6.14 3.49 11.28 15.93 1.41x
8192 3.84 11.19 7.43 15.36 29.90 1.95x
32768 13.10 34.09 21.98 52.41 91.15 1.74x

That is 3.0–3.4x faster than where this started, and the algorithm now behaves as advertised: capture reads the paged cache once instead of materialising an expanded FP32 row, and consume touches only the k retained keys instead of gathering K and V densely.

Engine level, tiny Ouro (2 layers), async self-speculation, prompt 32, max_tokens 16:

conc dense tok/s sparse r=1.0 tok/s match sparse r=0.1 tok/s match
1 68.15 12.99 1/1 12.95 0/1
8 398.27 38.82 8/8 58.31 2/8
32 1276.02 66.01 32/32 83.39 8/32

Full retention reproduces dense tokens exactly at every concurrency, which is the correctness result. Throughput is 5–19x worse, and the gap widens with concurrency, which is the diagnosis: the per-layer fixed cost dominates and it scales with rows, not with keys. With two layers and 48-token contexts there is no per-key work to amortise it against.

So the remaining work is a Python-overhead project, not an attention-kernel project: per-chunk index construction, per-row note_count, plan.groups()/_require_live walks, context_lengths.tolist() per layer, and q[row_ids]/block_tables[row_ids] gathers. A Triton launch in this environment is 0.015 ms and store views changed nothing measurable, so the cost is the Python around the kernels.

A meaningful serving number needs a trained Ouro-1.4B checkpoint (24 layers, long context), which this host does not have; the tiny model cannot answer that question.

Required human contributor answers

Answers below reuse the decision recorded by the human contributor @kunkunblueberry in issue #156, claim comment; they are referenced, not invented.

1. What outcome are you trying to achieve?

As recorded in the claim comment: contribute the asynchronous self-speculation integration for sparse decode attention, i.e. the lifecycle contract for per-token reference state across outstanding rounds.

2. What is the highest-leverage variable for that outcome?

The acceptance criteria of the RFC #156 §8 matrix cell for "Asynchronous self-speculation × Sparse decode attention": isolate references for outstanding rounds; GPU acceptance drives retirement, and allocation, conversion and invalidation dependencies across outstanding rounds. The state ownership rules of §5 (one slot per live token, incarnation after same-position replacement, reuse ordered after the last reader/writer) are the contract every later kernel change must respect.

3. How will this work move that variable rather than create busywork?

It implements that contract, plus the numerical contract that makes the result checkable: a runner-owned reference store, a per-row plan carried on the existing execution metadata (§4 "small named views rather than a new engine"), runner integration that reserves and retires slots per round, two fused kernels that make the arithmetic affordable, and focused tests for draft depths 1/2, rejection at every candidate position, same-position replacement, delayed completion and slot reuse.

4. Who is the customer or user, and what do they actually need?

The shared sparse decode attention feature and its reviewers: they need the lifetime/ownership layer to exist and be tested, and they need to know whether the feature can pay for itself. The second question now has a measured answer with a named remaining cost.

PR necessity and bugfix evidence

This is not a bugfix; it is a new RFC cell implementation. Scope follows the recorded claim in issue #156 and RFC §11 (cells contribute commits to the shared feature PR).

Purpose

Implements the asynchronous self-speculation × sparse decode attention cell of RFC #156:

  • attention/sparse.py — the §3.2 arithmetic contract (reference statistics + sparse correction, plain PyTorch, checked against the RFC's worked example), the per-row SparseDecodePlan named view, and the dispatch between the scalar oracle, the bitwise-exact batched path and the fused path.
  • worker/attention_reference.py — AttentionReferenceStore, the runner-owned per-token slot pool (§5): empty → source-ready → reusable → retired, one slot per live token, incarnation per (request, row) so a corrected token reusing a position rejects stale plans, device-side acceptance-driven invalidation with no host readback, metadata-only retirement, and host-side retained counts so a reader never syncs.
  • kernels/sparse_capture.py — one program per (row, kv head, position block) computes q @ K^T straight out of the paged cache: no expanded K, no per-row cuBLAS, one launch per layer.
  • kernels/sparse_consume.py — selected_scores and selected_values walk the retained indices and never materialise K or V.
  • kernels/sparse_decode.py — the opt-in fused implementation: fixed per-layer kernels for capture and consume with no host readbacks, selected with SparseDecodeAttentionConfig.implementation (batched by default). It changes consume-row numerics, so the bitwise contract stays with batched. The reference index buffer starts small (initial_key_capacity) and grows on the owner stream; the engine budgets the growth peak, not the initial allocation.
  • worker/attention_reference.py — retirement builds its index on the store's owner stream. Building it on the caller's stream let the host free and reuse that memory before the owner stream's index_fill_ ran, which surfaced as an index-out-of-bounds device assert once the host ran ahead under CUDA graphs.
  • worker/async_speculative.py — per round, the runner reserves one lease per verification row, publishes CAPTURE plans at depth index 1 and CONSUME plans at deeper loops, retires rejected candidates and their dependent descendants from the GPU acceptance counts, and releases slots on ticket retirement. Outstanding rounds hold disjoint slots.
  • config.py / entrypoints / stage_engine.py — the sparse_decode_attention group (enabled, default false; key_retention; key_bucket; max_context_tokens) with explicit capability checks and reference-state memory accounted before admission.

Numerical contract: dense and capture rows keep the selected backend's output, the actually rounded recurrence input is stored as O_ref, and full-retention consume rows use the same dense backend as the control. The fused capture pass computes the same FP32 dot product as the oracle with a different reduction order: max score delta 1.1e-5 on a magnitude of 65, retained mass delta 9e-8, attention output delta 2e-8, identical selected sets on 511/512 row-head pairs. batched_paged_attention without fused_scores still reproduces the scalar oracle bit for bit.

Explicitly out of scope: synchronous speculation, token pipelines, CUDA graph capture with sparse enabled, other matrix rows.

Test Plan

  • tests/test_sparse_decode_attention.py: plan and arithmetic contracts (RFC §3.2 worked example = 23.5, full retention ≡ dense, ceil retention rule, capacity), store lifecycle (one slot per live token, owner-checked release, acceptance-driven invalidation at accepted counts 0/1/2, root survival, stale-plan rejection after reincarnation, disjoint outstanding rounds, metadata-only retirement allocation bound, byte accounting, plan carriage), end-to-end decode through the real model + paged KV + attention, and engine-level full-retention parity.
  • tests/test_sparse_decode_batched.py (from @0z5a): the batched path must be scalar-exact.
  • tests/test_sparse_decode_sync.py (from @Levius-Fubuki): ordinary decode and synchronous speculation at draft depths 1/2.
  • benchmarks/sparse_decode_context_sweep.py: per-layer dense-baseline sweep and memory breakdown.
  • benchmarks/sparse_decode_engine_e2e.py: engine-level throughput, peak memory and token agreement across concurrency.

vllm-rlt Version / Commit: 88c8f96 (upstream/main, merged here) + this branch.

Test Result

A100 (SM80), torch 2.13.0+cu130:

  • CPU pytest tests/ -q: 981 passed, 652 skipped, 0 failed.
  • GPU (async speculative, async lifecycle, sync speculative, four sparse modules including the fused path): 219 passed, 4 skipped with -k "not flash_attn_4". The deselections are all flash_attn_4 (FlashAttention-4 is not supported on SM80).
  • Sparse modules alone on GPU: 87 passed.
  • Full-retention control, engine level: identical tokens at concurrency 1/8/32 on the tiny model, and 5904/5904 tokens with all exit depths on the real 1.4B model as reported by @Levius-Fubuki for the merged O_ref contract.
  • ruff check / ruff format --check clean on all touched files.

Not claimed: a serving speedup, TTFT under concurrency, or partial-retention task quality. The feature stays enabled=false by default.

…-speculation

Implements the asynchronous self-speculation cell of the sparse decode
attention column (RFC InferMatrixLab#156 section 8): the lifecycle contract for per-token
reference state across outstanding rounds.

- attention/sparse.py: the RFC 3.2 arithmetic contract and the per-row
  SparseDecodePlan named view (modes/slots/retention) that travels from the
  runner to the attention layer on the execution metadata.
- worker/attention_reference.py: AttentionReferenceStore, the runner-owned
  per-token slot pool. One slot per live token, empty -> source-ready ->
  reusable -> retired lifecycle, incarnation per (request, row) so a
  corrected token that reuses a position rejects stale plans, and
  acceptance-driven invalidation computed on the device without a host
  readback. Slots return to the pool only after the final reader completes
  (ticket retire).
- worker/async_speculative.py: the async speculative runner reserves one
  lease per verification row per round, publishes CAPTURE plans at depth 1
  (draft phase for draft depth >=2, verify phase for draft depth 1) and
  CONSUME plans at deeper loops, retires rejected candidates and their
  dependent descendants from the GPU acceptance counts, and releases slots
  on ticket retirement. Outstanding rounds hold disjoint slots.
- config/entrypoints/engine: the sparse_decode_attention group (enabled,
  key_retention, key_bucket), wired through LLM/CLI/StageEngine with
  explicit capability checks (async self-speculation, block mode) and
  reference-state memory accounted into the KV reserve.

The attention kernel consumption of the plan (dense source statistics and
the sparse correction) is the basic-implementation cell owned by
@Levius-Fubuki; the store publishes the slot views it will borrow.
@0z5a

0z5a commented Oct 27, 2026

Copy link
Copy Markdown
Contributor

https://github.com/Superone77/FlashLoop is an important reference, could you please add more e2e test result?

Copy link
Copy Markdown
Contributor

Hi @kunkunblueberry, I’ve reviewed #157 and would like to reuse its reference plan/store for the basic and synchronous paths. Before wiring attention into it, could we align on three points?

  1. Retained count. FlashLoop at aeaadee uses k = min(N, max(1, ceil(N * key_retention))), where N includes the current token. Sparse decode attention: per-token reference lifecycle for async self-speculation #157 currently rounds to nearest; for N=13, retention=0.1, these select 2 versus 1 key. I suggest following the reference’s ceil rule and adding boundary tests.

  2. key_bucket semantics. Could we separate buffer capacity from the algorithm’s retained count? Currently it caps both, so changing the bucket can change the algorithm, and retention=1 may cease to be a full-retention control. I suggest keeping k determined by retention, with capacity ≥ k and an explicit valid count/mask. Insufficient capacity should select a larger supported bucket or fail clearly, rather than silently truncate.

  3. State-pool capacity. Could we size key buffers from the supported serving context bound and retention, rather than always reserving model.max_position_embeddings entries per slot? With 24 layers, 16 heads and 65,536 entries, int32 indices alone cost 96 MiB per live-token slot. Slot count should separately cover the maximum simultaneously live input tokens—including verification rows and outstanding rounds—and include outputs, contributions, metadata and scratch in memory planning. The basic/synchronous paths can share this accounting with the async path.

If this sounds reasonable, I’ll build against the shared contract and cover retained-count boundaries, bucket overflow/full retention, and capacity accounting in the basic-path tests.

@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

https://github.com/Superone77/FlashLoop is an important reference, could you please add more e2e test result?

Thanks,i will check it and add it carefully.

…rate capacity

Aligns with the review on InferMatrixLab#157 and the FlashLoop reference (aeaadee
select_decode_keys):

- retained_count follows the reference rule k = min(N, max(1, ceil(N *
  key_retention))) instead of rounding to nearest; boundary tests include
  the N=13 / retention=0.1 case (2 keys, not 1).
- key_bucket is buffer capacity only, never an algorithm knob: the
  retained count is a pure function of key_retention, and a count above
  the capacity fails loudly instead of silently truncating retention.
- key_bucket=None sizes the index buffer as ceil(context x retention)
  rather than always max_position_embeddings, cutting per-slot index
  memory proportionally to retention (e.g. 24L/16H/65k at retention 0.1
  costs ~1/10 of the indices); explicit buckets below the worst-case
  count are rejected at startup.
- engine-level e2e: key buffer sizing, insufficient-bucket rejection,
  and a multi-request round with a mid-flight abort and full slot
  reclamation.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Thanks @Levius-Fubuki — this is exactly the alignment I was hoping for. All three points made sense, and I pushed ccf0aef adopting them:

  1. Retained count — retained_count now follows the FlashLoop rule verbatim, k = min(N, max(1, ceil(N * key_retention))) (verified against select_decode_keys at aeaadee), with boundary tests including your N=13, retention=0.1 → 2 example. One deliberate difference we keep: tie-breaking is stable by absolute position (RFC §9 "stable ties"), which plain torch.topk does not guarantee — happy to relax if you'd rather match the reference bit-for-bit.

  2. key_bucket semantics — agreed and separated. The retained count is now a pure function of key_retention; key_bucket is only the index-buffer capacity, with the store's existing per-(slot, layer, head) counts tensor serving as the explicit valid count. A count above the capacity now fails loudly (reference_statistics raises; the engine rejects key_bucket < ceil(context × retention) at startup) instead of silently truncating. Full retention needs no bucket headroom anymore.

  3. State-pool capacity — with key_bucket=None the index buffer is sized ceil(context × retention) instead of always max_position_embeddings, so at 24L/16H/65,536 with retention 0.1 the int32 indices cost ~1/10 of before; retention=1.0 honestly pays the full-retention control cost. Slot count separately covers the maximum simultaneously live tokens (verification rows × two outstanding rounds), and AttentionReferenceStore.bytes — indices, outputs, contributions and metadata — is added into the KV memory reserve at stage init. That property is the shared accounting entry point, so the basic/synchronous path can construct the same store and land in the same budget.

Please build against the shared contract: the store publishes the per-(slot, layer, head) selected/counts/mass/output/contribution views plus capture()/consume() for the attention side, and the per-row plan rides on DeviceKVMetadata.reference_plan. If the basic path wants the interface shaped differently (e.g. plan construction or capacity validation living elsewhere), say the word and I'll rework it here rather than fork the contract.

@Superone77

Superone77 commented Oct 28, 2026 •

Copy link
Copy Markdown

https://github.com/Superone77/FlashLoop is an important reference, could you please add more e2e test result?

For precision test, please use flashloop_torch implementation as gold result, thanks!

…dation

Makes the sparse decode column executable end to end so the optimization can
be measured instead of only asserted:

- attention/sparse.py: sparse_paged_attention, the RFC 3.2 reference path over
  the paged KV cache. Dense rows reproduce ordinary attention, capture rows
  publish (S, m, O_ref, C_ref) into the plan's store, consume rows correct the
  stored reference with the current loop's retained keys. Plain PyTorch, no
  backend dependency, so it is the numerical oracle for the optimized kernels.
- worker/attention_reference.py: write_reference/read_reference views for that
  path, with capacity enforced loudly.
- core/kv_cache_manager.py: a plan-carrying device metadata routes to the
  reference path; metadata without a plan keeps the existing backend dispatch
  unchanged.
- tests: end-to-end decode through the real model, paged cache and attention
  (full retention matches dense; divergence grows as retention shrinks; the
  consume path reads only the retained budget), plus an engine-level test that
  a sparse engine reproduces the dense engine's tokens exactly at retention 1.0.
- benchmarks/sparse_decode_attention_bench.py: reference-path measurement of
  the dense/capture/consume costs, keys touched, reference-state bytes and the
  same-loop identity vs later-loop divergence.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Done in db5ce44 — added the e2e evidence and a reference-path benchmark.

What I did: the lifecycle layer alone cannot show the optimization works, so I added sparse_paged_attention (in attention/sparse.py), the RFC §3.2 reference path over the paged KV cache: dense rows run ordinary attention, capture rows publish (S, m, O_ref, C_ref) into the plan's store, consume rows correct it with the current loop's retained keys. _attend_prepared routes to it only when the device metadata carries a plan; with sparse disabled the dispatch is byte-for-byte unchanged (full suite stays green).

How it validates:

  • model + paged KV + attention, end to end (CPU): full retention reproduces dense output (atol 1e-4), and the consume path is asserted to read only ceil(N * retention) keys per (token, layer, loop).
  • serving engine, end to end (GPU, async speculative + Triton): a sparse engine at retention=1.0 reproduces the dense engine's tokens 32/32 exactly; agreement then degrades monotonically with the knob — r=0.5 → 25/32, r=0.25 → 17/32.
  • benchmarks/sparse_decode_attention_bench.py, 24 layers / 16 heads / head_dim 128:
path 4-loop decode keys/row state/slot vs KV
GPU bf16, ctx 512, r=0.1 85.6 ms → 155.9 ms (0.55x) 512 → 52 5.7x
GPU bf16, ctx 2048, r=0.1 254.0 ms → 248.0 ms (1.02x) 2048 → 205 10.4x
CPU fp32, ctx 256, r=0.1 58.3 ms → 40.8 ms (1.43x) 256 → 26 4.4x

Honest reading: the algorithmic saving is real (10% of keys touched, consume cost scales with k), but the plain-PyTorch reference path is not a serving speedup on GPU — gather/argsort and per-step launch overhead eat it, and it only turns into wall-clock gain where launch overhead is negligible (CPU 1.43x). Per RFC §9 I am not claiming a runtime improvement from reduced work: the fused kernel (the basic-path cell) is where the GPU gain has to come from, and this reference path is the oracle to check it against. Same-loop O_ref - C_ref + m*C_new is an identity (7e-8); the approximation only appears across loops (1.2e-2 – 2.7e-2 at r=0.1). Reference state costs 4.4–10.4x the KV bytes per live token in fp32, which is a real budget item, not an optimization.

Regression: CPU pytest tests/ -q 891 passed; GPU async-speculative + lifecycle + sparse 113 passed / 1 skipped.

@Superone77

Copy link
Copy Markdown

Could you provide TTFT result? To evaluate the prefill. Thanks!

…tention

Prefill rows never take the decode correction (RFC InferMatrixLab#156 section 6.3), so the
feature must not move TTFT; this measures that plus the reference path's
inter-token cost against the dense Triton baseline.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

TTFT measurement added (benchmarks/ttft_bench.py, f53c35d) — medians of 5 sequential single-request decodes, tiny Ouro (2 layers, 128 ctx), A100.

prompt 64, max_tokens 32

case TTFT med TTFT min ITL e2e tok/s
dense 221.29 ms 219.31 ms 14.86 ms 682 ms 46.9
sparse r=1.0 222.59 ms (1.01x) 220.84 ms 123.06 ms 4040 ms 7.9
sparse r=0.5 222.40 ms (1.01x) 220.93 ms 114.02 ms 3757 ms 8.5

prompt 100, max_tokens 16 (prefill-dominated)

case TTFT med ITL e2e tok/s
dense 346.67 ms 15.33 ms 577 ms 27.7
sparse r=1.0 346.29 ms (1.00x) 119.62 ms 2140 ms 7.5
sparse r=0.5 346.76 ms (1.00x) 128.39 ms 2273 ms 7.0

Reading: TTFT is unchanged at 1.00–1.01x, i.e. within noise, which is what §6.3 predicts — prefill rows never take the decode correction, and the first token still comes from the prefill/coda path, so enabling the feature does not tax the latency users feel first. With a longer prompt the absolute TTFT doubles while the delta stays ~0.4 ms, so this is not a coincidence of a short prompt.

What it does cost today: ITL goes 15 ms → 120–128 ms (about 8x) because the reference path replaces the Triton kernel with plain PyTorch row loops (gather + argsort + per-step launch/sync). That is the reference implementation's cost, not the algorithm's — and it sets the bar for the fused kernel: it has to beat the 15 ms dense Triton ITL to be worth enabling. Note r=0.5 is not faster than r=1.0 here for the same reason: at this size the per-call overhead dominates the retained-count term, which is consistent with the standalone benchmark (GPU 0.55x / CPU 1.43x).

Until that kernel lands I would keep sparse_decode_attention.enabled=false by default (it already is) and treat these numbers as the reference baseline, not a serving configuration.

@0z5a 0z5a left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Requesting changes for the reproducible memory-capacity failure in the acceptance path. The current implementation can construct the engine and pass a short warmup, then OOM while processing an ordinary admitted concurrent batch. Please make acceptance/invalidation memory-bounded, account for the supported reference capacity before admitting work, and add coverage for the overlapping-round case below before merging.

The blocking inline finding is on AttentionReferenceStore.accept(). The performance observations and full-retention token divergence below are additional evidence: the token divergence has not been localized to a numerical or lifecycle cause, and this review does not require a plain-PyTorch reference to deliver an optimized-kernel speedup.

The sequential tiny-model TTFT observations in the latest comment also should not be generalized to concurrent serving. In the real-model B=64 case below, p95 engine first-token delivery changes from about 2.39 s (off) to 37.82/38.66 s (r=.1), even though prefill does not use sparse correction. These measurements include scheduling and output delivery and are not isolated prefill-kernel timings.


Supplementary real-model concurrency results for this PR.

I executed db5ce44642402102cc2a879196aac34d8a4026b5. While the campaign ran, HEAD advanced to f53c35d3bfa5b2412d4993018b1c07fd4bafd139; that commit only adds benchmarks/ttft_bench.py. The vllm_rlt/ and tests/ Git trees are identical between those two revisions. The new benchmark itself was not executed in this campaign.

Outcome: 10 measured cases completed, one measured case failed with CUDA OOM, and its repeat was skipped according to the frozen plan. The existing sparse test module passed all 28 tests, including its GPU tiny-model tests. The real-model results below expose additional capacity and output-equivalence limits.

Reproduction configuration

  • One NVIDIA GH200, 97871 MiB visible HBM, ARM CPU; driver 595.58.03, Python 3.12.3, PyTorch 2.14.0+cu130, Triton 3.8.0.
  • Native Ouro-1.4B, BF16 weights/activations/KV, unmodified max_position_embeddings=65536.
  • Model weight SHA256: 58872a72616c736595b8b7662079c5b12c5a162ec16eae94f21c348dfa9885af; config SHA256: ce9cc13da41591b8b4deca053d7dfee06424c0228628ee862ea86d725bc163f3.
  • Native StageEngine, not an HTTP server. For each batch size, enqueue all requests before calling step().
  • Requests i=0..B-1 use 256 deterministic prompt IDs: [1+(j*137+i*53)%49000 for j in range(256)].
  • Greedy sampling, max_tokens=16, max_loops=4, exit_threshold=1.0, ignore_eos=True.
  • Per case: fresh process/model/engine, same-concurrency 4-output-token warmup, then the 16-token measured batch. Startup and warmup are excluded from measurements. CPU/OMP threads=4; CPU affinity and GPU clocks were not pinned.
StageEngine(
    model,
    attention_backend="triton",
    cache_config=CacheConfig(
        num_blocks=2368 if B == 32 else 4672,
        block_size=16,
        incremental_allocation=True,
    ),
    scheduler_config=SchedulerConfig(
        max_num_seqs=B,
        max_num_batched_tokens=4096,
        prefill_chunk_size=256,
        admission_scan_limit=128,
    ),
    execution_config=ExecutionConfig(
        async_scheduling=True,
        cuda_graphs=False,
    ),
    speculative_config=SpeculativeConfig(3),  # draft 2 / target 4, block mode
    sparse_decode_attention_config=(
        None if mode == "off"
        else SparseDecodeAttentionConfig(enabled=True, key_retention=r)
    ),
)

key_bucket=None; prefix caching and preemption are disabled. The physical KV capacity covers all requests and four recurrent planes. Per batch size, the order was OFF1 / r=1 FULL1 / r=.1 SPARSE1 / SPARSE2 / FULL2 / OFF2. Each case had a 300-second wall limit including startup/warmup. A failed FULL1 was not repeated; independent configurations continued.

Measured observations

These are end-to-end native-engine output tokens/s, including enqueue, prefill, decode, output delivery and drain. They are not isolated attention-kernel timings. Both observations are retained because there is timing variability.

Requests Mode Output tokens/s, run 1 / 2 p95 engine TTFT seconds, run 1 / 2 Peak allocated / reserved GiB
32 Off 134.17 / 160.51 1.078 / 1.042 10.62 / 11.36
32 r=1 7.43 / 6.97 6.854 / 6.985 58.09 / 80.45
32 r=.1 14.45 / 10.44 5.339 / 8.070 14.89 / 17.46
64 Off 146.37 / 157.12 2.390 / 2.390 17.37 / 18.11
64 r=1 OOM / repeat not run — 82.94 allocated at OOM; another 18.00 requested
64 r=.1 9.40 / 9.94 37.819 / 38.664 26.45 / 35.02

Every successful case completed all B requests with exactly 16 output tokens. Both observed maximum resident requests and maximum requests in one speculative batch reached B=32/64. This was actual concurrency, not merely a larger pending queue. Sparse reference slots reached 256/256 at B=32 and 445/512 at B=64. Capture and consume path counters were nonzero. KV blocks and reference bindings drained to zero in every successful case, and same-mode repeats produced identical complete token/depth outputs.

The r=.1 observations do not show an acceleration on this short-context workload: throughput was approximately 6.3–6.4% of the paired dense controls at B=64. This is consistent with the PR's current scope as a plain PyTorch reference, not an optimized GPU serving kernel. Output changes also affect speculative acceptance, so the end-to-end ratio cannot isolate attention arithmetic cost.

Concrete B=64 full-retention failure

The 4-token warmup completed 64/64, but the 16-token measurement failed in AttentionReferenceStore.accept():

self.selected[slot_ids] = self.selected[slot_ids] * keep4.to(self.selected.dtype)

CUDA OOM requested 18.00 GiB, with 11.06 GiB free. PyTorch reported 82.94 GiB allocated and only 213.42 MiB reserved but unallocated. The last completed step observation had 64 resident requests and 0 completed measurement requests. Successful warmup is therefore not a successful measured case.

The reference pool alone is 24.09/48.19 GiB for B=32/64 at r=1, versus 2.49/4.99 GiB at r=.1. Pool sizing follows the model's 65,536-token limit even for a 256-token prompt, and the indexed assignment above introduces substantial transient allocation. No model-context reduction or allocator workaround was applied. Avoiding full-bucket indexed temporaries and making the service context capacity explicit are useful follow-up areas, but were not implemented or validated here.

Output equivalence

Full retention was not token-identical to dense on this larger BF16 workload:

Comparison to same-head dense Matching token positions Identical complete sequences
B=32, r=1, both repeats 503/512 31/32
B=32, r=.1, both repeats 259/512 10/32
B=64, r=.1, both repeats 450/1024 16/64

For r=1, request i=23 first differs at zero-based output position 7: dense token 30 versus full-retention token 2578. Both dense repeats are identical; both full-retention repeats are identical; exit depths match. This is reproducible autoregressive output divergence, not yet a localization of its numerical cause. No same-input logits-margin or FlashLoop-gold comparison was performed, so this observation alone does not identify an arithmetic or lifecycle bug.

These are synthetic-prompt generation comparisons, not benchmark task accuracy. Long contexts, long generation, HTTP serving, CUDA graphs and production performance remain outside this finite campaign.

All raw results, including the OOM, warmups, per-step traces, complete token/depth arrays and source/model/package provenance, were preserved in a 51-file archive and independently SHA-verified both remotely and after download. The original GPU lock was released, all owned children were joined, and the final GPU compute-process list was empty. Archive SHA256: f01759c21b454e936fc26bcc83ba39b172ccb2d3b202bf19e64f8d22df9ef8db.

Standalone native-engine reproduction harness used for the measurements

Run this against the pinned runtime revision with the local native Ouro-1.4B checkpoint above. For example, the failure case is python concurrency_probe.py --model /path/to/Ouro-1.4B --out b64-full1.json --batch 64 --retention 1.0 --context 256 --decode 16 --warmup 4. Use retention 0 for the off control or 0.1 for the sparse case. Run only within an allocated/exclusively owned GPU slot. The campaign wrapped each invocation in a 300-second timeout and isolated cases in separate processes.

"""Frozen native-engine concurrency probe; no HTTP or dataset accuracy claims."""
import argparse
import collections
import json
import math
import time
import traceback
from pathlib import Path

import torch
from vllm_rlt import CacheConfig, ExecutionConfig, SamplingParams, SchedulerConfig
from vllm_rlt import SparseDecodeAttentionConfig, SpeculativeConfig
from vllm_rlt.engine.stage_engine import StageEngine
from vllm_rlt.models import OuroForCausalLM
from vllm_rlt.request import RequestPhase
import vllm_rlt.core.kv_cache_manager as kv_module

p = argparse.ArgumentParser()
p.add_argument('--model', required=True)
p.add_argument('--out', required=True)
p.add_argument('--batch', type=int, required=True)
p.add_argument('--retention', type=float, required=True, help='0 means feature off')
p.add_argument('--context', type=int, default=256)
p.add_argument('--decode', type=int, default=16)
p.add_argument('--warmup', type=int, default=4)
a = p.parse_args()
torch.set_num_threads(4)
torch.manual_seed(157)
out = Path(a.out)
record = dict(status='STARTING', args=vars(a), head='db5ce44642402102cc2a879196aac34d8a4026b5',
              torch=torch.__version__, gpu=torch.cuda.get_device_name(),
              scope='native StageEngine synthetic prompts, not HTTP; timing includes prefill; fixed 4 loops')
def save():
    out.write_text(json.dumps(record, indent=2) + '\n')
save()
path_counts = collections.Counter()
original = kv_module.sparse_paged_attention
def counted(*args, **kwargs):
    plan = args[5]
    path_counts['calls'] += 1
    for mode, rows in plan.groups().items():
        path_counts[str(mode)] += len(rows)
    return original(*args, **kwargs)
kv_module.sparse_paged_attention = counted

def generate(engine, decode, label):
    path_counts.clear()
    start = time.perf_counter()
    done, timing, trace = {}, {}, []
    for i in range(a.batch):
        rid = str(i)
        prompt = [1 + (j * 137 + i * 53) % 49000 for j in range(a.context)]
        engine.add_request(rid, prompt, SamplingParams(max_tokens=decode, max_loops=4,
                           exit_threshold=1., temperature=0, ignore_eos=True))
        timing[rid] = dict(first=None, finish=None, chunks=[])
    max_resident = max_spec_batch = max_ref = max_blocks = 0
    with out.with_suffix('.' + label + '.jsonl').open('w') as stream:
        for step in range(10000):
            if not engine.has_unfinished_requests():
                break
            resident_before = len(engine.cache_manager._allocations)
            outputs = engine.step()
            now = time.perf_counter() - start
            resident = len(engine.cache_manager._allocations)
            max_resident = max(max_resident, resident_before, resident)
            schedule = engine.last_schedule
            phase = str(schedule.phase) if schedule else None
            scheduled = len(schedule.requests) if schedule else 0
            if schedule and schedule.phase == RequestPhase.SPECULATIVE:
                max_spec_batch = max(max_spec_batch, scheduled)
            ref = engine.reference_store.outstanding if engine.reference_store else 0
            max_ref = max(max_ref, ref)
            blocks = engine.cache_manager.num_used_blocks
            max_blocks = max(max_blocks, blocks)
            for item in outputs:
                t = timing[item.request_id]
                if item.token_ids:
                    if t['first'] is None:
                        t['first'] = now
                    t['chunks'].append([now, len(item.token_ids)])
                if item.finished:
                    if item.request_id in done:
                        raise AssertionError('duplicate completion')
                    done[item.request_id] = dict(tokens=list(item.token_ids), depths=list(item.exit_depths))
                    t['finish'] = now
            row = dict(step=step, seconds=now, phase=phase, scheduled=scheduled,
                       resident=resident, ref_slots=ref, used_blocks=blocks, completed=len(done))
            stream.write(json.dumps(row) + '\n')
            stream.flush()
        else:
            raise RuntimeError('bounded step limit')
    torch.cuda.synchronize()
    seconds = time.perf_counter() - start
    result = dict(seconds=seconds, tokens_s=sum(len(v['tokens']) for v in done.values())/seconds,
                  completed=len(done), outputs=done, timing=timing, steps=step,
                  max_resident=max_resident, max_speculative_batch=max_spec_batch,
                  max_reference_slots=max_ref, max_used_blocks=max_blocks,
                  drained_blocks=engine.cache_manager.num_used_blocks,
                  drained_references=engine.reference_store.outstanding if engine.reference_store else 0,
                  sparse_path_counts=dict(path_counts), peak_allocated=torch.cuda.max_memory_allocated(),
                  peak_reserved=torch.cuda.max_memory_reserved())
    assert len(done) == a.batch and all(len(v['tokens']) == decode for v in done.values()), 'incomplete batch'
    assert result['drained_blocks'] == result['drained_references'] == 0, 'leaked ownership'
    if a.retention:
        assert path_counts['1'] > 0 and path_counts['2'] > 0, 'sparse path did not execute'
    return result

engine = None
try:
    model = OuroForCausalLM.from_pretrained(a.model, device='cuda', dtype=torch.bfloat16)
    # Sufficient physical pages for all requests and all four recurrent planes.
    blocks = a.batch * math.ceil((a.context + a.decode + 4) / 16) * 4 + 64
    record['num_blocks'] = blocks
    record['model_max_position_embeddings'] = model.config.max_position_embeddings
    engine = StageEngine(model, attention_backend='triton',
        cache_config=CacheConfig(blocks, 16, incremental_allocation=True),
        scheduler_config=SchedulerConfig(max_num_seqs=a.batch, max_num_batched_tokens=4096,
                                         prefill_chunk_size=256, admission_scan_limit=128),
        execution_config=ExecutionConfig(async_scheduling=True, cuda_graphs=False),
        speculative_config=SpeculativeConfig(3),
        sparse_decode_attention_config=SparseDecodeAttentionConfig(enabled=True, key_retention=a.retention)
            if a.retention else None)
    record['reference_bytes'] = engine.reference_store.bytes if engine.reference_store else 0
    record['kv_bytes'] = engine.cache_manager.key_cache.nbytes + engine.cache_manager.value_cache.nbytes
    record['status'] = 'WARMUP'
    save()
    record['warmup_result'] = generate(engine, a.warmup, 'warmup')
    torch.cuda.synchronize()
    torch.cuda.reset_peak_memory_stats()
    record['status'] = 'MEASURING'
    save()
    record['result'] = generate(engine, a.decode, 'measure')
    record['status'] = 'COMPLETED'
except BaseException as exc:
    record['status'] = 'FAILED'
    record['error'] = repr(exc)
    record['traceback'] = traceback.format_exc()
    save()
    raise
finally:
    if engine is not None:
        engine.close()
        record['closed'] = True
    save()

The raw archive is retained locally; it is not publicly hosted by this review. The configuration, observed results, failure allocation sizes, source identity, and executable harness are included above so the failure can be independently reproduced without access to those private artifacts.

Comment thread vllm_rlt/worker/attention_reference.py Outdated
self.mass[slot_ids] = self.mass[slot_ids] * keep3
self.output[slot_ids] = self.output[slot_ids] * keep4
self.contribution[slot_ids] = self.contribution[slot_ids] * keep4
self.selected[slot_ids] = self.selected[slot_ids] * keep4.to(self.selected.dtype)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[P2] Avoid full-bucket temporaries when applying acceptance

The RHS materializes self.selected[slot_ids] and then another tensor for the masked result across the complete key_bucket, including unused positions. With native Ouro-1.4B BF16, its unmodified 65,536-token model capacity, max_num_seqs=64, K=3 async block speculation, r=1, 256-token prompts and 16 generated tokens, the engine constructs successfully and a four-token warmup completes, but this line raises CUDA OOM during the measured batch. The failed allocation is 18.00 GiB with only 11.06 GiB free (82.94 GiB already allocated by PyTorch); the last completed step has 64 resident requests and no completed measurement requests. The reference pool alone is 48.19 GiB, and the otherwise-matched dense cases complete 64/64. Please avoid temporaries proportional to the entire selected bucket here and make the configured memory budget cover the peak acceptance path (or reject an unsupported capacity before admitting requests). Add coverage that grows beyond the short warmup into overlapping rounds; complete reproduction and raw failure figures are in the review body.

…ence capacity

Addresses the review's reproducible OOM and capacity findings.

Retirement is memory-bounded (blocking finding):
- AttentionReferenceStore.accept() no longer reads-modifies-writes the retained
  index buffer. That indexed multiply materialised a bucket-sized temporary per
  batch (24 layers x 16 heads x 65,536 keys is tens of GiB), which is the OOM
  reported at B=64 full retention. Acceptance now clears only the per-row
  metadata that gates readers (valid, counts, mass) for rejected rows.
- _clear() likewise retires metadata only. Wiping the bucket on every release
  also cost tens of GiB of writes per round, which is part of why full
  retention measured 6.9 tokens/s.
- Readers gate on counts, and a capture overwrites the used prefix, so a stale
  payload is never observable.
- Regression test asserts retirement allocates less than an eighth of the
  bucket for a 512 MiB bucket (the old code allocated a full copy).

Capacity is explicit and accounted before admission:
- SparseDecodeAttentionConfig.max_context_tokens declares the supported serving
  context bound; the index buffer is sized from it instead of always from
  max_position_embeddings. At 24L/16H, B=32/64 and 512-token serving, the pool
  drops from 24.09/48.19 GiB to 0.28/0.56 GiB at retention 1.0.
- The engine now fails fast with an actionable message when the reference pool
  cannot fit in free device memory, and logs the resolved budget.
- Coverage for overlapping rounds: two submitted rounds of the same request must
  hold disjoint slots, and retiring the older round must not disturb the newer
  round's references.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Thanks — the OOM reproduction is solid and the finding is correct. Fixed in 048b9d7; both points are addressed at the root, not papered over.

Blocking finding: retirement is now memory-bounded. You are right that accept() was the problem, and the mechanism is what the trace shows: self.selected[slot_ids] * keep4 is a read-modify-write through advanced indexing, so it materialises a bucket-sized temporary per batch — 24L × 16H × 65,536 keys × 4 B is tens of GiB, which is the 18 GiB request against 11 GiB free at B=64. Acceptance now clears only the per-row metadata that gates readers (valid, counts, mass) for the rejected rows, with scalar fills that allocate nothing:

rejected = slot_ids[rows > limit]
if rejected.numel():
    self.valid[rejected] = False
    self.counts[rejected] = 0
    self.mass[rejected] = 0

While fixing it I found the same pattern in _clear(): wiping the bucket on every release did not allocate, but it wrote tens of GiB per round, which is a large part of why full retention measured 6.9 tokens/s. Retirement is now metadata-only there too. Readers gate on counts and a capture overwrites the used prefix, so a stale payload is never observable; the store docstring states that contract explicitly.

A regression test now asserts the property directly rather than the symptom: for a 512 MiB bucket, accept() + release() must allocate less than an eighth of it (test_accept_and_retire_do_not_allocate_bucket_sized_temporaries). The old code allocated a full copy and fails it.

Capacity is explicit and accounted before admission. New SparseDecodeAttentionConfig.max_context_tokens declares the supported serving bound, and the index buffer is sized from it instead of always from max_position_embeddings. With your configuration (24L/16H, B=32/64) the pool goes from the 24.09/48.19 GiB you measured to 0.28/0.56 GiB at r=1.0 with a 512-token bound (85x), and 0.11/0.23 GiB at r=0.1. The engine also fails fast with an actionable message when the pool cannot fit in free device memory, and logs the resolved budget, so an oversized configuration cannot reach a warmup-then-OOM path. Worth re-running your --batch 64 --retention 1.0 case with max_context_tokens=512 to confirm.

Overlapping rounds coverage added as requested: test_sparse_async_overlapping_rounds_keep_isolated_references submits two rounds of the same request before either is collected, asserts the rounds hold disjoint slots, and asserts that retiring the older round leaves the newer round's references bound and readable.

On the TTFT claim — you are right and I withdraw the generalization. What I measured was sequential, single-request, prefill-dominated latency; it says nothing about concurrent serving, and your p95 numbers (2.39 s → 37.8/38.7 s at B=64) are the relevant ones because they include scheduling and output delivery while the reference decode path is in the pipeline. The only defensible claim from my measurement is that prefill rows never take the correction (§6.3), not that enabling the feature preserves first-token latency under concurrency.

On the full-retention divergence — not localized, and I do not want to claim it is benign. Two data points I can add: on the tiny model the sparse engine matches dense exactly in both fp32 and bf16 (B=4 and B=8, 48-token prompts, 100% of positions), and the torch reference vs the Triton kernel agree to 2e-7 in fp32 and ~2e-9 in bf16 on identical paged inputs. That makes a structural arithmetic error unlikely but does not prove it, and it does not explain a 30 → 2578 flip at scale. The plausible cause is that the sparse path replaces Triton with fp32-accumulating PyTorch only for capture/consume rows, so hidden state is produced by mixed kernels across the four loops and small deltas compound; at 1.4B with sharper distributions a greedy argmax can flip on a near-tie. The localization you suggested is the right next step — capture the logit margin at the first divergent step for request 23 and compare against a same-input dense run — and I will do that before claiming full-retention equivalence at scale.

Defaults stay as they are: enabled=false, and the reference path is an oracle, not a serving configuration.

@0z5a

0z5a commented Oct 28, 2026

Copy link
Copy Markdown
Contributor

ok, I will add some high-concurrency e2e fix soon, just wait for my news.

@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

ok, I will add some high-concurrency e2e fix soon, just wait for my news.

Thanks,waitting for u.

@0z5a

0z5a commented Oct 28, 2026 •

Copy link
Copy Markdown
Contributor

@kunkunblueberry Please cherry-pick the high-concurrency follow-up from 0z5a/vllm-rlt#31. It lives in my fork and compares directly with your 048b9d7 head: only three incremental commits / two files, with no request to merge into upstream main.

git fetch https://github.com/0z5a/vllm-rlt.git codex/pr157-batched-attention
git cherry-pick e8f8261a447f778a280e20e25dec7d6ec3b4701f 90120f2116f1b09994b90b3a8578068d850e9e76 398333dd08fcb41e60d7a616846fb9d5738fa7d4

Please take all three in order. They batch equal-length gather/selection work, read consume counts once per layer, and prepare views once while retaining the scalar per-row cuBLAS shapes/strides. The first prototype's folded matmul batch changed BF16 rounding and generated tokens; the latter commits correct that regression. The rejected prototype's faster numbers are not used below. Your memory-bounded retirement fix is already in the baseline.

Completed GPU A/B results

  • A: 048b9d7c79a5711857ce3287812502ad8f274ad0 (your scalar reference after the retirement fix).
  • B: 398333dd08fcb41e60d7a616846fb9d5738fa7d4 (the complete cherry-pick range).
  • Output throughput in tok/s; primary r=0.1 entries are median (min–max) of three independent, unprofiled runs. Full-retention entries are single controls and are not three-repeat performance qualification.
Concurrency Retention Repetitions A: scalar tok/s B: candidate tok/s Improvement Tokens + exit depths vs A
32 0.1 3 each 10.40 (9.46–11.89) 43.93 (33.69–48.43) 4.23× / +322.6% All 3 exact
64 0.1 3 each 13.07 (11.54–13.47) 35.77 (35.05–38.26) 2.74× / +173.7% All 3 exact
32 1.0 1 each 8.43 26.08 3.09× / +209.4% Exact
64 1.0 1 each 6.30 26.74 4.24× / +324.4% Exact

B=32's three scalar observations were collected in the preceding validated phase. A separate fresh scalar drift check reached 12.15 tok/s; the candidate median is still 3.62× that observation. B=64 scalar/candidate repetitions were interleaved. Clocks and CPU affinity were not pinned, so the observed ranges are included rather than choosing the fastest run.

Correctness, memory and test status

  • 135 GPU tests passed, 1 skipped (the existing two-GPU reservation case). Final-head complete CPU pre-commit passed; the same-head GitHub lint/CPU check passed.
  • Real captured-input replay: FP32 bitwise exact, zero BF16 differences, for both capture and consume.
  • All 8 candidate model runs / 384 completed requests / 6,144 generated tokens matched their same-mode scalar oracle exactly, including exit depths. Actual resident/speculative concurrency reached 32/64; every run drained KV blocks and reference slots to zero. No OOM, failed request or timeout in these final unprofiled comparisons.
  • Peak allocated memory for full retention was 34.71 GiB at B=32 and 65.56 GiB at B=64, the same for A and B. The OOM correction belongs to your 048b9d7, not to this incremental patch.

Reproduction scope and remaining boundary

One GH200 (97,871 MiB HBM), real Ouro-1.4B BF16, PyTorch 2.14.0+cu130 / Triton 3.8, OMP=4, seed=157. Native StageEngine, Triton attention, async scheduling, default multistream, CUDA graphs off; 3 draft tokens / 2 draft loops / 4 target loops, block mode. Each fresh process warms up 4 output tokens at the target concurrency, then measures 256 input / 16 output tokens, greedy, ignore EOS, fixed four loops. Measurement includes enqueue, prefill, decode, output delivery, drain and CUDA synchronization.

Prompts for request i: [1 + (j * 137 + i * 53) % 49000 for j in range(256)]. Block size 16, KV blocks 2368/4672, max batched tokens 4096, prefill chunk 256, admission scan 128. key_bucket=None, max_context_tokens=None retains the original 65,536 context bound; no capacity reduction to 512 was used to obtain these gains. Weight SHA256: 58872a72616c736595b8b7662079c5b12c5a162ec16eae94f21c348dfa9885af.

These are short-burst native-engine measurements, not steady-state HTTP serving or task-quality results. The original full-retention-versus-dense discrepancy remains: current A matches dense on 503/512 tokens at B=32 and 1001/1024 at B=64. This patch preserves A exactly; it does not claim to resolve that separate numerical/quality issue.

Closeout completed: all four separate CPU+CUDA profiles (A/B × ops/full) passed, with matching layer-0 capture/consume windows (32/128 rows). Full captures contain readable shapes, memory and stack evidence. Capture CUDA launches fell 1,058 → 378; consume launches 3,588 → 441, while the corresponding bmm counts remained 96 / 256. Consume scalar readbacks fell 128 → 0 and stream synchronizations 130 → 9; capture synchronizations increased 3 → 15, so this is not a claim that all synchronization disappears. Acceptance launches stayed 46 → 46. These are work counts from separate profiles, not the timings used in the throughput table.

The complete output archive passed an independent local check of 126 files, and all 12 nested profile bundles passed their internal SHA256 inventories. No owned GPU children remain and the original GPU lock was released. Archive SHA256: d7d29e7eb6285acbe4cd3f77a4501c10c5ff4f2eb37e11c172022c01493782a7. The fork PR is ready for review/cherry-pick.

Copy link
Copy Markdown
Contributor

@kunkunblueberry — progress update for my attention/backend, ordinary decode and synchronous self-speculation work:

Draft collaboration PR: kunkunblueberry/vllm-rlt#1 targets kunkun/rfc156-sparse-decode-async, based on 048b9d7, with head fbd8dea (runtime/test fix: 1bbe6fd).

It reuses the shared reference plan/store, integrates ordinary synchronous decode and synchronous block speculation at draft depths 1/2, and covers per-input reference ownership, rejection/replacement, cancellation, capacity checks and stream-safe retirement.

Validation completed

  • Full pre-commit passed: 921 CPU passed, 479 skipped.
  • Applicable CUDA regression: 318 passed, 1 skipped, 16 FA4 deselected. The skip requires two GPUs; FA4 is unqualified on the tested SM89 device. These are local/reserved-server results.
  • Real Ouro-1.4B BF16, fixed four loops, ordinary + sync draft 1/2, arrival concurrency 1/8/32, retention off/1/.1: 27 configurations × 3 bursts = 81 bursts, 1,107 completed requests and 17,712 output tokens. Every run drained KV/reference state, and outputs/exits repeated exactly within each condition. Workload: 128 input / 16 output tokens, explicit sparse context bound 512.
  • Full-retention controls matched 5,904/5,904 dense tokens and all exit depths. The original C32/draft-1 case now matches 512/512 in all three repetitions, compared with 498/512 before the fix.

Numerical contract and remaining quality work

The original discrepancy was reproduced and traced to attention backend rounding: identical Q/K/V at the watched loop-2 physical layer produce a small BF16 difference between Triton and the Torch reference, which later recurrence amplifies. Frozen-input replay reproduces both outputs.

The fix preserves the selected backend for dense/capture outputs and stores the actual rounded recurrence input as O_ref. Retention 1 explicitly uses the same dense backend as a control, with reference lifecycle validation retained. The independent Torch paged capture/correction oracle remains separately tested, including full retention.

Partial-retention task-quality validation is still pending. Retention .1 completes the synthetic execution tests, but this does not establish task accuracy. Remaining work includes final-path flashloop_torch gold comparisons, real BF16 verifier-versus-serial checks under the same sparse policy, paired dense/sparse task evaluation, and long-context/perplexity/retrieval checks with a predeclared quality budget. Sparse stays disabled by default; no deployment-quality or serving-speed improvement is claimed.

Detailed diagnosis and validation report · Reproduction scripts, frozen tensors, raw results and SHA256 inventory. The complete archive's 322 files were independently verified locally, and GitHub asset digests match.

Could you please review the collaboration PR—especially the O_ref rounding and full-retention control contract—and merge it into the feature branch for #157 once acceptable, keeping partial-retention task-quality validation tracked as open work? The separate 0z5a/vllm-rlt#31 optimization is not included; their combined revision will need fresh regression validation.

kunkunblueberry and others added 2 commits October 29, 2026 13:18
…tion

Brings in the basic/synchronous cell from @Levius-Fubuki (PR #1 in this fork),
based on 048b9d7:
- ordinary decode and synchronous block speculation at draft depths 1/2,
  per-input reference ownership, rejection/replacement, cancellation,
  capacity checks and stream-safe retirement;
- numerical contract fix: dense/capture outputs keep the selected backend and
  the actually rounded recurrence input is stored as O_ref; retention 1.0 uses
  the same dense backend as the control, which removes the full-retention
  token divergence traced to Triton-vs-Torch BF16 rounding.
@0z5a

0z5a commented Oct 29, 2026

Copy link
Copy Markdown
Contributor

Thanks for the sweep and for converting to draft. I agree that this PR cannot meet the performance gate as scoped. Below are three end-to-end results that complement the kernel-level sweep, plus one point where I disagree: the lifecycle layer is not yet free of host readbacks, and I would fix that before freezing the store/plan surface.

All runs use the same engine settings for dense and sparse: async self-speculation with 3 draft tokens, fixed 4 loops, Triton, greedy, Ouro-1.4B BF16, key_retention=0.1.

1. The store still synchronizes with the host

With sparse enabled, a batch-1 decode step performed about 683 host synchronizations, versus 5.5 for dense. About 235 of them come from AttentionReferenceStore, although its docstring says nothing there reads device state back:

  • torch.tensor(list, device=...) in reserve, build_plan and release: pageable host-to-device copies wait for the stream.
  • plan.slot_ids[rows] with a Python list in capture, consume and bind_positions.
  • accept: slot_ids[rows > limit] followed by rejected.numel() reads the rejected count back.
  • reserve writes self.incarnation[slot] = count per row, and _clear / invalidate assign Python scalars through an index tensor; each of these copies a host scalar and blocks.

All of these can be fixed without changing any value written:

  • pinned non-blocking index copies (the idiom kv_cache_manager already uses);
  • reusing plan.slot_ids when a call covers the whole plan;
  • masked writes in accept instead of the boolean gather;
  • index_copy_ / index_fill_ for the scalar writes.

With the copy and acceptance changes, GSM8K tokens and exit depths were identical to e4f8988 on all 87 questions. The scalar-write change writes the same values and passes the GPU lifecycle suites. After all of them, the store's reserve/invalidate/release path performs no synchronization. The wall-time effect at batch 1 is small (about 1%). The point is the contract the fused kernel will build on, so I would land this before treating the surface as frozen.

2. A short-context quality data point for the third gate

Setup: the fixed 87-question GSM8K CoT subset the repo already uses (lm-eval strict-match), batch 1, on an RTX 5090.

dense sparse (e4f8988)
correct 59 60
engine steps 3229 3267
wall time 270 s 480 s

Accuracy holds at r=0.1 on this short-context set. Long-context retrieval is still missing, as you said.

3. Selection inside the kernel reaches parity end-to-end, not a win

I prototyped an opt-in fused path locally, with the default left as the current batched path and its bitwise contract:

  • Capture is the dense backend call plus one Triton kernel. That kernel scores the keys and selects the retained ones by bisection on an order-preserving key. It writes them in ascending positions with the oracle's tie rule, then publishes (S, m, O_ref, C_ref).
  • Consume is one Triton kernel.
  • There are no host readbacks. O_ref stays the backend's rounded output.

Gates were fixed before the runs, and dense and fused were interleaved three times on an RTX 5090:

dense fused
GSM8K correct (each of 3 runs) 59 60
wall time, mean of 3 264.9 s 276.6 s (1.044×; dense alone spans 254–277 s)
engine steps 3229 3235
peak allocated 4.79 GiB 4.79 GiB
32 × 256 synthetic, median tok/s (engine steps) 290 (14) 323 (13)

Two things limit the end-to-end gain. These may help size the follow-up kernel.

  • The eager engine is host-bound. On GH200 at 32 × 256 and 4 × 4096, the dense attention time that consume could replace is about 22 ms of 5.5 s and at most 34 ms of 8.2 s (about 0.4%). Once selection is fused, the per-step cost matches dense. The synthetic lead above comes from one fewer step on degenerate outputs, not from cheaper steps.
  • The speculative pass structure makes the kernel budget stricter than S + 2C < 2D. In each round, depths 0 and 1 run once per draft offset on B rows (8 passes), and capture happens at depth 1 for every offset. Only depths 2 and 3 run on 4B verify rows. Dense already reuses each request's KV in L2 across its 4 verify rows. In a kernel microbenchmark on an RTX 5090 at 8k–64k, the 4-row call cost about 1.1–1.6× the 1-row call, and 2.6–3.1× less than 4 rows with private KV. Per step, break-even is therefore roughly 4·S(B) + 2·C(4B) < 2·D(4B) ≈ 2.6·D(B). Even with consume free, the selector has to cost under about 0.65 of one dense call on the same rows.

4. Memory

Agreed that max_context_tokens is the right declaration. As an alternative that keeps the full 65,536 logical bound, I also tried growing the index pool on demand. On GH200 at 32 × 256, the reference state drops from 2.68 GB to 0.13 GB and the peak from 13.11 to 10.73 GiB (dense: 10.62), with tokens unchanged.

The commits are local; I can push them as a branch to build on if that is useful: the store sync fixes, the index-pool growth, and the fused prototype with its tests.

The capture step was CPU-bound and paid for an expanded FP32 copy of the whole
causal row. Measured at 32 rows / 24L / 16H / 128D / BF16, r=0.1, capture went
from 59.6 ms to 20.9 ms at 8k context and from 221.0 ms to 43.9 ms at 32k.

- kernels/sparse_capture.py: one program per (row, kv head, position block)
  computes q @ K^T for that KV head's query group straight out of the paged
  cache, writing FP32 scores. No expanded K, no per-row cuBLAS, one launch per
  layer. Accepts a 5-D cache or the per-layer view the cache manager hands the
  attention layer. Against the oracle at 8k: max score delta 1.1e-5 on a
  magnitude of 65, retained mass delta 9e-8, attention output delta 2e-8, and
  identical selected sets on 511/512 row-head pairs.
- attention/sparse.py: the capture branch takes that pass when Triton is
  available and a dense backend is present; the scalar oracle remains the
  definition of the arithmetic, and batched_paged_attention without
  fused_scores still reproduces it bit for bit (0z5a's test is unchanged).
- Chunk budgets follow the data each path actually materialises. The old 16M
  element budget was below a single 8k expanded row, so a group ran one row at
  a time: 32 small launches plus a host sync each. The fused path budgets the
  score matrix instead and keeps the whole group in one chunk.
- The capture contribution gathers retained values straight from the paged
  cache rather than from the expanded FP32 copy. BF16 -> FP32 is exact, so the
  bits are identical while the read no longer scatters through a row-length
  tensor.
- benchmarks/sparse_decode_context_sweep.py measures the production entry point
  instead of one implementation, so the sweep follows the dispatched path.

Consume is now the bottleneck and is untouched here: about 10 ms of per-layer
fixed cost (including a 2.45 ms host readback of the retained count) plus a
gather that costs 16 ms at 32k, which is why the four-loop total is still
2.5-4x dense even though capture alone is now cheaper than dense at 512.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Levius's basic path landed and is merged here (efe5bc7), so I took the next step on this branch: the fused selection pass. 25e15d7.

What the profile actually said

I had reported "the selector costs 25x per key". That was directionally right but the mechanism was wrong, and the correction matters:

  • capture's GPU time was only ~11 ms/call at 8k; the 98 ms was CPU — chunk = max(1, 16M // (extent * heads * dim)) fell below a single 8k expanded row, so a group ran one row per iteration: 32 launches, each with a host sync, each materialising a full expanded FP32 K row.
  • the retained-value gather read out of that expanded FP32 copy rather than the cache.

The fix

kernels/sparse_capture.py computes q @ K^T for a KV head's query group straight out of the paged cache: one program per (row, kv head, position block), FP32 scores out, no expanded K, no per-row cuBLAS, one launch per layer. The capture branch uses it when Triton and a dense backend are present; batched_paged_attention without fused_scores still reproduces the scalar oracle bit for bit, so @0z5a's exactness test is unchanged and still passes.

Numerics against the oracle at 8k: max score delta 1.1e-5 on a magnitude of 65, retained mass delta 9e-8, full attention output delta 2e-8, identical selected sets on 511/512 row-head pairs (the one difference is a near-tie). That is a new numerical reference, not a bitwise-identical one — same trade Levius made for O_ref, and it needs the same tolerance treatment in the acceptance tests.

Measured (32 rows, 24L/16H/128D, BF16, r=0.1, median of 10, dense = Triton)

ctx dense capture before capture now consume 4-loop dense 4-loop sparse ratio
512 2.49 13.52 1.97 12.84 9.94 30.13 3.03x
1024 2.65 17.91 17.91 12.95 10.61 46.46 4.38x
2048 2.80 20.48 18.41 13.26 11.20 47.83 4.27x
8192 3.84 59.55 20.88 17.28 15.34 59.28 3.86x
32768 13.08 221.00 43.85 36.15 52.31 129.23 2.47x

Capture is 2.9x faster at 8k and 5.0x at 32k, and at 512 it is now cheaper than dense attention itself — the algorithmic claim ("selecting 10% of keys can cost less than attending to all of them") holds once selection is fused. Total four-loop is still 2.5–4.3x dense because consume has not been touched.

Note the small-context rows: capture is ~18 ms flat from 1k to 4k while dense is 2.6–3.1 ms, and the 512 measurement (1.97 ms, below the 2.49 ms dense call it contains) is inside run-to-run noise. Those contexts are dominated by fixed per-layer overhead, not by the score pass.

Consume is the remaining bottleneck, and it is all overhead

At 512 with k=52 keys — a workload whose GPU work is trivial — consume still costs 12.8 ms per layer. Bisected:

  • store.counts[...].amax(...).tolist(): 2.45 ms, a host readback per layer. RFC §6.2 forbids exactly this; the retained count is host-known at capture time and should never be read back.
  • store.consume / store.capture: 3.0 / 5.7 ms for setting a handful of metadata words (the index is a Python list, so each call builds a CPU tensor and copies it to the device).
  • per-row bmm loops: 2.3 ms; the k-key K/V gathers: 0.6 ms at 512 but 16.3 ms at 32k.

Fixing the first two removes ~8 ms of the ~10 ms fixed cost at every context; that alone puts the four-loop total near or below dense at 512–2k. The 32k gather needs a layout fix (indexing [blocks, block_size, kv, dim] with three advanced indices runs at ~45 GB/s).

State of the PR

Still draft, enabled=false by default, and I am not claiming a serving win yet: the sweep is one layer at 32 rows, not a serving run, and consume alone currently costs more than dense. Next, in order: kill the consume readback and metadata overhead, then re-run this sweep and the engine-level parity before asking for another review.

Brings in InferMatrixLab#159 (scheduler wavefront selection / prefill progress) and InferMatrixLab#160
(native CUDA SiLU-and-Mul) so the branch is measured against current main.
…onsume

Consume was ~10 ms of fixed cost per layer at every context, so a 52-key
workload cost more than dense attention. Three measured causes:

- the retained count was read back from the device with
  counts[...].amax(...).tolist(), a host sync every layer that RFC InferMatrixLab#156
  section 6.2 forbids outright. The capture decides the count before it
  submits, so AttentionReferenceStore now keeps it host-side
  (note_count/host_count) and the reader sizes its gather without a readback.
- writing the source-ready flag used a Python scalar into a bool tensor through
  advanced indexing: 2.6 ms per call at 32 rows. Assigning a tensor instead is
  ~0.13 ms.
- torch.tensor(rows, device=cuda) built a device tensor from a Python list on
  every call, a blocking copy worth 2.4 ms. Grouped rows are contiguous, so
  they are derived with a device-side arange.

Host counts are dropped when a row is rejected, invalidated or retired, so a
stale count can never size a gather.

Measured at 32 rows / 24L / 16H / 128D / BF16, r=0.1, median of 10, dense
Triton baseline: four-loop sparse went from 306.9 ms to 100.8 ms at 32k
context (2.23x -> 1.92x dense), from 97.9 ms to 39.6 ms at 8k (6.39x -> 2.57x),
and from 41.7 ms to 26.4 ms at 512 (4.06x -> 2.57x).
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Consume is fixed and the branch is rebased on current main. db65440 on top of merge 61ac441 (upstream 88c8f96, i.e. #159 + #160).

Three measured causes, all of them overhead

At 512 context with k=52 keys, consume cost 12.8 ms per layer for work whose GPU time is negligible. Bisected:

  1. A host readback every layer. counts[slots, layer].amax(dim=-1).tolist() — 2.45 ms, and RFC §6.2 forbids it outright. The capture decides the retained count before it submits, so it is host-known by construction: AttentionReferenceStore now keeps it host-side (note_count / host_count) and the reader sizes its gather without touching the device. Host counts are dropped on rejection, invalidation and retirement, so a stale count can never size a gather.
  2. valid[...] = True — a Python scalar into a bool tensor through advanced indexing costs 2.6 ms at 32 rows, versus ~0.13 ms for a tensor assignment. Same for uses[...] += 1? No — that one was already 0.13 ms; the bool scalar store is the anomaly.
  3. torch.tensor(rows, device="cuda") — 2.4 ms per call to build a 32-element index from a Python list. Grouped rows are contiguous, so they come from a device-side arange now.

Measured (32 rows, 24L/16H/128D, BF16, r=0.1, median of 10, dense = Triton)

ctx dense capture consume 4-loop dense 4-loop sparse ratio this morning
512 2.57 12.88 5.47 10.28 26.39 2.57x 4.06x
8192 3.85 16.02 9.86 15.40 39.59 2.57x 6.39x
32768 13.10 38.96 24.36 52.40 100.78 1.92x 5.86x

Consume alone: 12.8 → 5.5 ms at 512, 17.3 → 9.9 ms at 8k, 36.2 → 24.4 ms at 32k. Together with the fused capture pass, the four-loop total is 3.0x faster than this morning at 32k (306.9 → 100.8 ms) and 2.5x at 8k (97.9 → 39.6 ms). Still 1.9–2.6x slower than dense, so the feature remains opt-in and unqualified for serving.

What is left, and it is now structural rather than silly

  • Capture carries ~12 ms of fixed cost per layer independent of context (12.9 at 512 vs 39.0 at 32k). Part is Triton's Python launch path (~3 ms per launch in this environment, against ~0.05 ms of actual work at 512), part is the bool-flag write and the argsort. A pre-bound kernel handle and a uint8 flag would take a few ms off every context.
  • Consume is one pass over k keys but does it as gather-then-matmul: 8.2 ms of K gather plus 8.1 ms of V gather at 32k (k=3277), through three advanced indices on a strided [blocks, block_size, kv, dim] at ~45 GB/s. That is the last big item, and the fix is the same shape as the capture fix: one fused kernel that walks the selected indices, accumulates p * V and writes the correction, with no materialised K/V.

Doing both would put the four-loop total at roughly capture ≈ dense + 2 ms and consume ≈ 2 ms at 32k, i.e. clearly below dense for the first time. I will take those two next unless you would rather I stop here and hand the numbers over.

Regression on the rebased branch: CPU pytest tests/ -q 969 passed, 614 skipped; GPU (async speculative, lifecycle, sync speculative, three sparse modules) 303 passed, 1 skipped, 13 deselected (FA4 on SM80).

@0z5a

0z5a commented Oct 31, 2026

Copy link
Copy Markdown
Contributor

Thanks for the two rounds today, and for keeping the sweep as the single harness. On your closing question: I already have both of your next items on a local branch. One is a selection with no fixed per-layer cost. The other is a fused consume that walks the selected indices without materialising K/V. I ran them against your head so we don't build the same kernels twice.

What the local branch does

The branch is your head db65440 merged with my fused path. The fused path is opt-in through SparseDecodeAttentionConfig(implementation="fused"), so the default batched path and its oracle contract are unchanged.

  • Capture: a copy of the dense Triton split kernel that also keeps the scores, so O_ref is bit-identical to the dense output.
    • The first radix histogram is computed in that kernel. Selection then runs as a 4×8-bit radix select over 16 parts per (row, head) and writes positions only.
    • The retained contribution is gathered over the compact list, and the last finishing part closes it with a ticket. Nothing is read back to the host.
  • Consume: split-K over the retained ranks with a fixed-order merge, one pass over the k retained keys.

Your sweep, both implementations on one GPU

I ran benchmarks/sparse_decode_context_sweep.py unchanged, with the store forced to batched (your head) or fused. Setup: RTX PRO 6000 Blackwell, 32 rows, 24L/16H/128D, BF16, r=0.1, median of 10, two repeats, kv_heads=4. Each cell is four-loop sparse time divided by four-loop dense time:

ctx head (batched) fused
512 16.3x / 15.6x 1.72x / 1.66x
2048 7.77x / 7.95x 1.11x / 1.11x
8192 4.02x / 4.00x 0.73x / 0.72x
16384 4.01x / 4.00x 0.68x / 0.68x
32768 4.09x / 4.11x 0.66x / 0.66x
  • At 32k, capture drops from 6.38 to 1.44 ms and consume from 4.99 to 0.155 ms, against 1.07 ms for one dense call.
  • kv_heads=16 gives the same picture: 0.69–0.72x at 8k and above.
  • All 32 rows of the sweep share one block table, so most of the KV sits in L2. That flatters the absolute numbers for both paths.

Eager GSM8K-87, batch 1, about 1k context

Three interleaved pairs on an RTX 5090, Ouro-1.4B, async self-speculation K=3, r=0.1:

  • At first, fused took 1.117x dense wall time (per pair 1.10 / 1.11 / 1.14), with 61 vs 59 correct.
  • Per step that was +10.3%, and all of it was host time. A fused capture is six Triton launches, about 112 µs of host time per call, against two launches (about 34 µs) for a dense call. Six launches is the minimum for this design, because each selection pass depends on the previous pass's global state.
  • So the launches themselves now skip Triton's Python dispatch once a specialization is compiled. This is limited to Triton 3.8, and the key mirrors Triton's specialization rules.
    • In the tests, every reused launch asserts that it is the kernel Triton's own dispatch would pick.
    • That check caught one real case: a positional tl.constexpr argument (table_stride) that Triton specializes on by value.
    • A separate test compares outputs and store state bit for bit against normal dispatch.
  • I also trimmed the call path's own Python. Launchers are looked up by kernel identity, since hashing a JITFunction recomputes its cache key. Arguments are no longer repacked, and the scratch lookup uses the raw stream query.
  • Capture host time went from 112 to 70 µs and consume from 39 to 25 µs, against 33 µs for a dense call.
  • The rerun fused GSM8K-87 at 1.064x dense (per pair 1.060 / 1.060 / 1.074), 61 vs 59 correct.

So in eager mode the fused path pays off from roughly 8k context. At short context with batch 1 it is about 6% slower than dense.

Two notes on db65440

  • _clear, invalidate and accept now call .tolist() on device slot tensors. Each call is a device-to-host sync on every retire or reject, which is the same RFC §6.2 class this commit removes elsewhere. release already holds the host list, so the host counts can be dropped from it directly.
  • In the merge I kept the readback-free masked accept. To preserve your invariant that a stale count never sizes a gather, accept now drops the host counts of every row of the round, using the host-known lease.slots. A later reader then sizes from the device counts, which are zero for rejected rows.

Proposal

The eager scope is ready as one commit on top of db65440: eeb07a0 on branch 0z5a:pr157-fused-eager. It leaves the graph and P/D files untouched, passes the CPU suite and ruff, and passes the GPU tests with the launch check on. I can open it as a PR against your branch, so your sweep becomes the acceptance check, or you can cherry-pick it yourself. Whichever is easier for you.

The same kernels also replay inside CUDA graphs, where the per-launch host cost disappears. I measured that end to end separately: Ouro-1.4B, 512 greedy tokens, interleaved dense/sparse engines. At 2 × 32k and 4 × 16k, throughput is +6% to +10% over dense at r=0.1/0.2. At 2 × 14.6k on an RTX 5090 it is +3% to +8% at r=0.1. At batch 1 it is parity. Graph replay and P/D wiring would come as their own PRs afterwards.

… store

Consume was the last item costing more than dense attention. Two changes:

- kernels/sparse_consume.py: selected_scores walks the retained indices and
  writes one FP32 score per (row, head, retained), selected_values accumulates
  p * V into one output row. Neither materialises K or V, which removes the two
  8 ms gathers and the per-row cuBLAS calls the Torch path made at 32k.
- The store no longer builds a device tensor from a Python list. capture,
  consume and bind_positions all did plan.slot_ids[rows], a blocking
  host-to-device copy worth 2.4 ms per call, and fused_capture_scores read the
  context bound back with .item() twice per layer, a stream sync that RFC InferMatrixLab#156
  section 6.2 forbids. The runner now passes the context and raggedness it
  already knows, and the slot index is derived on the device when the rows are
  contiguous.
- The capture and consume paths read and write the store through views when the
  chunk's slots are contiguous, falling back to the indexed form otherwise.

Measured at 32 rows / 24L / 16H / 128D / BF16, r=0.1, median of 10, dense Triton
baseline. Four-loop sparse vs dense:

  ctx      before today   after fused capture   now
  512        4.06x              2.57x          1.38x
  2048       4.43x              2.07x          1.41x
  8192       6.39x              2.57x          1.95x
  32768      5.86x              1.92x          1.74x

Capture alone is 221.0 -> 34.1 ms at 32k and 13.5 -> 5.5 ms at 512; consume is
36.2 -> 22.0 ms and 12.8 -> 3.0 ms. The remaining gap is per-layer fixed cost,
not per-key work.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Fused consume is in, and the whole column is now within 1.4–2.0x of dense instead of 4–6x. 4802af6.

Two changes

kernels/sparse_consume.py — the consume step only touches k retained keys, but the Torch path gathered them into a dense [rows, heads, k, dim] tensor twice (K, then V) and ran one cuBLAS call per row: 8.2 + 8.1 ms of gather at 32k through three advanced indices on a strided cache. Now selected_scores walks the retained indices and writes one FP32 score per (row, head, retained), and selected_values accumulates p * V into one output row. Neither materialises K or V; the softmax in between stays in Torch because k is small.

No host round-trips left in the store — three of them, all measurable:

  • plan.slot_ids[rows] in capture, consume and bind_positions built a device tensor from a Python list: 2.4 ms per call, a blocking host-to-device copy. Rows in a group are contiguous, so the index is derived on the device.
  • fused_capture_scores read the context bound back with .item() twice per layer — a stream sync, and exactly what RFC §6.2 forbids. The runner already knows the context and whether the rows are ragged, so it passes them.
  • store reads and writes go through views when the chunk's slots are contiguous, with the indexed form as the fallback.

Measured (32 rows, 24L/16H/128D, BF16, r=0.1, median of 10, dense = Triton)

ctx dense capture consume 4-loop dense 4-loop sparse ratio start of day
512 2.55 5.51 3.02 10.20 14.11 1.38x 4.06x
2048 2.82 6.14 3.49 11.28 15.93 1.41x 4.43x
8192 3.84 11.19 7.43 15.36 29.90 1.95x 6.39x
32768 13.10 34.09 21.98 52.41 91.15 1.74x 5.86x

Per step: capture 221.0 → 34.1 ms at 32k and 13.5 → 5.5 ms at 512; consume 36.2 → 22.0 ms and 12.8 → 3.0 ms. Cumulative over the day the four-loop total is 3.4x faster at 32k (306.9 → 91.2 ms), 3.3x at 8k, and 3.0x at 512.

What is left

The gap is now per-layer fixed cost, not per-key work. At 512 with k = 52 the GPU work is nothing, yet capture costs 5.5 ms against a 2.6 ms dense call and consume 3.0 ms. I ruled out the obvious suspects: a Triton launch in this environment is 0.015 ms, not the ~3 ms I assumed, and switching the store accesses to views changed nothing measurable. What remains is the Python per-chunk work (index construction, q[row_ids], block_tables[row_ids], the corrected-output scatter) plus the stable argsort, which is 5.9 ms of the 34 ms at 32k.

Closing the last 1.4x needs either a single kernel that does capture end to end (scores, selection, contribution, store write) or a CUDA-graph capture of the per-layer Python sequence — both are larger changes than anything in this branch so far, and the second one interacts with the graph support that is currently disabled for sparse.

Regression: CPU pytest tests/ -q 969 passed, 614 skipped; GPU (async speculative, lifecycle, sync speculative, three sparse modules) 303 passed, 1 skipped, 13 deselected. ruff check clean.

Next unless you would rather I stop: engine-level end-to-end (real Ouro-1.4B, concurrency 1/8/32) to confirm the per-layer gain survives scheduling, and updating the PR description, which still describes the branch as blocked.

Runs the real StageEngine with asynchronous self-speculation and compares dense
against sparse decode at full and reduced retention: throughput, peak memory and
per-request token agreement, across concurrency 1/8/32.

The result is a correctness pass and a performance diagnosis. Full retention
reproduces dense tokens exactly at every concurrency (1/1, 8/8, 32/32), so the
merged O_ref contract holds through scheduling, sampling and delivery. Throughput
is 5-19x worse on the tiny model and the gap widens with concurrency, which shows
the remaining cost is per-layer Python fixed cost scaling with rows, not per-key
attention work: with two layers and 48-token contexts there is nothing to
amortise it against. A serving answer needs a trained Ouro-1.4B checkpoint, which
this host does not have.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Engine-level e2e is in (fe83818, benchmarks/sparse_decode_engine_e2e.py) and the PR description is rewritten — it no longer says "blocked", it says what is measured and what is left.

Correctness: clean

Full retention reproduces the dense engine's tokens exactly at every concurrency: 1/1, 8/8, 32/32. That means the merged O_ref contract survives scheduling, sampling and delivery, not just a single attention call.

Throughput on tiny Ouro (2 layers, async self-speculation, prompt 32, max_tokens 16, median of 3)

conc dense tok/s sparse r=1.0 match sparse r=0.1 match peak GiB dense / r=1.0 / r=0.1
1 68.15 12.99 1/1 12.95 0/1 0.07 / 0.13 / 0.16
8 398.27 38.82 8/8 58.31 2/8 0.17 / 0.20 / 0.21
32 1276.02 66.01 32/32 83.39 8/32 0.07 / 0.13 / 0.17

5–19x slower, and the gap widens with concurrency. That is the useful part of the result: the per-layer sweep says 1.4–2.0x at 32 rows, but the engine says the cost grows with rows while the per-key work does not. With two layers and 48-token contexts there is nothing to amortise the fixed cost against, and dense batching keeps getting better while the sparse path adds per-row Python work.

So the remaining work is Python overhead, not attention math

Concretely, per layer: context_lengths.tolist() (a sync), the grouped-dict rebuild, q[row_ids] / block_tables[row_ids] gathers, plan.groups() and _require_live walking 32 slots in Python (2–3 times per call), per-row note_count, the corrected-output scatter, and the stable argsort (5.9 ms of the 34 ms at 32k). I ruled out the two candidates that would have been nicer: a Triton launch here is 0.015 ms, and switching store accesses to contiguous views changed nothing measurable.

Two ways out, both larger than anything in this branch: precompute the per-layer plan and chunk indices once per submission instead of per layer (the runner already knows every field), or capture the per-layer Python sequence in a CUDA graph — which currently conflicts with sparse being excluded from graph capture.

What I am not claiming

These are tiny-model numbers on a randomly initialised 2-layer architecture. They cannot answer whether 24 layers at 8k–32k amortise the fixed cost enough to win; the per-layer table suggests it is close (1.74x at 32k with 24 layers), but that needs a trained Ouro-1.4B checkpoint and this host has none. I would rather say that than extrapolate.

Regression unchanged: CPU 969 passed / 614 skipped, GPU 303 passed / 1 skipped / 13 deselected.

@0z5a

0z5a commented Nov 1, 2026

Copy link
Copy Markdown
Contributor

Thanks for 4802af6 and fe83818. Your diagnosis of where the remaining time goes matches mine: per-layer host work that scales with rows, not per-key attention work. Two things from my side that might save you a round.

Trained Ouro-1.4B through your e2e harness

I ran the structure of benchmarks/sparse_decode_engine_e2e.py with a trained Ouro-1.4B instead of the tiny model, on one GH200 (BF16, async self-speculation K=3, greedy, 256 tokens per request, eager). The prompts are real text, and three arms run on one build: dense, your default batched path at r=0.1, and the fused path at r=0.1. The build is your fe83818 merged with my fused path (see below).

context × concurrency dense tok/s batched (head) fused
1k × 1 5.36 0.60x 0.88x
1k × 8 59.7 0.19x 0.50x
8k × 1 4.81 0.57x 0.77x

So in eager mode, neither sparse path beats dense at these sizes, even with the real model. That is consistent with your tiny-model result.

On this host, eager decoding is itself dominated by host time. Dense with CUDA graphs is 322 tok/s at 1k × 8, against 59.7 in eager. That points at the second way out you mention.

With CUDA graphs, the same harness on the same GH200 puts the fused path at parity with dense: 0.92x at 1k × 1, 1.00x at 1k × 8, 0.98x at 8k × 1, and 1.02x at 8k × 4. The wall time includes prefill, which is the same for both arms. GH200 is the hardest case for sparse; the kernel breakdown below shows why.

CUDA graphs are where it pays off

Capturing the per-layer sequence in CUDA graphs works on a local branch. Selection counts are decided on the device, so there is no .tolist() per layer, and a static plan has its slot ids refilled before each replay. Measured end to end with Ouro-1.4B, 512 greedy tokens, interleaved dense/sparse engines, r=0.1:

setting GPU sparse vs dense
B=2 × 32k, block speculation RTX PRO 6000 +9.3%
B=4 × 16k, block speculation RTX PRO 6000 +9.6%
B=2 × 14.6k, block speculation RTX 5090 +3% to +8%
B=1 × 32k, block speculation L40S +7.3%
B=2 × 14.6k, P/D 1P1D (sparse on the decode stage) 2× L40S +6.5% / +8.2%
B=2 × 22.5k, token pipeline vs dense block L40S +13.5%

The gain shrinks on a GH200: block speculation is +7.7% at B=2 × 32k, and a token pipeline loses 1.3% at 14.6k. An nsys breakdown of the pipeline case (106 decode steps, graphs):

  • Attention is 47% of GPU time for both arms (641 ms dense, 652 ms sparse).
  • Sparse saves about 141 ms of dense attention on the consume rows, but consume plus selection add about 162 ms.
  • In a pipeline only the verify rows consume, about 22% of the attention work. On a GH200 that work is cheap enough (about 1.3 ms per step) that the fixed cost of consume and selection outweighs it. On lower-bandwidth GPUs the same rows cost more, which is where the gains above come from.

At batch 1 with short contexts, it is parity.

Merging with your head

My eager-scope fused path merges cleanly onto fe83818: ea2c58c on branch 0z5a:pr157-fused-eager-v2. The only conflicts were in the batched consume loop and the store's slot helpers:

  • The batched path keeps your fused consume.
  • The store ends up with one helper: the whole plan reuses its device slot ids, contiguous rows use a device arange, and other rows use a pinned copy.
  • The GPU suite passes on a GH200.

One numerical note: on sm_90, the split capture rounds FP32 outputs up to 1.8e-7 from the dense kernel. BF16 is bit for bit, and the FP32 test now allows one ulp.

If it helps, I can open the eager part against your branch as one commit, and the graph / P/D / token-pipeline part as a stacked follow-up.

@0z5a

0z5a commented Nov 1, 2026 •

Copy link
Copy Markdown
Contributor

@kunkunblueberry I think this PR finally has good result, and can be opened for review after such a long iteration.

@0z5a

0z5a commented Nov 1, 2026

Copy link
Copy Markdown
Contributor

Please resolve the conflicts

Token-id prompts drive a trained model into a fixed point, so every
implementation agrees and the comparison says nothing. This encodes real text
with the checkpoint's tokenizer and reports throughput, peak memory and
per-request token agreement against the dense engine.

Measured on Ouro-1.4B (24 layers, 16 KV heads, 200-token prompt, 16 tokens out,
median of 2, A100): concurrency 1 dense 4.77 tok/s, sparse r=1.0 3.74, r=0.1
3.87; concurrency 8 dense 42.75, r=1.0 17.27, r=0.1 22.01. Full retention
reproduces dense tokens exactly at both concurrencies.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Downloaded the trained checkpoint and re-ran the e2e on it. e128c04 adds the benchmark; the numbers below are the first ones from a real model.

Setup

ByteDance/Ouro-1.4B, revision 574fa66c, sha256 verified against the LFS digest (58872a72…), 24 layers, 16 attention heads, 16 KV heads (no GQA), 65,536 context, 4 UT steps, 1.435B parameters. A100, BF16, Triton attention, async self-speculation with 3 draft tokens, 200-token prompt from the checkpoint's own tokenizer, 16 tokens out, median of 2, 768 KV blocks.

Results

conc dense tok/s sparse r=1.0 match sparse r=0.1 match peak GiB (dense / r=1.0 / r=0.1)
1 4.77 3.74 (0.78x) 1/1 3.87 (0.81x) 1/1 7.21 / 9.51 / 14.05
8 42.75 17.27 (0.40x) 8/8 22.01 (0.51x) 8/8 18.60 / 14.24 / 18.79

This is much better than the tiny model, and it confirms the shape of the remaining cost: at concurrency 1 the feature is 20% slower than dense, not 5x, because 24 layers amortise the per-layer fixed cost. At concurrency 8 it is 2–2.5x slower, because that cost scales with rows while the dense Triton kernel keeps batching better. That is the same diagnosis as before, now measured on the model that matters.

Two findings worth recording

1. Token agreement is a weak quality signal, and I nearly reported a false pass. The first run used synthetic token-id prompts and the model collapsed to a fixed point (token 0 sixteen times), so every implementation "matched". After switching to real text I instrumented the retained probability mass:

retention k of 200 retained mass (mean / min) tokens vs dense
1.0 200 1.000000 / 1.000000 identical
0.1 51 0.920 / 0.773 identical
0.02 10 0.828 / 0.556 identical

So at 2% retention the correction discards 17% of the attention mass on average (44% at the worst row) and the 16 greedy tokens are still identical. The sparse path is definitely executing (capture 360, consume 192 calls per run) and the mass is definitely changing, so the tokens are simply robust over 16 steps. Any quality gate that compares tokens on a short generation will pass trivially; it has to look at retained mass, logit margins, or task accuracy over a real cohort.

2. A corrupted checkpoint looks like a broken implementation. My first download assembled to the right byte count but the wrong sha256, and the loader accepted it (it validates keys and shapes, not hashes). The model then emitted token 0 forever, which reads exactly like a broken attention path. Worth adding a hash check to the load path, or at least documenting that shape validation is not integrity validation.

What I am still not claiming

Peak memory at concurrency 8 is noisy across configurations (dense 18.60 vs r=1.0 14.24 GiB) because the KV pool and the reference pool are sized differently per run, so I would not read a memory conclusion from that column yet. Partial-retention task quality is still unmeasured: the mass table above is a mechanism check, not an accuracy result.

Adds --model-path, --cache-blocks and --sparse-context so the same harness can
run against Ouro-1.4B. The KV pool has to be sized to the workload: 4096 blocks
at 24 layers, 16 heads and 128 dims is 51 GiB of K and V, which OOMs; 768 blocks
fits a 32-request, 200-token burst.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

@0z5a conflicts resolved: 5c4aee6 merges upstream/main at 989921a (v0.2.0) into the branch, and the PR is MERGEABLE / CLEAN again. CI is green — lint and CPU unit tests passed in 2m12s (run 41328262422). The only conflict was README.md, where v0.2.0 restructured the feature list; I took upstream's structure and re-added the sparse-decode bullet under it, so docs/runtime.md#sparse-decode-attention-reference-path still resolves.

I reproduced the CI steps locally first (ruff 0.13.2 check + format --check on the whole tree, then OMP_NUM_THREADS=1 pytest -q): 969 passed, 614 skipped.

On your data — thank you, this is the piece I could not get. My e2e topped out at 0.78x at concurrency 1 and 0.40x at 8 on the A100, and I had no way to test the graph hypothesis from here. Your numbers make the ordering unambiguous:

  • eager is host-bound on every host: dense 59.7 tok/s at 1k × 8 against 322 with graphs;
  • with graphs the fused path reaches parity (0.92x at 1k × 1, 1.00x at 1k × 8, 0.98x at 8k × 1, 1.02x at 8k × 4);
  • the gains appear where attention is expensive per byte (RTX PRO 6000 +9.3/+9.6%, L40S +7.3%, token pipeline +13.5%) and vanish on GH200 (+7.7% block, −1.3% pipeline), which the nsys breakdown explains: 141 ms saved on consume rows, 162 ms added by consume plus selection, and only ~22% of attention work in a pipeline.

That also settles the argument about what this PR is: the per-layer work is done, and the remaining 2.5x at high concurrency is host time that graphs remove. Yes please open the eager part (ea2c58c) against kunkun/rfc156-sparse-decode-async as one commit — I will review and merge it here. The graph / P-D / token-pipeline part as a stacked follow-up is the right shape too, and given the GH200 result I would label it "GPU-dependent, not a universal speedup" in the description rather than quoting the L40S numbers alone.

Two notes on your conflicts description: I want the store to end up exactly as you describe — one helper, whole-plan device slot ids when they are contiguous, device arange for contiguous rows, pinned copy otherwise — that is strictly better than my _row_slots. And please include the sm_90 FP32 tolerance change (1.8e-7, one ulp) in that same commit; it belongs with the code that needs it, not in the graph follow-up.

One finding from my side that affects how we report quality: I instrumented the retained probability mass, and token agreement is a weak signal. At r=0.02 on Ouro-1.4B (10 of 200 keys) the correction discards 17% of attention mass on average and 44% on the worst row, yet the 16 greedy tokens are still bit-identical to dense, at every concurrency. Any quality gate that compares short generations will pass trivially. The mass table is in my previous comment; if your graph campaign can log retained mass per step, the two data sets would line up.

@0z5a

0z5a commented Nov 1, 2026

Copy link
Copy Markdown
Contributor

Opened as one commit against your branch: kunkunblueberry#2 (fb475f8 on 5c4aee6). It includes the single store helper and the sm_90 FP32 one-ulp tolerance. The graph / P-D / token-pipeline part will follow stacked on top. Its description will say the speedup depends on the GPU, with the GH200 rows next to the others.

On retained mass: agreed that short-generation token agreement says little. I'll log the retained mass per step in the graph runs, so the two data sets can be lined up.

@kunkunblueberry
kunkunblueberry marked this pull request as ready for review November 2, 2026 05:25
0z5a added 2 commits November 2, 2026 11:28
SparseDecodeAttentionConfig(implementation="fused") routes uniform capture or
consume submissions to Triton kernels with no per-layer host readbacks; the
default batched path and its oracle contract are unchanged, and it keeps the
fused consume from 4802af6.

- Capture is the dense Triton split kernel that also keeps the scores, so
  O_ref matches the dense output (bit for bit in BF16; within one ulp in FP32
  on sm_90). It builds the first radix histogram per part; selection is a
  4 x 8-bit radix select over 16 parts per (row, head), each pass closed in a
  fixed order by the last part to finish (integer tickets), writing positions
  only. The retained contribution is gathered over the compact list. Backends
  other than Triton keep their call plus one serial selection kernel.
- Consume is split-K over the retained ranks with a fixed-order merge. Its
  per-stream scratch follows rows, heads and the padded head dimension.
- Retained counts are decided on the device; a batched consume after a fused
  capture reads them back once.
- The store writes its metadata without host-scalar copies: one row-slot
  helper (the whole plan reuses its device slot ids, contiguous rows use a
  device arange, other rows a pinned copy), index_fill_ retirement, a
  readback-free masked accept that drops the round's host counts, and index
  storage that grows on demand within the logical bucket (the engine budgets
  the growth peak).
- Reused launches call the compiled Triton 3.8 kernel directly, skipping its
  Python dispatch. The key function is generated per argument-type layout and
  mirrors Triton's specialization; the tests assert on every reuse that it is
  the kernel Triton would pick.

Measured with benchmarks/sparse_decode_context_sweep.py (32 rows, r=0.1,
RTX PRO 6000): four-loop sparse/dense 0.66-0.73x at >= 8k, 1.1x at 2k,
1.6-1.7x at 512. Eager GSM8K-87 at batch 1 (RTX 5090): 1.064x dense wall time,
61 vs 59 correct.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
A speculative ticket releases its leases from the host's current stream, but
the store clears slot metadata on its owner stream. The slot index was copied
and allocated on the current stream and freed when release returned, so a
later allocation on that stream could overwrite it before the owner stream
ran index_fill_, clearing other slots or indexing out of range. With CUDA
graphs the host runs far enough ahead for this to surface as an index
out-of-bounds device assert. The index is now built inside the owner stream.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
@0z5a

0z5a commented Nov 2, 2026

Copy link
Copy Markdown
Contributor

Thanks for marking this ready, and for getting CI green on 5c4aee6. Before I change my review to an approval, could you merge kunkunblueberry#2 into this branch? There are two reasons.

  1. d6c52eb fixes a bug that is on this head.
    • AttentionReferenceStore.release builds the slot index on the caller's current stream (vllm_rlt/worker/attention_reference.py:171).
    • _clear reads that index on the store's owner stream (:385-386), but the index is freed as soon as release returns. The current stream can reuse that memory before the owner stream's indexing runs.
    • In eager mode this is rare. Under CUDA graphs the host runs far enough ahead that it shows up as an index-out-of-bounds device assert.
    • The commit builds the index on the owner stream and adds test_release_from_another_stream_builds_its_index_on_the_owner_stream.
  2. 21c2814 is the opt-in fused path that both our measurements used. Without it, this PR ships only the batched path. In eager mode on Ouro-1.4B, that path runs at 0.19–0.60x of dense in my GH200 runs and 0.40–0.81x in your A100 runs.

Once it is merged and CI is green, I'll approve with the scope as it stands:

  • sparse decode stays opt-in, and the default dense path is unchanged;
  • any speedup depends on the GPU and comes from the CUDA graph / P-D / token-pipeline follow-up, which I'll open stacked on this branch after it lands.

Brings in #2:
- kernels/sparse_decode.py: opt-in fused capture/consume kernels, the path both
  measurement campaigns used;
- attention_reference.py: build the retirement index on the store's owner
  stream, which removes a use-after-free that only showed up under CUDA graphs.
@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Merged: 43d9cc9 brings kunkunblueberry/vllm-rlt#2 into the branch (that PR is auto-marked MERGED), and CI is green again — lint and CPU unit tests passed in 2m34s (run 41454588044). The PR is MERGEABLE / CLEAN.

Both commits are in, and I read them rather than just taking the merge:

d6c52eb (the race). Confirmed on this head before merging: release built the slot index on the caller's current stream while _clear consumed it on the owner stream, and nothing kept the index alive across that gap. Eager mode rarely gets far enough ahead to hit it; under graphs the host runs ahead by design, which is exactly when an index-out-of-bounds assert would appear. Building it inside with torch.cuda.stream(self.stream) is the right fix, and the test is the right shape.

21c2814 (the fused path). implementation="batched" by default is the correct call: it keeps the bitwise contract on the path that has it and makes the numerically different path opt-in. Three details I want to keep in the record:

  • the reference index buffer now starts at initial_key_capacity=64 and grows on the owner stream, with peak_capacity_bytes budgeting every in-flight generation — that is a better answer to the memory gate than my "size it to ceil(context × retention)" approach, because it stops a short-prompt deployment from paying for a long one;
  • the mixed-submission fallback in batched_paged_attention (device readback only when a fused capture left no host count) is the right compromise: no sync on the common path, correctness on the rare one;
  • _row_slots now returns the whole plan's device slots when the rows cover the plan, which is strictly better than my contiguous-arange branch.

Verified locally on the merged head, in CI's own terms: ruff check . clean, ruff format --check . 139 files already formatted (ruff 0.13.2, the version CI pins), OMP_NUM_THREADS=1 pytest -q 981 passed, 652 skipped. GPU, with -k "not flash_attn_4": 219 passed, 4 skipped across the four sparse modules (including test_sparse_decode_fused.py), async speculative and the lifecycle suite.

The description now documents the implementation knob, the growth accounting and the stream fix, so a reviewer does not have to read the diff to learn that batched is the default.

Ready for your approval. After it lands I will take the stacked graph / P-D / token-pipeline follow-up, and I will label it GPU-dependent with the GH200 numbers next to the L40S ones rather than quoting the best case alone.

@0z5a 0z5a left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for merging kunkunblueberry#2. 43d9cc9 has the same tree as d6c52eb, and CI is green, so I'm switching my review to approve.

Scope, as agreed:

  • Sparse decode stays opt-in (SparseDecodeAttentionConfig), and the default dense path is unchanged.
  • In eager end-to-end runs it does not beat dense yet. The gains we measured come with CUDA graphs, P/D and token pipelines, and they depend on the GPU. That part comes in the stacked follow-up I'll open on top of this branch.

@kunkunblueberry

Copy link
Copy Markdown
Contributor Author

Before this goes to approval: thank you for staying with this one. The review has been unusually useful — the OOM reproduction, the dense baseline switch, the CUDA-graph measurements and now the retirement race each changed what this PR actually is, and none of them were things I would have found from my side. The race in particular only shows up on a host that runs far enough ahead, which is exactly the kind of bug a reviewer with different hardware catches and a single-machine author does not.

I appreciate the patience with the long iteration, and the willingness to keep the scope honest (opt-in, GPU-dependent, no serving claims) rather than let it ship on the best-looking number.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants