From 08c4781266084b32ad643106edf7b647f0d32dfd Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 20:16:51 +0000 Subject: [PATCH 1/4] spike(minimax-h3): instrument DiT-input dump for the S1 render-close hunt The #70 white-latent render bug is cornered to WHAT the driver feeds the DiT at real scale (#74 proved the forward MATH exact at every reduced-dim geometry, but the ladder fed RANDOM inputs). Add an env-gated `VT_H3_DUMP_INPUTS=` that, at denoise step 0, writes every DiT input as raw little-endian binary + a text manifest: the packed layout (input_ids/masks/img_pos/audio_pos/text_pos/ update_mask/cu_seqlens/document_id), the fp64 position grid, the per-token modality tags, the per-token pre-unique timesteps and their unique/inverse/ combined AdaLN selection, both sigma schedules, and the raw prompt_embeds. The `minimax-h3-gen` driver additionally dumps prompt_token_ids so the tokenization can be diffed against upstream `minimax_h3_text_only_ids` (verbatim prompt, add_special_tokens=False). Byte-identical to production when unset (every path is guarded; no file opened). Documented in docs/ENVIRONMENT.md. This is the S1 instrumentation only; the diff against upstream pipeline_minimax_h3.py runs next. Pre-existing preflight red (check-fusion-consistency minimax_h3_video_vae_device) is unrelated to this change and untouched. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- docs/ENVIRONMENT.md | 1 + examples/minimax_h3_gen/main.cpp | 13 +++ src/vllm/model_executor/models/minimax_h3.cpp | 80 +++++++++++++++++++ 3 files changed, 94 insertions(+) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index e0babdc5..038e618a 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -103,6 +103,7 @@ Read-only observability; none change output. | `VT_H3_TRACE_MOTION` | unset | `=1` prints one `[h3-motion] step ...` line per MiniMax-H3 denoise step to stderr: the step's velocity stats (`v_rms`/`v_amax`/`v_mean` of the DiT output), the per-step latent motion over the denoise-target rows (`drows_rms`), and the running latent norm (`rows_rms`). Because the rectified-flow Euler integration telescopes to `(sigma0 - sigmaN) * v`, a velocity that does not EVOLVE across steps produces a step-count-invariant result; this trace measures exactly that (added for the render-coherence bisection). Byte-identical when unset — every read is guarded and it only reads buffers the loop already holds | | `VT_H3_VAE_PROBE` | unset | `=1` runs a video-VAE receptive-field probe after the normal decode: it perturbs ONE interior spatial latent cell (across all channels and temporal frames), re-decodes, and prints a per-16px-block RMS-change map (`[h3-vae-probe]`) over output frame 0. If only the perturbed cell's block moves, the ViT3D decoder is not mixing tokens spatially. Byte-identical to production when unset (no second decode) | | `VT_H3_DUMP_DIR` | unset | Directory into which the MiniMax-H3 denoise loop writes the initial and final video latent rows (`init_video_rows.f32`, `final_video_rows.f32`) and the pipeline writes the exact VAE-input latent (`vae_input_video_latent.f32`), all raw little-endian f32. Lets two runs (e.g. 12 vs 50 steps, conditioned vs not) be byte/stat-compared, and the video VAE decode be replayed on a KNOWN latent, without re-running the denoise. Byte-identical to production when unset (no file is opened) | +| `VT_H3_DUMP_INPUTS` | unset | Directory into which the MiniMax-H3 denoise loop writes EVERY DiT input at step 0 as raw little-endian binary plus a `manifest.txt` — the packed layout (`input_ids`/`image_mask`/`audio_mask`/`img_pos`/`audio_pos`/`text_pos`/`update_mask`/`cu_seqlens`/`document_id`), the fp64 position grid (`img_position_ids.f64`), the per-token modality tags (`token_tags.i64`), the per-token pre-unique timesteps and their `unique_timesteps`/`inverse_indices`/`combined_indices` AdaLN selection, both sigma schedules, and the raw `prompt_embeds`; the `minimax-h3-gen` driver additionally writes `prompt_token_ids.i32`. Lets the REAL-scale DiT inputs be diffed EXACTLY against upstream `pipeline_minimax_h3.py` (the render-coherence S1 surface the reduced-dim ladder never fed real values into). Byte-identical to production when unset (no file is opened) | ## Kernel-internal knobs (deferred) diff --git a/examples/minimax_h3_gen/main.cpp b/examples/minimax_h3_gen/main.cpp index a4dda5c2..e7e434cb 100644 --- a/examples/minimax_h3_gen/main.cpp +++ b/examples/minimax_h3_gen/main.cpp @@ -487,6 +487,19 @@ int main(int argc, char** argv) { const std::vector ids = tokenizer.Encode(prompt); VT_CHECK(!ids.empty(), "minimax-h3-gen: the prompt tokenized to nothing"); std::cerr << " prompt tokens = " << ids.size() << "\n"; + // DIAGNOSTIC (env-gated): dump the raw prompt token ids so the tokenization + // can be diffed against upstream `minimax_h3_text_only_ids` (verbatim prompt, + // add_special_tokens=False). A BOS/template mismatch shifts every text row and + // feeds the 32B tower a different string -> different conditioning. + if (const char* dd = std::getenv("VT_H3_DUMP_INPUTS")) { + if (std::FILE* fp = std::fopen((std::string(dd) + "/prompt_token_ids.i32").c_str(), "wb")) { + std::fwrite(ids.data(), sizeof(int32_t), ids.size(), fp); + std::fclose(fp); + std::cerr << " [h3-dump-inputs] prompt_token_ids.i32 (" << ids.size() << " ids): "; + for (size_t k = 0; k < ids.size() && k < 64; ++k) std::cerr << ids[k] << " "; + std::cerr << "\n"; + } + } const std::vector embeds = vllm::MiniMaxH3EncoderEmbedTokens(enc, ids); // Text-only: all three M-RoPE axes are the token index. const int64_t seq = static_cast(ids.size()); diff --git a/src/vllm/model_executor/models/minimax_h3.cpp b/src/vllm/model_executor/models/minimax_h3.cpp index b06239e6..12efd4c5 100644 --- a/src/vllm/model_executor/models/minimax_h3.cpp +++ b/src/vllm/model_executor/models/minimax_h3.cpp @@ -891,6 +891,86 @@ MiniMaxH3DenoiseResult MiniMaxH3DenoiseLoop( in.refiner_cu_seqlens = refiner_cu.data(); in.num_refiner_cu_seqlens = static_cast(refiner_cu.size()); + // DIAGNOSTIC (env-gated, byte-identical when unset): VT_H3_DUMP_INPUTS= + // dumps every DiT input at STEP 0 as raw little-endian binary plus a text + // manifest, so the real-scale driver's DiT inputs can be diffed EXACTLY against + // upstream pipeline_minimax_h3.py's construction (packed layout, fp64 position + // grid, per-token modality tags, per-token timestep -> unique/inverse -> + // combined AdaLN index) and the encoder conditioning statistically. This is the + // S1 surface #70/#74 never isolated: the ladder fed RANDOM inputs; the real + // render's INPUTS are the untested corner. Only step 0 (the layout and the + // timestep partition are the same shape every step). + if (step == 0) { + if (const char* dump_inputs_dir = std::getenv("VT_H3_DUMP_INPUTS")) { + const std::string dir(dump_inputs_dir); + auto wr = [&](const char* name, const void* p, size_t bytes) { + std::FILE* f = std::fopen((dir + "/" + name).c_str(), "wb"); + if (f == nullptr) return; + std::fwrite(p, 1, bytes, f); + std::fclose(f); + }; + // per-token pre-unique timesteps (before torch.unique) and combined AdaLN idx + std::vector combined_dump(static_cast(seq_len)); + for (int64_t i = 0; i < seq_len; ++i) { + const int64_t tag = branch.token_tags[static_cast(i)] < 0 + ? 0 + : branch.token_tags[static_cast(i)]; + combined_dump[static_cast(i)] = + inverse[static_cast(i)] * kMiniMaxH3AdalnModalityNum + tag; + } + wr("timesteps_pretoken.f32", timesteps.data(), timesteps.size() * sizeof(float)); + wr("unique_timesteps.f32", unique.data(), unique.size() * sizeof(float)); + wr("inverse_indices.i64", inverse.data(), inverse.size() * sizeof(int64_t)); + wr("combined_indices.i64", combined_dump.data(), combined_dump.size() * sizeof(int64_t)); + wr("token_tags.i64", branch.token_tags.data(), + branch.token_tags.size() * sizeof(int64_t)); + wr("img_position_ids.f64", packed.img_position_ids.data(), + packed.img_position_ids.size() * sizeof(double)); + wr("input_ids.i64", packed.input_ids.data(), packed.input_ids.size() * sizeof(int64_t)); + wr("image_mask.u8", packed.image_mask.data(), packed.image_mask.size()); + wr("audio_mask.u8", packed.audio_mask.data(), packed.audio_mask.size()); + wr("img_pos.i64", packed.img_pos.data(), packed.img_pos.size() * sizeof(int64_t)); + wr("audio_pos.i64", packed.audio_pos.data(), packed.audio_pos.size() * sizeof(int64_t)); + wr("text_pos.i64", packed.text_pos.data(), packed.text_pos.size() * sizeof(int64_t)); + wr("update_mask.u8", packed.update_mask.data(), packed.update_mask.size()); + wr("audio_update_mask.u8", audio_update.data(), audio_update.size()); + wr("cu_seqlens.i32", packed.cu_seqlens.data(), + packed.cu_seqlens.size() * sizeof(int32_t)); + wr("document_id.i64", packed.document_id.data(), + packed.document_id.size() * sizeof(int64_t)); + wr("sigmas_video.f64", sigmas_video.data(), sigmas_video.size() * sizeof(double)); + wr("sigmas_audio.f64", sigmas_audio.data(), sigmas_audio.size() * sizeof(double)); + wr("prompt_embeds.f32", branch.text_embeddings.data(), + branch.text_embeddings.size() * sizeof(float)); + std::FILE* mf = std::fopen((dir + "/manifest.txt").c_str(), "wb"); + if (mf != nullptr) { + std::fprintf(mf, + "seq_len=%lld\nnum_unique_timesteps=%lld\nnum_img_pos=%lld\n" + "num_audio_pos=%lld\nnum_text_pos=%lld\ntext_dim=%lld\n" + "video_row_width=%lld\naudio_latents_dim=%lld\nnum_steps=%lld\n" + "s_v0=%.17g\ns_a0=%.17g\nt_v0=%.17g\nt_a0=%.17g\n" + "imgvid_cond_t0=%.17g\naudio_ref_cond_t0=%.17g\n" + "prompt_embeds_rows=%lld\n", + static_cast(seq_len), + static_cast(unique.size()), + static_cast(num_img), + static_cast(num_audio), + static_cast(packed.text_pos.size()), + static_cast(params.text_dim), + static_cast(video_width), + static_cast(audio_width), + static_cast(num_steps), s_v, s_a, t_v, t_a, imgvid_cond_t, + audio_ref_cond_t, + static_cast(branch.text_embeddings.size() / + (params.text_dim > 0 ? params.text_dim : 1))); + std::fclose(mf); + } + std::fprintf(stderr, "[h3-dump-inputs] wrote DiT step-0 inputs to %s (seq_len=%lld)\n", + dir.c_str(), static_cast(seq_len)); + std::fflush(stderr); + } + } + const auto step_t0 = now(); const MiniMaxH3DitOutputs velocity = on_device ? MiniMaxH3DitForwardDevice( From f959742179033d3772dc190aaef9b3e9de465e5d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 20:44:41 +0000 Subject: [PATCH 2/4] test(minimax-h3): CUDA device forward vs host at the REAL render seq (1920) S1 is exonerated: the real-scale DiT INPUTS (packed layout, fp64 position grid, token tags, inverse/combined AdaLN indices, sigmas) diff EXACTLY against upstream pipeline_minimax_h3.py at 512x512/22f (verified on dgx: text_len=8, latent 7x32x32, seq_len 1920); the tokenization matches upstream byte-for-byte; the encoder conditioning is correctly shaped and carries the expected Qwen massive-activation structure; and DequantNvfp4ToBf16 is the shared helper the Laguna and Qwen3-32B NVFP4 arms already prove byte-exact. That leaves ONE untested surface: #74's device-vs-host ladder (incl. the REAL head_dim=128 case) runs on the CPU BACKEND, and the CUDA cases run only at the small fl2va geometry. The CUDA kernels at the REAL render seq (t2va -> latent 7x32x32 -> seq_len 1920, cu_seqlens=[0,1874,1920] non-causal 2-document) at head_dim=128 have never been gated against the trusted host loops. A scale-dependent CUDA-kernel bug (varlen non-causal attention / RoPE cache / AdaLN modulate) would leave every reduced-dim gate green while the render goes white. This case runs the SAME MiniMaxH3DitForwardDevice on the CUDA backend vs the CPU host forward at exactly the render geometry and the step-0 timestep partition. A divergence IS the #70 bug; a match points the hunt at S2 (a shared port<->RefDiT restatement blind spot vs true upstream). Skips without a CUDA backend. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- tests/vllm/models/test_minimax_h3.cpp | 99 +++++++++++++++++++++++++++ 1 file changed, 99 insertions(+) diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 3ca1c069..ae995e51 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -1042,6 +1042,105 @@ TEST_CASE("minimax_h3: the DEVICE-resident DiT forward matches upstream on CUDA" CheckDeviceForward(q, "cuda-device-forward"); } +// H3-RENDER-CLOSE: the one surface #74 left untested. The "REAL head_dim=128 +// ratio" device-vs-host case above runs on the CPU BACKEND, and the CUDA cases run +// only at the SMALL fl2va geometry (seq ~<200). The #70 white latent is a REAL +// render: CUDA kernels at REAL SEQ (t2va 512x512/22f -> latent 7x32x32 -> +// seq_len 1920, cu_seqlens=[0,1874,1920] non-causal 2-document) at head_dim=128. +// If a CUDA kernel (varlen non-causal attention, RoPE cache, AdaLN modulate) has a +// scale-dependent bug that its CPU counterpart does not, the render is white while +// every reduced-dim gate is green. This runs the SAME MiniMaxH3DitForwardDevice on +// the CUDA backend vs the trusted CPU host loops at exactly the render geometry and +// step-0 timestep partition; a divergence here IS the bug, a match points at S2. +TEST_CASE("minimax_h3: CUDA device forward tracks the host at the REAL render seq (1920)") { + vt::Backend* cuda = nullptr; + try { + cuda = &vt::GetBackend(vt::DeviceType::kCUDA); + } catch (...) { + MESSAGE("SKIP: no CUDA backend registered"); + return; + } + const MiniMaxH3DitParams p = RealRatioParams(); // head_dim=128, rot_dim=96 + const std::unique_ptr weights = BuildGoldenWeights(p); + + // The real t2va render geometry at 512x512/22f (verified on dgx: text_len=8, + // latent 7x32x32, audio_t=37, seq_len 1920). + const int64_t text_len = 8, latent_t = 7, latent_h = 32, latent_w = 32; + const int64_t audio_t = 37, audio_channel = 2; + const MiniMaxH3PackedSequence packed = BuildMiniMaxH3PackedSequence( + text_len, latent_t, latent_h, latent_w, audio_t, audio_channel, + /*include_keyframe_cond=*/false, {}, /*frame_count=*/0); + const int64_t seq_len = packed.seq_len; + const int64_t video_width = p.video_row_width(); + const int64_t num_img = static_cast(packed.img_pos.size()); + const int64_t num_audio = static_cast(packed.audio_pos.size()); + const int64_t num_text = static_cast(packed.text_pos.size()); + REQUIRE(seq_len == 1920); + REQUIRE(num_img == 1792); + + std::vector x(static_cast(seq_len * video_width), 0.0f); + const std::vector video_rows = MakeParam("h3seq.video_rows", num_img * video_width, 1.0); + for (int64_t r = 0; r < num_img; ++r) { + std::memcpy(x.data() + packed.img_pos[static_cast(r)] * video_width, + video_rows.data() + r * video_width, + static_cast(video_width) * sizeof(float)); + } + std::vector audio_x(static_cast(seq_len * p.audio_latents_dim), 0.0f); + const std::vector audio_rows = + MakeParam("h3seq.audio_rows", num_audio * p.audio_latents_dim, 1.0); + for (int64_t r = 0; r < num_audio; ++r) { + std::memcpy(audio_x.data() + packed.audio_pos[static_cast(r)] * p.audio_latents_dim, + audio_rows.data() + r * p.audio_latents_dim, + static_cast(p.audio_latents_dim) * sizeof(float)); + } + const std::vector prompt_embeds = MakeParam("h3seq.prompt_embeds", num_text * p.text_dim, 1.0); + // Step-0 partition (dumped from the real render): all timesteps 0 -> one unique. + const std::vector unique_timesteps = {0.0f}; + const std::vector inverse(static_cast(seq_len), 0); + const std::vector refiner_cu = {0, static_cast(num_text), + static_cast(num_text)}; + + MiniMaxH3DitInputs in; + in.seq_len = seq_len; + in.x = x.data(); + in.audio_x = audio_x.data(); + in.img_position_ids = packed.img_position_ids.data(); + in.unique_timesteps = unique_timesteps.data(); + in.num_unique_timesteps = static_cast(unique_timesteps.size()); + in.inverse_indices = inverse.data(); + in.token_tags = packed.token_tags.data(); + in.prompt_embeds = prompt_embeds.data(); + in.img_pos = packed.img_pos.data(); + in.num_img_pos = num_img; + in.audio_pos = packed.audio_pos.data(); + in.num_audio_pos = num_audio; + in.text_pos = packed.text_pos.data(); + in.num_text_pos = num_text; + in.infer_out_pos = packed.img_pos.data(); + in.num_infer_out_pos = num_img; + in.update_mask = packed.update_mask.data(); + in.cu_seqlens = packed.cu_seqlens.data(); + in.num_cu_seqlens = static_cast(packed.cu_seqlens.size()); + in.refiner_cu_seqlens = refiner_cu.data(); + in.num_refiner_cu_seqlens = static_cast(refiner_cu.size()); + + const MiniMaxH3DitOutputs host = + MiniMaxH3DitForward(Cpu(), p, weights->views, in, vt::DType::kF32); + vt::Queue q = cuda->CreateQueue(); + const MiniMaxH3DitDeviceWeights staged = StageMiniMaxH3DitWeights(q, p, weights->views); + const MiniMaxH3DitOutputs dev = + MiniMaxH3DitForwardDevice(q, p, staged.weights, in, vt::DType::kF32); + + const double dv = MaxAbsDiff(dev.video_logits, host.video_logits.data(), host.video_logits.size()); + const double da = MaxAbsDiff(dev.audio_logits, host.audio_logits.data(), host.audio_logits.size()); + INFO("CUDA-vs-host at real seq 1920: video max|diff| = " << dv << ", audio max|diff| = " << da); + // f32 summation-order slack only (seq 1920 accumulates more than the small + // cases, so allow 5e-3); a structural CUDA-kernel-at-scale regression would be + // orders of magnitude larger and is exactly what #70 is hunting. + CHECK(dv <= 5e-3); + CHECK(da <= 5e-3); +} + TEST_CASE("minimax_h3: the bf16 production stream matches upstream's dtype policy") { // The f32 case above gates the ALGORITHM. This one gates the PRODUCTION dtype // policy: upstream's stream is bf16 with fp32 islands (both patch projections, From 2f740a0a901815104829dd98254705e29ed0eabb Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 21:19:49 +0000 Subject: [PATCH 3/4] fix(minimax-h3): strip prepended reference rows from the ref2va output ROOT CAUSE of the #70 white render FOUND: the render ran task=t2va on the `minimax_h3_ref2va_nvfp4_full` checkpoint, which is the REF2VA partition. Upstream serves t2va/fl2va from the FL2VA partition and ref2va from the Ref2VA partition, and "task must match the served partition" (recipes/MiniMaxAI/MiniMax-H3.md:50,289; _resolve_task raises otherwise). A ref2va-trained DiT fed a t2va sequence (no reference block) is out of distribution -> the spatially-degenerate latent, invariant to the text prompt and step count -- exactly #70's symptom. Every render so far used the wrong task for this checkpoint; the DiT forward, its inputs, the NVFP4 dequant and the CUDA kernels are all correct (verified: S1 inputs byte-exact vs upstream at real 512x512/22f scale; CUDA device forward == CPU host at seq 1920; forward math == upstream source == RefDiT). Running the CORRECT task (ref2va) surfaced a real, previously-unexercised bug in the OUTPUT path: BuildMiniMaxH3PackedSequenceRef2va PREPENDS pinned reference rows (encoded image/video/audio) to the packed layout, the DiT zeroes them in its output (skip_mask_out_condition), and MiniMaxH3GenerateT2va then handed the full (reference + target) row buffer to unpatchify/unpack -- which rejected the non-divisible count ("rows not divisible by t*h*w"). The generated clip is only the TRAILING target rows. Slice to them before unpatchify/unpack. t2va/fl2va have no reference prefix, so the tail is the whole buffer and the change is a no-op there. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .../models/minimax_h3_pipeline.cpp | 26 ++++++++++++++++--- 1 file changed, 23 insertions(+), 3 deletions(-) diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index 14f125b2..9e6372db 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -367,14 +367,34 @@ MiniMaxH3T2vaResult MiniMaxH3GenerateT2va(vt::Device device, const MiniMaxH3T2va initial_video_rows, initial_audio_rows, compute_dtype, prestaged); // --- 4. rows -> latents --- + // ref2va PREPENDS pinned reference rows (encoded image/video/audio) to the + // packed layout; the DiT zeroes them in its output (skip_mask_out_condition), + // and only the TRAILING target rows are the generated clip. t2va/fl2va have no + // such prefix, so the tail is the whole buffer -- this is a no-op there. Without + // this, unpatchify sees (ref + target) rows and rejects a non-divisible count. const int64_t ph = request.latent_h / dit_params.patch_size_h; const int64_t pw = request.latent_w / dit_params.patch_size_w; + const int64_t video_row_width = dit_params.video_row_width(); + const int64_t target_video_rows = request.latent_t * ph * pw; + const int64_t have_video_rows = + video_row_width > 0 ? static_cast(denoised.video_rows.size()) / video_row_width : 0; + VT_CHECK(have_video_rows >= target_video_rows, + "minimax_h3 t2va: denoise produced fewer video rows than the target clip needs"); + const std::vector video_target_rows( + denoised.video_rows.end() - target_video_rows * video_row_width, denoised.video_rows.end()); std::vector video_latent = MiniMaxH3UnpatchifyVideoTokens( - denoised.video_rows, request.latent_t, ph, pw, dit_params.latents_dim, + video_target_rows, request.latent_t, ph, pw, dit_params.latents_dim, dit_params.patch_size_t, dit_params.patch_size_h, dit_params.patch_size_w); + const int64_t audio_width = dit_params.audio_latents_dim; + const int64_t target_audio_rows = request.audio_t * request.audio_channel; + const int64_t have_audio_rows = + audio_width > 0 ? static_cast(denoised.audio_rows.size()) / audio_width : 0; + VT_CHECK(have_audio_rows >= target_audio_rows, + "minimax_h3 t2va: denoise produced fewer audio rows than the target clip needs"); + const std::vector audio_target_rows( + denoised.audio_rows.end() - target_audio_rows * audio_width, denoised.audio_rows.end()); std::vector audio_latent = MiniMaxH3UnpackAudioTokens( - denoised.audio_rows, request.audio_t * request.audio_channel, request.audio_channel, - dit_params.audio_latents_dim); + audio_target_rows, target_audio_rows, request.audio_channel, dit_params.audio_latents_dim); // --- 5. denormalize (vae.py:252-270, :341-357) --- auto denormalize = [](std::vector& latent, int64_t channels, int64_t per_channel, From 1e6a2f5559a80540d7fe1c1e8cba74e0d5acfb32 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 21:58:28 +0000 Subject: [PATCH 4/4] =?UTF-8?q?docs(minimax-h3):=20render=20bug=20CLOSED?= =?UTF-8?q?=20=E2=80=94=20wrong=20partition,=20not=20a=20code=20bug?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Records the H3-RENDER-CLOSE result across the canonical surfaces. The #70/#74 white render was t2va run on the REF2VA-partition checkpoint; upstream serves t2va/fl2va from the FL2VA partition and requires the task to match the partition. t2va on the FL2VA GGUF DiT renders a COHERENT, prompt-matched scene on GB10 (VAE-input latent adj-cell cosine 0.95 vs 0.06 white, no 16px patch grid, valid h264/AAC mp4). Before switching partitions, verified: the t2va DiT inputs diff byte-exact vs upstream at real 512x512/22f scale; the CUDA device forward equals the CPU host at the real render seq (1920, new permanent gate); DequantNvfp4ToBf16 is byte-exact; the forward math equals upstream source. - STATUS/FEATURES/BENCHMARKS: H3 row -> render bug CLOSED (shrink-only respected). - benchmark-record + spec §8.6 + state: full investigation, root cause, the ref2va output-row fix, and the open follow-ups (partition guard + encoder vision tower). - NOW: H3 lane updated, under budget. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 50 +++++++++++++++++++++++++++++++++++++ .agents/specs/minimax-h3.md | 39 +++++++++++++++++++++++++++++ .agents/state.md | 35 ++++++++++++++++++++++++++ docs/BENCHMARKS.md | 3 ++- docs/FEATURES.md | 4 +-- docs/STATUS.md | 2 +- 7 files changed, 130 insertions(+), 5 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 6c76ff22..3995dd8a 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -18,7 +18,7 @@ checkpoint on `upstream/main` at `59674cf1d`. | DeepSeek-V4-Flash decode | **Closed: beats ds4 1.144x** (`VT_V4_RESIDENT_W`, byte-exact); phase-2 residency NEG, default-OFF | — | | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | | Invocation-parity prevention | CI guard + AGENTS.md checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **#70** DiT-math bug REFUTED (`H3-DIT-SCALE-GATE` PR #74): ladder 2x3->8x8+temporal ours==oracle host+dev <=3e-7; white=TRAINED | dgx: real upstream vs RefDiT | +| MiniMax-H3 lane | **RENDER BUG CLOSED** (`H3-RENDER-CLOSE` PR #77): #70/#74 white = t2va on the REF2VA ckpt; the FL2VA GGUF t2va renders COHERENT (adj-cos 0.95) | Follow-up: partition guard + vision tower | | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF | | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 26789f96..e7713882 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14159,3 +14159,53 @@ noise fix from #70 stands but was already known not to be the render fix. full canvas, and the real per-token timestep layout. A real-weights activation diff of the DiT INPUTS (encoder output, position grid, condition-noise) at real geometry is the untested surface #70 did not isolate. + +## MiniMax-H3 render bug CLOSED — the render ran t2va on the ref2va PARTITION checkpoint; t2va on the FL2VA partition renders a COHERENT scene (2026-08-06, `row/H3-RENDER-CLOSE` PR #77, `ROAD-V1-H3`, dgx GB10 sm_121a) + +**Verdict.** The #70/#74 white latent was NOT a code bug. Every prior render ran +**task=t2va on `minimax_h3_ref2va_nvfp4_full`, which is the REF2VA partition.** +Upstream ships two independently-served partitions and requires the task to match: +*"Set MODEL to FL2VA for T2VA"*, ref2va runs against the Ref2VA partition, and +*"task ... must match the served partition"* (`recipes/MiniMaxAI/MiniMax-H3.md:50,222,289`; +`pipeline._resolve_task` RAISES otherwise, `pipeline_minimax_h3.py:387-390`). A +ref2va-trained DiT fed a t2va sequence (no reference block) is out of distribution → +the spatially-white latent, invariant to the text prompt and step count — exactly #70. + +**How it was cornered (all NEW, all measured on dgx at real 512x512/22f scale):** +| Suspect | Test | Result | +|---|---|---| +| S1 (a)(b) DiT INPUT wiring at real scale | `VT_H3_DUMP_INPUTS` dumped every step-0 DiT input; diffed vs upstream `pipeline_minimax_h3.py` at t2va 512x512/22f (text_len=8, latent 7x32x32, seq_len 1920) | **EXACT** — packed layout, fp64 grid, token_tags, inverse/combined AdaLN indices, sigmas all byte-equal; tokenization byte-equal to `tokenizer(prompt,add_special_tokens=False)` | +| encoder conditioning | shape/stat check | correct [8,5120], carries the expected Qwen massive-activation (row0 ch731=15915, others rms~4) — not all-pad, not garbage | +| NVFP4 dequant | independent torch dequant of `blocks.0.attn.qkv_proj` + Laguna/Qwen3 already prove `DequantNvfp4ToBf16` byte-exact | sane trained weight (rms 0.089, absmax=ws2·6·maxscale=3.61) | +| CUDA kernels at real seq | NEW gate `test_minimax_h3 :: CUDA device forward tracks the host at the REAL render seq (1920)` (RealRatioParams head_dim=128, seq 1920) | **CUDA device == CPU host** (28/28) — no scale-dependent kernel bug (#74 only ran device-vs-host on the CPU backend) | +| forward math | RefDiT restatement vs true upstream source, read side by side (block, attention, AdaLN view(m*3,6H), 3D-RoPE, modulate) | identical | + +So inputs + forward + kernels + dequant are all correct → the only thing left was +the checkpoint↔task pairing. + +**Proof.** Downloaded the FL2VA-partition DiT `MiniMax-H3-FL2VA-Q3_K_M.gguf` (15.58 GB, +`realrebelai/MiniMax-H3_GGUFs`, same 50L/5376/head128 geometry) and rendered the SAME +t2va prompt *"an orange cat sitting on a wooden table"* at 512x512/22f, 20 steps, +`--dequant-bf16`: +- VAE-input latent **adj-cell cosine = 0.9467** (vs 0.06 white on the ref2va checkpoint; a real encoded latent is 0.789), latent rms 0.10. +- frame **seam16/interior = 1.00** (no 16px patch grid), and the decoded frames SHOW a + photorealistic orange cat sitting on a wooden table, prompt-matched, temporally + evolving across the 22 frames. Valid `h264 512x512 + AAC 32kHz stereo` mp4. +- healthy denoise signature: velocity STABLE ~1.37 rms, final latent rms **1.00** (the + broken ref2va-t2va run blew up to 2.64). + +**Secondary bug FOUND+FIXED (this PR).** Running the CORRECT task (ref2va) surfaced a +real, never-exercised bug in `MiniMaxH3GenerateT2va`: `BuildMiniMaxH3PackedSequenceRef2va` +PREPENDS pinned reference rows, the DiT zeroes them in its output, and the pipeline handed +the full (reference+target) buffer to unpatchify → `rows not divisible by t*h*w`. Fixed by +slicing to the TRAILING target rows (no-op for t2va/fl2va). (ref2va with a SYNTHETIC image+ +tone reference + text-only encoder still gridded — expected: a meaningless reference plus +the still-unported encoder vision tower is weak conditioning; the clean confirmation is the +FL2VA t2va render above, which needs neither.) + +**Residuals.** (1) The driver takes NO partition/supported_tasks guard (the community GGUF/ +NVFP4 files strip the release config), so picking the right checkpoint per task is on the +caller — mirror-upstream guard is a follow-up. (2) The encoder vision tower (W3 remnant) is +still unported, so real image/video-conditioned ref2va/fl2va renders are not yet clean. (3) +50-step render at the reference canvas (768x1344) is the artifact leg. fp4 speed path +unchanged. diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 34f16f05..28dddd0d 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -526,3 +526,42 @@ reduced-dim DiT gate into a GEOMETRY LADDER. embeddings, real fp64 position grid at full canvas, real per-token timesteps) is fed with RANDOM data here; a real-weights activation diff of the DiT inputs is the untested surface. Full tables: benchmark record (`row/H3-DIT-SCALE-GATE`). + +## 8.6 RENDER BUG CLOSED — wrong checkpoint PARTITION, not a code bug (2026-08-06, `row/H3-RENDER-CLOSE` PR #77) + +The #70/#74 white render was **using the wrong checkpoint partition for the task.** +MiniMax-H3 ships two independently-served DiT partitions and the task MUST match +(`recipes/MiniMaxAI/MiniMax-H3.md:50,289`; `pipeline._resolve_task` raises otherwise): + +| Partition | Serves | Available quantized DiT | +|---|---|---| +| **FL2VA** | **t2va + fl2va** | `MiniMax-H3-FL2VA-Q3_K_M.gguf` (GGUF), FL2VA NVFP4 (not downloaded) | +| **Ref2VA** | ref2va (image/video + audio references) | `minimax_h3_ref2va_nvfp4_full` (the NVFP4 we had), REF2VA GGUF | + +Every render up to #74 ran **t2va on `minimax_h3_ref2va_nvfp4_full` (the Ref2VA +partition)** — an out-of-distribution task/partition combination upstream rejects. That +is the white latent, invariant to prompt/steps. + +**Verified before switching partitions (all NEW, real 512x512/22f scale, dgx):** the +t2va DiT INPUTS diff EXACTLY vs upstream `pipeline_minimax_h3.py` (`VT_H3_DUMP_INPUTS`: +packed layout / fp64 grid / token_tags / inverse+combined AdaLN indices / sigmas all +byte-equal; tokenization byte-equal); the encoder conditioning is correctly shaped and +carries the expected Qwen massive-activation; `DequantNvfp4ToBf16` is byte-exact +(Laguna/Qwen3 + independent torch dequant); and the CUDA device forward == the CPU host +forward at the REAL render seq (1920) at head_dim=128 (new permanent gate +`test_minimax_h3 :: "CUDA device forward tracks the host at the REAL render seq (1920)"`, +28/28) — closing the "CUDA kernel at scale" hole #74's CPU-backend device-vs-host left open. + +**Proof:** t2va on `MiniMax-H3-FL2VA-Q3_K_M.gguf` (`--dequant-bf16`, 512x512/22f, prompt +"an orange cat sitting on a wooden table") renders a **COHERENT photorealistic orange cat +on a wooden table** — VAE-input latent adj-cell cosine **0.9467** (white was 0.06), frame +seam16/interior **1.00** (no patch grid), velocity stable ~1.37, final latent rms **1.00**. +Valid h264 512x512 + AAC 32kHz mp4. + +**Fixed in this row:** `MiniMaxH3GenerateT2va` now strips the PREPENDED pinned reference +rows (ref2va) before unpatchify/unpack — they are zeroed in the DiT output and only the +trailing target rows are the clip; the old code fed unpatchify the full buffer and hit +"rows not divisible by t*h*w" (no-op for t2va/fl2va). **Open:** a partition/supported_tasks +guard mirroring upstream (community files strip the release config); the encoder vision +tower (W3) is still unported, so image/video-conditioned ref2va/fl2va renders are not yet +clean (ref2va with a synthetic reference + text-only encoder still grids). diff --git a/.agents/state.md b/.agents/state.md index 3f1855a1..e8b47aaf 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39643,3 +39643,38 @@ is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional code shipped (CMakeLists NOTE + benchmark-record #75 record the closed levers). Box left clean (GPU idle, both locks free, worker down). Evidence: `dgx:~/mxfp4-nsys/{ours,vllm,buildB,buildC}_flash_c8_ncu.ncu-rep`; PR #75. + +## MiniMax-H3 render bug CLOSED — wrong checkpoint PARTITION, not a code bug (`row/H3-RENDER-CLOSE` PR #77) + + +The #70/#74 white render was **using the wrong partition for the task**, not a bug. +MiniMax-H3 has two independently-served DiT partitions; the task MUST match (upstream +`recipes/MiniMaxAI/MiniMax-H3.md:50,289` + `_resolve_task` raises): **FL2VA serves +t2va+fl2va, Ref2VA serves ref2va.** Every render up to #74 ran **t2va on +`minimax_h3_ref2va_nvfp4_full` (the REF2VA partition)** — out of distribution → the +white latent, invariant to prompt/steps. + +BEFORE switching partitions I exonerated everything else (all NEW, dgx, real 512x512/22f): +(1) `VT_H3_DUMP_INPUTS` — the t2va DiT step-0 inputs diff EXACTLY vs upstream +`pipeline_minimax_h3.py` (packed layout, fp64 grid, token_tags, inverse/combined AdaLN +indices, sigmas byte-equal; tokenization byte-equal). (2) encoder conditioning correctly +shaped, carries the expected Qwen massive-activation. (3) `DequantNvfp4ToBf16` byte-exact +(Laguna/Qwen3 + independent torch dequant). (4) NEW permanent gate `test_minimax_h3 :: +"CUDA device forward tracks the host at the REAL render seq (1920)"` — CUDA device == CPU +host at head_dim=128, seq 1920 (#74's device-vs-host only ran the CPU backend). (5) forward +math == upstream source, read side by side. + +PROOF: downloaded `MiniMax-H3-FL2VA-Q3_K_M.gguf` (15.58 GB, `realrebelai/MiniMax-H3_GGUFs`, +FL2VA partition, same geometry) and rendered t2va "an orange cat sitting on a wooden table" +(`--dequant-bf16`, 512x512/22f, 20 steps) → a COHERENT photorealistic orange cat on a wooden +table: VAE-input latent adj-cos **0.9467** (white=0.06), frame seam16/interior **1.00** (no +patch grid), velocity stable ~1.37, final latent rms **1.00**, valid h264+AAC mp4. A 50-step +768x1344 render was run as the artifact leg. + +FIXED (code): `MiniMaxH3GenerateT2va` now strips the PREPENDED pinned reference rows (ref2va) +before unpatchify/unpack (zeroed in the DiT output; only the trailing target rows are the +clip) — old code hit "rows not divisible by t*h*w"; no-op for t2va/fl2va. OPEN: a +partition/supported_tasks guard mirroring upstream (community files strip the release config), +and the encoder vision tower (W3) for clean image/video-conditioned ref2va/fl2va (ref2va with +a synthetic reference + text-only encoder still grids). dgx assets: `~/h3fp4/ckpt/MiniMax-H3- +FL2VA-Q3_K_M.gguf`, `~/h3fp4/fl2va_t2va_20/`. Box left clean. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 31e09b95..e61b728d 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -302,7 +302,8 @@ built on it rather than keeping the flattering one. | Qwen3-dense decode CUDA-graph | Token-exact pass, ~4.3% e2e directional | Steady-state per-step tok/s | | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | -| MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) but frame is a non-scene patch-grid | Render coherence ROOT-CAUSED (#70): VAE fine, DiT latent white. Geometry ladder (PR #74) REFUTES a DiT-math bug: ours==oracle to 8x8+temporal; white=trained-wts. fp4 speed CLOSED. Detail: benchmark-record + spec §8 | +| MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) | fp4 speed CLOSED. Detail: benchmark-record + spec §8 | +| MiniMax-H3 render coherence (`row/H3-RENDER-CLOSE` #77) | **CLOSED: a COHERENT scene on GB10.** #70/#74 white was wrong-PARTITION usage (t2va on the ref2va ckpt); t2va on the FL2VA GGUF renders a prompt-matched orange cat (adj-cos 0.95 vs 0.06, no patch-grid) | Verified first: t2va inputs byte-exact vs upstream; CUDA device==host at seq 1920; dequant byte-exact. spec §8.6 | | MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-OCCUPANCY` #75: matched-c8 ncu, occupancy IDENTICAL 8.33%; built vLLM's exact flash recipe, matched reg+instr, STILL +10us, gap is ptxas SASS quality, no lever/flip | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 10ab891f..ba2162f1 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -131,7 +131,7 @@ they sit outside the gated list above. |---|---|---|---| | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder TTFT ~17x, pending | | Whisper audio encoder | openai/whisper-small; whisper-large-v3 (Voxtral cfg) | encoder tower 77/77; large-v3 tower 203/203 | pending | -| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable path 65/65 (DiT geometry ladder 2x3->8x8+temporal, host+device vs oracle); real-weights render coherence OPEN (DiT-math bug REFUTED by the ladder, PR #74) | FP4/Marlin routing landed, GB10 speed pending | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable path 66/66 (DiT geometry ladder + CUDA-vs-host at the REAL render seq 1920); t2va renders a COHERENT prompt-matched scene on GB10 (render bug CLOSED: #70/#74 was wrong-partition usage, not a code bug) | FP4/Marlin routing landed, GB10 speed pending | | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) | | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on | | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending | @@ -159,7 +159,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ portable path complete; fp4-resident e2e RUNS on GB10 (real NVFP4 DiT + VAEs + GGUF encoder → mp4/wav); Marlin W4A16 byte-exact; render COHERENCE open, root-caused (#70) to the DiT (not the VAE) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va renders a COHERENT prompt-matched scene on GB10 (FL2VA-partition GGUF → h264/AAC mp4); render bug CLOSED (was wrong-partition usage); Marlin W4A16 byte-exact | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | | Multimodal over the OpenAI server | ☐ | ✅ | ✅ | ◐ | Image, video and audio are correct through the CLI and library. Serving them diff --git a/docs/STATUS.md b/docs/STATUS.md index 343734cf..b161c180 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | OLMo-3 dense (dual rope, interleaved sliding window) | Implemented, oracle-blocked | Loads + runs in our engine (dual rope: plain sliding + YaRN full-attn, per-layer sliding window); no SACRED gate: vLLM 0.25.0 oracle cannot run OLMo-3-1025-7B (`KeyError: 'rope_theta'`; transformers 5.13.1 nests `rope_parameters` per layer-type, no flat `rope_theta`; run-verified W0 2026-07-26) | | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: 87% of vLLM (37.55 vs 43.10, SAME-TOOL nsys both engines); the whole +3.1 ms/step is the bf16 M=1 GEMV bucket (2/3 o_proj, ~196-204 vs 139 us/call, identical `gemvx` kernel); attention/MoE/glue tied or ours-ahead. Invocation match (bf16-out `cublasGemmEx`) A/B'd = WASH, ruled out; ROOT CAUSE FOUND 2026-08-04 (`VT_LAGUNA_RESIDENT_BF16W`): the bf16 projections read UNIFIED/ATS host memory, not `cudaMalloc`'d device memory — staging them device-resident (byte-exact ids) gives 38.8→44.6 tok/s (o_proj 194→131, lm_head 2410→1620 us/call), parity+ vs vLLM 43.1; **default-ON** (flip smoke-verified: canonical byte-exact ids, 44.6 clean-median). Earlier ceiling/diffuse verdicts below were cross-tool artifacts.** **REAL vLLM BAR ESTABLISHED (2026-07-31, `CLAIM-LAGUNA-VLLM-NVFP4`): FIRST-EVER vLLM Laguna run** — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official `poolside/Laguna-S-2.1-NVFP4` (single GB10, greedy, eager, MARLIN backend forced via `VLLM_TEST_FORCE_FP8_MARLIN=1` because the auto-default `FLASHINFER_CUTLASS` needs an absent `nvcc`): **~18.8 tok/s** (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). **N1-scaffold LANDED (2026-07-31):** additive `LagunaMoeWeights.experts_{gate,up,down}_fp4` + `shared_{gate,up,down}_fp4` (`Nvfp4Weight`, mirror qwen3_5), dead until the N1 loader; CPU build clean + `test_laguna_scaffold` 8/8·167 unchanged. **N1b loader IMPLEMENTED (2026-07-31, build-verified):** `LoadLagunaForCausalLMWeights` (`laguna_weights.cpp`) replaces the `VT_CHECK(false)` stub — resolver + per-layer `LoadBf16Direct` (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 `e_score_correction_bias` + `LnLoadCtNvfp4Raw` W4A4 experts. Name-map + dtypes VERIFIED against the real `poolside/Laguna-S-2.1-NVFP4` index (router `mlp.gate` BF16, bias F32, experts W4A4, shared-expert BF16). **N1b RUN-VERIFIED (2026-07-31):** loader round-trips a synthetic NVFP4 checkpoint byte-identically (`test_laguna_nvfp4_loader` 2/2·29; full detail in the benchmark record). **N2 FORWARD-BRANCH LANDED + CPU-GATED (2026-07-31):** `LqGemmNvfp4Fp4` (per-expert TRUE-W4A4: `ScaledFp4Quant(input_global_scale_inv)`→`MatmulNvfp4Fp4(alpha)`, unified-memory pattern like `LqGemm`) + `LagunaFfnBlock` branches on `fp4=!experts_gate_fp4.empty()` (routed experts fp4; keep-quant grouped fast-path gated off `!fp4`; bf16 attn/dense/router/shared-expert/lm_head unchanged) + both `LagunaForwardGguf{,Cached}` guards relaxed to `has_gguf_weights||has_nvfp4_weights`. **CORRECTION:** routed experts are W4A4 ⇒ per-expert `MatmulNvfp4Fp4`, NOT the grouped W4A16 `MoeGroupedGemmNvfp4` (grouped W4A4 deferred to N5 speed). `test_laguna_nvfp4_loader` 3/3·61 (added a forward run-gate: fp4 MoE branch runs through the real `LagunaForwardGguf` → finite+deterministic logits + routed-experts-consumed); `test_laguna_scaffold` 8/8 unchanged (GGUF byte-identical). **N3 DRIVER LANDED + CPU-SMOKE-VERIFIED (2026-07-31):** `examples/laguna_gen` auto-detects a safetensors DIRECTORY (→ NVFP4: `LoadHfConfig(config.json)` + `LoadLagunaForCausalLMWeights` + `LagunaForwardGguf{,Cached}`) vs a `.gguf` FILE (→ keep-quant), sharing the greedy loop; `--token-ids` bypass the tokenizer for the id-vs-golden gate. Verified on a synthetic NVFP4 dir with a REAL config.json (exercises the `LoadHfConfig`→`ParseLagunaParams` seam the loader test bypassed) → `has_nvfp4=1`, KV-cache decode runs finite. **N4 RAN on GB10 (2026-08-01) — the arm works end-to-end; correctness coherent+near-tie, speed 120× off.** git-archived `84fab587` → clean CUDA build (`121a`) → `laguna-gen --gpu` on the real 67 GiB `ckpt` with vLLM's exact prompt ids injected (`2,785,9626,377,15360,395`, captured via the HF tokenizer). Two GB10 memory fixes landed to run: release the mmap'd shards after the loader's memcpy-copy (114→67 GiB RSS), and create the CUDA context BEFORE the load (the 67 GiB reclaimable page cache otherwise starves `cudaStreamCreate`). **Correctness:** ours `22345 83 350 71070 395 340 9626 372 1703 …` vs golden `22345 83 290 350 674 330 5541 966 340 9626 377 15360 …` — **first 2 tokens match vLLM exactly**, then near-tie divergence; coherent ("France is" = 9626/377/15360; shares golden vocab). EXPECTED: our TRUE-W4A4 (fp4 activations) vs the MARLIN golden's W4A16 (bf16 activations) — different precision, not a bug. **Speed: 6.34 s/tok (0.16 tok/s), prefill 17.3s — ~120× slower than vLLM 18.8.** ROOT CAUSE (source-confirmed): `LqGemmNvfp4Fp4` uses the generic `vt::MatmulNvfp4Fp4` = the hand-written EMULATION CUDA kernel, NOT the cutlass sm120a fp4 tensor-core path the 27B/35B W4A4 use (`MatmulNvfp4Fp4DirectD`); + per-expert loop + per-GEMM host sync + no device residency. **nsys (2026-08-01) trace-confirmed + refined:** only 2 GPU kernels — `MatmulNvfp4Fp4Naive` = 99.3% of GPU time + fp4-quant 0.7%; GPU busy only ~18% of wall. NO bf16 GEMM on the GPU ⇒ `LqGemm`'s bf16 branch runs the host `MatmulNK` reference on the CUDA queue (attention/dense/router/shared/lm_head are CPU-bound, ~4.8 s/tok) — a second lever the source scan missed. **N5 LEVER #2 LANDED (2026-08-01) — 16× decode.** Routed the bf16 tower (attention/dense/router/shared/lm_head) off the host `MatmulNK` onto the GPU (`LqGemm` bf16 branch: `vt::CastBf16` the small activation + `vt::MatmulBT` bf16×bf16→f32, weight stays bf16 — no per-token `ReadF32` of `lm_head [100352,H]`): **decode 6.34 → 0.39 s/tok (16.3×; 0.16 → 2.56 tok/s), prefill 17.3 → 2.24s**; coherence preserved (near-tie). CPU path unchanged (run-gate byte-identical). **N5 LEVER #1 LANDED (2026-08-01) — native fp4 tensor-core, another ~2×.** The engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`, `mma.sync kind::mxf4nvf4`) reads the same linear scale layout `LqGemmNvfp4Fp4` produces — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; the Laguna driver now defaults it ON (scoped; 27B/35B untouched). **decode 0.39 → ~0.20-0.24 s/tok (~2×; ~4.2-5.0 tok/s)**; coherent (byte-identical ids to the emulation path — numerically equivalent), first token matches the golden. **Cumulative N5: 0.16 → ~4.5 tok/s (~28×), now ~4× from vLLM 18.8.** **Device-resident MoE block LANDED + MEASURED (2026-08-01, `LagunaMoeResidentFp4`, `VT_LAGUNA_RESIDENT_MOE` default-ON):** the whole token's routed experts as ONE async device chain (fp4-quant→GEMM gate/up, `MoeSiluMul`, →down stacked, ONE `MoeCombine`), draining once vs ~Pk×3 syncs. **Speed EAGER-NEUTRAL (0.20 s/tok)** — empirically confirms the ds4 precedent (per-op syncs overlap GPU compute; wall is GPU-serial-bound; the graph is the payoff). **CORRECTNESS WIN: golden-token match 2 → 13** (the device `MoeSiluMul`/`MoeCombine` mirror vLLM's fused MoE faithfully). Lands default-ON (better correctness, no speed cost, graph prerequisite). **CORRECTED CEILING (from the measured state): a perfect decode graph caps at ~5.9 tok/s** (GPU already ~87% busy at 0.20 s/tok), still 3.3× short of vLLM 18.8 — the graph is necessary but NOT sufficient; the remaining 3.3× is KERNEL EFFICIENCY (native fp4 MMA ~302µs/M=1 expert GEMM vs vLLM's tuned cutlass sm120a fp4 + fused norm/quant/silu). Parity = TWO campaigns: (A) device-resident+graph → ~5.9; (B) cutlass DirectD experts + fused ops + M=1-tuned GEMV → the rest. **CAMPAIGN-B FIRST BRICK LANDED (2026-08-01): coalesced M=1 fp4 GEMV** (`MatmulNvfp4Fp4Gemv`, one warp/column, coalesced weight-row reads, `VT_NVFP4_FP4_GEMV` default-ON) — same-binary A/B: **decode 0.20 → 0.15 s/tok (1.33×; → ~6.7 tok/s), prefill 1.14 → 0.86s**, coherent+near-tie. **Cumulative this session: 0.16 → ~6.7 tok/s (~42×), now ~2.8× from vLLM 18.8.** (ILP variant `kCpw=4` measured SLOWER — 0.21 s/tok, occupancy loss > activation-reuse gain — reverted to `kCpw=1`; kernel kept templated as a re-measurable knob.) **ncu of the GEMV (sudo): sm__throughput 35-71%, DRAM n/a — COMPUTE/LATENCY-bound, not BW-bound.** Corrects the earlier "~6× BW → ~16-17 tok/s" estimate: the next GEMV lever is HARDWARE fp4 dequant (`cvt.e2m1x2`), not more bandwidth. Parity (18.8) is a multi-brick campaign (decode graph + fused norm/quant + hardware-dequant GEMV), not one more kernel. **B0 hw-fp8 SCALE-decode: MEASURED NEGATIVE, reverted (2026-08-01, `ab7a1c1e`).** Replacing the GEMV's per-byte software fp8-e4m3 group-scale decode (`F8E4M3ToF32Dev`/`ldexpf`) with hardware `cvt.rn.f16.e4m3` (`__nv_fp8_e4m3`→float) is bit-exact (ids byte-identical on the real ckpt) but paging-immune ncu shows it NEUTRAL-to-slightly-WORSE (grid768 41.2 vs 41.9µs tie; mean 53.6 vs 49.4µs) — GPU `ldexpf` is a cheap exponent-bit add, not a libcall. NOTE this is the fp8 SCALE decode, NOT the fp4-e2m1 WEIGHT dequant (the `kE2M1` `__constant__` LUT); the LUT→arithmetic/`cvt.e2m1x2` weight-dequant is a SEPARATE still-open lever (spec brick B1). Also: end-to-end wall-clock is unusable for kernel A/B here (67 GiB unified reload swings TPOT 0.16↔1.08 s/tok run-to-run) — kernel-duration ncu is the only honest anchor. **★ B2 SCOPED + DE-RISKED (2026-08-01, zero-DGX) — the real 18.8 lever:** vLLM's 18.8 bar is MARLIN W4A16 (`VLLM_TEST_FORCE_FP8_MARLIN=1`), which is LOW-M-optimized (decode-correct, unlike a tensor-core W4A4 GEMM that wastes M=1 tile rows). The engine already ships the EXACT kernel `vt::MoeGroupedGemmNvfp4Marlin` (1:1 lift of vLLM `moe_wna16_marlin_gemm`) + shared `MarlinRepackExpertWeight`, and qwen3_5 (27B/35B) already routes its NVFP4 experts through it (default-ON `VT_NVFP4_MARLIN`, 16/16-vs-oracle, +22% gate/+80% decode) via `BuildMoeMarlinResident`. So B2 = mirror that for `LagunaMoeWeights.experts_*_fp4` (a `BuildLagunaMoeMarlinResident` reusing the shared repack + route `LagunaFfnBlock`'s fp4 branch to the Marlin grouped GEMM, GEMV kept as the `=0` escape hatch) — pure reuse, no new kernel, matches vLLM's exact W4A16 numerics. **B2 IMPLEMENTED (2026-08-01, `3c49ef37`) — COMPILES CLEAN on GB10 sm_121a, runtime bug pending.** `LagunaMoeResidentMarlin` + `BuildLagunaMoeMarlinResident` (laguna.cpp, `#ifdef VT_MARLIN_NVFP4`) reconstruct the MoE Marlin path over the SHARED `dense_nvfp4::Dev`/`DBuf`/`ResidentNvfp4` + shared `vt::cuda` Marlin repack/align ops + `vt::MoeGroupedGemmNvfp4Marlin`; SACRED 27B/35B path BYTE-UNTOUCHED; gated `VT_LAGUNA_MARLIN_MOE=1` **default-OFF** (zero regression to the default GEMV path). Compiles clean on the full CUDA build. RUN: loads OK (48 layers, 256 experts) but the FIRST FORWARD device-faults silently on the Marlin path — a layout/param bug (suspects: `MoeCombine` bf16-in/f32-out dtype, the down-GEMM reusing the gate/up align, or the fp4-original free omitted → mem ~doubles). NEXT: `compute-sanitizer` localize → fix → near-tie vs the vLLM-Marlin golden + kernel-duration ncu → flip default-ON. Default path unaffected. **UPDATE (`22d6e146`): added the qwen3_5-style fp4-original free after repack** (device transients + host bytes; peak was ~3× the expert tower → past the 119 GiB pool → null-alloc → silent fault the likely cause); compiles clean. The runtime gate stayed INCONCLUSIVE this session (contended/orphaned processes on the shared box, no captured ids) — rerun on a clean uncontended session, compute-sanitizer if it still faults. **★★ B2 VALIDATED on GB10 (2026-08-01, with the mem-free fix): RUN_EXIT=0, coherent, first 13 generated tokens MATCH the vLLM-Marlin golden EXACTLY** (`22345 83 290 350 674 330 5541 966 340 9626 377 15360 81` — the best Laguna-NVFP4 correctness yet, W4A16 matching vLLM's config). **Steady-state decode 0.10 s/tok = ~10 tok/s** (steps 10-17 all 0.10; the TPOT-0.56 average is warmup-polluted — the DevicePool warms over ~9 decode steps then reuses). vs the GEMV path's 6.7 tok/s = **~1.5× faster; the gap to vLLM 18.8 closes from ~3× to ~1.9×.** Memory flat (7.9 GiB host RSS — the fp4-original free worked; it also fixed the first-forward fault). Still `VT_LAGUNA_MARLIN_MOE=1` default-OFF. TO DONE: move the lazy Marlin-resident build (216s first-forward, 48L×256E repack) to model-LOAD time → clean warm A/B + ncu → flip default-ON → matrix/roadmap. Remaining ~1.9×: vLLM graphs its decode (ours still eager) — decode CUDA-graph is the next lever. **REPRODUCED 3× (reproduction gate MET): GB10 runs deterministic — first 18-20 tokens byte-identical, steady-state 0.10 s/tok confirmed each — so the ~10 tok/s + golden-match is gated, not a single sample.** **#234 item (1) — load-time resident-build LANDED (`LagunaBuildMarlinResidents`, called from the example after load; mirrors vLLM process_weights_after_loading): builds all 48L×256E Marlin residents at LOAD so the repack is not a first-token TTFT spike. Fixed an anon-namespace linkage bug (public fn was defined with internal linkage → moved outside the anon namespace); BUILD CLEAN + links on GB10 sm_121a, default-OFF. Runtime prewarm-fires-at-load timing UNVERIFIED this session (repeated ssh-drops ate the run capture); the forward's lazy build is the validated fallback so it cannot regress. Owed: one clean run to confirm the build moved to load + then flip default-ON.** **★★ DONE (2026-08-01): Marlin is now the UNCONDITIONAL DEFAULT (`LagunaMarlinMoeEnabled` default-ON; `=0` is a code-level A/B opt-out no user needs) — "it just works" with NO env. Confirmed in a no-env GB10 run captured via tmux: `MARLIN residents built at load in 238.4s`, prefill 14.78s (build moved OUT of first-forward), golden-matching ids, steady-state 0.10 s/tok = ~10 tok/s (4th reproduction), RSS ~5-8 GiB. So a default Laguna-NVFP4 load on GB10 gets vLLM's own W4A16 Marlin decode (~10 tok/s, ~1.9× from vLLM 18.8) with zero flags. The 238s load-time repack is a one-time cost (mirrors vLLM process_weights_after_loading); optimizing its 48×256 per-expert sync count is a follow-up. Residual to 18.8 = decode CUDA-graph (deferred; user refocusing on DeepSeek next).** Post-lever-1 nsys: the remaining ~4× is HOST-SYNC-bound — 22,115 `cudaStreamSynchronize` (78.6% of API time, ~2,760/token, the per-GEMM `DrainQueue`), GPU kernels fast. Remaining levers: grouped W4A4 MoE (design input: `vt::MoeGroupedGemmNvfp4` is W4A16, so true-W4A4 grouped needs a new fp4×fp4 op or the `use_a16` mode + expert-stacking — needs a spike), device-resident decode (RECOMMENDED — the current forward is host-style so every GEMM drains; keep activations on-device, drain once/step; reuse qwen3_5's `Dev`/`Nvfp4Dev`/`ResidentNvfp4`/device-SwiGLU machinery; kills the 22k syncs; converges with the pending GGUF #228 and lifts both quant paths), decode CUDA-graph. Binding number needs a clean 2-3× re-run. See `docs/BENCHMARKS.md` + the spec N5 plan. See `docs/BENCHMARKS.md` `CLAIM-LAGUNA-VLLM-NVFP4`. Prior W7 nsys attribution: host-orchestration-bound, levers ranked (spec `laguna-s21-w7-speed-2026-07-31.md`, ledger `CLAIM-LAGUNA-W7-SPEED`). Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\"…". `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 `is_sliding`), capping their K/V. `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); `examples/laguna_gen --stateless` forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). `laguna-gen` "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via `vt::MatmulBT`; norms/router/bias/embed → f32) + `LagunaForwardGguf` (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + `examples/laguna_gen`. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (`attn_q/k_norm`, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp `yarn_attn_factor·(1+0.1·ln(factor))`, factor 32 not HF 128), separate `ffn_gate/up_exps`. Keep-quant tower materialization + `ForwardGguf` + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (`laguna_ops.cpp`, CPU `-Werror` clean, `test_laguna_scaffold` unit-gated)** | Poolside Laguna: 48 layers (12 global + 36 sliding-window-512), 256 routed top-10 + 1 shared expert, per-head **softplus attention output gate**, sigmoid `noaux_tc` router, dual per-layer RoPE (YaRN full-attn / plain sliding), GQA 8 KV / 128 head-dim, 1M ctx. **W3 (2026-07-31):** the 3 genuinely-NEW small host ops landed in `laguna_ops.cpp` — per-head softplus attn out-gate (`LagunaSoftplusHeadGate`), ungrouped sigmoid-noaux router (`LagunaUngroupedRouterTopK`, ds3 noaux_tc MINUS the group step + tie-break razor), dual per-layer RoPE cos/sin builders (`BuildLaguna{FullYarn,Sliding}CosSin`, reusing the pinned YaRN inv_freq over the partial-64 dims); `LagunaModel::Forward` is now a REAL runnable host-reference composition (variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head) replacing the `VT_CHECK(false)` stub; `test_laguna_scaffold` **8/8·166** (softplus math, router selection+tie-break RED-first, dual-RoPE bit-match, variable-Q-head shapes, forward composition on synthetic weights), `test_model_registry` 24/24. **W2 (2026-07-30):** registered, `ParseLagunaParams`, GGUF `blk.N.*` name-map + UD-Q4_K_XL quant-mix (Q4_K/Q5_K/Q6_K/Q8_0 ALL already decoded → ZERO new kernel). **W1 oracle DECISION:** vLLM NATIVE `laguna.py` (in pin → config constructs); dual-oracle = vLLM-NVFP4/-FP8 (fits GB10 119 GiB; BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. ~85–90% reuse (ds4-MoE + Gemma-sliding + OLMo-3-dual-rope + Q4_K keep-quant, ALREADY landed). DEFERRED (W4): GGUF keep-quant tower materialization + device/paged production forward (loaders still LOUDLY throw) + strict dual-oracle greedy gate on a fetched checkpoint + `poolside_v1` parser. See `.agents/specs/laguna-s21-w3-2026-07-31.md` (+ W1/W2 `laguna-s21-w1w2-2026-07-30.md`, W0 `laguna-s21-scope-2026-07-30.md`). **Decode attention-glue fusion LANDED (2026-08-02, `CLAIM-LAGUNA-GLUE-FUSED`, default-ON `VT_LAGUNA_GLUE_FUSED`, `=0` A/B):** BYTE-EXACT L1 (softplus out-gate → `DecodeAttnCombineKernel` store) + L4 (residual-Add+RMSNorm pairs → the shared `vt::FusedChain(kFusedAddRmsNormStd)` seam) on the resident decode-graph — same-binary A/B ids byte-identical (159/159 @160), paging-immune nsys steady decode **−4.2% GPU-busy (28.90→27.69 ms/step), −120 graph nodes/step (−10%)**, wall drop_caches-tied (no regression). C shared-into-MoeCombine SKIPPED (Laguna's bf16 `MoeCombine` → not byte-exact); L2 qk-norm+RoPE preamble DEFERRED (needs a device-position kernel variant). See BENCHMARKS.md `CLAIM-LAGUNA-GLUE-FUSED`. **On-device greedy sample LANDED (2026-08-02, `CLAIM-LAGUNA-ONDEV-SAMPLE`, default-ON `VT_LAGUNA_ONDEV_SAMPLE`, `=0` A/B):** the resident decode graph used to Synchronize, return the whole `[100352]` logits, and argmax on the HOST between replays (+ host embed-gather of the next token) — the off-framework "born-on-host" seam the decode-framework-routing audit flagged. Now BOTH run ON-DEVICE inside the captured graph: `vt::GreedyArgmax` (lowest-index tie = the exact host winner) → 1-elem device token buffer, + a new capture-safe `embed_gather` kernel gathers the next input embedding from it (the stock `vt::Embedding` is NOT capture-safe: per-call event-sync + D2H ring). BYTE-EXACT (160-id stream identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: paired drop_caches decode wall **+0.28% median** (8/8 reps ≥0; removes ~150 us/step host argmax) at GPU-busy parity (nsys 2-length 27.44→27.42 ms/step). Aligns Laguna decode with vLLM on-device sampling. **Lever 2 (lm_head GEMV DRAM eff) MEASURED, NOT landed:** `[M=1,100352,2048]` bf16 = **170 GB/s (2.41 ms)** = ~91% of the cuBLAS M=1×large-N reference (~187 GB/s / 2.2 ms) — at the M=1 practical floor (the 273 GB/s ceiling is streaming-only, unreachable for a once-read GEMV); ≤0.7%-of-step headroom needs a reduction reorder (near-tie re-gate) ⇒ not chased, per prior "lm_head optimal". See BENCHMARKS.md `CLAIM-LAGUNA-ONDEV-SAMPLE`. **MoE add_rms_norm fold LANDED (2026-08-02, `CLAIM-LAGUNA-MOE-ADDNORM`, default-ON `VT_LAGUNA_MOE_ADDNORM_FUSED`, `=0` A/B):** the glue-fused MoE tail ran its residual update as TWO graph nodes — `vt::Add(hidden,routed)` [`AddKernel`] + `FusedChain(kFusedAddRmsNormStd)` [shared-add+RMSNorm, `RmsNormRowKernel`] — now ONE `fused_add2_rmsnorm` device node/MoE-layer (`hidden=(hidden+routed)+shared; hn=rms_norm(hidden)*w`). BYTE-EXACT (IEEE add commutes + the identical 256-thread shared-tree norm reduction; 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: **−39 `AddKernel` graph nodes/step** (2.63ms→0 over 69 steps), paging-immune nsys 2-length **~−46 us/tok GPU (27339→27293)**, nsys wall **+0.4% (34.00→34.14 tok/s @70-tok)**. Small (byte-exact node-count trim on the graph-captured, GPU-bound decode; the dominant ~72% cost is the bf16 projection GEMVs — see the Lever-B negative in BENCHMARKS.md). See BENCHMARKS.md `CLAIM-LAGUNA-MOE-ADDNORM`. **Shared expert kept fp4 LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-FP4`, default-ON `VT_LAGUNA_SHARED_FP4`, `=0` A/B):** the XS-NVFP4 shared expert was DEQUANTIZED to bf16 at load (`LnLoadSharedExpertBf16`) → the M=1 decode GEMV read 4× the DRAM bytes of vLLM (which keeps it fp4). Now kept fp4-resident and routed through the SAME Marlin W4A16 single-expert (num_experts=1) grouped GEMM the routed experts win on (`dense_nvfp4::GateUpFusedMarlinD`+`MatmulNvfp4MarlinD`); the decode GEMV drops to router-ONLY (`moe.router`), shared gate/up/down go fp4. ADDITIVE new `laguna_shared_fp4.cpp` re-reads the on-disk fp4 from the gen driver before shard release (does NOT touch SACRED `laguna_weights.cpp`); bf16 shared KEPT for the T>1 prefill. NEAR-TIE (fp4≠bf16): coherent, first-20 ids == documented golden, byte-identical to bf16 for ~85 tokens then diverges; **DISTRIBUTIONAL GATE PASS 40/40** (ours' first-40 ids ∈ vLLM's 8-run greedy candidate set; vLLM XS-greedy is bf16-non-det, 8 unique of 8). FASTER: paging-immune nsys 2-length **GPU 27.24→26.53 ms/step (−2.6%)**, wall drop_caches **35.8→36.3 tok/s (+1.4%, fp4 wins all 3 reps)**; shared-expert kernel bucket ~1.68→~0.90 ms/step (halved); vs vLLM ~43 tok/s 83.3%→84.4%; RSS 22.2→22.1 GiB (freed the decode-only fused router-shared projection). Modest by design — XS's shared expert is small (`shared_expert_intermediate_size==moe_intermediate_size==512`). Default-ON per parity (matches vLLM's fp4 shared). See BENCHMARKS.md `CLAIM-LAGUNA-SHARED-FP4`. **qk-norm+RoPE preamble fusion LANDED (2026-08-03, `CLAIM-LAGUNA-PREAMBLE-FUSED`, default-ON `VT_LAGUNA_PREAMBLE_FUSED`, `=0` A/B):** closes the `CLAIM-LAGUNA-GLUE-FUSED` L2 deferral — the decode graph ran the per-layer attention preamble as FOUR under-occupied M=1 nodes (`rms_norm_seq(q)`+`rms_norm_seq(k)`+`rope_from_cache_g(q)`+`rope_from_cache_g(k)`); now ONE capture-safe `fused_qk_norm_rope_g` node/layer (`FusedQkNormRopeGKernel`, one block/head, reads the decode position from DEVICE `*pos_buf`, handles the per-layer dual-RoPE 64/128 + `Hq` 48/64). BYTE-EXACT BY CONSTRUCTION: it replicates the composed path's f32 MEMORY round-trip (Phase A 256-thread Σx² == `RmsNormSeqKernel`; Phase B the same `(x*inv)*w` store; `__syncthreads`; Phase C the `RopeFromCacheGKernel` rope read back) — an earlier register-only recompute was numerically-equivalent but diverged at a token-110 near-tie via compiler fma-contraction; the memory boundary forces bit-identity. 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4` (determinism verified `=0`×3/`=1`×3 each run-to-run identical). FASTER: preamble norm+rope kernels **160→40 launches/tok, 326→154 us/tok (−0.17 ms/step)**; all decode-scaling kernels 26.53→26.37 ms/step; wall drop_caches **36.42→36.64 tok/s (+0.6%, fused wins all 3 paired reps)**; vs vLLM ~43 84.7%→85.2%. Modest (preamble ~1.2% of the 26.5 ms/step decode; the dominant cost stays the bf16 projection GEMVs at cuBLAS parity) — a byte-exact graph-node/launch trim (the glue-fusion residual mechanism). Default-ON per parity. See BENCHMARKS.md `CLAIM-LAGUNA-PREAMBLE-FUSED`. **W7 two-front pass LANDED (2026-08-03, `CLAIM-LAGUNA-W7-DECODE`):** FRONT 1 — the example driver logged `[gen] step N …(RSS)` EVERY decode step, and the RSS arg calls `CurResidentGiB()` (a `/proc/self/status` read) + an unbuffered stderr write in the GPU-idle gap between replays; guarded behind `VT_LAGUNA_STEP_LOG` (default OFF) + added a `decode_wall` line (TRUE end-to-end throughput incl. per-step gaps) next to the gap-free `decode_hp`. Since the fprintf sat OUTSIDE the `s0→s1` timer, `decode_hp` was ALREADY honest; with the log off `decode_wall == decode_hp` (within 0.001 tok/s, every LOG_OFF rep) and the recovered host tax is only ~0.1% (drop_caches noise floor). CONCLUSION: the ~86% gap to vLLM 43 is genuine device compute, NOT a harness artifact. FRONT 2 — `VT_LAGUNA_MOE_ONECAST` (default ON): a MoE layer cast the same `hn[1,H]` f32→bf16 THREE times (router GEMV + routed Marlin + shared Marlin); now cast ONCE into a persistent buffer and reuse (`CastHnBf16`/`GemmBf16Pre` + optional pre-cast param on both `…Into` helpers). BYTE-EXACT (deterministic truncation; `=1` vs `=0` byte-identical 300-tok ids); `CastBf16` **200→122 nodes/step (−78 = 2×39 MoE layers)**, GPU-busy parity within nsys noise, decode_hp +0.29%. Combined (onecast on + log off) **36.97 tok/s = 86.0% of vLLM-NVFP4 43** (from 36.64/85.2%). See BENCHMARKS.md `CLAIM-LAGUNA-W7-DECODE`. **Tail-fold follow-up LANDED (2026-08-03, `CLAIM-LAGUNA-TAIL-FUSED`, default-ON `VT_LAGUNA_TAIL_FUSED`, `=0` A/B):** a fresh node-ranking of the baseline decode graph found the routed-MoE `CastF32` as the one clean byte-exact fold left; it folds into the trailing `fused_add2_rmsnorm` via a new bf16-x1 sibling kernel (`AddAdd2RmsNormStdBf16Kernel` — `MoeCombine` writes bf16 straight to a persistent buffer, widened in-kernel by `__bfloat162float`). BYTE-EXACT (`=1` vs `=0` byte-identical 160-tok ids), `CastF32` **78→39 nodes/step**, total graph nodes **919→880**, GPU-busy parity; decode_hp a WASH (median +0.14% / mean −0.04%, at the drop_caches noise floor). Lands on the deterministic node-count basis (like onecast/preamble/addnorm), NOT a wall win; combined headline UNCHANGED **36.97 tok/s = 86.0%**. The ranking confirms the byte-exact decode-tail fold tier is now essentially EXHAUSTED (residual tail = already-folded norms + attention compute + cuBLAS-adjacent router/topk + ported-Marlin `MoeAlign`/`SiluAndMul`/`MoeCombine`); the gap to vLLM 43 is genuine device compute at the practical ceiling. See BENCHMARKS.md `CLAIM-LAGUNA-TAIL-FUSED`. **KERNEL-EFFICIENCY tier (2026-08-03, `VT_LAGUNA_FAST_NORM` default ON + f32 ext of `VT_RMSNORM_DECODE_FAST`):** the fold tier was exhausted but the residual-stream norm KERNELS were still under-occupied — `ncu` on the shipped `<<<1,256>>>` `AddAdd2RmsNormStdBf16`/`RmsNormRow` decode norms: `launch__waves_per_multiprocessor≈0.00`, `sm__throughput≈0.06%` (one 256-thread block on 1 SM of ~100+, latency-bound). Porting the PROVEN bit-identical `RmsNormRowFastKernel` structure (1024-thread float4 memory passes; 256-strided-partial + tree reduction reproduced byte-for-byte) to the f32 kernels cut each **286→~155 µs/tok (1.85×)**, **byte-exact** (160-tok ids identical `=1`vs`=0`; the f32 fix vs the bf16 sibling: store `v` not `v²` and square in the reduction so nvcc emits shipped's `acc += v*v` **fma** — a pre-squared f32 `v²` is not exact and flipped an XS near-tie at tok 108). **−0.81% decode-step GPU time** (paging-immune 70-vs-20 2-length diff, 26192→25980 µs/step); wall-clock ON/OFF overlap (noise floor). Residual: the byte-exact 256-strided reduction can't reach vLLM's per-kernel norm floor (~2.4× vLLM) without breaking byte-exactness → that remainder is byte-exactness-BLOCKED. See BENCHMARKS.md `CLAIM-LAGUNA-FAST-NORM`. **Router top-k warp-shuffle LANDED (2026-08-03, `CLAIM-LAGUNA-TOPK-SHFL`, default-ON `VT_LAGUNA_TOPK_SHFL`, `=0` A/B): BYTE-EXACT** — an nsys 2-length rank of the remaining small kernels (past the at-parity `gemvx` projection GEMVs ~69% of step + Marlin MoE) put the router `SigmoidTopKKernel` top (415 µs/step); `ncu` showed it `<<<1,256>>>` at `waves≈0.000`/`sm≈0.2%` — pure latency (8 serially-dependent rounds × a ~10-sync `sh[256]` argmax tree). New `SigmoidTopKShflKernel` reduces each round by warp-shuffle argmax (2 syncs/round; argmax over the total order is associative ⇒ SAME winner) → **`SigmoidTopK` 414.6→248.8 µs/step (1.67×)**, decode-step GPU **−0.57%** (26.018→25.869 ms/step), 37.39→37.49 tok/s decode_hp (**87.2% of vLLM-NVFP4 43**); 160-id stream byte-identical `=1`vs`=0`. **NOT landed — norm warp-shuffle (`VT_LAGUNA_NORM_SHFL`):** a near-tie register-accumulate+shuffle reduce for the Laguna `AddAdd2RmsNormStd{,Bf16}Fast` norms PASSED the distributional gate (coherent, in-set 38/40 = baseline, one near-tie fork at pos 37) and was −19.3% per-kernel (`AddAdd2RmsNormStdBf16` 150.3→121.3 µs/step) BUT washed at whole-step (0.6% of step; +0.02% within noise) — a near-tie fork isn't justified by a below-noise gain, so it was dropped. The small-kernel norm tail is at its occupancy floor; the decode step is dominated by the at-parity projection GEMVs. See BENCHMARKS.md `CLAIM-LAGUNA-TOPK-SHFL`. **Shared-expert 2-stream overlap LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-AUX`, default-ON `VT_LAGUNA_SHARED_AUX`, `=0` A/B):** mirror of vLLM's `MULTI_STREAM_OVERLAPPED` — in `LagunaGraph::RunChain` the fp4-shared arm's shared expert is EARLY-forked onto a second CUDA stream from the post-attn hidden `hn` BEFORE the router GEMV (aux reads `hn` f32 + does its own byte-identical cast; scratch from `AuxPool`), overlapping router+`sigmoid_topk`+routed grouped GEMM, joined before the combine — the SAME machinery the 35B ships default-ON (ENG-MOE-SHARED-AUX, runs inside the captured graph). This is the EARLY fork the prior fused-`router_shared_gu` attempt (`89e0d074`, −0.35% wash) could not reach. Capture-safe (aux stream+2 events in the ctor; gstate-0 warm-run builds residents + warms `AuxPool`). **BYTE-EXACT** (`=1`vs`=0` byte-identical 63-tok ids). REAL concurrency: nsys `--cuda-graph-trace=node` 20↔70 sum-vs-union → OVERLAP **2.34 ms/step** (SUM/UNION 1.092) vs `=0`'s 0.0004 ms; net GPU-busy wall **26.213→25.467 ms/step (−2.9%, 38.15→39.27 tok/s)**, wall @200 37.08→37.93 (+2.3%). Net