Skip to content

feat(loader): native U8 NVFP4 loading for pre-quantized checkpoints#257

Merged
tbraun96 merged 3 commits into
Avarok-Cybersecurity:mainfrom
Sujimoshi:feat/nvfp4-native-u8-loading
Jul 10, 2026
Merged

feat(loader): native U8 NVFP4 loading for pre-quantized checkpoints#257
tbraun96 merged 3 commits into
Avarok-Cybersecurity:mainfrom
Sujimoshi:feat/nvfp4-native-u8-loading

Conversation

@Sujimoshi

@Sujimoshi Sujimoshi commented Jul 5, 2026

Copy link
Copy Markdown
Contributor

Summary

Third in a planned split of #218 (DFlash speculative decoding) into smaller, independently reviewable PRs (first two: #255 merged, #256 closed as unneeded).

Some pre-quantized checkpoints (Standard NVFP4 convention, e.g. sakamakismile) ship weights already packed as U8 NVFP4 on disk at the .weight key (rather than .weight_packed, the compressed-tensors convention). This adds a direct-load path that skips the BF16-dequant-then-requantize roundtrip:

  • FullAttention Q/K/V/O (TP-sharded path): detect UInt8 dtype, load via quantized_auto(Nvfp4Variant::Standard) instead of dense_auto (which bails on UInt8). TP sharding of pre-quantized NVFP4 isn't supported yet (tp_size=1 only).
  • LinearAttention (SSM/GDN): new native_nvfp4 fast path — when in_proj_qkv/in_proj_z/out_proj are all UInt8, load+concat+transpose them as QuantizedWeight directly. load_ssm_proj also gained a UInt8 branch for partial-U8 checkpoints that fall through to the legacy path. in_proj_a/in_proj_b now route through load_ssm_proj instead of raw dense() so a U8-packed A/B doesn't get read as BF16 garbage.
  • quant_helpers.rs: quantized_v2 now sets weight_scale_2_vec (a field added to QuantizedWeight since this code was originally written).

Note on provenance: re-derived against current main rather than patched mechanically. Main has since added a native FP8 GDN fast path (fixing a ~7pt BFCL-ST regression) and an ATLAS_GDN_BF16_WEIGHTS opt-in path — both sit ahead of this new native-NVFP4 branch in priority and are untouched. Common SSM loads (A/B, conv1d, A_log, dt_bias, norm, ba_dense, qkvz_size) are hoisted once above all branches so the new fast path can share them.

Test plan

  • cargo check -p spark-model (ATLAS_SKIP_BUILD=1, CUDA stubbed) — clean
  • cargo fmt --check -p spark-model — clean
  • cargo clippy -p spark-model --tests (ATLAS_SKIP_BUILD=1) — no warnings
  • Runtime verification with an actual pre-quantized Standard-NVFP4 checkpoint on GB10 — see comment below

🤖 Generated with Claude Code

Some pre-quantized checkpoints (Standard NVFP4 convention, e.g.
sakamakismile) ship weights already packed as U8 NVFP4 on disk at the
`.weight` key rather than `.weight_packed` (the compressed-tensors
convention). Load these directly into QuantizedWeight, skipping the
BF16-dequant-then-requantize roundtrip:

- FullAttention Q/K/V/O (TP-sharded path): detect UInt8 dtype, load via
  quantized_auto(Nvfp4Variant::Standard) instead of dense_auto (which
  bails on UInt8). TP sharding of pre-quantized NVFP4 isn't supported yet
  (tp_size=1 only).
- LinearAttention (SSM/GDN): new native_nvfp4 fast path — when
  in_proj_qkv/in_proj_z/out_proj are all UInt8, load+concat+transpose them
  as QuantizedWeight directly. load_ssm_proj also gained a UInt8 branch
  for partial-U8 checkpoints that fall through to the legacy path.
  in_proj_a/b now route through load_ssm_proj instead of raw dense() so a
  U8-packed A/B doesn't get read as BF16 garbage.
- quant_helpers.rs: quantized_v2 now sets weight_scale_2_vec (field added
  since this code was originally written).

