Conversation
DeepSeek-V4.1-Flash stores each KV token as 528B: 512B of FP8 E4M3 covering the full 512-wide K (rope lanes quantized, no BF16 rope segment) plus a 16B footer of 16 UE8M0 scales over 32-wide groups (the deepseek-ai/FlashMLA#221 layout, SM100-only there). Add ModelType::DSV4_1 on top of the ScaleSpec groundwork: GLM53_NOPE-style geometry (D_ROPE=0) composed with a 32-wide UE8M0 footer. Kernel side the format is simpler than DSV4 (one 512B bulk per row, a single unified block-scaled FP8 QK pass, no rope MMA paths). The two new mechanics: - the 16B footer row generalizes the scale gather from one uint64_t to a width-selected load (uint4) in both the prefill IO helper and the decode kernel's staged gather; - the 32-wide groups constrain the XV W-fold (one scale group per W buffer), so an 8-warp XV split would floor NT_PER_WARP_XV to 0: decode runs a 4-math-warp tile at BI=64 (the DOTS3_SWA halving precedent) and prefill is SG-only on the BI=32 producer/consumer tile. The ComputeTraits/SmemLayout aliases now take the XV warp count so the NT_PER_WARP_XV assert evaluates at the warp count the XV MMA actually runs at. Selection is explicit everywhere (d_qk=512 collides with DSV4, the 528B payload with GLM53_NOPE): the functional API gains kv_scale_format= "ue8m0_g32", the TRTLLM-compat entry kv_cache_format="fp8_dsv41", and the decode-dsv4 FFI takes an explicit model_type (-1 keeps the legacy width inference, matching the paged-entry kAuto convention). Dual-cache decode is supported like DSV4; dual-cache prefill (MG_DUAL) is intentionally not instantiated. Calibration keys a new dsv4_1 family at topk 512. Tests: decode matrix (dedicated/runtime-H heads, partial-tail topk, sinks), SG prefill, dual-cache through the TRTLLM entry, poisoned-slot-0 NaN safety on both gather paths, dispatch/plan/envelope coverage. 666 passed on RTX PRO 6000 (SM120). Signed-off-by: lucifer1004 <13583761+lucifer1004@users.noreply.github.com>
…ile) The XV W-fold binds one FP8 weight buffer to one V scale group, so a QUANT_TILE-wide chunk feeds at most QUANT_TILE/8 warps. DSV4_1's 32-wide groups would floor NT_PER_WARP_XV to 0 at 8 math warps; instead of halving the warp count, fold W once per scale group within a 2-chunk group (XV_FOLD=2): warp groups of N_WARPS/XV_FOLD share one chunk and tile its dims. The mapping degrades to the existing per-chunk loop at XV_FOLD=1, so DSV4/DOTS3_SWA are unchanged. DSV4_1 decode returns to the standard 8-warp BI=64 tile: halves per-warp QK/XV MMA latency and acc registers vs the 4-warp fallback, at +2.5KB smem (a second W buffer per parity). Signed-off-by: lucifer1004 <13583761+lucifer1004@users.noreply.github.com>
Bring the dedicated eight-head instantiation and compact/masked-row tests from flashinfer-ai#5022 onto the canonical-payload branch.
Honor flat-cache page spacing, enforce aligned gather addresses and packed footer rows, and cover padded RoPE paths and eight-head scratch bounds.
00e7865 to
b0e7bac
Compare
…decode, stale-cache rereads Carried from ormandj's review branch (#3), adapted to the pre-DSV4_1 tree: - parse_paged_kv_layout derived the flat-2D block advance from the dim-1 size instead of stride(0); a cache view with gaps between pages read the wrong page. The physical block stride is now honored and validated. - Cache origins and block strides must be 16B-aligned (cp.async.bulk), and footer-scale 3D/4D views must keep token rows packed — previously only the row width was checked. - GLM53_NOPE H=8 decode gets a dedicated instantiation: the runtime-H fallback strides split-K scratch by the 16-head tile while the public allocator reserves 8 rows (the flashinfer-ai#5022 eight-head fix, carried with its regression coverage). - A stale-schema calibration file's mtime is remembered so repeated lookups no longer reparse it; replacement tuning loads when the file changes. The DSV4_1-parametrized coverage in the original branch is deferred to the DSV4.1 follow-up PR (the type does not exist on this branch). Co-authored-by: David Orman <ormandj@corenode.com> Signed-off-by: lucifer1004 <13583761+lucifer1004@users.noreply.github.com>
|
Thanks @ormandj — this is a great catch set, and the timing needs one note: flashinfer-ai#5075 was force-pushed back today to its pre-DSV4.1 state (the DSV4.1 cache-format support moved to a follow-up PR so the NaN/row-stride fixes can land first). Your branch was based on the pre-rewind head, so I integrated your changes directly rather than merging this PR as-is. Now in flashinfer-ai#5075 (single squashed commit,
Deliberately not taken (they belong with the DSV4.1 follow-up PR, since the rewound branch has no DSV4_1 type): the two 528B-format disambiguation / DSV4.1 scale-group commits. If you are willing, rebase just those two onto the follow-up once it is open — they are exactly the coverage it needs. During integration I also had to adapt Two replies to your notes:
Closing this PR since the content is incorporated — thank you again. (Developed-with-AI disclosure noted; same here per repository policy.) |
… NaN-safe masked gathers in sparse MLA (#5075) ## 📌 Description Three changes to the SM120 sparse-MLA cache handling, motivated by the compact-row discussion in #5022. This is an alternative design for the same goals. It also carries follow-up fixes from community review: cache-view validation hardening and the GLM NoPE eight-head decode instantiation from ormandj's branch (lucifer1004#3, retaining the #5022 implementation), and the calibration-cache invalidation reported by qsang-nv. **Runtime gmem row stride ("stride is data, not format").** GLM53_NOPE's packed payload is 528B/token (512 FP8 values + 4 inline FP32 scales); the 656B row came from inheriting vLLM's fp8_ds_mla ABI, whose trailing 128B are RoPE bytes the NoPE model never uses. Instead of adding a second layout selected by shape sniffing, this PR makes the payload canonical (`bytes_per_token = 528`) and takes the gmem row advance as a runtime stride everywhere: - decode already honored `stride_kv_row`; prefill now derives the advance from `stride_kv_block / page_block_size` (no new kernel parameter), and swapAB decouples its smem row stride (528B payload) from the gmem advance. - A legacy 656B pool (vLLM fp8_ds_mla) and a compact 528B pool are the same kernel with a different stride; the payload prefix is identical, so existing callers keep working unchanged. Per-token KV storage drops 19.5% for callers that switch to 528B rows. - Binding validation: inline-scale models accept padded rows (advance >= width >= payload, 16B-aligned, blocks contiguous); footer-scale models keep the packed-rows requirement. Flat 2D caches must be packed at 528B for GLM53_NOPE (a flat block carries no stride to infer); GLM53_NOPE has no released downstream, so this is the last moment the canonical width can flip. - The cpb calibration store bumps to schema v2 for the flip: a pre-PR cache entry persists `bytes_per_chunk = 41984` for glm53_nope and would silently overstate the L2 footprint in the cpb guard by ~24%, potentially excluding valid candidates. Stale files count as absent and recalibrate on the next tuning-mode pass (review feedback from qsang-nv on this PR), and a stale file's mtime is remembered so repeated lookups no longer reread it. **Cache-view validation hardening** (carried from ormandj's review branch, lucifer1004#3): - Flat 2D caches derived the block advance from the dim-1 *size* instead of the real `stride(0)`, so a view with gaps between pages read the wrong page. The physical block stride is now honored (and must be >= the packed block width). - Cache origins and block strides must be 16B-aligned (cp.async.bulk); footer-scale 3D/4D views must keep token rows packed (previously only the row *width* was checked, not the stride). - GLM53_NOPE H=8 decode now has a dedicated instantiation: the runtime-H fallback strides split-K scratch by the 16-head tile while the public allocator reserves only 8 rows for H=8 (the #5022 eight-head fix, carried). H=16 joins the dedicated grid for symmetry with DSV3_2/GLM_NSA — it was already correct via runtime-H (its scratch stride equals the true head count there), this just drops the runtime loop bounds. **Universal NaN-safe masked gathers.** Masked (-1) candidates were clamped to mutable cache slot 0 in seven places (prefill MG/SG gather + scale gather, swapAB gather, decode-v32, decode-dsv4, and the two NVFP4 kernels from #4955). A NaN in that slot leaks into valid outputs through `0 * NaN` in the value MMA (and `0 * inf` from a poisoned 0xFF footer scale on footer-scale models). All seven now gather a shared zero-initialized row (`sparse_mla_zero_row`), which is mbar-compatible and adds no synchronization. The NVFP4 case is narrower but real: the masked scale load there is explicitly zeroed, which covers the E2M1 NoPE values, but the 128B BF16 rope tail rides the same clamped pointer with no scale in its path (review feedback from qsang-nv on this PR). The property now holds unconditionally for every model type — no opt-in and no contract version to negotiate. This is a live bug on main for every family, not only the new one: this PR's seven poison tests each fail against a build without the fix (glm53_nope, dsv4 decode, dsv4 prefill, dsv3_2 prefill verified on main fb96128; nvfp4 verified against the pre-fix kernels, which are main's). **Unconditional topk bounds in prefill index staging.** The prefill `load_idx` lambdas (MG, MG dual-cache, swapAB) bounded lanes by `actual_ni * BI` but not by `topk_len`, so lanes in the last partial tile read caller padding. Callers that leave stale values there (instead of -1) would gather wild gmem addresses — the new `bounds_stale_padding` tests crash a main build with a CUDA illegal memory access on both the MG and the swapAB route. Lanes past `topk_len` now stage -1 for every model type. The math side normalizes the same way before forming gmem addresses from the raw index row: the QK rope operand loads (SG/MG/producer-consumer) are real loads, and DSV4's XV rope MMA reads rope straight from gmem, so stale positive padding past `topk_len` could form a wild address there as well (review feedback from coderabbit on this PR; the initial stale-padding test only covered rope-free GLM53_NOPE). **ScaleSpec extraction (layout groundwork).** Each model's scale configuration (numeric format, group size, inline-vs-footer placement) is now a composed `ScaleSpec` row in `model/scale_spec.cuh` instead of hand-spelled constants per `KVCacheTraits` specialization, with the conversion helpers shared through `ScaleConvert`. Consumers keep reading the forwarded members (`QUANT_TILE`/`NUM_SCALES`/`SCALE_FORMAT`/`SCALE_INLINE`/...) unchanged, and static asserts pin every derived constant to its pre-refactor value. A future DSv4-MXFP8 (group-32 UE8M0) or GLM53-NVFP4 cache becomes one traits row plus its dispatch entry, not another parallel traits/kernel copy; combined with the runtime row stride, the container side needs no kernel change at all. Relation to #5022: same goals (compact rows, masked-read correction); different mechanism — one canonical layout + runtime stride instead of dual layouts + `shape[-1] == 528` sniffing + `compact_bytes_per_token`/`glm53_nope_contract_version` advertisement fields. #5022 constexpr-gates its masked-read and bounds fixes to GLM53_NOPE; this PR applies them to every family, and the on-main reproduction above shows the hazards are not glm53-specific. If this PR lands, it is intended to supersede #5022's compact-row and masked-read changes; the 8-head decode instantiation from #5022 is now carried here (via ormandj's branch, with its regression coverage). ## 🧪 Tests RTX PRO 6000 (SM120): ```bash python -m pytest -q tests/attention/test_sparse_mla_sm120.py \ tests/attention/test_sparse_mla_sm120_dispatch.py \ tests/attention/test_sparse_mla_sm120_cpb_model.py \ tests/attention/test_sparse_mla_nvfp4_sm120.py \ tests/attention/test_sparse_mla_nvfp4_sm120_plan.py # 782 passed ``` Performance (RTX PRO 6000; paired same-session A/B vs upstream main fb96128, per-arm forced JIT rebuilds, isolated autotune caches, CUDA-graph replay timing for the runner suite plus direct-binding timing with explicit cpb for the DRAM-pool rows): - Decode (dsv4/dsv3_2/glm53_nope, T=1..64, single and dual cache): flat within ±1% on dsv3_2 and glm53_nope; dsv4 dual-cache decode carries +2-3% at small T, cleanly attributable to the zero-row select in the issue loop (a probe with main's clamp instead matches main exactly). At 11-16us per call this is ~0.3us of added predication in the latency-bound regime. - Prefill: dsv3_2/glm_nsa MG+swapAB flat across a 57-row matrix (0.986-1.013x); dsv4 MG prefill (H=128, topk=1024) on the pre-review build was +0.7% at T=128 and +2.4% at T=2048 (3855us vs main's 3761-3776us). An earlier revision spilled 8B to stack in the DSV4 MG instantiation and lost +4.6% there; restructuring the zero-row selects to clamp-first form removed the spill (register/stack usage is now identical to main). The residual resists micro-optimization: five formulations of the bounds/zero-row predication — including a fused per-lane limit that is strictly fewer instructions per tile than main's original guard — all measure identically (+2.2-2.5%). On this shape the IO warps sit exactly at a per-tile issue edge where a handful of added instructions cost the full delta and the cost saturates: spreading the bulk gather from 64 to 128 IO threads (3870us vs 3855us), staging fully resolved row pointers and footer scale values a tile ahead so the issue point is a bare cp.async.bulk (3880us, +2.6% at T=128), address-clamped unconditional index loads (3862-3866us), and a full-tile fast branch whose executed loop body differs from main only by the three zero-row selects (3855-3864us) all measure the same; only a loop body with zero added instructions — no protection at all — recovers main's 3761-3776us. NCU stall attribution puts the delta on long-scoreboard (gmem dependency) cycles in the gather/index path: +0.19 cycles per issued instruction on +0.7% more issued instructions, with barrier, MIO, and short-scoreboard stalls unchanged. Static compute/MMA/LSU SASS opcode counts are identical to main. Two closing controls (separate event-timed harness, same GPU): (a) adding a single always-true per-lane select to main's staging loop reproduces nearly the whole delta by itself (3895us vs main's 3795us; the PR adds only ~10-25us on top of that), so the loop has no slack for any added instruction — this is not specific to the protection logic; (b) the +2.2-2.5% is a property of the 10MB L2-resident pool used above, not of production shapes: with a 4.9GB DRAM-resident pool the same change costs +0.4% (3940us vs 3924us), because the longer per-tile gather latency restores the slack the L2-resident microbenchmark lacks. The math-side topk_len normalization added during review shifted the same equilibrium the other way — the final build measures 3808-3815us at T=2048 (+1.3% vs main) and 250.9-251.1us at T=128 (-0.9%, i.e. slightly faster than main) — consistent with the reading that small instruction changes anywhere in this kernel move this L2-resident shape by ~1% in either direction. - GLM53_NOPE 528B rows are time-neutral on the measured shapes: at the production decode shape (T=64, H=64, topk=2176, 640MiB pool, cpb=16) both row widths run ~220us at a ~415 GB/s unique-byte rate — not DRAM-byte-bound on this GPU. The 528B row's benefit is the -19.5% KV footprint and the layout compatibility, not local kernel time. New tests: - `test_sparse_mla_sm120_glm53_nope_compact_rows`: 656B pool vs packed 528B vs a 528B slice of the 656B pool (row stride 656) are bitwise-identical across decode, prefill-MG, and prefill-swapAB shapes. - `test_sparse_mla_sm120_{glm53_nope,decode_dsv4}_masked_rows_ignore_poisoned_slot_zero`: slot 0 poisoned with NaN values and scales; masked candidates keep outputs finite and on-reference (inline and footer gather paths). - `test_sparse_mla_sm120_prefill_glm53_nope_bounds_stale_padding`: garbage (huge positive) indices past `topk_length` are never gathered (MG + swapAB). - `test_sparse_mla_sm120_prefill_{dsv3_2,dsv4,dots3_swa}_bounds_stale_padding`: the same stale-padding discipline for rope models — SG, MG, MG dual-cache, and the BI=32 producer-consumer route — covering the math-side rope address formation. Each crashes the pre-fix build of this PR with an illegal memory access and passes after. Carried from ormandj's branch (lucifer1004#3): - `test_sparse_mla_sm120_footer_flat_block_stride` / `test_glm53_decode_flat_block_stride`: flat 2D caches with aligned gaps between pages decode bit-identically to packed pages (the size-vs-stride regression). - `test_sparse_mla_sm120_cache_alignment_rejected`, `test_sparse_mla_sm120_{decode,prefill}_footer_row_gap_rejected`, `test_sparse_mla_sm120_inline_scale_rejects_padded_block_stride`: the new binding rejections (misaligned origin/block stride, footer row padding, inter-block gaps). - `test_glm53_eight_head_decode_preserves_scratch_guards`: the H=8 scratch ABI, with guard regions. - `test_glm53_compact_rows_match_padded_rows`, `test_glm53_masked_cache_rows_ignore_poisoned_slot_zero`, `test_glm53_canonical_payload_is_scoped_to_model`: #5022's compact-row and masked-row regression coverage on the canonical-payload design. - `test_legacy_glm_layout_calibration_is_invalidated` / `test_stale_schema_is_read_once_until_file_changes`: schema-v1 caches are retired and a stale file is not reparsed until it changes. Updated tests: - `test_sparse_mla_sm120_inline_scale_prefill_accepts_padded_rows` (was `..._rejects_padded_rows`): padded-row inline caches now work in prefill; the inter-block-gap rejection test keeps its semantics with the new message. - `test_sparse_mla_sm120_envelope_consistency`: 2D fixtures pack GLM53_NOPE at its 528B payload. pre-commit (clang-format, ruff, mypy) passes on the changed files. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit - **New Features** - Added support for GLM-5.3 NoPE sparse MLA KV caches with compact 528-byte payloads and padded rows. - Added runtime row-stride handling for inline-scale KV caches. - Masked or out-of-range KV entries now use safe zero values instead of reading cache slot 0. - **Bug Fixes** - Prevented stale padding and invalid cache indices from producing unintended memory reads or non-finite attention results. - Improved support for padded KV rows across prefill, decode, and dual-cache workflows. - **Tests** - Added coverage for compact rows, padded layouts, masked entries, and multiple sparse MLA model variants. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Signed-off-by: lucifer1004 <13583761+lucifer1004@users.noreply.github.com> Co-authored-by: David Orman <ormandj@corenode.com>
📌 Description
Follow-up to flashinfer-ai/flashinfer#5075, targeting its branch. Preserve its canonical payload and runtime-stride design while tightening cache-view validation and bringing over the remaining GLM NoPE fix and regression coverage from #5022.
Flat 2D cache parsing currently substitutes the packed width for the actual block stride. A view with gaps between pages therefore reads the wrong page. The bindings also accept misaligned cache origins and block strides, and prefill accepts padded token strides for footer-scale caches whose kernels require packed rows. Preserve the actual block stride and reject these unsupported layouts before launching kernels.
GLM NoPE with eight query heads currently selects a kernel that uses sixteen scratch rows, while the caller allocates eight. Add the dedicated eight-head instantiation. Carry over compact/padded, masked-row, and CUDA graph replay tests; extend coverage to aligned storage offsets, additional padded row widths, RoPE prefill routes, and scratch guard regions.
These binding and scratch gaps also exist before flashinfer-ai#5075. This contribution retains Gabriel Wu's broader masking and layout implementation.
Invalidate the earlier calibration-cache schema: it can retain the former 656-byte GLM footprint and measured tuning choices after the payload becomes 528 bytes, as reported by @qsang-nv. This retires version-1 tuning files for all sparse MLA families. Normal untuned dispatch remains available until they are regenerated. Remember a stale file’s timestamp so repeated misses do not reread it; load replacement tuning when the file changes.
Add explicit coverage distinguishing GLM's 512 FP8 values plus four FP32 scales from DeepSeek-V4.1's 512 FP8 values plus sixteen UE8M0 scales. Both have a 528-byte payload, so byte size alone must not select the format. A deterministic per-group test checks all sixteen DeepSeek-V4.1 scale groups across decode chunk boundaries.
🔍 Related Issues
🚀 Pull Request Checklist
Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete.
✅ Pre-commit Checks
pre-commitby runningpip install pre-commit(or used your preferred method).pre-commit install.pre-commit run --all-filesand fixed any reported issues.🧪 Tests
unittest, etc.).Configured hooks pass on the changed files. The full-repository hook run is not claimed.
Author-run CPU validation: 64 passed, 17 GPU cases skipped across
test_sparse_mla_sm120_dispatch.py,test_sparse_mla_sm120_cpb_model.py, andtest_sparse_mla_nvfp4_sm120_plan.py. The new calibration regression fails when the module uses schema 1 and passes with schema 2, covering old constants, measured overrides and crossover entries, plus a current-schema save/reload.Author-run GPU validation on RTX PRO 6000 Blackwell Max-Q (SM120):
tests/attention/test_sparse_mla_sm120.pymemcheck, layout/masking/scratch/scale-group subsetThese GPU runs used the final changed files from this branch over an integration containing flashinfer-ai#5075, including a fresh run of the SGLang adapter tests. Source hashes were checked before testing. This is author-run validation of that integration; the complete repository suite and other GPU architectures are not claimed.
The regression matrix covers compact and padded GLM rows, physical page gaps, aligned storage offsets, rejected misalignment/footer-row gaps, eight-head scratch guards, poisoned unused cache entries, and CUDA graph replay with changed inputs. The standalone decode tests also cover its existing packed-row checks for both main and extra caches. Whole-model serving checks passed with the combined integration, but do not isolate this PR's effect; no serving-speedup or exact-output-equivalence claim is made.
🔬 Experimental Track
flashinfer/experimental/and/or an@flashinfer_experimental_api. Tracking issue: #tests/experimental/and were validated on the intended hardware; a runnable example is included.flashinfer/aot.py, and no experimental backend is reachable frombackend="auto"withoutFLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1. (Calling an@flashinfer_experimental_apior naming a backend explicitly is itself the opt-in and needs no environment variable.)Reviewer Notes
Review physical block strides, packed footer rows, the eight-head scratch contract, and calibration invalidation. Flat 2D GLM caches use packed 528-byte rows; existing padded 656-byte pools need an explicit token axis in a 3D/4D view.
The separate NVFP4 slot-zero masking issue reported on flashinfer-ai#5075 is outside this FP8 change. DeepSeek-V4.1 dual-cache prefill above 64 query tokens remains unsupported by the parent.
Please preserve the contributed commits' authorship when incorporating them, including co-author credit if squash-merging.
Developed with AI assistance.