Conversation
…3 NaN The SM70 grouped-page4 QSA prefill route can emit NaN for every query row and head on a fixed set of head dims when the KV cache is E4M3 and a hybrid (GDN/Mamba) layout co-locates non-attention state in the paged KV pool. The same batch is correct through the XQA route. Root cause: the grouped planner pads each category to a multiple of eight with the null block -- (physical microblock 0, mask 0) -- and counts the padding in seq_len. The forward loads page 0's K/V for those padded rows and masks them multiplicatively (P = 0). But 0 * NaN = NaN survives the P@V tensor-core reduction, and under E4M3 the null block's bytes decode to NaN. Every group's every tile reads the same page-0 microblock, so the whole tile's rows and heads go NaN on identical dims. XQA never reads the null block, which is why it stays clean on the same inputs. Fix (defense in depth, both default-on): - Kernel: in flash_attention_grouped_verify_e5m2_partial_kernel under SPARSE_PAGE4, after loading each K/V panel, zero every row no query attends (unattended iff (mask & (0x11111111u << (row & 3))) == 0) so the reduction sees 0 * 0 = 0. A leading __syncthreads() before the V-panel zeroing keeps the panel load from racing past the zero and resurrecting the NaN; the trailing one publishes the zeros before P@V. - Python (VLLM_SM70_QSA_GROUPED_PAD_FIX, default 1): after the planner, repoint every mask==0 padding entry at the group's first real microblock. torch.where + copy_ has no host sync, so it stays CUDA-graph capturable. This backstops binaries built before the kernel fix. The grouped route stays default-on (VLLM_SM70_QSA_GROUPED_PAGE4=1); only what the padded rows contribute changes. Adds tests/models/qwen4_exp/test_qsa_e4m3.py::test_qsa_grouped_page4_null_block_padding_no_nan, an SM70-only regression that poisons the null block's value plane with E4M3 NaN and asserts the grouped route stays finite and matches XQA, parametrized over the Python pad-fix off (kernel-only) and on. It reproduces the NaN on an unfixed kernel and passes on either fix. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…adding test The regression added in the previous commit used the XQA page4 route as its clean reference. On Flash-V100 builds whose XQA E4M3 page4 kernel requires an fp32 partition buffer, that reference raises "XQA decode tmp_out must be fp32 for E4M3 KV" before the grouped route is exercised, so the test cannot run against such a build. Use a kernel-independent einsum ground truth instead: decode the E4M3 K/V to fp32 and attend only the selected tokens (which map to physical pages >= 1, so the null block is excluded). This is .so-agnostic and directly checks that the grouped route stays finite and numerically correct. The grouped route and the parametrized pad-fix coverage are unchanged. Verified on V100 (SM70): the grouped null-block NaN reproduces with the stock kernel and pad-fix off (test fails), and passes with either the kernel fix or the Python pad-fix. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Keep the downstream integration branch clean under the fork's current ruff and clang-format hooks. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Staging branch for the separate E4M3 + MTP draft-scale enablement; NOT for merge and no PR. Ports two work-in-progress patch sets on top of the null-block-padding fix branch: - phase2: opt-in gate (VLLM_QWEN4EXP_QSA_E4M3_MTP) that relaxes the QSA E4M3 MTP0 requirement, an extracted _verify_e4m3_kv_requirements, the envs entry, an MTP weight-remap hook, calibration-overlay tooling, and CPU/GPU tests. - phase2b: a CUDA-graph-capture-safe XQA workspace (torch.full on-device instead of a host torch.tensor copy) with capture/envelope tests. CPU check (CUDA hidden): 27 passed, 19 skipped (GPU-only), 2 failed. The two failures (test_validate_scale_overlay_lists_missing_names, test_finalize_qsa_scale_load_success_and_missing) exercise draft-side q-scale visibility whose model.py scale-load changes are not part of these two patches; that piece is out of scope for this staging branch. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Finish the staged draft-scale path without changing the existing main-model fallback contract. Add a repository-owned, revision-checked target/draft scale overlay materializer and make its CPU tests portable. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Bind generated target and draft scale overlays to the source manifest, draft calibration report, base checkpoint, and merged index hashes. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Normalize generated scale, index, and provenance files to mode 0644 so overlays created in a root-owned container remain loadable by non-root runtimes. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
The 1CatAI#647 overlay kept the stricter native E4M3 ABI but lost the matching Python workspace dtype selection. Key the cached workspace by KV dtype and allocate FP32 temporary output for E4M3 while retaining FP16 for the existing path. Co-authored-by: OpenAI Codex <codex@openai.com> Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
This was referenced Sep 19, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Purpose
Rebuild the downstream E4M3 QSA + MTP qualification lane from the current fork
main, upstream 1Cat PR 1CatAI#647, and areslp's separate phase-2/phase-2b stagingcommit. PR #16 remains unchanged as the first failed experiment and evidence
record.
Current scope
0fda07e: opt-in E4M3+MTP admission, draft scale-name remap,calibration support, graph-capture-safe XQA workspace, and focused tests.
E4M3 MTP draft loading requires calibrated 2/2 K/V scales and fails closed
instead of using unit scales. The existing main-model fallback contract is
unchanged.
materializer. It records the base, input manifest, draft calibration report,
and merged-index hashes and emits non-root-readable artifacts. The imported
tests no longer depend on the author's private
/home/l/...tool path.the native E4M3 ABI; the pre-existing FP16 path is unchanged.
The calibrated CT252 values remain an external checkpoint artifact; they are
not hard-coded into runtime source.
Duplicate-work check
feat/qsa-e4m3-mtp-draft-scalesexplicitly declares itself astaging branch with no PR; its original authorship is preserved here.
lane without closing or rewriting it.
implementation of [Bugfix][SM70] Stop grouped-page4 null-block padding from leaking E4M3 NaN 1CatAI/1Cat-vLLM#647.
Tests
28 passed, 20 skipped(the skips are GPU-only SM70/runtime tests).8 passed.forbidden-import, and torch CUDA API checks.
remap to the draft layer.
b86e6414a7:19 passed, 15 warnings.24/24, MTP draft2/2.variants, and grouped/XQA warmups captured successfully.
4 concurrent exact responses, and the 804-token boundary regression passed.
drafted 53 and accepted 51 tokens (96.2%), with 32/32 in the concurrent arm.
Qualification boundary
This is a functional GPU admission, not production promotion. First-use shapes
still produced Triton JIT latency warnings, so production warmup coverage needs
expansion. CPU KV offload is not admitted: the inherited hybrid-offload
configuration was rejected by the scheduler layout assertion and was removed
from the successful GPU-only arm.
AI assistance
OpenAI Codex assisted with repository synchronization, integration, test
repair, the overlay materializer, static validation, GPU qualification,
evidence capture, and this draft description. The human submitter must review
and understand every changed line before promotion.