Split out of Avarok-Cybersecurity#218. Re-derived against current main rather than patched:
main has since added a native FP8 GDN fast path (fixing a ~7pt BFCL-ST
regression) and an ATLAS_GDN_BF16_WEIGHTS opt-in path, both of which sit
ahead of this new native-NVFP4 branch in priority and are left untouched.
Common SSM loads (A/B, conv1d, A_log, dt_bias, norm, ba_dense, qkvz_size)
are hoisted once above all branches so the new fast path can share them
without re-deriving the legacy dense-path ordering.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@Sujimoshi

Copy link
Copy Markdown
Contributor Author

Runtime verification on GB10 — done ✅

Checkpoint: Qwen3.6-27B-NVFP4 (local, modelopt-produced, confirmed Standard NVFP4 convention: dtype=U8 at .weight + separate weight_scale/weight_scale_2/input_scale, for both GDN/SSM projections and FullAttention Q/K/V/O — checked directly against the safetensors headers, not just key names).

Build: cargo build --release -p spark-server with ATLAS_TARGET_MODEL=qwen3.6-27b ATLAS_TARGET_QUANT=nvfp4 (needed explicitly — default kernel target doesn't match this checkpoint's hidden_size/model_type) and LIBRARY_PATH=/tmp/nccl-stubs (no system NCCL on this box). Both environment gaps, unrelated to this PR's code.

Server: spark serve --model-from-path .../Qwen3.6-27B-NVFP4 --kv-cache-dtype nvfp4 --gpu-memory-utilization 0.75 --max-batch-size 1 --max-seq-len 4096, TP=1. Weights loaded in ~7s, 19.15 GB.

Native-path confirmation:

  • GDN: all 48 SSM layers logged SSM[layers.N] native NVFP4 GDN: qkvz+out_proj loaded pre-quantized (U8-packed on disk; no BF16 dequant/requant roundtrip) — exactly 48/48, no fallbacks.
  • FullAttention Q/K/V/O: fast path confirmed in code (qwen35_dense.rs:291-309) and indirectly in logs — the only dense_auto/quantize_to_nvfp4 log hit in the entire run was the LM head (n=248077,k=5120, i.e. vocab×hidden), not any attention projection. If attention had fallen back to BF16-dequant, each of the 16 attention layers would've logged similarly — none did.

Smoke test: two prompts via /v1/chat/completions ("capital of France", "haiku about a GPU") — both coherent, correct, no repetition/degradation. ~14 tok/s single-batch (expected, no speculative decoding active), TTFT ~128ms.

Finding (not a bug, logging gap): the GDN branch has an explicit tracing::info! confirming the native path; the attention branch doesn't. Recommend adding an analogous log line in a follow-up so the fast path is directly log-visible instead of inferred from absence of fallback logs. Not fixed here since this PR was scoped to testing only.

Server was shut down cleanly after the test, no leftover process.

@rrstesiak

Copy link
Copy Markdown
Contributor

1. qwen35_dense.rs - documented constraint with no code behind it.
The FullAttention U8 fast path comment says "TP sharding of pre-quantized NVFP4
isn't supported yet (tp_size=1 only)" - but there is no tp_size check. At
tp_size>1 every rank silently loads the full unsharded weight and serves wrong
results. This needs && tp_size == 1 (falling through to the dequant path) or
a hard bail.

2. Same path - NULL dense weight handed to a caller that frees it.
The early return produces DenseWeight { weight: NULL }, and the call site
unconditionally gpu.free()s the dense intermediates after quantization.
Please confirm gpu.free(NULL) is a no-op on every backend this compiles for,
or skip the free for the U8 path.

3. quantized.rs::concat_rows - bit-exact f32 equality as a load gate.
ensure!(self.weight_scale_2 == other.weight_scale_2) means any checkpoint
whose quantizer emitted independent per-tensor scales for in_proj_qkv vs
in_proj_z fails at load. It fails loud (good), but this is a checkpoint-
convention assumption, not an invariant - worth a comment naming which
quantizer guarantees it, and an error message saying "re-quantize with X".

4. loaders_fp8.rs::quantize_to_nvfp4_per_row - dead code, wrong comment.
#[allow(dead_code)], no caller in this PR, and the comment says "use max for
compatibility diagnostics" while the code reads row 0. Belongs in the PR that
calls it, with the comment matching the code.

@Sujimoshi

Copy link
Copy Markdown
Contributor Author

Thanks for the thorough review — going through each:

