[Model] Qualify calibrated E4M3 and disk PLE with QSA DCP2 - #702
Draft
yangzhuxinyzx wants to merge 49 commits into
Draft
yangzhuxinyzx wants to merge 49 commits into
yangzhuxinyzx wants to merge 49 commits into
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> (cherry picked from commit cab8358)
…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> (cherry picked from commit 31d137c)
Adapt vllm-project/vllm#51865 (0a94d85a66499cad8297ead86a470967de5c0212) to the current target and Eagle/MTP draft paths. A K+1-token prompt tail must not replay a decode graph merely because its shape matches. Preserve dummy capture and genuine decode classification. Co-authored-by: Rahul Chalamala <22563365+rchalamala@users.noreply.github.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Janelle Cai <janelle.cai@modal.com> Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com> (cherry picked from commit 1892fb2)
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> (cherry picked from commit 464f538)
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> (cherry picked from commit bb4bc78)
The #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> (cherry picked from commit b86e641)
ruff-format and clang-format only; no behaviour change. Keeps the native tree byte-identical to the already-qualified SM70 build so its compiled binaries can be reused. Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
test_qsa_e4m3_loader_requires_all_24_scales asserted that an incomplete overlay raises, but left VLLM_QWEN4EXP_QSA_E4M3_STRICT_SCALES unset, so _validate_qsa_e4m3_scale_load took the non-strict warning path and the test failed on main. The gate it claimed to cover was never exercised. Set the env explicitly and assert both contracts: non-strict returns the missing scale names so the caller can fall back to unit scales, strict raises with the loaded/required count. Also assert the non-FP8 cache short circuit returns an empty set rather than only that it does not raise. Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
… integration/ct252-e4m3-offload-c4-20260919
Include the existing graph-padded token-count regression before changing QSA metadata for DCP2. Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Prepare the experimental target-sharded, draft-replicated layout without enabling QSA DCP serving. Track per-layer physical pages and pad shared Mamba views to the actual tensor stride. CPU allocation and real worker reshape tests cover the mixed layout. Ownership design adapted from vllm-project/vllm PRs #56723 and #57431. Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Preserve the token-ID/-1 selection ABI and the original MTP buffer. Expose optional base-2 LSE and FP32 partial output for cross-rank merging, with the output gate deferred until after the merge. V100 tests pass for localization, graph replay, FP16/E4M3 G12 attention and the original DCP1 call. This remains an operator prototype; serving DCP is still gated off. Index localization adapted from vllm-project/vllm PR #57431. Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Scale target attention blocks by DCP while keeping replicated draft and recurrent owners on their own token grids. Preserve worker physical specs and normalize only prefix lookup specs. Match grouped CPU budgets between worker and scheduler layouts. CPU tests exercise prefix reuse, shared page strides and the full target-only capacity projection. Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Assisted-by: OpenAI Codex Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com>
qsa_dcp_block_geometry declares a replicated selector page as the global token span it covers, but the group's common block table enumerates the kernel blocks this rank owns. _canonical_block_table divided the global span by the local kernel block, doubling the virtual expansion under DCP2: the canonical table collapsed from 11 columns to 6 and every page ID was divided by twice the correct factor, mapping the table onto page 0. The QSA slot kernel then found no column for the newest compressed groups and wrote PAD_SLOT_ID for all 392 group boundaries of a 1568-token prefill chunk, so the selector cache never received the current keys and the indexer scored an identical query against a cache missing them. A 30,030-token prompt answered "blue" under DCP2 and "Blue" under DCP1. Condition the existing division on whether the table counts local tokens, which is true for a sharded page and equally for a replicated selector sharing a sharded main owner's table. A standalone MTP draft replicates every QSA cache and reads a global table, so it is unchanged, and DCP1 is unaffected because the guard still requires dcp_world_size > 1. This supersedes cc026e7, which cleared a startup failure by removing the division rather than correcting its span, and replaces the test that encoded that behaviour. Verified on TP4/V100: DCP2 now returns Blue for the fixed 30,030-token prompt, the canonical table is 11 columns, and the valid compressed-slot positions match DCP1 exactly (392/1568, stride 4) where they were 0/1568 before. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
get_dcp_local_seq_lens built its rank offset with torch.tensor([[dcp_rank]], device=seq_lens.device). Creating a CUDA tensor from a Python list is a pageable H2D copy that synchronizes the stream, and every DCP metadata build calls this function, so each call waited for the previous step's graph replay to drain. With MTP3 that is four builds per decode step, and the stall broke CPU/GPU overlap: the GPU sat idle while the CPU prepared the rest of the step, and the slowest rank's delay propagated to every rank through the tensor-parallel all-reduce. A TP4/DCP2 V100 profile attributed 99% of blocked host time (52 of 305 synchronizations, about 375 ms per rank over one 38-token request) to this call site. Use the Python scalar directly in the broadcast, as upstream vLLM already does, while keeping this tree's return shapes. Measured on the same Qwen4Exp TP4/DCP2 configuration, median of five streamed samples: output speed 57.5 -> 69.2 tok/s on a short prompt and 55.9 -> 68.5 tok/s at 8K (DCP1: 73.1 and 72.4). TTFT is unchanged. The fixed-sample token and MTP acceptance gates remain identical to DCP1. The new test checks every rank against a per-position reference and runs the rank-scalar path under CUDA sync debug mode, which fails on the previous implementation. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Three changes to the DCP path of the QSA target layers, measured on TP4/DCP2 V100 decode (per layer, 4 query rows): - Narrow the localized selection. With interleave 1 every compressed group of 4 positions splits evenly over the DCP ranks, so a rank owns at most compress_ratio // dcp positions of every group and of the open tail. Compaction never writes past (topk // ratio + 1) * (ratio // dcp) columns (1026 of 2051 at DCP2), and the sparse kernel iterates every column it is given. Passing only that prefix: 96.9 -> 61.2 us. - Apply the output gate with the Triton kernel the DCP1 page4 route uses after rounding the merged output to its dtype, instead of a chain of five elementwise kernels: 12.2 -> 1.4 us. - Default Qwen4Exp DCP to the a2a combine (one all-to-all instead of all-gather LSE + reduce-scatter output): 31.9 -> 19.9 us. The option is now unset by default so a model can supply its preference through ParallelConfig.set_dcp_defaults(); an explicit user choice still wins and every other model keeps ag_rs. End to end (C1, MTP3, E4M3, graphs, 32K), decode relative to DCP1 moves from -5.4% to -3.4%. Token outputs and MTP acceptance of the fixed accuracy gate (six MTP prompts, five code repeats, 30K marked prompt with prefix hit/changed suffix) stay identical to DCP1. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
prepare_attn reads input_batch.is_dummy_batch since the new runner marks dummy QSA batches (bdbe40a), but this fixture's input batch never gained the field, so the test failed with an AttributeError on every run. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Every GDN cache group built its spec-decode metadata on every step by uploading the pageable CPU row mask (a stream sync per upload) and indexing device tensors with it (a nonzero sync per index): about a dozen host syncs per group and step, plus an unused query-length upload and a device-side assert. In steady MTP decode the speculative rows lead the batch, with graph padding and any non-speculative rows after them, so: - the state contract selects rows by slicing (cloned, so still fresh contiguous tensors) when they lead, and by mask otherwise; - the device row mask is built from an arange when they lead; - query lengths are uploaded only in the mixed branch that reads them; - the device total check joins the opt-in state-contract assertions (VLLM_SM70_GDN_STATE_CONTRACT_ASSERT=1). A GPU test compares the contract with mask indexing for align, all-cache and block-table state sources under all-spec, spec-then-padding, no-spec and interleaved rows, and runs the leading cases under CUDA sync debug mode. V100 TP4, MTP3, E4M3, C1 decode: DCP1 73.10 -> 77.15 tok/s (+5.5%); with DCP2 packed pages (six GDN groups instead of three) 66.01 -> 70.58 tok/s. Token outputs and MTP acceptance of the accuracy gate are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Grouped (per-cache-group) CPU offload rows are chosen whenever every prefix-cacheable group has one global block span, which is the DCP1 hybrid layout and now also the DCP2 layout with packed target pages. Tiering then rejected any decode context parallelism, so such a server failed at startup with "Grouped tiering currently requires single-node TP". Nothing in the grouped rows depends on DCP: every rank registers and moves only its own pages, one row slice per worker, and the group spans are already global block sizes. Keep rejecting pipeline and prefill context parallelism and multi-node setups. On V100 TP4/DCP2 the offload regression (eviction under pressure, then CPU->GPU recovery with external prefix hits and unchanged output) passes on the grouped rows, matching the DCP1 control. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Under DCP2 each rank holds half of a target QSA layer's main K/V. The shared CSA+linear pool sizes a block so that one physical page holds one recurrent (GDN) state, so keeping a whole state-sized page per sharded layer doubled the global block span (3,200 tokens instead of DCP1's 1,600 on the E4M3/MTP3 checkpoint). Replicated selector, draft and every request's recurrent-state blocks paid for that doubled span, short requests fit worse than on DCP1, and a prefix shorter than 3,200 tokens could never be reused. A block now spans cache_config.block_size global tokens at every DCP size: a sharded target layer holds block_size // dcp slots per rank, and the allocator packs sharded owners whose pages cannot hold a state, in layer order, as few per physical page as hold one (two at DCP2). A packed page interleaves its members one kernel block at a time, so each member is a strided view of the shared tensor and every kernel that honours the block stride (the cache write and the QSA sparse attention) works unchanged. - qsa_dcp_block_geometry returns span = block_size; the model asks the platform to align the block size to kernel block x dcp so the rank-local share stays kernel aligned (an lcm with any user block size). - The allocator groups owners (_pack_csa_linear_main_kv), places recurrent states per physical owner, and emits one KVCacheTensor per physical page with packed_members. - The v2 GPU runner lays out packed members as interleaved strided views; the v1 runner rejects packed configs. - The offload worker registers a packed page once, by its first member, covering every member's kernel blocks of that pool block. - The replicated draft expands its table from the target's rank-local page. Fewer physical owners per block (7 instead of 13) means six GDN state groups instead of three; the previous commit removes the per-group host syncs that otherwise cost DCP2 decode 6.5%. V100 TP4/DCP2, MTP3, E4M3, graphs, grouped CPU offload: - block size 1,600 (DCP1 1,600, previous DCP2 span 3,200) - 32K/C2 capacity 439,296 -> 493,244 tokens (DCP1 396,336) - 262,144/C4/GMU 0.96 capacity 1,142,744 -> 1,178,337 tokens, 4.50x concurrency (DCP1 775,096, 2.96x); allocator projection 1,177,026 - repeated 2K/8K prefix TTFT 834/818 -> 270/256 ms (DCP1 226/235) - accuracy gate identical to DCP1 (tokens and MTP acceptance); offload eviction+recovery, mixed prefill/decode and 5x30K admission regressions pass with DCP1-identical outputs Unit tests cover the geometry, allocator packing and capacity projection, interleaved views (NHD/HND), offload registration, and a V100 test that writes packed members through the real cache writer and reads them with QSA sparse attention bit-identically to unpacked caches (FP16 and E4M3). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Under DCP the FlashAttention metadata builder computes every rank's context lengths for FlashAttention's own DCP attention, a dozen small kernels per build. The QSA builders inherit it, and QSA's DCP attention localizes its own selections and never reads those lengths, yet paid for them on every target and draft build (four per decode step with MTP3). Add builds_dcp_context_lens to FlashAttentionMetadataBuilder (default True) and turn it off for Qwen4ExpQSAMetadataBuilder. A test builds the same DCP2 batch with both and checks that FlashAttention still produces rank 1's lengths while QSA metadata carries none and matches otherwise. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Replace the prototype-stage status with what now serves: the packed target layout at DCP1's block span, the attention path, the accuracy gate and regressions, capacity, prefix reuse, decode and TTFT against DCP1 in one measurement window, and the remaining gaps (XQA-only prefill route, six GDN state groups, packed pages without pipeline parallelism or the v1 runner). The prototype checkpoint stays as history. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Hybrid models build GDN attention metadata once per cache group per step. In steady MTP3 decode one build costs about 0.65 ms of host time on V100 (measured without the profiler), and Qwen4Exp serves its 36 GDN layers in three cache groups at DCP1 and six at DCP2, so the builds alone cost 2-4 ms of every ~40 ms step while the GPU waits. The GDN backend already has the fix, used by the SM70 DFlash2 verifier route: request classification, token indices and offsets are computed once per step (compute_common_gdn_attn_metadata), and for a pure speculative decode batch one Triton launch prepares every group's state indices (prepare_dflash2_gdn_group_metadata); any other batch falls back to per-group builds. Native MTP verifies a linear draft chain exactly as DFlash2 does; only the DFlash selector-engine gate kept it out. Let native MTP use the shared metadata, and for SM70 Qwen4Exp native MTP (a narrow contract, like the DFlash2 verifier defaults) default VLLM_SM70_DFLASH2_VERIFY_FASTPATH and VLLM_SM70_DFLASH2_FUSED_GDN_METADATA to 1; an explicit value still wins and 0 restores per-group builds. V100 TP4, MTP3, E4M3, graphs, same measurement window: C1 decode DCP1 72.98 -> 77.92 tok/s (+6.8%), DCP2 71.21 -> 75.76 tok/s. Accuracy gate (tokens and MTP acceptance) identical to the DCP1 reference on both; offload eviction/recovery, mixed prefill/decode and five 30K admissions pass with DCP1-identical outputs, offload bytes and peak memory. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Leon <tomsel1010.us2@gmail.com>
Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
…ation Reuse PR #696 at 9e4236c. Preserve main GDN CPU-index selection and per-group MTP prefix semantics. Add guarded E4M3 defaults, PLE admission and four-rank graph gates; model acceptance remains pending. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Keep 71c7de1 intact as an integration dependency; E4M3 reduces main KV storage only, not compute reductions. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
…ults Co-authored-by: Codex <noreply@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Record the installed runtime version and resolve benchmark provenance without shadowing the clean wheel. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Co-authored-by: Codex <codex@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Co-authored-by: Codex <codex@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Co-authored-by: Codex <codex@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
Co-authored-by: Codex <codex@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
…ntegration-20260927 Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
4 tasks
Contributor
|
Heads-up for this integration: #696 is now rebased onto current main ( The DCP2 code is the same as what you imported, with three differences that matter here:
Rebasing this PR onto main after #696 and #664 should leave only your own changes and #701. #664 (E4M3 KV with MTP) merges with #696 without conflicts. |
Co-authored-by: Codex <codex@openai.com> Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com>
This was referenced Sep 26, 2026
Contributor
|
Update for whoever (human or agent) picks up this PR next. The pieces #702 imported are now separate, rebased PRs on current main (
The three merge onto main without conflicts, in any order; #705 first is simplest because #664 carries the same fix and drops it on rebase. The old heads this PR imported are kept at Suggested next steps for #702
Why #705 matters for #702's default: on current main an E4M3 server on the QSA DCP1 path stops during startup ( What we measured on main
|
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
3 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
Follow-up qualification/integration of #696 (DCP2) and #701 (high-precision
reductions / host PLE placement). Draft: paired dataset acceptance and DCP2
long-context model checks pending.
This does not implement a competing DCP layout and does not supersede #696.
Both original histories are preserved. The distinct changes are safe E4M3
defaults for SM70 Qwen3.8 NVFP4 TP4/MTP0, admitting the existing disk-PLE
producer under DCP2, reconciling current-main GDN/prefix semantics, real
four-rank collective/graph regression tests, and a source-installed bounded
model/256K quality driver. The original PRs remain dependencies/review scopes.
b034648012244ab712df05b93e6d8fff877a6f2f.9e4236c2a6f5a15c44c1a0542ab53a779eade730.71c7de134e9107655ad01a1e329f0813c8b0c1b8.fd4b545e0c,32acc8bda7,f6c62234e9.dd7624462fae2b7016aaa8843527e9fc3af85b67.f6c62234e9959edb5ea2fda2b386ee0cd71d276b.1e90d17f2cmerged cleanly atd6ff5be488; normal new native build andinstalled wheel prepared. Current-main full-model dataset acceptance pending.
Automatic E4M3 applies only to the admitted no-MTP route; explicit FP16 wins.
Automatic selection turns on strict calibration checks unless overridden.
The published revision-bound 24-scale payload matches our base index and has
been materialized into an independent model view; original weights unchanged.
Compute reductions remain FP32 / high-precision. FP8 is main KV storage only.
Test Plan
communication.
private DSOs or
LD_PRELOAD.token differences as diagnostics, 8K/32K/near-256K retrieval and exact 262143+1 boundary,
allocator/worker memory, separate prefill/decode metrics. Shut workers down.
permitted if paired task quality does not regress. First screen: fixed
128 GSM8K + 32 LongBench multifieldqa_en + 32 multifieldqa_zh; identical model,
official sampling and per-question seeds; report per-dataset scores and paired
wins/losses, truncation, empty output and nonfinite logprobs. This finite subset
is not proof of universally lossless quality. Keep high-precision reductions.
Test Result
graph-padding metadata tests.
three-tuple unpacking of the new four-field
SpecGroup; targeted rerun passed.tests/models/qwen4_exp/test_qsa_dcp_collectives.py: passed on four V100s;FP16/E4M3, rows1/5/33, AG+RS and A2A, eager and two changed-input graph
replays, checked against an independent FP32 oracle.
passed after [Core] Keep CUDA reductions and PLE memory on safe defaults #701, including environment-isolation repair.
all 192 selected questions (100–13,601 input tokens); local GSM8K rows exactly
match a pinned official source after JSON normalization.
verifies each worker's sampler setting. The previous full-model short health
cases checked natural EOS/answers but did not opt into full-logit NaN scanning;
their exact-boundary sampled logprob finite check remains valid. NaN auditing
disables the greedy-only shortcut, so these health timings are not production
throughput results.
CUDA/Torch dependencies only; no private DSO, preload or source overlay.
Wheel SHA256:
e1cb69363ff7e1ff66c2fc04f25564f6d52bb17cc9f8d278177f9e9fd29460be.natural EOS, including 261,632-token input. Exact 262,143+1 boundary returned
finite logprob. This is health evidence, not a dataset accuracy result.
reasoning first diverged at zero-based token 27 and stopped the old parity
gate before long tests. Under the updated criterion this is diagnostic, not
by itself a dataset-quality regression.
difference 2.44e-4; short-prefill 9.77e-4. DCP2 was not farther from the FP64
oracle in three prefill seeds. This does not establish the exact cause of
the full-model token difference.
Same 4 GiB/card physical KV budget, TP4, E4M3:
Allocator capacity increases 76.98%; four simultaneous 256K requests have
not been tested. PLE logs confirm GPU table rows=0, disk mmap prefill and
pinned-host/UVA decode. Runtime arithmetic precision flags remain disabled
for reduced-precision reductions/accumulation.
Commands and the bounded model contract are in
docs/design/qwen38_calibrated_dcp_integration.md.This PR does not claim restored 97 tok/s, paired dataset acceptance or
accepted DCP2 256K model quality yet. #696 author's measurements are not ours.
Raw local evidence is retained in the owned worktree's
.artifacts/:model_dcp1.json,model_dcp2.json,dcp_numerics.log,dataset_preflight.json,dataset_driver_cpu_tests.logandWORKLOG.md.Dataset screen correction: the first DCP1 192-question run accidentally
selected source-checkout workers: the reused GSM8K benchmark indirectly
prepended source paths to
sys.path. Concurrent source/build changes invalidateit as a frozen-artifact comparison. Its scores are retained as diagnostic only;
they are not promotion evidence. Earlier non-dataset health/capacity worker
manifests confirm installed-wheel paths and remain valid.
The scorer no longer imports benchmark entrypoints. A real CPU spawn now
verifies package/native-extension paths and version before model load, then
actual worker manifests must match. 20 CPU regression cases and the actual
192-question/tokenizer/scorer/spawn preflight pass with the current-main wheel.
The first DCP2 dataset attempt failed the pre-load memory check (28.15 GiB free,
28.56 GiB requested), before output evaluation. Exact transient memory ownership
was not recorded; do not attribute this to another task or DCP numerics.
All owned GPU processes exited. Other tasks now hold both four-card groups.
The corrected pair uses the same utilization0.85 on both sides with the
unchanged explicit4GiB/card KV budget and preserved precision/context.
Updated source-installed wheel SHA256:
3687562a695ccc2052059eb503aefce705e6a72555a324173cc8c9feac0ec091.65 targeted CPU warmup/graph tests passed after current-main integration.
Corrected full-model dataset pair has not started; PR remains Draft.
Review / Risks
AI assistance was used (Codex). Human line-by-line review and relevant reruns
are still required before promotion/merge. No direct main push or running
service is part of this change. DCP4, automatic MTP enablement, and a lower
arithmetic-precision speed path are out of scope. The default now requires a
calibrated model view;
--kv-cache-dtype float16is the explicit control.