From bda3d72cb652e7980ac7d953f898379183680233 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 16:21:01 +0000 Subject: [PATCH 1/6] diag(minimax-h3): env-gated denoise motion trace + latent dumps for render-coherence bisection The #64 real-checkpoint NVFP4 t2va render produces a structured non-scene (a grid of 16-px multicolour blocks) IDENTICALLY at 12/20/50 steps whether conditioned or not. The rectified-flow Euler integration telescopes to (sigma0 - sigmaN) * v = 1.0 * v, so step-count invariance is the exact signature of a velocity that does NOT evolve across the denoise trajectory. This adds, env-gated and byte-identical when unset: - VT_H3_TRACE_MOTION: per-step [h3-motion] velocity + latent-motion stats (v_rms/v_amax/v_mean, per-step drows_rms, running rows_rms) plus INIT/FINAL displacement, to see whether the loop moves the latent and whether the velocity is frozen at real geometry. - VT_H3_DUMP_DIR: dump init/final video latent rows and the exact VAE-input latent as raw f32, so 12-vs-50-step and cond-vs-uncond runs can be byte/stat-compared and the VAE decode replayed on a known latent. Both documented in docs/ENVIRONMENT.md (check-env-doc green). No production path changes; the pre-existing check-fusion-consistency red (minimax_h3_video_vae_device) is untouched. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] Row: row/H3-RENDER-COHERENCE (#70) --- docs/ENVIRONMENT.md | 2 + src/vllm/model_executor/models/minimax_h3.cpp | 92 +++++++++++++++++++ .../models/minimax_h3_pipeline.cpp | 18 ++++ 3 files changed, 112 insertions(+) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 244d498e..a736ebd4 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -99,6 +99,8 @@ Read-only observability; none change output. | `VT_POOL_BYPASS` | off | `=1` makes every device-scratch pool allocation an exact-size driver `Alloc` and every release a real `Free`, so `compute-sanitizer` can see tensor boundaries and use-after-free that the caching, size-class-rounding pool hides. DEBUGGING ONLY: it reinstates the per-op `cudaMalloc`/`cudaFree` device-sync storm the pool exists to remove, so it is never a timing configuration | | `VT_TTFT_DUMP` | unset | `=1` prints one `TTFTSPLIT rid=... intake=.. queued=.. prefill=.. decode=.. e2e=..` line per finished request to stderr, reconstructing the per-request timing split from the event-populated `req_state` timestamps. The async serving frontend otherwise tracks no per-request stats (passes `iteration_stats=nullptr` and never stamps `EngineCoreOutputs.timestamp`); under this flag both are wired so a serving TTFT attribution can read the queue-vs-execution split against vLLM's own `request_{queue,prefill,decode}_time_seconds`. Generation is byte-identical when unset (the default path is instruction-identical to production); the durable replacement is the async `/metrics` stat logger | | `VT_LOOP_TRACE` | unset | `=1` prints one `LOOPTRACE ...` line per ~1 s window to stderr from the engine busy loop: the full-iteration cadence (`interval`), `process_engine_step` wall (`step`), per-window admits, input-queue residence (`resid` = enqueue-to-drain, the same endpoints `VT_TTFT_DUMP`'s `intake` measures), per-drain admit max and max backlog depth. Diagnoses whether the admission wait is one busy-loop iteration or the input queue is backing up (it attributed the 35B INTAKE deficit to bursty arrival during long prefill steps). Byte-identical when unset: every trace read is guarded, and the enqueue timestamp is stamped only under the flag | +| `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_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) | ## Kernel-internal knobs (deferred) diff --git a/src/vllm/model_executor/models/minimax_h3.cpp b/src/vllm/model_executor/models/minimax_h3.cpp index 62e684c4..b06239e6 100644 --- a/src/vllm/model_executor/models/minimax_h3.cpp +++ b/src/vllm/model_executor/models/minimax_h3.cpp @@ -777,6 +777,44 @@ MiniMaxH3DenoiseResult MiniMaxH3DenoiseLoop( ++audio_ref_index; } + // DIAGNOSTIC (env-gated, byte-identical when unset): VT_H3_TRACE_MOTION prints + // per-step velocity + latent-motion stats to stderr; VT_H3_DUMP_DIR writes the + // initial and final video latent rows as raw f32 so runs can be byte-compared. + // The rectified-flow Euler integration telescopes to (sigma0 - sigmaN) * v ~= v, + // so a velocity that does NOT evolve across steps yields a step-count-INVARIANT + // result -- this trace measures exactly that (see H3 render-coherence bisection). + const bool trace_motion = std::getenv("VT_H3_TRACE_MOTION") != nullptr; + const char* dump_dir = std::getenv("VT_H3_DUMP_DIR"); + auto rms_absmax = [](const float* p, int64_t n, double* rms, double* absmax, double* mean) { + double s2 = 0.0, amax = 0.0, sum = 0.0; + for (int64_t i = 0; i < n; ++i) { + const double v = static_cast(p[i]); + s2 += v * v; + sum += v; + if (std::fabs(v) > amax) amax = std::fabs(v); + } + const double denom = n > 0 ? static_cast(n) : 1.0; + *rms = std::sqrt(s2 / denom); + *absmax = amax; + *mean = sum / denom; + }; + auto dump_rows = [&](const char* name, const std::vector& v) { + if (dump_dir == nullptr) return; + std::string path = std::string(dump_dir) + "/" + name; + std::FILE* f = std::fopen(path.c_str(), "wb"); + if (f == nullptr) return; + std::fwrite(v.data(), sizeof(float), v.size(), f); + std::fclose(f); + }; + dump_rows("init_video_rows.f32", initial_video_rows); + if (trace_motion) { + double r = 0, a = 0, mn = 0; + rms_absmax(initial_video_rows.data(), static_cast(initial_video_rows.size()), &r, &a, + &mn); + std::fprintf(stderr, "[h3-motion] INIT video_rows n=%lld rms=%.6g absmax=%.6g mean=%.6g\n", + static_cast(initial_video_rows.size()), r, a, mn); + } + const int64_t num_steps = static_cast(sigmas_video.size()) - 1; for (int64_t step = 0; step < num_steps; ++step) { const double s_v = sigmas_video[static_cast(step)]; @@ -892,6 +930,8 @@ MiniMaxH3DenoiseResult MiniMaxH3DenoiseLoop( } }; + std::vector motion_prev; + if (trace_motion) motion_prev = result.video_rows; advance(result.video_rows, velocity.video_logits, packed.update_mask, num_img, video_width, t_v, s_v, s_v_next); cond_index = 0; @@ -913,6 +953,58 @@ MiniMaxH3DenoiseResult MiniMaxH3DenoiseLoop( static_cast(audio_width) * sizeof(float)); ++audio_ref_index; } + + if (trace_motion) { + double v_rms = 0, v_amax = 0, v_mean = 0; + rms_absmax(velocity.video_logits.data(), + static_cast(velocity.video_logits.size()), &v_rms, &v_amax, &v_mean); + // Motion this step, over the UPDATE (denoise-target) rows only. + double d2 = 0.0; + int64_t dn = 0; + for (int64_t rr = 0; rr < num_img; ++rr) { + if (!packed.update_mask[static_cast(rr)]) continue; + for (int64_t c = 0; c < video_width; ++c) { + const size_t idx = static_cast(rr * video_width + c); + const double diff = + static_cast(result.video_rows[idx]) - static_cast(motion_prev[idx]); + d2 += diff * diff; + ++dn; + } + } + const double drows_rms = std::sqrt(d2 / (dn > 0 ? static_cast(dn) : 1.0)); + double rows_rms = 0, rows_amax = 0, rows_mean = 0; + rms_absmax(result.video_rows.data(), static_cast(result.video_rows.size()), + &rows_rms, &rows_amax, &rows_mean); + std::fprintf(stderr, + "[h3-motion] step %lld/%lld sig_v=%.5f->%.5f v_rms=%.6g v_amax=%.6g " + "v_mean=%.6g drows_rms=%.6g rows_rms=%.6g rows_amax=%.6g\n", + static_cast(step + 1), static_cast(num_steps), + sigmas_video[static_cast(step)], + sigmas_video[static_cast(step + 1)], v_rms, v_amax, v_mean, drows_rms, + rows_rms, rows_amax); + std::fflush(stderr); + } + } + + dump_rows("final_video_rows.f32", result.video_rows); + if (trace_motion) { + double r = 0, a = 0, mn = 0; + rms_absmax(result.video_rows.data(), static_cast(result.video_rows.size()), &r, &a, + &mn); + // Total displacement from the initial noise, over all rows. + double d2 = 0.0; + const size_t n = std::min(result.video_rows.size(), initial_video_rows.size()); + for (size_t i = 0; i < n; ++i) { + const double diff = + static_cast(result.video_rows[i]) - static_cast(initial_video_rows[i]); + d2 += diff * diff; + } + const double disp_rms = std::sqrt(d2 / (n > 0 ? static_cast(n) : 1.0)); + std::fprintf(stderr, + "[h3-motion] FINAL video_rows rms=%.6g absmax=%.6g mean=%.6g " + "disp_from_init_rms=%.6g\n", + r, a, mn, disp_rms); + std::fflush(stderr); } return result; } diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index 2f79cc2a..42d8a6f1 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -27,6 +27,8 @@ #include #include #include +#include +#include #include #include @@ -406,6 +408,22 @@ MiniMaxH3T2vaResult MiniMaxH3GenerateT2va(vt::Device device, const MiniMaxH3T2va video_latent = MiniMaxH3VideoVaePostQuantConv(video_weights, video_latent, dit_params.latents_dim, video_per_channel); } + // DIAGNOSTIC (env-gated, byte-identical when unset): dump the exact latent that + // enters the video VAE (post unpatchify + denormalize + post_quant_conv) as raw + // f32, so the VAE decode can be replayed on a KNOWN latent and runs byte-compared + // (H3 render-coherence bisection at the VAE boundary). + if (const char* dump_dir = std::getenv("VT_H3_DUMP_DIR")) { + std::string path = std::string(dump_dir) + "/vae_input_video_latent.f32"; + if (std::FILE* f = std::fopen(path.c_str(), "wb")) { + std::fwrite(video_latent.data(), sizeof(float), video_latent.size(), f); + std::fclose(f); + std::fprintf(stderr, "[h3-dump] wrote %s (%zu floats, channels=%lld per_channel=%lld)\n", + path.c_str(), video_latent.size(), + static_cast(dit_params.latents_dim), + static_cast(video_per_channel)); + } + } + // On a device, run the ViT3D decoder device-resident. The portable decoder is a // scalar reference; at real resolutions it is the stage that does not finish. It // stays the CPU path, and stays the thing the device path is gated against. From e295c9cba9c3d57781b2176cfef46189f662d1be Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 16:39:16 +0000 Subject: [PATCH 2/6] diag(minimax-h3): env-gated video-VAE receptive-field probe (VT_H3_VAE_PROBE) The render-coherence bisection localized the bug BELOW the denoise loop: final latents differ byte-for-byte at 3/12/50 steps (loop moves the latent; velocity evolves 1.4->5.5 rms), yet the frame is a hard grid of independent 16px blocks at the VAE latent-cell scale. To isolate the video VAE decoder, this perturbs one interior spatial latent cell, re-decodes, and prints the per-16px-block RMS-change map: a single hot block proves the decoder does not mix tokens spatially. Documented; byte-identical when unset. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] Row: row/H3-RENDER-COHERENCE (#70) --- docs/ENVIRONMENT.md | 1 + .../models/minimax_h3_pipeline.cpp | 52 +++++++++++++++++++ 2 files changed, 53 insertions(+) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index a736ebd4..e4d548cd 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -100,6 +100,7 @@ Read-only observability; none change output. | `VT_TTFT_DUMP` | unset | `=1` prints one `TTFTSPLIT rid=... intake=.. queued=.. prefill=.. decode=.. e2e=..` line per finished request to stderr, reconstructing the per-request timing split from the event-populated `req_state` timestamps. The async serving frontend otherwise tracks no per-request stats (passes `iteration_stats=nullptr` and never stamps `EngineCoreOutputs.timestamp`); under this flag both are wired so a serving TTFT attribution can read the queue-vs-execution split against vLLM's own `request_{queue,prefill,decode}_time_seconds`. Generation is byte-identical when unset (the default path is instruction-identical to production); the durable replacement is the async `/metrics` stat logger | | `VT_LOOP_TRACE` | unset | `=1` prints one `LOOPTRACE ...` line per ~1 s window to stderr from the engine busy loop: the full-iteration cadence (`interval`), `process_engine_step` wall (`step`), per-window admits, input-queue residence (`resid` = enqueue-to-drain, the same endpoints `VT_TTFT_DUMP`'s `intake` measures), per-drain admit max and max backlog depth. Diagnoses whether the admission wait is one busy-loop iteration or the input queue is backing up (it attributed the 35B INTAKE deficit to bursty arrival during long prefill steps). Byte-identical when unset: every trace read is guarded, and the enqueue timestamp is stamped only under the flag | | `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) | ## Kernel-internal knobs (deferred) diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index 42d8a6f1..14f125b2 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -437,6 +437,58 @@ MiniMaxH3T2vaResult MiniMaxH3GenerateT2va(vt::Device device, const MiniMaxH3T2va result.frames = MiniMaxH3VideoVaeDecodeTemporalDevice( device, video_config, staged_vae, video_latent, request.latent_t, request.latent_h, request.latent_w, request.num_frames, &result.frame_shape); + + // DIAGNOSTIC (env-gated): VAE receptive-field probe. Perturb ONE interior spatial + // latent cell across all channels + temporal frames, re-decode, and report the + // per-16px-block RMS change on output frame 0. If ONLY the perturbed cell's block + // moves, the decoder is not mixing tokens spatially (render-coherence bisection). + if (std::getenv("VT_H3_VAE_PROBE")) { + const int64_t lh = request.latent_h, lw = request.latent_w, lt = request.latent_t; + const int64_t C = dit_params.latents_dim; + const int64_t per = lt * lh * lw; + const int64_t ch = lh / 2, cw = lw / 2; // center latent cell + std::vector pert = video_latent; + for (int64_t c = 0; c < C; ++c) { + for (int64_t t = 0; t < lt; ++t) { + const int64_t idx = c * per + (t * lh + ch) * lw + cw; + pert[static_cast(idx)] += 8.0f; // large, unambiguous impulse + } + } + MiniMaxH3VideoFrameShape ps{}; + std::vector pf = MiniMaxH3VideoVaeDecodeTemporalDevice( + device, video_config, staged_vae, pert, request.latent_t, request.latent_h, + request.latent_w, request.num_frames, &ps); + const int64_t oh = result.frame_shape.h, ow = result.frame_shape.w, oc = result.frame_shape.channels; + const int64_t ratio = oh / lh; // pixels per latent cell (== vae spatial ratio) + // block-diff map over the lh x lw grid, output frame 0 + std::fprintf(stderr, "[h3-vae-probe] latent %lldx%lldx%lld -> frame %lldx%lld, ratio=%lld, " + "perturbed cell (h=%lld,w=%lld). Per-block RMS |delta| (x1000):\n", + (long long)lt, (long long)lh, (long long)lw, (long long)oh, (long long)ow, + (long long)ratio, (long long)ch, (long long)cw); + const int64_t plane = oh * ow; + for (int64_t bh = 0; bh < lh; ++bh) { + std::string line; + for (int64_t bw = 0; bw < lw; ++bw) { + double s2 = 0.0; int64_t n = 0; + for (int64_t c = 0; c < oc; ++c) { + for (int64_t py = 0; py < ratio; ++py) { + for (int64_t px = 0; px < ratio; ++px) { + const int64_t oy = bh * ratio + py, ox = bw * ratio + px; + if (oy >= oh || ox >= ow) continue; + const int64_t k = c * plane + oy * ow + ox; + const double dd = static_cast(pf[static_cast(k)]) - + static_cast(result.frames[static_cast(k)]); + s2 += dd * dd; ++n; + } + } + } + const int v = static_cast(1000.0 * std::sqrt(s2 / (n > 0 ? n : 1))); + char buf[16]; std::snprintf(buf, sizeof(buf), "%5d", v); line += buf; + } + std::fprintf(stderr, "[h3-vae-probe] %s\n", line.c_str()); + } + std::fflush(stderr); + } } else { result.frames = MiniMaxH3VideoVaeDecode(video_config, video_weights, video_latent, request.latent_t, request.latent_h, request.latent_w, From 70a3fd4673572951f05bba7ba4ca3de86cf790f7 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 16:49:12 +0000 Subject: [PATCH 3/6] diag(minimax-h3): --decode-latent mode to decode a dumped VAE-input latent (no DiT) The VAE-branch oracle test: decode the SAME real latent (dumped via VT_H3_DUMP_DIR) through the device ViT3D path (cuda) OR the scalar CPU reference (cpu, gated vs upstream at 8.9e-8), with no DiT and no conditioning loaded. Divergence between the two localizes a device-kernel bug at real seq; agreement means the grid is inherent to the latent. Self-contained early branch; unused paths and validation unchanged. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] Row: row/H3-RENDER-COHERENCE (#70) --- examples/minimax_h3_gen/main.cpp | 70 +++++++++++++++++++++++++++++--- 1 file changed, 65 insertions(+), 5 deletions(-) diff --git a/examples/minimax_h3_gen/main.cpp b/examples/minimax_h3_gen/main.cpp index 675aca74..552bd9a9 100644 --- a/examples/minimax_h3_gen/main.cpp +++ b/examples/minimax_h3_gen/main.cpp @@ -164,6 +164,7 @@ int main(int argc, char** argv) { std::string device_name = "cpu"; std::string encoder_path, prompt, tokenizer_path, save_embeds_path; std::string first_frame_path, last_frame_path; + std::string decode_latent_path; // diagnostic: decode a dumped VAE-input latent std::vector ref_image_paths; std::string ref_video_prefix, ref_audio_path; double imgvid_noise_aug = 1.0; @@ -188,6 +189,7 @@ int main(int argc, char** argv) { else if (f == "--dry-run") dry_run = true; else if (f == "--denoise-only") denoise_only = true; else if (f == "--dump-params") dump_params = true; + else if (f == "--decode-latent") decode_latent_path = Need(argc, argv, ++i, f); else if (f == "--device") device_name = Need(argc, argv, ++i, f); else if (f == "--encoder") encoder_path = Need(argc, argv, ++i, f); else if (f == "--prompt") prompt = Need(argc, argv, ++i, f); @@ -213,11 +215,14 @@ int main(int argc, char** argv) { // no VAEs, no conditioning, no output path. Requiring them would make the one // tool that works on a checkpoint too large to load unusable on exactly that // checkpoint. - const bool need_vaes = !denoise_only && !dump_params; - const bool need_cond = !dump_params; - if (dit_path.empty() || (need_vaes && (video_vae_path.empty() || audio_vae_path.empty())) || - (need_vaes && out_path.empty()) || - (need_cond && embeds_path.empty() && (encoder_path.empty() || prompt.empty()))) { + const bool need_vaes = !denoise_only && !dump_params && decode_latent_path.empty(); + const bool need_cond = !dump_params && decode_latent_path.empty(); + // --decode-latent needs NO DiT and NO conditioning (its own block validates its + // inputs); the shared check below would otherwise reject the missing --dit. + if (decode_latent_path.empty() && + (dit_path.empty() || (need_vaes && (video_vae_path.empty() || audio_vae_path.empty())) || + (need_vaes && out_path.empty()) || + (need_cond && embeds_path.empty() && (encoder_path.empty() || prompt.empty())))) { std::cerr << "usage: minimax-h3-gen --dit --video-vae --audio-vae " "--prompt-embeds --out [--video-vae-config ] " "[--audio-vae-config ] [--keep-quant] [--steps N] [--frames N] " @@ -278,6 +283,61 @@ int main(int argc, char** argv) { return 0; } + // --decode-latent DIAGNOSTIC: decode a dumped VAE-input latent + // (VT_H3_DUMP_DIR/vae_input_video_latent.f32) directly, with NO DiT and NO + // conditioning, on either device. Lets the device ViT3D decoder be compared + // against the scalar CPU reference (gated vs upstream at 8.9e-8) on the SAME + // real latent -- the VAE-branch oracle test for the render-coherence bisection. + if (!decode_latent_path.empty()) { + if (video_cfg_path.empty() || video_vae_path.empty() || out_path.empty()) { + throw std::runtime_error( + "--decode-latent needs --video-vae, --video-vae-config, --out and --width/--height/--frames"); + } + vllm::MiniMaxH3LatentStats vstats; + vllm::MiniMaxH3VideoVaeDecoderConfig vcfg = + vllm::ParseMiniMaxH3VideoVaeDecoderConfig(ReadJson(video_cfg_path), &vstats); + vllm::SafetensorsFile vfile = vllm::SafetensorsFile::Open(video_vae_path); + vllm::MiniMaxH3AudioVaeWeights vweights = vllm::LoadMiniMaxH3VideoVaeDecoderWeights(vfile); + const vllm::MiniMaxH3ShapePlan plan = vllm::MiniMaxH3ResolveShape( + "t2va", 0.0, frames, height, width, 0, 0); + const int64_t lt = plan.latent_t, lh = plan.height / vllm::kMiniMaxH3VaeRatio, + lw = plan.width / vllm::kMiniMaxH3VaeRatio, ch = vcfg.in_channels; + const int64_t need = ch * lt * lh * lw; + std::ifstream lf(decode_latent_path, std::ios::binary); + if (!lf) throw std::runtime_error("cannot open --decode-latent file"); + std::vector latent(static_cast(need)); + lf.read(reinterpret_cast(latent.data()), need * static_cast(sizeof(float))); + if (!lf) throw std::runtime_error("--decode-latent file too small for [C,T,H,W]"); + std::cerr << "decode-latent: [" << ch << "," << lt << "," << lh << "," << lw << "] on " + << device_name << "\n"; + vllm::MiniMaxH3T2vaResult result; + if (device_name == "cuda") { + vt::Device dev = vt::GetBackend(vt::DeviceType::kCUDA).CreateQueue().device; + vt::Queue vq = vt::GetBackend(dev.type).CreateQueue(); + const vllm::MiniMaxH3VideoVaeDeviceWeights staged = + vllm::StageMiniMaxH3VideoVaeWeights(vq, vcfg, vweights); + result.frames = vllm::MiniMaxH3VideoVaeDecodeTemporalDevice( + dev, vcfg, staged, latent, lt, lh, lw, plan.num_frames, &result.frame_shape); + } else { + result.frames = vllm::MiniMaxH3VideoVaeDecode(vcfg, vweights, latent, lt, lh, lw, + &result.frame_shape); + } + vllm::MiniMaxH3VideoDenormalizePixels( + result.frames, result.frame_shape.channels, + result.frame_shape.t * result.frame_shape.h * result.frame_shape.w); + std::string mkc = "mkdir -p '" + workdir + "'"; + if (std::system(mkc.c_str()) != 0) throw std::runtime_error("cannot create " + workdir); + for (int64_t fr = 0; fr < result.frame_shape.t; ++fr) { + char nm[512]; + std::snprintf(nm, sizeof(nm), "%s/frame_%06lld.ppm", workdir.c_str(), + static_cast(fr)); + WriteFile(nm, vllm::MiniMaxH3WritePpmFrame(result.frames, result.frame_shape, fr)); + } + std::cerr << "decode-latent: wrote " << result.frame_shape.t << " frames to " << workdir + << "\n"; + return 0; + } + // --- 1. DiT --- std::cerr << "loading DiT " << dit_path << (keep_quant ? " (keep-quant)" : "") << "\n"; vllm::MiniMaxH3GgufDit dit; From 9843254165fa5f1d398662897d161e63b7c51a03 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 17:27:25 +0000 Subject: [PATCH 4/6] diag(minimax-h3): --roundtrip VAE encode->decode gold-standard test (+encoder-latent dump) The DiT-independent test that exonerated the VAE: encode a real image through the video VAE encoder, apply post_quant_conv, decode. On the real checkpoint a structured test pattern round-trips to a COHERENT image (same bars, timecode, diagonal) with no grid -- proving the ViT3D decoder, post_quant_conv and the device path are correct, and localizing the render bug to the DiT-produced latent (which is spatially white). VT_H3_DUMP_DIR also dumps the encoder's (in-distribution) latent for spatial-structure comparison against the DiT's. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] Row: row/H3-RENDER-COHERENCE (#70) --- examples/minimax_h3_gen/main.cpp | 77 +++++++++++++++++++++++++++++--- 1 file changed, 72 insertions(+), 5 deletions(-) diff --git a/examples/minimax_h3_gen/main.cpp b/examples/minimax_h3_gen/main.cpp index 552bd9a9..8dc078c6 100644 --- a/examples/minimax_h3_gen/main.cpp +++ b/examples/minimax_h3_gen/main.cpp @@ -30,6 +30,7 @@ #include #include #include +#include #include #include #include @@ -165,6 +166,7 @@ int main(int argc, char** argv) { std::string encoder_path, prompt, tokenizer_path, save_embeds_path; std::string first_frame_path, last_frame_path; std::string decode_latent_path; // diagnostic: decode a dumped VAE-input latent + std::string roundtrip_path; // diagnostic: encode->decode a real image std::vector ref_image_paths; std::string ref_video_prefix, ref_audio_path; double imgvid_noise_aug = 1.0; @@ -190,6 +192,7 @@ int main(int argc, char** argv) { else if (f == "--denoise-only") denoise_only = true; else if (f == "--dump-params") dump_params = true; else if (f == "--decode-latent") decode_latent_path = Need(argc, argv, ++i, f); + else if (f == "--roundtrip") roundtrip_path = Need(argc, argv, ++i, f); else if (f == "--device") device_name = Need(argc, argv, ++i, f); else if (f == "--encoder") encoder_path = Need(argc, argv, ++i, f); else if (f == "--prompt") prompt = Need(argc, argv, ++i, f); @@ -215,11 +218,12 @@ int main(int argc, char** argv) { // no VAEs, no conditioning, no output path. Requiring them would make the one // tool that works on a checkpoint too large to load unusable on exactly that // checkpoint. - const bool need_vaes = !denoise_only && !dump_params && decode_latent_path.empty(); - const bool need_cond = !dump_params && decode_latent_path.empty(); - // --decode-latent needs NO DiT and NO conditioning (its own block validates its - // inputs); the shared check below would otherwise reject the missing --dit. - if (decode_latent_path.empty() && + const bool diag_vae_only = !decode_latent_path.empty() || !roundtrip_path.empty(); + const bool need_vaes = !denoise_only && !dump_params && !diag_vae_only; + const bool need_cond = !dump_params && !diag_vae_only; + // --decode-latent / --roundtrip need NO DiT and NO conditioning (their own blocks + // validate their inputs); the shared check below would otherwise reject --dit. + if (!diag_vae_only && (dit_path.empty() || (need_vaes && (video_vae_path.empty() || audio_vae_path.empty())) || (need_vaes && out_path.empty()) || (need_cond && embeds_path.empty() && (encoder_path.empty() || prompt.empty())))) { @@ -338,6 +342,69 @@ int main(int argc, char** argv) { return 0; } + // --roundtrip DIAGNOSTIC: encode a real image through the video VAE encoder, + // apply post_quant_conv, and decode -- a DiT-independent gold-standard test of + // the decoder. A coherent round-trip proves the decoder works and localizes the + // render bug to the DiT-produced latent; a grid proves the decoder itself. + if (!roundtrip_path.empty()) { + if (video_cfg_path.empty() || video_vae_path.empty()) { + throw std::runtime_error("--roundtrip needs --video-vae and --video-vae-config"); + } + vllm::MiniMaxH3LatentStats vstats; + vllm::MiniMaxH3VideoVaeDecoderConfig vcfg = + vllm::ParseMiniMaxH3VideoVaeDecoderConfig(ReadJson(video_cfg_path), &vstats); + vllm::SafetensorsFile vfile = vllm::SafetensorsFile::Open(video_vae_path); + vllm::MiniMaxH3AudioVaeWeights dec_w = vllm::LoadMiniMaxH3VideoVaeDecoderWeights(vfile); + vllm::MiniMaxH3AudioVaeWeights enc_w = vllm::LoadMiniMaxH3VideoVaeEncoderWeights(vfile); + int64_t ih = 0, iw = 0; + std::vector chw = ReadPpmAsChw(roundtrip_path, &ih, &iw); // [3,H,W] in [0,1] + vllm::MiniMaxH3VideoNormalizePixels(chw, 3, ih * iw); // -> imagenet space + vllm::MiniMaxH3EncoderFcn3dConfig enc_cfg; + enc_cfg.z_channels = 2 * vcfg.in_channels; // moments (mean|logvar) + enc_cfg.t = 1; enc_cfg.h = ih; enc_cfg.w = iw; + vllm::MiniMaxH3VideoFrameShape ls{}; + std::vector z = vllm::MiniMaxH3VideoVaeEncodeToLatent(enc_cfg, enc_w, chw, &ls); + std::cerr << "roundtrip: encoded [" << vcfg.in_channels << "," << ls.t << "," << ls.h << "," + << ls.w << "]\n"; + const int64_t per = ls.t * ls.h * ls.w; + // per-channel stats of the ENCODED latent (the in-distribution reference) + { double s2 = 0; for (float v : z) s2 += double(v) * v; + std::cerr << "roundtrip: encoded-latent rms=" << std::sqrt(s2 / z.size()) << "\n"; } + if (const char* dd = std::getenv("VT_H3_DUMP_DIR")) { + std::string p = std::string(dd) + "/encoder_latent.f32"; + if (std::FILE* fp = std::fopen(p.c_str(), "wb")) { + std::fwrite(z.data(), sizeof(float), z.size(), fp); std::fclose(fp); + std::cerr << "roundtrip: dumped encoder latent [" << vcfg.in_channels << "," << ls.t + << "," << ls.h << "," << ls.w << "] to " << p << "\n"; + } + } + if (dec_w.Has("post_quant_conv.weight")) { + z = vllm::MiniMaxH3VideoVaePostQuantConv(dec_w, z, vcfg.in_channels, per); + } + vllm::MiniMaxH3T2vaResult result; + if (device_name == "cuda") { + vt::Device dev = vt::GetBackend(vt::DeviceType::kCUDA).CreateQueue().device; + vt::Queue vq = vt::GetBackend(dev.type).CreateQueue(); + const vllm::MiniMaxH3VideoVaeDeviceWeights staged = + vllm::StageMiniMaxH3VideoVaeWeights(vq, vcfg, dec_w); + result.frames = vllm::MiniMaxH3VideoVaeDecodeTemporalDevice( + dev, vcfg, staged, z, ls.t, ls.h, ls.w, 1, &result.frame_shape); + } else { + result.frames = vllm::MiniMaxH3VideoVaeDecode(vcfg, dec_w, z, ls.t, ls.h, ls.w, + &result.frame_shape); + } + vllm::MiniMaxH3VideoDenormalizePixels( + result.frames, result.frame_shape.channels, + result.frame_shape.t * result.frame_shape.h * result.frame_shape.w); + std::string mkc = "mkdir -p '" + workdir + "'"; + if (std::system(mkc.c_str()) != 0) throw std::runtime_error("cannot create " + workdir); + WriteFile(workdir + "/roundtrip.ppm", + vllm::MiniMaxH3WritePpmFrame(result.frames, result.frame_shape, 0)); + std::cerr << "roundtrip: wrote " << workdir << "/roundtrip.ppm (" + << result.frame_shape.w << "x" << result.frame_shape.h << ")\n"; + return 0; + } + // --- 1. DiT --- std::cerr << "loading DiT " << dit_path << (keep_quant ? " (keep-quant)" : "") << "\n"; vllm::MiniMaxH3GgufDit dit; From a35cbabbfdff9f0964f788fbd5e61ae2aac13051 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 17:54:25 +0000 Subject: [PATCH 5/6] =?UTF-8?q?diag(minimax-h3):=20VT=5FH3=5FGAUSSIAN=5FNO?= =?UTF-8?q?ISE=20toggle=20=E2=80=94=20flow=20noise=20should=20be=20N(0,1),?= =?UTF-8?q?=20not=20uniform?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Found a real secondary fidelity bug while bisecting: the driver seeded the diffusion initial noise from uniform[-1,1] (std 0.577), but a flow-matching model is trained on Gaussian N(0,1) (torch.randn). The old comment claiming the noise distribution "does not matter" is wrong -- only the exact VALUES (RNG identity) don't matter, the DISTRIBUTION does. VT_H3_GAUSSIAN_NOISE=1 draws Box-Muller Gaussians (INIT rms 0.58 -> 1.0). A/B verdict: it is NOT the render bug (latent adjacent-cos 0.057 uniform -> 0.077 Gaussian, still white vs the 0.789 of a real encoded latent) but is a correctness deviation from upstream worth flipping. Left as a toggle pending the operator's default call. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] Row: row/H3-RENDER-COHERENCE (#70) --- docs/ENVIRONMENT.md | 1 + examples/minimax_h3_gen/main.cpp | 22 +++++++++++++++++++--- 2 files changed, 20 insertions(+), 3 deletions(-) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index e4d548cd..e0babdc5 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -99,6 +99,7 @@ Read-only observability; none change output. | `VT_POOL_BYPASS` | off | `=1` makes every device-scratch pool allocation an exact-size driver `Alloc` and every release a real `Free`, so `compute-sanitizer` can see tensor boundaries and use-after-free that the caching, size-class-rounding pool hides. DEBUGGING ONLY: it reinstates the per-op `cudaMalloc`/`cudaFree` device-sync storm the pool exists to remove, so it is never a timing configuration | | `VT_TTFT_DUMP` | unset | `=1` prints one `TTFTSPLIT rid=... intake=.. queued=.. prefill=.. decode=.. e2e=..` line per finished request to stderr, reconstructing the per-request timing split from the event-populated `req_state` timestamps. The async serving frontend otherwise tracks no per-request stats (passes `iteration_stats=nullptr` and never stamps `EngineCoreOutputs.timestamp`); under this flag both are wired so a serving TTFT attribution can read the queue-vs-execution split against vLLM's own `request_{queue,prefill,decode}_time_seconds`. Generation is byte-identical when unset (the default path is instruction-identical to production); the durable replacement is the async `/metrics` stat logger | | `VT_LOOP_TRACE` | unset | `=1` prints one `LOOPTRACE ...` line per ~1 s window to stderr from the engine busy loop: the full-iteration cadence (`interval`), `process_engine_step` wall (`step`), per-window admits, input-queue residence (`resid` = enqueue-to-drain, the same endpoints `VT_TTFT_DUMP`'s `intake` measures), per-drain admit max and max backlog depth. Diagnoses whether the admission wait is one busy-loop iteration or the input queue is backing up (it attributed the 35B INTAKE deficit to bursty arrival during long prefill steps). Byte-identical when unset: every trace read is guarded, and the enqueue timestamp is stamped only under the flag | +| `VT_H3_GAUSSIAN_NOISE` | unset | `minimax-h3-gen` only: `=1` seeds the diffusion initial noise from Box-Muller GAUSSIAN N(0,1) (what a flow-matching model is trained on) instead of the historical uniform[-1,1] draw. A/B knob for the render-coherence investigation; the exact values still do not match torch's RNG (that only selects WHICH sample), but the DISTRIBUTION does | | `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) | diff --git a/examples/minimax_h3_gen/main.cpp b/examples/minimax_h3_gen/main.cpp index 8dc078c6..e75a0cad 100644 --- a/examples/minimax_h3_gen/main.cpp +++ b/examples/minimax_h3_gen/main.cpp @@ -29,6 +29,7 @@ #include #include +#include #include #include #include @@ -749,15 +750,30 @@ int main(int argc, char** argv) { // NOT reproduce torch's RNG: matching it bit-for-bit decides WHICH sample you // get, not whether the pipeline is correct, and pretending otherwise would // invite comparing our sample against upstream's as if they should match. - auto fill = [](std::vector& out, uint64_t seed) { + // A flow-matching model is trained with GAUSSIAN N(0,1) noise at sigma=1 + // (torch.randn); feeding uniform[-1,1] (std 0.577) is out-of-distribution. + // VT_H3_GAUSSIAN_NOISE=1 draws Box-Muller Gaussians from the same stream for the + // render-coherence A/B; default stays the historical uniform draw. + const bool gaussian = std::getenv("VT_H3_GAUSSIAN_NOISE") != nullptr; + auto fill = [gaussian](std::vector& out, uint64_t seed) { uint64_t x = seed; - for (float& v : out) { + auto u01 = [&x]() { x += 0x9E3779B97F4A7C15ULL; uint64_t z = x; z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ULL; z = (z ^ (z >> 27)) * 0x94D049BB133111EBULL; z ^= z >> 31; - v = static_cast((z >> 11) * 0x1.0p-53 * 2.0 - 1.0); + return (z >> 11) * 0x1.0p-53; // [0,1) + }; + for (size_t i = 0; i < out.size(); ++i) { + if (gaussian) { + double u1 = u01(), u2 = u01(); + if (u1 < 1e-12) u1 = 1e-12; + out[i] = static_cast(std::sqrt(-2.0 * std::log(u1)) * + std::cos(2.0 * 3.14159265358979323846 * u2)); + } else { + out[i] = static_cast(u01() * 2.0 - 1.0); // uniform [-1,1] + } } }; std::vector noise_video( From cb0dbde46766c313fa09a86773b809a75d084634 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 18:03:31 +0000 Subject: [PATCH 6/6] =?UTF-8?q?docs(minimax-h3):=20render=20bug=20ROOT-CAU?= =?UTF-8?q?SED=20=E2=80=94=20re-localize=20from=20VAE=20to=20the=20DiT=20(?= =?UTF-8?q?#70)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Records the latent-bisection verdict: the VAE decoder is CORRECT (a real image encode->decode round-trip is coherent), the denoise loop moves the latent step-dependently, and the DiT emits a spatially-WHITE latent at real geometry (adjacent-cell cosine 0.06 vs 0.789 for a real encoded latent) — which the VAE faithfully renders as one independent patch per token = the grid. Not fp4 (bf16 equally white), not the attention kernel (MMA==chunk, chunk==warp==keylane), not the init noise. Overturns the #64 "device VAE decode / denoise convergence" framing. Secondary: driver init noise was uniform, not Gaussian. Updated: state.md (anchored), benchmark-record.md, spec minimax-h3 §8.4, NOW.md, STATUS/BENCHMARKS/FEATURES H3 rows. row/H3-RENDER-COHERENCE. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 46 ++++++++++++++++++++++++++++++ .agents/specs/minimax-h3.md | 15 +++++++++- .agents/state.md | 56 +++++++++++++++++++++++++++++++++++++ docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- 7 files changed, 120 insertions(+), 5 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 9117103d..92328e3b 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 | **GB10 leg LANDED** (#64): Marlin W4A16 byte-exact, fp4=MEMORY win; e2e RUNS, frame=non-scene | **OPEN: render bug**; fp4 CLOSED | +| MiniMax-H3 lane | **Render bug ROOT-CAUSED** (#70): VAE FINE, **DiT latent spatially WHITE** (0.06 vs 0.79) | Pin DiT line; fp4 CLOSED | | 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 d52b6c1f..c01dabeb 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -13925,3 +13925,49 @@ and we are now at ~281 on `131,2048,512`, so we are **ahead of ggml's own kernel by ~1.3x**; ggml WITH llamafile is ~419, so ~1.5x still separates us and that residue is the FMA plus K-vectorised-hsum structure we are deliberately not adopting. + +## MiniMax-H3 render-coherence ROOT-CAUSED — VAE decoder is CORRECT (encode->decode round-trip is coherent); the DiT emits a spatially-WHITE latent at real geometry (2026-08-06, `row/H3-RENDER-COHERENCE` PR #70, `ROAD-V1-H3`, dgx GB10 sm_121a) + +**Setup.** dgx.casa GB10, dual-lock (`$HOME/gpu.lock`+`/tmp/gpu`), worker down, +`drop_caches` before loads. Real ~39 GB NVFP4 arm cached at `~/h3fp4/ckpt` +(DiT `minimax_h3_ref2va_nvfp4_full.safetensors` 18.75 GB + both VAEs + GGUF +encoder). Diagnostics built on `7d05aee9`: env-gated `VT_H3_TRACE_MOTION` +(per-step velocity + latent-motion), `VT_H3_DUMP_DIR` (raw f32 latents), +`VT_H3_VAE_PROBE` (receptive field), `VT_H3_GAUSSIAN_NOISE`, and two driver +modes `--decode-latent` / `--roundtrip`. Spatial coherence metric = mean cosine +of adjacent latent-cell channel-vectors (16x16 grid) vs random-pair cosine. + +**Latent bisection (each rung MEASURED at 256x256/22f):** +| Rung | Measurement | Verdict | +|---|---|---| +| Denoise loop | final latents byte-DIFFER at 3/12/50 steps (md5; s12-s50 corr 0.25); velocity 1.40->5.52 rms; disp-from-init ~1.9 | loop MOVES latent, not frozen -> not upstream-loop | +| **VAE decoder** | `--roundtrip`: real test pattern encode->post_quant_conv->decode returns a COHERENT image (bars/timecode/diagonal), no grid | **DECODER CORRECT** (overturns #64 "device VAE decode" framing) | +| DiT latent | adjacent-cos: encoded (coherent) **0.789** vs DiT **0.06**; laplacian 0.93 vs 4.2; 8x8-token adj-cos ~0 (== random) | DiT latent is spatially WHITE -> the grid | +| fp4 vs bf16 | both white (cos 0.057 / 0.040); NOT byte-exact on real weights (max|diff| 11.2) | not the residency/precision path | +| DiT attn kernel | MMA vs chunk (`VT_DFLASH_ATTN_MMA=0`): both white (0.057/0.056) | not the attention kernel | +| VAE attn kernel | f32 chunk/warp/keylane identical (max|diff| 1/255) | not the VAE attention kernel | +| init noise | Gaussian vs uniform: 0.077 vs 0.057 | marginal, not the fix | + +**Why the gates missed it.** DiT-forward gate runs latent 4x6 (spatial 2x3 = 6 +tokens) and matches upstream at **1.6e-7** there, so spatial mixing is correct +at small scale; the divergence is real-geometry-specific, appearing only between +2x3 and the real 8x8 (64 tokens) — same shape as the earlier temporal-chunking +and VAE-tiling misses (every gate sits below the regime that breaks). + +**The mechanism of the grid.** The ViT3D decoder's `proj_out` maps EACH latent +token to its own 16-px pixel patch; cross-patch coherence comes only from the +latent's spatial structure. Given a real (spatially-coherent) latent it renders +a coherent image (round-trip proves this); given the DiT's spatially-white +latent it renders one independent smooth block per cell = the observed grid, +identically at any step count (the loop only decides WHICH white latent). + +**Secondary real bug.** Driver seeded uniform[-1,1] init noise (std 0.577); a +flow model needs Gaussian N(0,1) (torch.randn). `VT_H3_GAUSSIAN_NOISE=1` fixes +INIT rms 0.58->1.0; correctness deviation, not the render fix. + +**Residual / handoff.** Bug RE-LOCALIZED to the DiT forward's spatial mixing at +real geometry (video tokens not mixed). Exact line needs an upstream (vllm-omni) +DiT-activation diff at real geometry — impractical on one GB10 (no quantized +vllm-omni H3 arm; bf16 is 4xB300). Diagnostics committed on PR #70; latents at +`dgx:~/h3fp4/diag`. fp4 speed path unchanged/CLOSED. No source/kernel/model/gate +mark changed (all additions are env-gated diagnostics, byte-identical when off). diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 4a7d4b29..2e14610d 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -463,7 +463,20 @@ video flow_shift 12 / audio 3, no CFG; default canvas **768×1344**, default fra fp4-resident DiT → both VAEs → ffmpeg, producing a valid `h264 256×256 + AAC 32 kHz` mp4 + wav. **But the decoded frame is a structured multicolour patch-grid at the latent-cell scale, NOT a coherent scene — identically at 12/20/50 steps, conditioned - or not.** So the composed path is proven to RUN e2e on the real checkpoint, but a + or not.** + - **ROOT-CAUSED (2026-08-06, `row/H3-RENDER-COHERENCE` PR #70) by latent + bisection:** the VAE decoder is **CORRECT** — a real image encode→post_quant_conv→ + decode round-trip (`--roundtrip`) returns a coherent frame — and the denoise loop + moves the latent step-dependently (byte-different finals at 3/12/50 steps). The bug + is the **DiT forward emitting a spatially-WHITE latent** at real geometry: adjacent + latent-cell cosine is **0.06** vs **0.789** for a real encoded latent, so every VAE + token decodes an independent patch = the grid. NOT fp4 (bf16 equally white), NOT the + attention kernel (MMA≡chunk, VAE chunk≡warp≡keylane), NOT the init noise. The DiT + gate runs spatial 2×3 (matches upstream 1.6e-7); the divergence is real-geometry + only (2×3→8×8). Secondary: driver used uniform init noise, not Gaussian + (`VT_H3_GAUSSIAN_NOISE`). Exact DiT line pends an upstream-oracle diff at real + geometry. See the benchmark record + state entry. + - So the composed path is proven to RUN e2e on the real checkpoint, but a coherent render is an OPEN bug (device video-VAE decode and/or denoise convergence at real geometry), independent of the fp4 speed work. **DiT s/step (full 50-layer fp4-resident, per forward):** 5.45 s @512×512/22f, 20.03 s @768×768/61f, 209.09 s diff --git a/.agents/state.md b/.agents/state.md index 212e4488..130e121d 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39404,3 +39404,59 @@ convergence at real geometry) — the frame-sanity gate caught it (all unit gate green + valid mp4, yet a non-scene). fp4 speed path is CLOSED. Box left clean (GPU idle, locks free, worker down, disk ≥15 G; checkpoint cached for reruns). Benchmark record + spec §8 + STATUS/BENCHMARKS/FEATURES + model-matrix/roadmap updated. + +## 2026-08-06T20:30 - H3 render-coherence ROOT-CAUSED by latent bisection: VAE is FINE (round-trip coherent), the DiT emits a spatially-WHITE latent at real geometry + + +`row/H3-RENDER-COHERENCE` (helper, DRAFT PR #70, off `7d05aee9`). The #64 grid +(regular 16-px multicolour blocks, IDENTICAL 12/20/50 steps, cond or not) is +NOT the VAE and NOT the loop. Bisected the pipeline at the VAE boundary on the +real ~39 GB NVFP4 arm (dgx GB10, dual-lock, worker down) with env-gated +instrumentation (`VT_H3_TRACE_MOTION`, `VT_H3_DUMP_DIR`, `VT_H3_VAE_PROBE`) and +two new driver diagnostics (`--decode-latent`, `--roundtrip`). + +**The bisection ladder (each rung a measurement, not an argument):** +1. **Loop MOVES the latent, step-dependent -> upstream is NOT frozen.** Final + latents at 3/12/50 steps are byte-DIFFERENT (distinct md5, s12-vs-s50 corr + 0.25); velocity evolves 1.40->5.52 rms as sigma 1->0; disp-from-init ~1.9. + The Euler integral telescopes to `(sig0-sigN)*v ~= v`, so "identical across + steps" would have meant a frozen velocity; it is not frozen. +2. **The VAE DECODER IS CORRECT (the decisive rung).** `--roundtrip` encodes a + real test pattern through the video-VAE encoder, applies `post_quant_conv`, + and decodes: the frame comes back COHERENT (same colour bars, timecode, + diagonal), no grid. So ViT3D decoder + post_quant_conv + temporal decode + + pixel-denormalize all work on an in-distribution latent. This OVERTURNS the + #64 "device VAE decode and/or denoise convergence" framing. +3. **The DiT emits a spatially-WHITE latent.** Adjacent-cell cosine at the 16x16 + VAE-token scale: a real ENCODED latent (coherent decode) = **0.789** + (laplacian 0.93); the DiT-produced latent = **0.06** (laplacian 4.2). At the + 8x8 DiT-patch scale adjacent tokens are as uncorrelated as random pairs. The + VAE faithfully renders that white latent as one independent 16-px patch per + token (`proj_out` maps each token to its own patch), which IS the grid. +4. **Not fp4, not the attention kernel, not the noise.** fp4-resident vs bf16: + BOTH white (cos 0.057 vs 0.040), so it is not the residency/precision path + (they are NOT byte-exact on real weights, max|diff| 11.2, but both wrong). + DiT bf16 MMA kernel vs the CUDA-core chunk kernel (`VT_DFLASH_ATTN_MMA=0`): + BOTH white (0.057 vs 0.056). VAE f32 chunk/warp/keylane kernels: identical + (max|diff| 1.0/255). Gaussian vs uniform init noise: 0.077 vs 0.057 (marginal). +5. **The gate blind spot.** The DiT forward gate runs at latent 4x6 (spatial + 2x3 = 6 tokens) and matches upstream at 1.6e-7 there, so spatial mixing IS + exercised at small scale and is correct; the divergence appears only between + 2x3 and the real 8x8 (64 tokens). A real-geometry-scaling divergence in the + DiT forward's spatial mixing -- the same shape as the temporal-chunking miss + (every gate below one chunk/tile) and the VAE-tiling miss. + +**Secondary real bug found:** the driver seeded uniform[-1,1] init noise +(std 0.577); a flow model wants Gaussian N(0,1) (torch.randn). `VT_H3_GAUSSIAN_NOISE=1` +fixes the distribution (INIT rms 0.58->1.0) but not the coherence. The "noise +distribution doesn't matter" comment was wrong (only the RNG identity doesn't). + +**Verdict / handoff.** Render bug RE-LOCALIZED: **VAE exonerated, bug is in the +DiT forward's spatial mixing at real geometry** (video tokens not spatially +mixed). Exact line needs an upstream-oracle diff at real geometry (vllm-omni +DiT activations), which is impractical on one GB10 (the quantized arm does not +run in vllm-omni; bf16 is 4xB300). Diagnostics + `--decode-latent`/`--roundtrip` +are committed on the row for the next session. fp4 speed path stays CLOSED. +Box left clean (GPU idle, both locks free, worker down, ~39 G ckpt cached at +`dgx:~/h3fp4/ckpt`, diagnostic latents at `dgx:~/h3fp4/diag`). Evidence: +`dgx:~/h3fp4/{diag,rt_out,out_small,out_cond}`; PR #70. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 5324fd6b..6ec8a77d 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -299,7 +299,7 @@ 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 (OPEN bug) | Root-cause H3 render coherence (VAE/denoise); 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) but frame is a non-scene patch-grid | Render coherence ROOT-CAUSED (#70): VAE fine, DiT latent spatially white. fp4 speed CLOSED. Detail: benchmark-record + spec §8 | | 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** banks +951us marlin. `FUSED-GLUE` W0: glue-into-marlin REFUTED (vLLM doesn't fuse into extern marlin); c8 residual FLASH-dominant, not glue. Flash same-tool audit OWED; see record | | 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 5c96d595..f8caf608 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -115,7 +115,7 @@ oversight. | 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 18.75 GB NVFP4 DiT + VAEs + GGUF encoder → valid mp4/wav); Marlin W4A16 GB10-verified byte-exact; render COHERENCE open (non-scene patch-grid) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| 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) | ☐ | ☐ | | 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 eb1e7750..152607d3 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -73,7 +73,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