1. TP-sharding gap on the U8 fast path — confirmed, real bug, you're right. There's no tp_size guard at all on that branch; at tp_size>1 it'd silently load the full unsharded weight per rank. Fixing now (loud bail on tp_size>1 rather than a silent wrong result, since sharding pre-quantized NVFP4 isn't implemented yet).

2. gpu.free(NULL) — checked all three backends (cuda_backend/gpu_impl.rs, metal_backend.rs, gpu/mock.rs): free() explicitly no-ops on a null pointer in all of them (CUDA and Metal check ptr.is_null() up front; the mock just no-ops removing a nonexistent key). So this is safe as-is, no code change needed there.

3. Bit-exact weight_scale_2 equality gate — agreed it's a checkpoint-convention assumption rather than a hard invariant. Adding a comment naming the quantizer this holds for and a clearer error message.

4. Dead quantize_to_nvfp4_per_row with a stale comment — confirmed it's new in this PR with no caller yet, and the comment contradicts the code (reads row 0, not max). Removing it from this PR — it belongs in whichever future PR actually wires up the per-row-scale GEMV path.

Pushing a fix for 1/3/4 shortly.

…ode, comments)

Per review from @rrstesiak:

- qwen35_dense.rs: the FullAttention U8 pre-quantized fast path had no
  tp_size check despite its comment claiming tp_size=1-only — at
  tp_size>1 every rank would have silently loaded the full unsharded
  weight (duplicated, not sharded) with no error. Now bails loudly with
  a clear message instead of producing silently wrong results.
- quantized.rs: documented why concat_rows' bit-exact weight_scale_2
  equality check is valid (ModelOpt/Standard NVFP4's single global
  per-tensor scale, not a general invariant) and improved the error
  message to say what checkpoint/quantizer this expects.
- loaders_fp8.rs: removed quantize_to_nvfp4_per_row — added in this PR
  with no caller anywhere in the codebase, and its comment ("use max
  for compatibility diagnostics") didn't match what the code actually
  did (reads row 0). Belongs in whichever future PR wires up the
  per-row-scale GEMV path.

(gpu.free(NULL) being a no-op on all three backends — the review's
other finding — was confirmed and answered directly in the PR thread,
no code change needed.)

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@Sujimoshi

Copy link
Copy Markdown
Contributor Author

Pushed fixes for 1/3/4 in `b3469ae`:

  1. TP-sharding guard: the U8 fast path now bails loudly (`anyhow::ensure!`) if `tp_size > 1`, instead of silently loading the full unsharded weight per rank.
  2. `concat_rows` comment/error message: clarified that the bit-exact `weight_scale_2` equality is valid specifically for ModelOpt/Standard NVFP4 (single global per-tensor scalar), not a general invariant across all conventions.
  3. Dead code removed: `quantize_to_nvfp4_per_row` deleted — confirmed no callers anywhere in the codebase and new to this PR; will come back in whichever future PR actually wires up the per-row-scale GEMV path.

`cargo check`/`fmt --check`/`clippy --tests -p spark-model` all clean after the changes. Thanks again for catching the TP bug — that one was real.

@pragmaxim

Copy link
Copy Markdown
Contributor

@Sujimoshi I just tested it with :

./target/release/spark serve morosystems/ThinkingCap-Qwen3.6-27B-NVFP4 \
      --port 8888 \
      --host 0.0.0.0 \
      --max-seq-len 65536 \
      --kv-cache-dtype bf16 \
      --gpu-memory-utilization 0.9 \
      --scheduling-policy slai \
      --enable-prefix-caching \
      --tool-call-parser qwen3_coder

And it fixed the error 👍

rsafier added a commit to MonumentalSystems/atlas that referenced this pull request Jul 8, 2026
…ision 256K PASS) + Nemotron-Puzzle-75B bring-up scope

ThinkingCap-Qwen3.6-27B-NVFP4: Avarok-Cybersecurity#257 native U8-NVFP4 load, prefill ~1.1-1.4K
tok/s, decode ~12-14 tok/s (dense hybrid), agentic 3/3, Mona-Lisa vision PASS
at --vision-max-pixels 262144.

Nemotron-3-Puzzle-75B-A9B-NVFP4: assessed not-loadable (heterogeneous
nemotron_h_puzzle) — dispatch + per-block MoE + null layer-count + shared-expert
gaps documented; assessment-only, deferred.

Claude-Session: https://claude.ai/code/session_01Hyvzj7tixGc2o6jpefko8x
@SeedSource

Copy link
Copy Markdown
Collaborator

[BETA] Automated PR review — SeedSource AI pipeline

Advisory only, not a merge gate. AI-generated (internal risk-triage, Atlas serve gate, GLM-5.2 review), attributed per Atlas convention.

Triage

  • Risk: high - serve-test-required (triage advisory): True - path-heuristic (authority): True
  • Rationale: Adds NVFP4 pre-quantized U8 handling and enforces tp_size=1 to prevent silent TP sharding failures.
  • Serving-path files: crates/spark-model/src/mistral_loader/loader_impl/phase_assemble.rs, crates/spark-model/src/tp_shard/quant_shard.rs, crates/spark-model/src/weight_loader/qwen35_dense.rs, crates/spark-model/src/weight_loader/step3p7.rs, crates/spark-model/src/weight_loader/step3p7/load_layers.rs, crates/spark-model/src/weight_map/loaders_fp8.rs

Serve gate

  • Build: 19s - Qwen3.6-35B-A3B-FP8 on GB10 node1
  • Top-k KL vs main: 0.0002 (no shift vs main) - decode delta -0.1%
  • Coherence probe: coherence=pass ping=10/10 (note: our test_coherence.py is flaky here; KL is the reliable serve signal)

GLM-5.2 code review

Correctness concerns:

  1. SSM out_proj field mismatch: In the native NVFP4 path, SsmWeights.out_proj is assigned out_proj_nvfp4 (a QuantizedWeight), but the legacy path assigns out_proj_dense (a DenseWeight). If SsmWeights.out_proj is typed DenseWeight, this won't compile; if it's a union/enum, the downstream consumer must handle both — not visible in this diff. Verify the type and that Qwen3SsmLayer::new_sequential + predequant_for_prefill correctly handle the quantized out_proj path.

  2. weight_scale_2_vec propagation: Added to struct construction in two places but no sharding logic for it. For tp_size=1 this is fine (pass-through), but the field is new — confirm it's initialized everywhere else or defaults to NULL.

  3. in_proj_a/in_proj_b now routed through load_ssm_proj with nv/h dims instead of raw dense(). If A/B were previously BF16 and are now potentially U8-dequantized, confirm interleave_ba still receives BF16. The comment claims this is correct but no assertion enforces it.

  4. quantized_auto for SSM projections — no tp_size guard here unlike the attention path. If tp_size > 1 reaches this code, quantized_auto may load full unsharded weights silently. Add the same anyhow::ensure!(tp_size == 1) guard.

Risk: Medium. The SSM native path is substantial new code with no visible tests. Partial-U8 checkpoints fall through to legacy path — verify load_ssm_proj's new U8 branch produces correct shapes for concat.

Quality: Comments are excessive (multi-paragraph essays). Trim to essential context.

Pipeline: build+serve on GB10, GLM-5.2 via NVIDIA NIM, internal risk-triage classifier. External diff treated as untrusted input. Public repo.

@tbraun96
tbraun96 merged commit e220d40 into Avarok-Cybersecurity:main Jul 10, 2026
12 checks passed
AzeezIsh pushed a commit that referenced this pull request Jul 10, 2026
…s, end-to-end on GB10 (#293)

* feat(deepseek-v4): native MXFP4 routed-expert loader (transcode-free) — ARM-2 Phase L

Serve DeepSeek-V4-Flash native MXFP4 experts without the two lossy 4-bit
conversions the old UInt8 arm did (dequant_nvfp4_e8m0_to_bf16 -> quantize_to_nvfp4
= MXFP4 E8M0 -> BF16 -> standard-NVFP4). Land packed I8 nibbles + F8_E8M0
per-32 scales device-resident UNCHANGED; tag the MoE layer Mxfp4E8m0 so the
Phase-K E8M0 GEMM variants dispatch on it.

- WeightQuantFormat::Mxfp4E8m0 variant (quantized.rs)
- quantized_mxfp4_e8m0() passthrough loader fn, asserts GROUP_SIZE=32 (model_a.rs)
- load_expert_proj UInt8 arm -> native passthrough (assemble.rs)
- detect_routed_scale_kind() + MoeLayer.experts_scale_kind tag (assemble.rs/moe/mod.rs/init.rs)
- retain dequant_nvfp4_e8m0_to_bf16 as Phase-K Leg-2 host-ref (allow dead_code)

cargo check -p spark-model clean. Kernel dispatch (consuming the tag) = Phase K.
Leg-1 byte-check: device bytes == on-disk MXFP4 by construction; disk ref
sha256 banked (spark-bench .planning/ARM-2-LEG1-BYTE-CHECK.md). No eval fired.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* ARM-2 Phase-K STEP-0: ATLAS_DUMP_EXPERT0 device byte-check hook (Leg-1 DEVICE-VERIFIED)

One-shot copy_d2h dump of layers.0.ffn.experts.0.w1 {weight,scale} device
buffers, gated by ATLAS_DUMP_EXPERT0=1, pre-transpose. Device sha256 matched
on-disk MXFP4 targets 177ac128.../ea4ac989... byte-for-byte on EP=2 n3/n4 —
loader is device-proven transcode-free. Kept env-gated for re-verification.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* ARM-2 Phase-K Family A (decode): E8M0 dequant variants for transposed MoE GEMV

moe_shared_expert_fused_t.cu: extract gate_up + silu_down decode bodies into
template<int GS,bool E8M0> __device__ impls; emit _e8m0 extern-C wrappers
(<32,true>) alongside the NVFP4 defaults (<GROUP_SIZE,false>). mx_block_scale<E8M0>
bit-constructs 2^(sb-127) (from_bits(sb<<23), NaN/zero guard) — byte-exact to Rust
fp8_e8m0_to_f32, no exp2f drift. GROUP_SIZE 16->32, scale2 dropped on E8M0 path.
Routing/index/combine untouched. PTX-verified: shl<<23 present, no e4m3 on e8m0
path, guard = ((sb+1)&255)<2.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* ARM-2 Phase-K Family B (prefill): E8M0 dequant variants for W4A16 MoE GEMM

RIDER 1: extract mx_block_scale<E8M0> + atlas_dec_e4m3 to shared header
kernels/gb10/common/mx_block_scale.cuh — ONE E8M0 scale primitive across
Family A (decode) and Family B (prefill). Family A refactored to include it;
PTX byte-identical to d951661 (pure no-op, verified).

Family B: all 5 W4A16 entries in moe_w4a16_grouped_gemm.cu templated
<int GS, bool E8M0>, extern-C wrappers <GROUP_SIZE,false> (NVFP4) +
<32,true> (_e8m0, native MXFP4 per-32, no global):
  1 moe_w4a16_grouped_gemm_ptrtable      (BF16-MMA, inline; compile-time
    if(E8M0) split keeps NVFP4 (LUT*fp8)*scale2 order byte-identical)
  2 moe_w4a16_grouped_gemm_ptrtable_t    (K32, 2->1 group)
  3 moe_w4a16_grouped_gemm_ptrtable_t_k64(K64, 4->2 group)
  4 moe_w4a16_fused_gate_up_t_k64        (K64, 4->2 group; V4-hit primary)
  5 moe_w4a16_fused_gate_up_t            (K32, 2->1 group)
Hand-unrolled per-quarter dequant macros replaced by GS-generic loop
(sv[K_STEP/GS] = mx_block_scale<E8M0>; s = sv[kp/(GS/2)]). For the FP8-MMA
entries mx_block_scale<false> returns exactly the old (float)f*scale2, so the
NVFP4 path is numerically identical (PTX differs only in reg-alloc/scheduling;
entries 1-2 diffs confirmed = local smem labels + BB index + reg renumber).

SKIPPED (out of scope): moe_fp8_grouped_gemm_ptrtable_t (FP8 experts),
*_fp4 (W4A4 fp4-activation) — V4 routes W4A16 not A4.

Grep-gate PASS on all 5 _e8m0 entries: shl<<23 present, cvt.f32.e4m3
(NVFP4 scale decode) = 0, ex2 = 0. nvcc --ptx -arch=sm_121a clean.

Leg-2 host-ref + cross-family (RIDER 2) + host dispatch = next.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Phase-K: move Family B E8M0 edits to the correct V4 build target

PROVENANCE CORRECTION. The Phase-K spec pinned the Family B kernel at
kernels/gb10/qwen3.6-35b-a3b/nvfp4/moe_w4a16_grouped_gemm.cu, but the
deepseek-v4-flash build compiles a MODEL-SPECIFIC OVERRIDE copy
(kernels/gb10/deepseek-v4-flash/nvfp4/moe_w4a16_grouped_gemm.cu) —
collect_cu_files() shadows common/qwen by stem, model dir wins. The qwen
copy is the 35B target, never compiled for the V4 serve.

- Ported the 5-entry E8M0 templating (commit 5516e0e) verbatim to the
  deepseek-v4-flash override (entries 1-5 were byte-identical to the qwen
  baseline; git patch applied clean, 25 hunks). Include path
  ../../common/mx_block_scale.cuh resolves (same dir depth).
- Reverted the qwen3.6-35b-a3b copy to baseline (wrong target; the e8m0
  variants are V4-only dead weight there).
- deepseek copy has no FP4 (W4A4) entries — only the 10 W4A16 + FP8.

Family A (moe_shared_expert_fused_t.cu) is common-only, no deepseek
override -> the V4 build already compiles the edited file. Correct.

nvcc --ptx -arch=sm_121a clean on the deepseek copy; grep-gate PASS on
all 5 _e8m0 entries (shl<<23 present, cvt.f32.e4m3=0, ex2=0).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Phase-K: dual-format fused decode kernel (Mike ruling, Option A + riders)

The native V4 ckpt is heterogeneous: routed experts MXFP4-E8M0, shared expert
FP8->NVFP4. The fused decode kernel moe_expert_*_shared_t serves BOTH in one
launch, so a single <GS,E8M0> template mis-decodes one. Ruling: keep the fused
kernel (execution shape == 15/40 baseline, one variable), make it dual-format.

- Template <int GS_R, bool E8M0_R, int GS_S, bool E8M0_S>: routed region uses
  (GS_R,E8M0_R), shared region (GS_S,E8M0_S). Block-uniform on is_shared
  (grid.y = expert_slot, one expert/block — RIDER A2).
- Accumulation extracted to a single parameterized macro (GATEUP_ACCUM /
  SILUDOWN_ACCUM) so routed and shared run byte-identical logic differing only
  in (GS,E8M0) — no drift (RIDER A1/A3).
- `if constexpr (GS_R==GS_S && E8M0_R==E8M0_S)` collapses same-format wrappers
  to a SINGLE loop: NVFP4 wrappers <GROUP_SIZE,false,GROUP_SIZE,false> emit PTX
  BYTE-IDENTICAL to baseline d951661 (RIDER A3, verified). Only the e8m0
  wrapper <32,true,GROUP_SIZE,false> (routed E8M0, shared NVFP4) emits the branch.

Verified (nvcc --ptx sm_121a): NVFP4 gate_up + silu_down PTX-identical to
baseline; e8m0 routed = shl<<23 bit-construct, zero scale-path ex2 (the 5 ex2
in silu_down e8m0 == NVFP4 silu_down = SiLU __expf, unchanged); e8m0 shared =
NVFP4 e4m3 decode (correct heterogeneous). RIDER A1 (tag-keyed, not positional)
+ RIDER A4 (mixed-fusion Leg-2) land in the Rust dispatch + Leg-2 next.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Phase-K: host dispatch (E8M0 select + straggler net) + Family C

Wires the native-MXFP4 E8M0 kernels into MoE dispatch. cargo check -p
spark-model clean under #![deny(warnings)].

Handles (mod.rs + init.rs): 7 _e8m0 KernelHandle fields (5 prefill W4A16 +
2 dual-format decode) resolved via try_kernel (0 on targets that don't ship
them). e8m0_or(nvfp4, e8m0, site) select helper — panics if E8M0 selected but
the variant is absent (straggler net, not silent NVFP4-on-E8M0).

RIDER A1 (tag-keyed): shared_experts_scale_kind field + detect_shared_scale_kind
(assemble.rs) on ffn.shared_experts.w1. Native V4 shared expert is FP8->NVFP4 =>
tag Nvfp4. Decode dispatch asserts shared==Nvfp4 via `expect` before using the
_e8m0 kernel (which is <32,true,GROUP_SIZE,false> = routed E8M0 / shared NVFP4).

Prefill (forward_prefill_routed.rs): leading `if experts_scale_kind==Mxfp4E8m0`
branch at the gate_up + down transposed chokepoints -> _e8m0 handles; this
structurally bars E8M0 from the NVFP4-only cutlass/fp4/m128/fp8_down sub-paths.
`expect(Nvfp4)` straggler net on the two non-transposed else fallbacks.

Decode (forward_batched.rs + forward_phase.rs): e8m0_or select at gate_up +
silu_down launches + the shared-format expect guard.

Family C (quantized.rs + helpers_a.rs): transpose_for_gemm_gs(.., group_size)
(transpose_for_gemm delegates 16). Routed experts transpose per-32 for E8M0
(both prefill-transpose paths); shared expert stays per-16 NVFP4. Without this
the E8M0 kernels read a mis-shaped ([N,K/16] vs [N,K/32]) scale table.

Kernel registration is automatic (per-module cuModuleGetFunction) — no manifest.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Phase-K Leg-2: numeric gate harness — ALL CHECKS PASS (bit-exact)

Leg-2 (spec §SESSION-3, Mike-blessed method) run on n3 GB10:
- CHECK 1 Family A decode e8m0 vs host f32 GEMV (full-range E8M0): 256/256 exact, maxULP 0
- CHECK 2 RIDER A3 NVFP4 shared branch baseline-vs-e8m0: bit-identical
- CHECK 3 RIDER A4 mixed-fusion (routed-E8M0 + shared-NVFP4): no cross-contamination
- CHECK 4 Family B 5 W4A16 entries x K{64,128,448}: bit-exact vs NVFP4-equivalent
- RIDER 2: two independent anchors (host f32 GEMV + shipping-NVFP4 kernel)

Method: prefill kernels are FP8-MMA (both operands -> E4M3) so a plain host GEMM
cannot match within bf16 tol; kernel-vs-NVFP4-equivalent (power-of-2 scale) bit-exact
isolates the one changed line (K/16->K/32 group-count threading). Unique-per-group
scales (rider-1) so a group-index slip yields a bit diff.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Leg-2 CHECK 5: real-shape multi-expert integration gate — memcheck CLEAN (0 errors)

Routed E8M0 GEMM prefill path (fused_gate_up_t_k64_e8m0 + ptrtable_t_k64_e8m0 down)
at real V4 dims (dim=4096, inter=2048, 256 experts, top_k=6, 1536 rows, real
expert_offsets/sorted_token_ids/A-gather) is memory-safe under compute-sanitizer.
The new gate the CELL-BLOCKED scar demanded (numeric bit-exact != integration safety).

Tree-split: the layer-1 OOB is NOT in the routed GEMM path (dispatch E8M0-guarded,
kernels index scales via template GS not GROUP_SIZE, scale2_vals full [ne]). Crash
geometry grid=[11] block=[256] shmem=0 != GEMM shape confirms. Fault is upstream/
adjacent (routing/HC/attention/index-hash/EP all-to-all).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 CHECK 5: add mt>=2 imbalance case (hash-layer condition) — memcheck CLEAN (max/expert=202, mt=4)

Routed E8M0 GEMM exonerated on the multi-m-tile axis too. Hash-MoE layers
(layer_idx<num_hash_layers=3) route via static tid2eid = unbalanced → one expert
>M_TILE=64 → mt>=2; that path is memory-safe (0 memcheck errors). Routed GEMM now
cleared 4 ways (dispatch, scale indexing, real-shape mt=1, real-shape mt=4).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2: wire deepseek_v4 into the MoE prefill-transpose gate

DeepSeek-V4 native-MXFP4 routed experts were never getting transposed
prefill tables (gate_ptrs_t/up_ptrs_t/down_ptrs_t): the only model-level
caller of transpose_moe_for_prefill* (factory/m2_setup) gated on
model_type in {minimax_m2, step3p7}, and the deepseek_v4 loader (unlike
qwen35/qwen3/mistral) never calls transpose_for_prefill. Result:
gate_ptrs_t=None -> run_routed_grouped_gemm takes the non-transposed
fallback, which reads the E8M0 [N,K/32] scales as NVFP4 [N,K/16] -> OOB
on a stale binary, or panics at the expect(Nvfp4) straggler-net at HEAD.

Allow deepseek_v4 through the gate so it runs transpose_for_prefill_unified
(routed_gs=32 for Mxfp4E8m0, shared expert NVFP4) and routed prefill takes
the validated moe_w4a16_fused_gate_up_t_k64_e8m0 (Leg-2 bit-exact,
CHECK-5 memcheck-clean). Requires ATLAS_UNIFIED_MOE_LAYOUT=1 (unified frees
originals to fit the EP=2 budget; use_t_layout_for_decode() routes decode to
the _e8m0 _t kernels). Additive -- minimax_m2/step3p7 match earlier in the
gate, dispatch byte-identical.

Verified: EP=2 n3/n4 serve loads clean, q1 prefills coherent (no OOB/panic);
core-8 discriminator runs all 8 (serve alive throughout).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Leg-2: relocate 845-LoC example harness → #[ignore]d integration tests

The native-MXFP4 (E8M0) numeric gate lived as examples/arm2_leg2.rs (845 LoC),
which trips the ≤500-LoC .rs cap (file-size-cap.yml). Relocate it as the proper
regression net for the E8M0 format path, split under the cap:

  tests/arm2_common/support.rs   (391) — shared PRNG/quant/upload/compare/launch
  tests/arm2_leg2_decode.rs      (269) — Family A: CHECK 1 host-GEMV, 2 RIDER-A3, 3 RIDER-A4
  tests/arm2_leg2_prefill.rs     (256) — Family B: CHECK 4 bit-exact, 5 multi-expert memcheck

Each CHECK is a #[test] #[ignore] fn (Requires GB10 GPU). CI builds+links them
against the libcuda stubs (catches kernel-signature drift) but skips execution;
run on a GB10 host with `cargo test -p spark-model --test arm2_leg2_* -- --ignored`.
Check bodies transcribed verbatim (same kernels, arithmetic, launches).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2: rebase fixup — set QuantizedWeight.weight_scale_2_vec (NULL) in native-MXFP4 loader

avarok/main #257 (native U8 NVFP4) added a per-output-row `weight_scale_2_vec:
DevicePtr` field to QuantizedWeight. Our transcode-free deepseek_v4 loader
constructor (model_a.rs native_mxfp4) predates it → E0063 missing-field after
rebase. Native MXFP4 uses the scalar `weight_scale_2` (E8M0 per-group), not the
per-row vector, so NULL — matching every sibling constructor (#257 convention).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2 Leg-2: one #[test] per binary — fix per-test CUDA-context death

The AtlasRegistry singleton's CUDA context is current only on the thread that
first initialized it; cargo runs each #[test] on its own thread, so 5 separate
backend-init tests failed after the first (cuCtxGetCurrent → ctx 0x0). Mirror the
original single-`main` harness: one #[test] per binary (leg2_family_a_decode,
leg2_family_b_prefill) inits the backend once and runs its checks sequentially on
that thread. Checks are now plain (gpu, st) helpers. GB10 verified:
  decode  — CHECK 1 256/256 maxULP 0; CHECK 2/3 bit-identical 0-diff
  prefill — CHECK 4 all 5 entries bit-exact; CHECK 5 memcheck ERROR SUMMARY 0

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

* ARM-2: curate for CI green — split model_a/moe under 500-LoC cap + fmt

The native-MXFP4 feature pushed two files over the ≤500-LoC cap (main→ours:
model_a.rs 492→534, moe/mod.rs 490→546; neither allow-listed). Split each into
a cohesive sibling module, re-exported so all call sites are unchanged:
  weight_map/fp8_dequant.rs   (187) — FP8-E4M3→BF16 per-tensor dequant helpers
  layers/moe/ptr_table_build.rs (135) — NVFP4/BF16/FP8 expert ptr-table builders
Leaves model_a.rs 364, moe/mod.rs 425. Also `cargo fmt` (feature files carried
unformatted `==` line-wraps from prior sessions) + drop now-unused imports.

CI mirror GREEN: cargo fmt --check, cargo clippy --workspace --tests, ≤500-LoC
gate. GB10 re-verify after the split (behavior-preserving): Leg-2 CHECK 1 256/256
maxULP 0; CHECK 2/3 bit-identical; CHECK 4 all 5 entries bit-exact; CHECK 5 clean.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XPm7ArAwQKXEx6dkjVBuKx

---------

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
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.

5 participants