From e48c862533512c8094fbf9f8646dcefd90dc028b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 00:09:37 +0000 Subject: [PATCH 1/2] port(ltx-2.5): L9C -- the connector on the render path, and a pool that is drained FOLLOWING_AGENTS_PROTOCOL Two gaps stood between the L9B render and a real scene at a real resolution. This closes the second one and makes the first REACHABLE rather than closing it. THE CONNECTOR WAS NEVER CALLED. `Ltx2ConnectorForward` landed at L5, gated against upstream on five arms, and was called by NOTHING outside its own test. The render path handed the prompt-embeds file straight to the DiT's cross-attention, and the two `*_embeddings_connector` families the shipped DiT carries -- 129 tensors each, in the file -- were refused as "unported" and stepped over. Upstream never does that: `EmbeddingsProcessor` runs an 8-layer 1-D transformer over the caption projections before the DiT sees them (embeddings_processor.py:70-95), and those weights ship INSIDE the DiT file because upstream's key ops rewrite them into the text encoder (encoder_configurator.py:331-346). So they load here, beside the DiT and not inside it, and the render applies them. PRESENCE DECIDES, not a flag: the connector runs exactly when the checkpoint carries it, which is the rule `DetectLtx2Video` already uses on the DiT. Half a connector is refused -- conditioning the picture on eight transformer layers and the soundtrack on none is not a degraded render, it is two different requests. THE CONFIG IS READ, AND ONE OF ITS VALUES IS NOT NEAR ITS DEFAULT. LTX-2.5 declares `connector_positional_embedding_max_pos = [4096]` where `Embeddings1DConnector`'s class default is `[1]`, and `get_fractional_positions` DIVIDES the token index by it (rope.py:132-141). No shape can see the difference; the default is every RoPE angle wrong. `Ltx2ParseConnectorConfig` mirrors both configurators key for key, including the audio one's fallback onto the VIDEO spellings, and deliberately does NOT read the DiT's `positional_embedding_theta`, because neither configurator passes it. THE PROCESSOR IS A SEPARATE PORT, and it carries the thing the connector's own goldens cannot see: upstream RIGHT-PAD-SORTS the features before the connector ("Connectors expect right-padded input"), because the register table is indexed by ABSOLUTE position. A left-padded batch handed straight in puts registers where caption tokens belong and is finite, correctly shaped and conditioned on nothing. It also zeroes the VIDEO encoding at masked positions and NOT the audio one; that asymmetry is upstream's and is mirrored rather than tidied. WHERE THE TWO REFERENCES DISAGREE, reported rather than silently picked: `diffusers` folds that sort INTO the connector (pipelines/ltx2/connectors.py) and its comment claims that matches "the original LTX implementation" -- true only because `ltx_core` sorts one level up, in the processor. They compose to the same function and differ in which module owns it. This follows `ltx_core`, so `Ltx2ConnectorForward` stays a faithful port of `Embeddings1DConnector`. THE POOL WAS NEVER DRAINED. MiniMax-H3 calls `ActivePool()->Drain()` at the denoise-to-decode boundary and records why (minimax_h3_pipeline.cpp:546-563); the LTX path called it nowhere, and GB10's `device_pool_cap_bytes` is 0, so every activation size class a phase touched was retained and unreusable by the next phase at twice the resolution. On unified memory those are host bytes and this box reboots instead of OOM-killing. The drain lands at every phase boundary, which is also the denoise-to-decode boundary for the last phase. `VLLM_LTX2_POOL_DRAIN=0` exists ONLY so the A/B that measures it runs on one binary; it is not a configuration. WHAT IS STILL SYNTHETIC, stated so it cannot be read as closed: the Gemma-4 tower is not ported, so what ENTERS the connector is still whatever the caller put in the prompt-embeds file. `prompt_embeds_valid_rows` exists because a file carries no tokenizer mask and padding is not inert -- it is what the register table replaces. Issue: #435 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/benchmark-record.md | 150 ++++++++++++ .agents/porting-inventory.md | 60 +++++ docs/BENCHMARKS.md | 2 +- docs/ENVIRONMENT.md | 1 + docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 58 +++-- examples/ltx2_gen/main.cpp | 12 +- .../model_executor/models/ltx2_connector.h | 53 ++++ .../vllm/model_executor/models/ltx2_loader.h | 64 +++++ include/vllm/multimodal/ltx2_video.h | 23 ++ .../model_executor/models/ltx2_connector.cpp | 101 ++++++++ .../model_executor/models/ltx2_loader.cpp | 227 +++++++++++++++++- src/vllm/multimodal/ltx2_video.cpp | 189 ++++++++++++++- tests/vllm/models/test_ltx2_loader.cpp | 43 +++- tests/vllm/models/test_ltx2_pipeline.cpp | 190 +++++++++++++++ tests/vllm/multimodal/ltx2_video_fixture.h | 109 ++++++++- tests/vllm/multimodal/test_ltx2_video.cpp | 189 +++++++++++++++ 18 files changed, 1430 insertions(+), 45 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 298ce2526..546bfce9b 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -19405,3 +19405,153 @@ streams are byte-identical, so upstream's DSpark is lossless on that lane. Evidence: `dgx:~/work/dspark-w6/pinned_{35b,27b}_{on,off}.json`, `xengine.log`, `xengine_ours.log`, `xengine_pinned.log`. + +## LTX-2.5 L9c — the per-phase pool drain is worth 0.11 GiB, and L9B's 58 GB runaway does not reproduce (2026-08-13, `row/LTX25-L9C-CONNECTOR-DRAIN`, issue [#435](https://github.com/mudler/vllm.cpp/issues/435)) + +**No speed number is claimed and none is implied.** LTX-2.5's speed axis is +structurally `PENDING` (spec [ltx-2-5.md](specs/ltx-2-5.md) §0): vLLM-Omni has no +native 2.5 and its diffusers adapter is a black box, so no production-configuration +denominator exists. Every wall clock below is SIZING — how long a render takes on +this box — and is not comparable to anything. + +### What was measured, and on what + +dgx.casa (GB10, sm_121a, 119 GiB unified), one `flock $HOME/gpu.lock` hold, +`local-ai-worker` down at both ends of every arm. Build: Release, CUDA `121a`, +`VLLM_CPP_CUTLASS_DIR=$HOME/cutlass-4.5.0`, `VLLM_CPP_TRITON=ON`, configure log +verified to print `CUTLASS found ... sm120a NVFP4 cutlass GEMM`, +`FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a]` and +`Triton AOT: ... sm_121a`. + +Artifacts, named per spec §3.1 because a render is only a statement about the +files that produced it: + +| | | +|---|---| +| DiT | `vonkaiser/LTX-2.5-FP8-NVFP4` `ltx-2.5-22b-distilled-fp8.safetensors` (21.00B, FP8, 6124 tensors) | +| Video VAE | `Lightricks/LTX-2.5` `ltx-2.5-video-vae-conv-bf16.safetensors` (Conv arm) | +| Audio VAE | `Lightricks/LTX-2.5` `ltx-2.5-audio-vae-bf16.safetensors` + its BWE vocoder | +| Upsampler | `ltx-2.5-latent-spatial-upscaler-x2-bf16-1.0.safetensors` | +| Config | LTX-2.5's **DECLARED** config, `sha256 30d08fad…4b21` — `frequencies_precision=float64`, `av_ca_timestep_scale_multiplier=1000`, `connector_positional_embedding_max_pos=[4096]`. NOT the manifest defaults | +| Conditioning | 128 rows, `--prompt-valid-rows 24`, `sha256 c7dff715…11e8` / `b7ff5ff3…ef0f`. **Synthetic** N(0, 0.2): the Gemma-4 tower is not ported | + +### The drain A/B — same binary, `VLLM_LTX2_POOL_DRAIN`, 320x192 / 25 frames + +| arm | drain | wall (sizing only) | user CPU | peak host RSS | lowest MemAvailable | result | +|---|---|---|---|---|---|---| +| E0 | OFF | 23:39.30 | 1265.63 s | 32.84 GB | **68.23 GiB** | 25 frames, h264 320x192 + AAC | +| E1 | ON | 23:40.13 | 1265.82 s | 32.84 GB | **68.14 GiB** | 25 frames, h264 320x192 + AAC | + +0.8 s of wall and 0.09 GiB of floor separate them, and the floor moves the WRONG +way, so both are noise. The drain's own report says why: + +| geometry | after `generate_lowres` | after `refine` | +|---|---|---| +| 128x128 / 9f | 0.01 GiB | 0.02 GiB | +| 320x192 / 25f | 0.03 GiB | 0.08 GiB | + +**The retained scratch at an LTX phase boundary is 0.11 GiB, not 58 GB.** The +drain is correct, costs one free per retained block, and matches what MiniMax-H3 +does at the same boundary — and it is NOT what makes a bigger render possible. + +**And it is numerically INERT, proven rather than argued.** `diff -r -q` over the +two arms' output directories returns 0: all 25 frames, the WAV and the MP4 are +BYTE-IDENTICAL between drain-off and drain-on (rolled-up md5 `2eba29bf…e656` on +both sides). Draining a free list cannot change arithmetic, and this is what +saying so looks like when it is measured instead of asserted. + +### The ladder, and where L9B's ~58 GB actually lives + +| rung | geometry | denoise | outcome | +|---|---|---|---| +| D0 | 128x128 / 9f | drains 0.01 + 0.02 GiB | **completes**, 9 frames, 6:31 | +| E0/E1 | 320x192 / 25f | drains 0.03 + 0.08 GiB | **completes**, 25 frames, 23:39, floor 68.2 GiB | +| F1 | 448x256 / 25f | drains 0.04 + 0.14 GiB | **STOPPED by the watchdog**, 0 frames | + +**The highest rung that completes is 320x192 / 25 frames.** That is a measurement, +not a ceiling — see the next hypothesis below. + +F1 is where L9B's number turns up, and it is not where L9B put it. Both denoise +phases finished and drained normally, and MemAvailable was flat at **75.2 GiB** +through all of it. Then, AFTER the last drain: + +``` +03:47:11 avail_kB=73014000 rss_kB=4972520 +03:47:21 avail_kB=57800944 rss_kB=4972520 +03:47:31 avail_kB=27711644 rss_kB=4899272 +03:47:35 WATCHDOG_KILL avail_kB=13774472 floor=18000000 +``` + +**~59 GB in 24 seconds with the process's own RSS flat at 4.9 GB.** That is L9B's +~58 GB and L9B's "host RSS flat", reproduced exactly — on the DECODE side of the +denoise-to-decode boundary, not inside the denoise loop L9B attributed it to. No +drain can shrink it: the drain runs before it, and what it returns is 0.14 GiB. + +At 320x192/25f the same code path completes with the drain OFF and MemAvailable +flat, so L9B's own arm — which was 320x192 at `--max-phase 0`, a SMALLER decode +than E0's — remains unreproduced at its own geometry. Both statements are true +and they are about different rungs. + +**Two obvious attributions are each contradicted by a measurement, so neither is +claimed.** `Ltx2ConvVideoDecode` is pure host C++ — no `vt::` op, no queue, no +device pointer anywhere in `ltx2_video_vae.cpp` — and GPU utilization is **0%** +throughout, so it is not a device pool. But `ps -o rss=` reports 4.9 GB across all +twelve samples of the decline, so it is not a plain resident host allocation +either. **Next traceable step: instrument `Ltx2ConvVideoDecode`'s own allocations +directly.** It is ONE function; inferring its footprint from a system-wide counter +is what produced L9B's mis-attribution and would produce another. + +### The instrument, stated because it is weaker than it looks + +`nvidia-smi --query-gpu=memory.used` returns **`[N/A]`** on GB10. There is no +per-process device-memory reading on this box, so "device usage" is only +observable as unified-memory pressure — `MemAvailable` — which moves for anything +on the machine. That is what L9B had too, and it is why L9B's attribution of a +MemAvailable fall to the denoise loop could not have been checked at the time. + +### What bounds the ladder is the DECODE, in both of its costs + +At 320x192/25f the process ran at **0% GPU utilization and ~110% CPU** for most of +the 23:39, and at 448x256/25f it is the decode that takes the box to the floor. So +the decode is the wall twice over — it is the time and it is the memory — +and `Ltx2VideoDecode` is a HOST path: it takes `std::vector` and the host +VAE weights, and `ltx2_video_vae.cpp` contains no `vt::` op at all. MiniMax-H3 has +the device analogue (`MiniMaxH3VideoVaeDecodeTemporalDevice`, +minimax_h3_pipeline.cpp) — and, just as relevantly, H3's decode is CHUNKED in time +and TILED in space by default, so it never materializes the whole canvas at once. +LTX-2.5's does neither. + +**No ceiling is declared.** The next traceable steps, in order: measure +`Ltx2ConvVideoDecode`'s own allocation footprint; then temporal chunking + spatial +tiling, which is what H3 already needed at a real canvas; then the device decode. + +### The frames ARE a scene + +Measured with L9B's OWN analyzer, unchanged, so the numbers are comparable: +frame-to-frame mean |diff| 0.500-1.465 (avg 0.980, so not a still); neighbour +|dx| / whole-image sd **0.093** where white noise gives ~1.13; 8-px block-mean +aligned/offset ratio 1.012 and 32-px 1.193, so not H3's patch grid. + +**Those statistics are almost identical to L9B's, and L9B's frames were not a +scene while these are.** The 25 frames are a temporally coherent photorealistic +clip: one subject, consistent identity and background, frame-to-frame motion. The +statistics could not separate "smooth colour field" from "photograph", and only +looking did. That is a finding about the INSTRUMENT and it belongs beside the +result. + +It is not a depiction of a prompt. With `--prompt-valid-rows 24`, 104 of the 128 +conditioning rows are the connector's own trained `learnable_registers` — which is +exactly what upstream substitutes at padded positions — and the other 24 are +synthetic noise. So what conditions the render is the checkpoint's own learned +default, reached through the real connector, and not anything a caller asked for. + +**The audio is NOT claimed to be anything.** 1.0100 s, 48 kHz stereo, ch0/ch1 rms +26.82/26.83, peak 217/218, zero-fraction 0.024 — so it is not silence, and it is +FAINTER than L9B's at the same settings (rms 131.28/127.30). Whether it is +speech-shaped, or matches the mouth movements in the frames, is not something +these numbers answer and is not asserted. Owed: a spectral check against the +video, which is the audio half of the question the frames just answered. + +Evidence: `dgx:~/work/ltx25-l9c/{hold.log,baselines.log,render-*.log,mem-*.log}`, +renders under `dgx:~/work/ltx25-l9c/render/`, contact sheets under +`dgx:~/work/ltx25-l9c/contact/`. diff --git a/.agents/porting-inventory.md b/.agents/porting-inventory.md index db65f3f9e..08984f416 100644 --- a/.agents/porting-inventory.md +++ b/.agents/porting-inventory.md @@ -1480,6 +1480,66 @@ Examples: `examples/cli` ✅ (C-API client), `examples/server` ✅ (OpenAI serve * **Spec:** [ltx-2.5 spec](specs/ltx-2-5.md) §6 (L8). Lifecycle: shipped (CPU + CUDA, bf16 and f32 streams). Owner: the LTX-2.5 row. +21. **LTX-2.5 phase L9c — the EMBEDDINGS CONNECTOR reaches the render path, and the + two references disagree about which module owns the sort.** + 2026-08-13, `MODEL-DIFFUSION-ltx-2-5-ltx2-video-transformer-3d-model` phase L9c, + issue [#435](https://github.com/mudler/vllm.cpp/issues/435). + * **Upstream source:** Lightricks/LTX-2 @ `fd4ded7f`, + `packages/ltx-core/src/ltx_core/text_encoders/gemma/embeddings_processor.py:70-95` + (`EmbeddingsProcessor.create_embeddings`) and `:23-48` (its three helpers), plus + `embeddings_connector.py:194-256` (the two configurators) and + `encoders/encoder_configurator.py:331-346` (`EMBEDDINGS_PROCESSOR_KEY_OPS`, + which is what says these weights belong to the TEXT ENCODER even though they + ship inside the DiT file). + * **Local anchor:** `Ltx2ConnectorCreateEmbeddings` + (`src/vllm/model_executor/models/ltx2_connector.cpp`), `Ltx2ParseConnectorConfig` + and `Ltx2LoadConnectorWeights` (`ltx2_loader.cpp`), and the call in + `Ltx2VideoEngine::Load` (`src/vllm/multimodal/ltx2_video.cpp`). + * **What was actually wrong before it.** `Ltx2ConnectorForward` landed at L5 and + was gated against upstream on five arms — and its ONLY caller was its own test. + The render handed the prompt-embeds file straight to cross-attention, and the + two `*_embeddings_connector` families (129 tensors each, present in both shipped + DiTs) were reported as unported and stepped over. A brick with a golden and no + caller is indistinguishable, from the outside, from a brick that is wired. + * **THE TWO REFERENCES DISAGREE, and it is recorded rather than resolved by + preference.** `ltx_core` right-pad-sorts the features in the PROCESSOR before + calling the connector ("Connectors expect right-padded input", + `embeddings_processor.py:80-84`); `diffusers` folds that sort INTO the connector + (`src/diffusers/pipelines/ltx2/connectors.py`, the + `torch.argsort(1 - binary_attn_mask, stable=True)` branch) and its comment + claims that matches "the original LTX implementation" — true only because + `ltx_core` does it one level up. They compose to the same function and differ in + which module owns it. This port follows `ltx_core`, so `Ltx2ConnectorForward` + stays a faithful port of `Embeddings1DConnector`. + * **WHERE THEY AGREE AND IT LOOKS LIKE A BUG.** `_to_binary_mask` is + `encoded_mask < 0.000001` (`embeddings_processor.py:46-48`). An additive mask + holds `0.0` for KEPT and `-finfo(f32).max` for PADDED, and BOTH satisfy it — so + the mask handed to the DiT is one at every position, and the video-only multiply + that follows is an identity on every reachable path. `diffusers` writes the + identical `(video_attn_mask < 1e-6)` and the identical multiply. Checked on both + BEFORE mirroring it, because the reading a port arrives at by reasoning about + intent (`>= 0`) is the opposite at padded positions. Gated as the surprising + behaviour, so "fixing" it REDs (`test_ltx2_pipeline`, mutation M3). + * **A CONFIG VALUE THAT IS NOT NEAR ITS DEFAULT.** LTX-2.5 declares + `connector_positional_embedding_max_pos = [4096]` where `Embeddings1DConnector`'s + class default is `[1]`, and `get_fractional_positions` DIVIDES the token index by + it (`rope.py:132-141`). No shape can see the difference; the default is every + RoPE angle wrong. `positional_embedding_theta` is deliberately NOT read from the + DiT config even though one is declared, because neither configurator passes it — + reading it would be a re-invention rather than a port. + * **The record changed, not just the code.** The two connector families are no + longer reported as unported: this port reads them, so naming them would say + something untrue about the tree and would demand `allow_unported_modules` from a + caller whose checkpoint is read completely. Asserted as an ABSENCE in + `test_ltx2_loader`, so restoring the old behaviour REDs (mutation M4). + * **OWED, and precisely:** the Gemma-4 TOWER, so what ENTERS the connector is still + whatever the caller put in the prompt-embeds file — the link below the tower is + real, the tower is not. `prompt_embeds_valid_rows` exists because a file carries + no tokenizer mask and the connector REPLACES padding with learned registers. + * **Spec:** [ltx-2.5 spec](specs/ltx-2-5.md) §6. Lifecycle: shipped (CPU host + module; f32, an annotated escape — it runs once per load and its output is + narrowed to the stream dtype on upload). Owner: the LTX-2.5 row. + ## 10. E2E test suites (T0 deliverable) 1. **Op parity**: golden dumps from upstream vLLM (Python, test-time only) → diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index f6243713c..be956e94d 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -414,7 +414,7 @@ built on it rather than keeping the flattering one. | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | **RUNNER FOLD LANDS (ROW 7, §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (near-tie profile); FA2 MLA default-ON; SACRED green.** Server 19.0 tok/s wall; CLI 18.93 reproduced | vLLM ~21 (#111 floor; in-session re-measure ABORTED by GB10 reboot at util 0.82, §21): **~0.90×**, >= vLLM NOT met; residual = KDA host islands + grouped MoE + decode graph | | 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) | fp4 speed CLOSED. bf16-vs-quant A/B: ENCODER half MEASURED (§8.15), DiT half NOT (no bf16 render exists). Detail: benchmark-record + spec §8 | -| LTX-2.5 axes | BOTH `PENDING`. Speed: no native vLLM-Omni 2.5 (vllm-omni#6066), its diffusers adapter is a black box, so no production denominator. Correctness vs that BINDING oracle also pending | Bricks gate vs the ltx-core cross-check; max abs diff in the spec | +| LTX-2.5 axes | Speed `PENDING` (vllm-omni#6066 has no native 2.5; its adapter is a black box), binding oracle too. **SIZE: 320x192/25f completes on GB10, 448x256/25f does not** | The wall is the HOST VAE decode, not the pool: the drain returns 0.11 GiB, byte-inert | | 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. Follow-up `H3-TASK-PARTITION-GUARD`: the task/partition mismatch now RAISES 1:1 with `_resolve_task` (spec §8.6-8.7) | | MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`) | **fl2va COHERENT; ref2va assembly bug FIXED+gated.** vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the **NVFP4 checkpoint/loader**, NOT assembly/fp4 (§8.10) | | MiniMax-H3 Thor render speed (sm_110, no FA2) | **34.6 s/step** at 864x480/124f/50 steps on Q4_K_M, **16.6x** off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 `mma.sync` (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head's K+V is 3.9 MB against 32 MB of L2) | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index ca51e6222..d610f3988 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -157,6 +157,7 @@ Read-only observability; none change output. | `VT_FP4_AUTOTUNE_VERBOSE` | off | Log the NVFP4 GEMM autotuner's tactic selection | | `VT_H3_PROGRESS` | unset | Trace the MiniMax-H3 denoise loop's phases to stderr: which forward path was taken (device vs the CPU reference), how long the ONE-TIME device weight staging took, and per-step forward seconds with the sequence length. A real-checkpoint run spends its minutes in exactly one of those phases, and this says which without guessing — it was added after GPU-utilization counters proved unreliable on Tegra-class boards | | `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 | +| `VLLM_LTX2_POOL_DRAIN` | on | `=0` suppresses the LTX-2.5 render's per-phase device scratch-pool drain. The drain is on by default and there is no supported reason to turn it off: a denoise phase leaves the pool holding every activation size class it touched, the next phase (twice the resolution) and the VAE decode allocate different classes, and on an uncapped pool those blocks are headroom the next stage never gets. This exists so the A/B that MEASURES what the drain is worth runs on ONE binary, which is what a same-binary A/B requires. Not a 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 | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 8677af152..42487871a 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -149,7 +149,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 79/79; all three modalities COHERENT on Q4_K_M (§8.20); PRUNED ckpts run, Q8_0 seam 0.9941 (§8.21); ref2va grid was NVFP4 quant error, §8.9 REFUTED; GGUF/NVFP4/bf16 shards stream | FP4/Marlin landed; speed pending; no bf16 render yet. Render from the Q4_K_M GGUF, not the NVFP4 arm. Krea 2 text-to-image (roadmap C11) is scoped to reuse these DiT seams | -| LTX-2.5 DiT (`LTX2VideoTransformer3DModel`, Lightricks lane) | LTX-2.5 (21.00B video+audio) | `SPIKE`. DiT, both VAE decoders, the pipeline and the quantized loaders gated vs upstream at reduced dims: max\|diff\| 1.2e-07 DiT, 1.8e-06 VAE, 2.8e-06 pipeline. SHIPPED weights render e2e at 128x128/9f, both distilled phases, valid MP4+WAV. Frames are NOT a scene | Family `ltx-2.5`, driven by `ltx2-gen`. FP8 DiT stages on GB10 (~44 GB); 320x192/25f exhausts unified memory. First-party NVFP4 does NOT load. Text tower, connector, DiffVAE, LoRA OWED, so conditioning is incomplete | +| LTX-2.5 DiT (`LTX2VideoTransformer3DModel`, Lightricks lane) | LTX-2.5 (21.00B video+audio) | `SPIKE`. DiT, both VAE decoders, the pipeline and the quantized loaders gated vs upstream at reduced dims: max\|diff\| 1.2e-07 DiT, 1.8e-06 VAE, 2.8e-06 pipeline. SHIPPED weights render e2e at 320x192/25f, both distilled phases, valid MP4+WAV. With the connector wired the frames ARE a scene: a temporally coherent photorealistic clip | Family `ltx-2.5`, driven by `ltx2-gen`. FP8 DiT stages on GB10 (~44 GB). The embeddings connector RUNS with the checkpoint's own weights; scratch drains per phase. NVFP4 does NOT load. Text tower, DiffVAE, LoRA OWED | | 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 | diff --git a/docs/STATUS.md b/docs/STATUS.md index b9cc7b033..dce6457cd 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -106,7 +106,7 @@ token-for-token correctness against the pinned oracle. | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: was 87% of vLLM (37.55 vs 43.10, same-tool nsys)**; root cause was bf16 projections on UNIFIED/ATS host memory, and device-resident staging (byte-exact) gives 44.6, parity+ vs 43.1, default-ON | 48 layers (12 global + 36 SWA-512), 256 routed top-10 + 1 shared expert, per-head softplus attn out-gate, sigmoid `noaux_tc` router, dual per-layer RoPE, GQA 8 KV / 128 head-dim, 1M ctx. History: benchmark-record | | InternLM2 dense (fused-`wqkv` interleaved split) | Correctness-complete, speed-pending | Token-exact 16/16 (internlm2-chat-1_8b): 12/16 strict + 4/16 bf16 near-tie (max gap 0.0 nats), 0 divergent; first InternLM model; ZERO new compute kernel (reuses the Llama dense forward; the only delta is a loader-side de-interleave of the fused `wqkv`, which packs q/k/v interleaved by KV-group) | | MiniMax-H3 (`MiniMaxH3DiTModel`, video+audio DIFFUSION) | **ABI v12 ONE SURFACE; device selector uses generic `DeviceType`; DSR 32.** t2va+fl2va COHERENT; bf16 shards STREAM | ref2va ckpt fidelity §8.12; encoder A/B §8.15; GB10 re-verify residual; CPU fold 6/137 (one queue + device provenance mutation-gated) | -| LTX-2.5 (`LTX2VideoTransformer3DModel`, video+audio DIFFUSION) | **L1-L8 landed (#435).** 21.00B / 48 blocks. `VideoEngine` seam + ABI **v18**, DiT forward (CPU f32 parity, bf16 device-resident), Gemma-4 TE, both VAEs, pipeline, NVFP4/FP8 arms, `/v1/videos` | A shipped 21.00B FP8 DiT stages and RUNS device-resident on GB10, all ops `vt-native`, zero reference-tier hits. Speed axis `PENDING` (spec section 0: no production-config denominator); binding-oracle parity `PENDING` | +| LTX-2.5 (`LTX2VideoTransformer3DModel`, video+audio DIFFUSION) | **L1-L9c landed (#435).** 21.00B / 48 blocks. `VideoEngine` seam + ABI **v18**, DiT forward (CPU f32 parity, bf16 device-resident), Gemma-4 TE, both VAEs, the embeddings connector, pipeline, NVFP4/FP8 arms, `/v1/videos` | A shipped 21.00B FP8 DiT stages and RUNS device-resident on GB10, all ops `vt-native`, zero reference-tier hits. With the connector wired the frames ARE a scene at 320x192/25f; 448x256 stops in the HOST VAE decode. Speed axis `PENDING` (spec section 0: no production-config denominator); binding-oracle parity `PENDING`; the Gemma-4 TOWER still owed, so nothing encodes a prompt | | Command-R / Cohere dense (`CohereForCausalLM`) | Implemented, gate-blocked | ZERO-new-kernel port grounded in vLLM `commandr.py`: weight-only Cohere LayerNorm + GPT-J full-width RoPE + PARALLEL residual + `logit_scale` + tied embeddings, all reuse; compiles, links, self-registers. No SACRED gate yet (real checkpoints HF-gated, ungated ones tiny-random, GPU box disk-full); oracle run-verified at W0. See docs/BENCHMARKS.md | | Phi-1 / Phi-2 dense (`PhiForCausalLM`, parallel residual) | Correctness-complete, speed-pending | Token-exact 16/16 (microsoft/phi-2): 9/16 strict + 7/16 bf16 near-ties (max gap 0.25 nats), 0 forward-divergent; the OLDER Microsoft Phi arch, DISTINCT from Phi-3/Phi-4; ZERO new compute kernel (GPT-J parallel residual, LayerNorm-with-bias, biased qkv/dense, partial NeoX rope 32/80, non-gated NewGELU MLP reusing `vt::GeluTanh`, untied biased lm_head); F16 dtype-aware loader | | MiniCPM dense (`MiniCPMForCausalLM`, three scalars) | Correctness-complete, speed-pending | Token-exact 16/16 (openbmb/MiniCPM-2B-sft-bf16): 10/16 strict + 6/16 bf16 near-ties (max gap 0.0 nats), 0 forward-divergent; first OpenBMB MiniCPM model; ZERO new compute kernel (the Llama/Granite dense forward plus three scalars: scale_emb, scale_depth/sqrt(layers) residual, dim_model_base logit scaling), tied lm_head; `.bin`-only weights converted to safetensors via trusted torch | diff --git a/docs/USAGE.md b/docs/USAGE.md index 42b0a64ca..05e15d71e 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -331,6 +331,17 @@ with the same row count in both. Supplying a `--prompt` is refused, and supplyin only one of the two files is refused, because a stream left unconditioned renders instead of failing. +**Those rows go through the embeddings connector.** Both shipped LTX-2.5 DiTs +carry two `*_embeddings_connector` families, 129 tensors each, and they are the +8-layer 1-D transformer upstream runs between the caption projections and the +DiT's cross-attention. The render applies it with the checkpoint's own weights, +under the checkpoint's own `connector_*` configuration. Two consequences for the +command line: the row count must be a multiple of the connector's learnable +register count (128 on the shipped files), and `--prompt-valid-rows N` says how +many of those rows are real tokens. The rest are padding, and padding is not +inert here: the connector REPLACES it with its learnable register table, so a +run that leaves the default renders as if every supplied row were caption. + **The DiT config is required when the checkpoint does not carry one.** The shipped `vonkaiser` FP8 transformer has no `__metadata__` at all, and the values a config decides are ones no tensor shape encodes: `frequencies_precision` and @@ -369,21 +380,27 @@ knobs the flags above map onto. Both are described under **Two things about that command are worth knowing before you run it.** -*It is bounded by memory well below the recipe's own defaults.* Staging the -21.00B FP8 transformer costs about 44 GB on a 119 GB GB10, and the render then -runs a forward per denoise step on top of that. A 128x128 clip at 9 frames -completes comfortably through both distilled phases; a 320x192 clip at 25 frames -consumed a further 58 GB and had to be stopped. Unified memory makes those host -bytes, and this class of box reboots rather than OOM-killing, so start small and -grow. The recipe default (1024x1536 at 121 frames) is far beyond what one GB10 -holds today. - -*It cannot yet render a scene, and that is conditioning, not the model.* The -text tower is not ported and neither is the embeddings connector that sits -between it and the DiT's cross-attention, so whatever you supply as prompt -embeds reaches the transformer unprocessed. A completed render at the sizes -above is a structurally valid MP4 of smooth colour fields, not a depiction of -anything. Treat it as an end-to-end plumbing check. +*It is bounded by the VIDEO DECODE, well below the recipe's own defaults.* +Staging the 21.00B FP8 transformer costs about 44 GB on a 119 GB GB10. **320x192 +at 25 frames completes** through both distilled phases; 448x256 at 25 frames +finishes its denoise and then loses about 59 GB in 24 seconds inside the decode +and has to be stopped. The denoise itself is flat at either size. Unified memory +makes those host bytes and this class of box reboots rather than OOM-killing, so +start small and grow, and put a memory watchdog in front of anything larger. The +recipe default (1024x1536 at 121 frames) is far beyond what one GB10 holds today. +Expect minutes, not seconds: most of a 320x192/25f render is spent single-threaded +in the host VAE decode at 0% GPU. + +*It renders a scene, and it does not render YOUR scene.* With the connector on +the path the shipped 21.00B FP8 transformer produces a temporally coherent +photorealistic clip at 320x192 / 25 frames: consistent subject, consistent +background, frame-to-frame motion. Before the connector was wired the same +weights at the same settings produced smooth colour fields. What conditions it, +though, is mostly the connector's own trained `learnable_registers` table, which +is what upstream substitutes at PADDED positions — so the render is the model's +own default, not a depiction of anything you asked for. The Gemma-4 text tower is +still not ported, so the rows you supply as prompt embeds are not an encoded +prompt. Ask for a subject and you will not get it. LTX-2.5 ships two video decoders behind one checkpoint field. The convolutional one is implemented; the higher quality diffusion one (`NADiffusionDecoder`) is @@ -1349,8 +1366,15 @@ knobs from `extras`. H3 takes `partition`. LTX-2.5 takes seam's `prompt_embeds_path`, which carries the video stream), `pipeline_kind` (default `distilled_two_stage`), `model_version` (only for a checkpoint that declares none), `dit_config_path`, `allow_unported_modules`, `max_phase`, -`upsampler_path` and `duration_head_path`. An extra a family does not define is -refused, never ignored. +`prompt_embeds_valid_rows`, `upsampler_path` and `duration_head_path`. An extra a +family does not define is refused, never ignored. + +`prompt_embeds_valid_rows` is how many of the supplied conditioning rows are real +tokens; absent, every row is. It matters because the embeddings connector +substitutes its learnable register table at PADDED positions, so padding decides +which of the connector's inputs are learned constants rather than caption +features. Upstream always knows this because its tokenizer produced the mask; +this seam reads conditioning from a file, which carries none. `dit_config_path` names a JSON file holding the DiT's `{"transformer": {...}}` configuration, and it exists because only one of the two shipped LTX-2.5 DiTs diff --git a/examples/ltx2_gen/main.cpp b/examples/ltx2_gen/main.cpp index 1da8931dc..863190d30 100644 --- a/examples/ltx2_gen/main.cpp +++ b/examples/ltx2_gen/main.cpp @@ -85,6 +85,7 @@ const char* Need(int argc, char** argv, int i, const char* flag) { " [--model-version 2.5] [--pipeline-kind distilled_two_stage]\n" " [--upsampler ] phase 2 needs it\n" " [--max-phase N] [--allow-unported]\n" + " [--prompt-valid-rows N] how many embed rows are real tokens\n" " [--frames N] [--width N] [--height N] [--seed N]\n" " [--device cpu|cuda]\n\n" "Renders LTX-2.5 (family \"ltx-2.5\") through vllm_video_engine_load +\n" @@ -92,7 +93,14 @@ const char* Need(int argc, char** argv, int i, const char* flag) { "little-endian f32, the video one 4096 wide and the audio one 2048, with the\n" "SAME row count. There is no --prompt: the Gemma-4 text tower is not ported and\n" "the engine refuses a prompt rather than silently rendering these embeddings\n" - "as if they were it.\n"); + "as if they were it.\n\n" + "Those rows are the EMBEDDINGS CONNECTOR's input, not the DiT's: when the\n" + "checkpoint carries the two *_embeddings_connector families (both shipped\n" + "LTX-2.5 DiTs do) they run through it first, with the checkpoint's own\n" + "weights. The row count must then be a multiple of the connector's learnable\n" + "register count (128 on the shipped files), and --prompt-valid-rows says how\n" + "many of them are real: the rest are padding, and padding is REPLACED by the\n" + "learnable register table rather than ignored.\n"); std::exit(code); } @@ -135,6 +143,8 @@ int main(int argc, char** argv) { else if (f == "--duration-head") SetExtra("duration_head_path", Need(argc, argv, ++i, f.c_str())); else if (f == "--max-phase") SetExtra("max_phase", Need(argc, argv, ++i, f.c_str())); + else if (f == "--prompt-valid-rows") + SetExtra("prompt_embeds_valid_rows", Need(argc, argv, ++i, f.c_str())); else if (f == "--allow-unported") SetExtra("allow_unported_modules", "1"); else if (f == "--device") device = Need(argc, argv, ++i, "--device"); else if (f == "--frames") vp.num_frames = std::atoi(Need(argc, argv, ++i, "--frames")); diff --git a/include/vllm/model_executor/models/ltx2_connector.h b/include/vllm/model_executor/models/ltx2_connector.h index ba86c901d..eb2a36193 100644 --- a/include/vllm/model_executor/models/ltx2_connector.h +++ b/include/vllm/model_executor/models/ltx2_connector.h @@ -129,4 +129,57 @@ Ltx2ConnectorOutput Ltx2ConnectorForward(const Ltx2ConnectorConfig& config, const float* additive_attention_mask, int64_t batch, int64_t seq); +// ─── THE PROCESSOR AROUND THE TWO CONNECTORS ───────────────────────────────── +// +// Upstream: `EmbeddingsProcessor.create_embeddings` +// (text_encoders/gemma/embeddings_processor.py:70-95). It is a SEPARATE module +// from the connector and it does three things the connector does not, each of +// which changes the conditioning silently when it is skipped: +// +// * IT RIGHT-PAD-SORTS THE FEATURES FIRST (:82-84, `_compute_right_pad_order` / +// `_apply_right_pad_order`). Upstream's own comment is "Connectors expect +// right-padded input ([valid, pad])", because the register table is indexed +// by ABSOLUTE position (`s % num_registers`) rather than by which positions +// were padded. A LEFT-padded batch handed straight to the connector puts +// real tokens where registers belong and registers where tokens belong, and +// the result is finite, correctly shaped and conditioned on nothing. +// * IT MULTIPLIES THE VIDEO ENCODING BY A BINARY MASK (:86-87) and does NOT do +// the same to the audio one (:91-93) — an asymmetry that reads like an +// oversight, is upstream's behaviour, and is mirrored rather than tidied. +// * IT RETURNS THE BINARY MASK THE DiT CONSUMES (:89), which is the connector's +// OUTPUT mask, not the caller's input one. That mask is derived with +// `encoded_mask < 0.000001` (:46-48), and BOTH values an additive mask can +// hold — 0.0 and -finfo(f32).max — satisfy it, so it is all ones for every +// input either reference can produce. See the implementation: the direction +// is surprising, `diffusers` writes the identical comparison, and the +// multiply is consequently an identity on every reachable path. +// +// THE TWO REFERENCES DISAGREE ABOUT WHERE THE SORT LIVES, and the disagreement is +// recorded rather than resolved by preference. `diffusers`' +// `LTX2ConnectorTransformer1d.forward` folds the sort INTO the connector +// (`src/diffusers/pipelines/ltx2/connectors.py`, the `torch.argsort(1 - +// binary_attn_mask, stable=True)` branch) and its comment claims that matches +// "the original LTX implementation" — which is true only because `ltx_core` +// sorts one level up, in the processor. The two compose to the same function; +// they differ in which module owns it. This port follows `ltx_core`, so +// `Ltx2ConnectorForward` stays a faithful port of `Embeddings1DConnector` and +// the sort lands here, at the processor. +struct Ltx2ConnectorEmbeddings { + std::vector video; // [batch, seq, video inner_dim] + std::vector audio; // [batch, seq, audio inner_dim] + // `binary_mask.squeeze(-1)` (:89): [batch, seq], 1.0 for a position the DiT's + // cross-attention may attend to and 0.0 for one it may not. With registers + // enabled every position is attendable, which is the whole point of them. + std::vector mask; +}; + +// `additive_attention_mask` is [batch, seq] with 0 for a kept token and +// -finfo(f32).max for a padded one, and it is required: it is what decides which +// positions become registers. +Ltx2ConnectorEmbeddings Ltx2ConnectorCreateEmbeddings( + const Ltx2ConnectorConfig& video_config, const Ltx2VaeWeights& video_weights, + const float* video_features, const Ltx2ConnectorConfig& audio_config, + const Ltx2VaeWeights& audio_weights, const float* audio_features, + const float* additive_attention_mask, int64_t batch, int64_t seq); + } // namespace vllm diff --git a/include/vllm/model_executor/models/ltx2_loader.h b/include/vllm/model_executor/models/ltx2_loader.h index 5ec3064d9..ff357d4bd 100644 --- a/include/vllm/model_executor/models/ltx2_loader.h +++ b/include/vllm/model_executor/models/ltx2_loader.h @@ -114,6 +114,7 @@ #include "vllm/model_executor/models/ltx2.h" #include "vllm/model_executor/models/ltx2_audio_vae.h" +#include "vllm/model_executor/models/ltx2_connector.h" #include "vllm/model_executor/models/ltx2_text_encoder.h" #include "vllm/model_executor/models/ltx2_upsampler.h" #include "vllm/model_executor/models/ltx2_video_vae.h" @@ -388,6 +389,69 @@ Ltx2DitParams Ltx2AdoptDeclaredDitParams(const nlohmann::json& config, bool allow_unported_modules, const std::string& source); +// --------------------------------------------------------------------------- +// The embeddings connector (phase L9c) +// --------------------------------------------------------------------------- +// +// The two `*_embeddings_connector` families live in the DiT FILE but not in the +// DiT's own weight contract: upstream loads them into the TEXT ENCODER's +// `EmbeddingsProcessor`, through `EMBEDDINGS_PROCESSOR_KEY_OPS`, which rewrites +// `model.diffusion_model.video_embeddings_connector.` to `video_connector.` +// (text_encoders/gemma/encoders/encoder_configurator.py:331-346). They are +// therefore loaded HERE, beside the DiT and not inside it, and they stay outside +// `EnumerateLtx2DitTensors`. + +enum class Ltx2ConnectorStream { kVideo, kAudio }; + +// The tensor-name prefix each stream's family carries in the checkpoint, WITHOUT +// the ComfyUI `model.diffusion_model.` prefix (`PlanDit` strips that). +const char* Ltx2ConnectorCheckpointPrefix(Ltx2ConnectorStream stream); + +// Does this DiT file carry a connector at all? Keyed on `learnable_registers`, +// the one tensor the family always has and the only one that is not per-block. +bool Ltx2CheckpointHasConnector(const SafetensorsFile& file, Ltx2ConnectorStream stream); + +// `Embeddings1DConnectorConfigurator.from_metadata` and its audio twin +// (embeddings_connector.py:194-256) applied to a `{"transformer": {...}}` object. +// +// FOUR VALUES NO SHAPE ENCODES, and one of them is not close to its default: +// LTX-2.5 declares `connector_positional_embedding_max_pos = [4096]` where the +// class default is `[1]`. `get_fractional_positions` divides the token index by +// it (rope.py:132-141), so the default turns every position into a fractional +// position 4096x too large and every RoPE angle with it. The others are +// `rope_type`, `frequencies_precision` (-> double-precision frequencies) and +// `connector_apply_gated_attention`. +// +// `positional_embedding_theta` is deliberately NOT read from the config even +// though the DiT declares one: neither configurator passes it, so upstream's +// connector always runs at the class default of 10000.0. Reading the DiT's key +// would be a re-invention, not a port — and the shipped file declares 10000.0 +// anyway, so the two agree and only the RULE differs. +Ltx2ConnectorConfig Ltx2ParseConnectorConfig(const nlohmann::json& config, + Ltx2ConnectorStream stream); + +// Materialize one connector family out of the DiT checkpoint, widened to f32 — +// which is `Ltx2ConnectorForward`'s declared parity dtype. +// +// IT IS NOT CHEAP AND THE CALLER MUST TREAT IT AS EXPENSIVE. 129 tensors is 8 +// blocks of four dim x dim projections plus a 4x-wide feed-forward, so at the +// shipped widths the video family is ~1.61G parameters and the audio family +// ~0.40G: **about 8 GB of f32 together**. That is small against the DiT's 21 GB +// on disk and 44 GB staged, and it is NOT small on a 119 GB unified-memory box +// that reboots rather than OOM-killing. The video engine therefore loads these, +// runs the connector once, and drops them inside one scope — the conditioning is +// resolved at load time, so the weights never outlive their single use. +// +// It is also the CONTRACT CHECK on the config. `Ltx2ParseConnectorConfig` reads +// values that mostly cannot be seen in a shape — but `connector_num_layers`, +// `connector_apply_gated_attention`, `connector_ff_bias` and the head geometry +// all CAN, so every enumerated tensor must exist at its enumerated shape AND no +// tensor of the family may be left over. A config that says 2 layers against a +// file carrying 8 is refused by name here rather than binding the first two and +// rendering. +Ltx2VaeWeights Ltx2LoadConnectorWeights(const SafetensorsFile& file, + const Ltx2ConnectorConfig& config); + // `__metadata__["model_version"]` ("2.5.0"), which is what // `detect_model_version` reads to pick a recipe (ltx-pipelines // utils/constants.py:161) and what `should_use_ancestral_sampler` keys on diff --git a/include/vllm/multimodal/ltx2_video.h b/include/vllm/multimodal/ltx2_video.h index 8019e8679..99ab43729 100644 --- a/include/vllm/multimodal/ltx2_video.h +++ b/include/vllm/multimodal/ltx2_video.h @@ -51,6 +51,15 @@ // request is refused BY NAME. Conditioning comes from prompt-embeds, which is // the seam's own documented fallback (video_engine.h:55-57). // +// WHAT PHASE L9c CHANGED, AND WHAT IT DID NOT. The link BELOW the tower is now +// real: when the checkpoint carries the two `*_embeddings_connector` families +// — both shipped LTX-2.5 DiTs do, 129 tensors each — the supplied prompt +// embeds are run through the connector before the DiT sees them, with those +// weights, under the checkpoint's own `connector_*` configuration. Before +// L9c they went to cross-attention verbatim and the connector was reachable +// only from a test. The TOWER is still owed, so what enters the connector is +// still whatever the caller put in the file rather than an encoded prompt. +// // 3. Any pipeline kind / model version the recipe table does not carry. // `ResolveLtx2PipelineRecipe` already throws rather than defaulting // (ltx2_pipeline.h:543-562); this engine passes the checkpoint's OWN @@ -141,6 +150,20 @@ inline constexpr char kLtx2AllowUnportedExtra[] = "allow_unported_modules"; // say so rather than skipping the phase silently. inline constexpr char kLtx2MaxPhaseExtra[] = "max_phase"; +// How many of the supplied prompt-embeds rows are REAL tokens; the rest are +// padding. Absent means every row is real. +// +// WHY A SEAM WITH NO TOKENIZER NEEDS THIS. The embeddings connector substitutes +// its `learnable_registers` table at PADDED positions +// (embeddings_connector.py:139-152), so the padding is not inert — it is what +// decides which of the connector's inputs are learned constants rather than +// caption features. Upstream always knows this, because the tokenizer produced +// the mask. This seam takes prompt embeds from a FILE, which carries no mask, so +// without this extra the padded tail would be conditioned on as if it were text +// and every register would go unused. Recorded as a knob rather than assumed, +// and it is the field the Gemma-4 tower will supply when it lands. +inline constexpr char kLtx2PromptValidRowsExtra[] = "prompt_embeds_valid_rows"; + // A loaded LTX-2.5 checkpoint set. Construct through // `vllm::multimodal::LoadVideoEngine` (detection) or by declaring // `family = kLtx2VideoFamily`; this type is exposed so a test can name it. diff --git a/src/vllm/model_executor/models/ltx2_connector.cpp b/src/vllm/model_executor/models/ltx2_connector.cpp index 21fd85211..cda141711 100644 --- a/src/vllm/model_executor/models/ltx2_connector.cpp +++ b/src/vllm/model_executor/models/ltx2_connector.cpp @@ -282,4 +282,105 @@ Ltx2ConnectorOutput Ltx2ConnectorForward(const Ltx2ConnectorConfig& config, return state; } +Ltx2ConnectorEmbeddings Ltx2ConnectorCreateEmbeddings( + const Ltx2ConnectorConfig& video_config, const Ltx2VaeWeights& video_weights, + const float* video_features, const Ltx2ConnectorConfig& audio_config, + const Ltx2VaeWeights& audio_weights, const float* audio_features, + const float* additive_attention_mask, int64_t batch, int64_t seq) { + Require(video_features != nullptr && audio_features != nullptr, + "ltx2 connector: both modality feature streams are required " + "(embeddings_processor.py:76-79 refuses one without the other)"); + Require(additive_attention_mask != nullptr, + "ltx2 connector: the processor requires the additive attention mask; it is what " + "decides which positions become learnable registers " + "(embeddings_processor.py:82)"); + const int64_t vdim = video_config.inner_dim(); + const int64_t adim = audio_config.inner_dim(); + + // `_compute_right_pad_order` (:23-38): a STABLE descending argsort of the + // binary mask, so valid tokens move to the front keeping their relative order + // and padded ones follow. Written as a stable partition because that is what a + // stable argsort of a 0/1 key IS, and it needs no comparator. + std::vector order(static_cast(batch * seq)); + std::vector reordered_mask(static_cast(batch * seq)); + for (int64_t b = 0; b < batch; ++b) { + int64_t write = 0; + for (int64_t s = 0; s < seq; ++s) { + if (additive_attention_mask[b * seq + s] >= 0.0f) order[static_cast(b * seq + write++)] = s; + } + const int64_t valid = write; + for (int64_t s = 0; s < seq; ++s) { + if (additive_attention_mask[b * seq + s] < 0.0f) order[static_cast(b * seq + write++)] = s; + } + // `new_additive = (new_binary - 1) * finfo.max` (:37): 0 for the valid + // prefix, -finfo.max for the padded tail. + for (int64_t s = 0; s < seq; ++s) { + reordered_mask[static_cast(b * seq + s)] = + s < valid ? 0.0f : -std::numeric_limits::max(); + } + } + + // `_apply_right_pad_order` (:41-43): gather the feature rows into that order. + auto gather = [&](const float* src, int64_t width) { + std::vector out(static_cast(batch * seq * width)); + for (int64_t b = 0; b < batch; ++b) { + for (int64_t s = 0; s < seq; ++s) { + const int64_t from = order[static_cast(b * seq + s)]; + std::memcpy(out.data() + static_cast((b * seq + s) * width), + src + static_cast((b * seq + from) * width), + static_cast(width) * sizeof(float)); + } + } + return out; + }; + const std::vector video_sorted = gather(video_features, vdim); + const std::vector audio_sorted = gather(audio_features, adim); + + const Ltx2ConnectorOutput video_out = Ltx2ConnectorForward( + video_config, video_weights, video_sorted.data(), reordered_mask.data(), batch, seq); + const Ltx2ConnectorOutput audio_out = Ltx2ConnectorForward( + audio_config, audio_weights, audio_sorted.data(), reordered_mask.data(), batch, seq); + + Ltx2ConnectorEmbeddings out; + out.mask.resize(static_cast(batch * seq)); + // `_to_binary_mask` (:46-48): `encoded_mask < 0.000001`. + // + // THE COMPARISON DIRECTION IS SURPRISING AND IT IS MIRRORED EXACTLY. An + // additive mask holds 0.0 for a kept position and -finfo(f32).max for a padded + // one, and BOTH are `< 0.000001` — so this returns 1 at every position, for + // every mask upstream can produce. The reading a port would arrive at by + // reasoning about intent (`>= 0`, "keep the unmasked ones") is the OPPOSITE at + // padded positions, and it is not what either reference does. + // + // Checked against both, because a line this surprising is exactly where one + // implementation being wrong would show: `diffusers` + // `LTX2TextConnectors.forward` writes `(video_attn_mask < 1e-6).to(int64)` and + // then the same video-only multiply. They agree, down to the constant. + // + // The consequence is that the multiply below is a NO-OP for every mask value + // that can reach it, and the mask this returns is all ones. That is not a + // reason to drop either: with registers enabled the connector zeroes the mask + // itself, so the shipped configuration reaches this line with all-zeros and + // the identity is CORRECT rather than incidental — and removing upstream's + // line would silently change a future connector whose output mask is not one + // of those two values. + for (int64_t i = 0; i < batch * seq; ++i) { + out.mask[static_cast(i)] = video_out.mask[static_cast(i)] < 0.000001f ? 1.0f : 0.0f; + } + // :86-87 — the VIDEO encoding is multiplied by that binary mask. The AUDIO + // encoding is not (:91-93). Mirrored, not tidied. + out.video = video_out.hidden_states; + for (int64_t b = 0; b < batch; ++b) { + for (int64_t s = 0; s < seq; ++s) { + const float m = out.mask[static_cast(b * seq + s)]; + if (m != 0.0f) continue; + for (int64_t i = 0; i < vdim; ++i) { + out.video[static_cast((b * seq + s) * vdim + i)] = 0.0f; + } + } + } + out.audio = audio_out.hidden_states; + return out; +} + } // namespace vllm diff --git a/src/vllm/model_executor/models/ltx2_loader.cpp b/src/vllm/model_executor/models/ltx2_loader.cpp index cc44f2a9c..93e58df63 100644 --- a/src/vllm/model_executor/models/ltx2_loader.cpp +++ b/src/vllm/model_executor/models/ltx2_loader.cpp @@ -411,7 +411,22 @@ vt::DType MaterializeDitTensor(const SafetensorsFile& file, const DitPlan& plan, Fail("'" + spec.name + "' has dtype " + t.dtype + ", which this loader does not read"); } -// The families the file carries and the L2 contract does not. +// The families the file carries and NOTHING IN THIS PORT reads. +// +// The two `*_embeddings_connector` families are outside the DiT contract and are +// NOT unported: phase L9c materializes them through `Ltx2LoadConnectorWeights` +// and the video engine runs them, which is where upstream puts them too — its +// key ops rewrite `model.diffusion_model.video_embeddings_connector.` into the +// TEXT ENCODER's `EmbeddingsProcessor` (encoder_configurator.py:331-346). Naming +// them here would make the refusal say something untrue about the tree, and +// would demand `allow_unported_modules` from a caller whose checkpoint this port +// reads completely. Their contract is checked where it can be: against a parsed +// connector configuration, in `Ltx2LoadConnectorWeights`, which refuses a +// missing tensor, a wrong shape, and a leftover tensor by name. +bool LoadedElsewhere(const std::string& family) { + return family == "video_embeddings_connector" || family == "audio_embeddings_connector"; +} + std::vector UnportedFamilies(const DitPlan& plan, const std::vector& contract) { std::set known; @@ -421,6 +436,7 @@ std::vector UnportedFamilies(const DitPlan& plan, for (const Ltx2TensorSpec& spec : plan.manifest) { if (known.count(spec.name) != 0) continue; const std::string family = FamilyOf(spec.name); + if (LoadedElsewhere(family)) continue; if (seen.insert(family).second) families.push_back(family); } return families; @@ -441,9 +457,10 @@ std::vector ContractOf(const Ltx2DitParams& params) { ". They are not dropped silently: prompt_adaln_single / " "audio_prompt_adaln_single mean use_prompt_adaln_single is TRUE, which " "contradicts .agents/specs/ltx-2-5.md section 1.2 and voids the prompt-K/V " - "cache's premise; keyframes_abs_pos_embedding contradicts ltx2.h:47-49; the " - "two *_embeddings_connector families are the Embeddings1DConnector " - "ltx2_text_encoder.h:319-324 already records as owed. Pass " + "cache's premise; keyframes_abs_pos_embedding contradicts ltx2.h:47-49. The " + "two *_embeddings_connector families are NOT in this list and never will be " + "— they are outside the DiT contract by design and are loaded by " + "Ltx2LoadConnectorWeights, which is what the video engine calls. Pass " "Ltx2DitLoadOptions::allow_unported_modules to load the ported SUBSET, which " "still reports every one of them."); } @@ -910,6 +927,208 @@ nlohmann::json Ltx2ReadCheckpointConfig(const SafetensorsFile& file) { return parsed; } +// --------------------------------------------------------------------------- +// The embeddings connector (phase L9c) — see ltx2_loader.h for why it loads here +// --------------------------------------------------------------------------- + +const char* Ltx2ConnectorCheckpointPrefix(Ltx2ConnectorStream stream) { + return stream == Ltx2ConnectorStream::kVideo ? "video_embeddings_connector." + : "audio_embeddings_connector."; +} + +bool Ltx2CheckpointHasConnector(const SafetensorsFile& file, Ltx2ConnectorStream stream) { + const std::string bare = + std::string(Ltx2ConnectorCheckpointPrefix(stream)) + "learnable_registers"; + const std::string prefixed = std::string(kLtx2DitCheckpointPrefix) + bare; + for (const std::string& n : file.Names()) { + if (n == bare || n == prefixed) return true; + } + return false; +} + +namespace { + +// `transformer_config.get(key, fallback)` for the four scalar shapes the two +// connector configurators read. Kept local because a config value that is +// present but of the WRONG JSON type must refuse rather than fall back — a +// silent fallback is exactly how a declared 8-layer connector becomes a 2-layer +// one. +int64_t ConnectorInt(const nlohmann::json& t, const std::string& key, int64_t fallback) { + const auto it = t.find(key); + if (it == t.end() || it->is_null()) return fallback; + if (!it->is_number_integer() && !it->is_number_unsigned()) { + Fail("the connector config key '" + key + "' is " + it->dump() + ", not an integer"); + } + return it->get(); +} + +bool ConnectorBool(const nlohmann::json& t, const std::string& key, bool fallback) { + const auto it = t.find(key); + if (it == t.end() || it->is_null()) return fallback; + if (!it->is_boolean()) { + Fail("the connector config key '" + key + "' is " + it->dump() + ", not a boolean"); + } + return it->get(); +} + +} // namespace + +Ltx2ConnectorConfig Ltx2ParseConnectorConfig(const nlohmann::json& config, + Ltx2ConnectorStream stream) { + const auto tit = config.find("transformer"); + if (tit == config.end() || !tit->is_object()) { + Fail("the connector configuration needs a `{\"transformer\": {...}}` object; the two " + "configurators read the DiT's own transformer config " + "(embeddings_connector.py:196, :227)"); + } + const nlohmann::json& t = *tit; + + Ltx2ConnectorConfig out; + out.prefix = Ltx2ConnectorCheckpointPrefix(stream); + + // :198, :226 — `LTXRopeType(transformer_config.get("rope_type", "split"))`. + const auto rope = t.find("rope_type"); + std::string rope_type = "split"; + if (rope != t.end() && !rope->is_null()) { + if (!rope->is_string()) Fail("the connector config key 'rope_type' is not a string"); + rope_type = rope->get(); + } + if (rope_type == "split") { + out.rope_type = Ltx2RopeType::kSplit; + } else if (rope_type == "interleaved") { + out.rope_type = Ltx2RopeType::kInterleaved; + } else { + Fail("the connector config declares rope_type '" + rope_type + + "', which is neither 'split' nor 'interleaved' (rope.py:11-13)"); + } + + // :199, :228 — `transformer_config.get("frequencies_precision", False) == "float64"`. + // Absent or false means SINGLE precision, so a non-string value is not an + // error upstream; it simply is not "float64". + const auto freq = t.find("frequencies_precision"); + out.double_precision_rope = + freq != t.end() && freq->is_string() && freq->get() == "float64"; + + // :200, :229 — `connector_positional_embedding_max_pos`, default [1]. The one + // value the shipped config moves off its default, and it scales every position. + const auto max_pos = t.find("connector_positional_embedding_max_pos"); + if (max_pos == t.end() || max_pos->is_null()) { + out.positional_embedding_max_pos = {1}; + } else { + if (!max_pos->is_array()) { + Fail("the connector config key 'connector_positional_embedding_max_pos' is " + + max_pos->dump() + ", not an array"); + } + out.positional_embedding_max_pos.clear(); + for (const nlohmann::json& e : *max_pos) { + if (!e.is_number()) { + Fail("'connector_positional_embedding_max_pos' holds a non-numeric entry " + e.dump()); + } + out.positional_embedding_max_pos.push_back(e.get()); + } + if (out.positional_embedding_max_pos.size() != 1) { + // `get_fractional_positions` asserts the grid's position-dimension count + // equals `len(max_pos)` (rope.py:130-134), and the connector's grid is 1-D. + Fail("'connector_positional_embedding_max_pos' has " + + std::to_string(out.positional_embedding_max_pos.size()) + + " entries; the connector's position grid is 1-D (embeddings_connector.py:172-174) " + "so exactly one is required"); + } + } + + // :203-206 for video; :232-241 for audio, which falls back to the VIDEO keys + // "for backwards compatibility" rather than to the class defaults. + if (stream == Ltx2ConnectorStream::kVideo) { + out.num_attention_heads = ConnectorInt(t, "connector_num_attention_heads", 30); + out.attention_head_dim = ConnectorInt(t, "connector_attention_head_dim", 128); + out.num_layers = ConnectorInt(t, "connector_num_layers", 2); + } else { + out.num_attention_heads = ConnectorInt( + t, "audio_connector_num_attention_heads", ConnectorInt(t, "connector_num_attention_heads", 30)); + out.attention_head_dim = ConnectorInt( + t, "audio_connector_attention_head_dim", ConnectorInt(t, "connector_attention_head_dim", 128)); + out.num_layers = + ConnectorInt(t, "audio_connector_num_layers", ConnectorInt(t, "connector_num_layers", 2)); + } + // :212-213, :253-254 — BOTH configurators read the VIDEO spellings of these two. + out.apply_gated_attention = ConnectorBool(t, "connector_apply_gated_attention", false); + out.ff_bias = ConnectorBool(t, "connector_ff_bias", true); + + // NOT read by either configurator: `Embeddings1DConnector`'s own default of 128 + // is what upstream always runs (:102, :131). The shipped config declares + // `connector_num_learnable_registers` and declares 128, so the two agree; it is + // read here so a checkpoint that declares something else is not silently run at + // 128, and `Ltx2LoadConnectorWeights` checks it against the stored table. + out.num_learnable_registers = ConnectorInt(t, "connector_num_learnable_registers", 128); + + if (out.num_attention_heads < 1 || out.attention_head_dim < 1 || out.num_layers < 1) { + Fail("the connector config resolves " + std::to_string(out.num_layers) + " layers of " + + std::to_string(out.num_attention_heads) + " x " + + std::to_string(out.attention_head_dim) + " heads, which is not a module"); + } + return out; +} + +Ltx2VaeWeights Ltx2LoadConnectorWeights(const SafetensorsFile& file, + const Ltx2ConnectorConfig& config) { + const DitPlan plan = PlanDit(file); + const std::vector specs = EnumerateLtx2ConnectorTensors(config); + + Ltx2VaeWeights out; + for (const Ltx2ConnectorTensorSpec& spec : specs) { + const auto shape_it = plan.logical.find(spec.name); + if (shape_it == plan.logical.end()) { + Fail("the checkpoint is missing '" + spec.name + "' " + ShapeText(spec.shape) + + ", which the connector configuration resolved from " + + std::to_string(config.num_layers) + " layers of " + + std::to_string(config.num_attention_heads) + " x " + + std::to_string(config.attention_head_dim) + + " heads. The config and the file describe different modules; refusing rather " + "than binding the subset that happens to be present."); + } + if (shape_it->second != spec.shape) { + Fail("'" + spec.name + "' is " + ShapeText(shape_it->second) + + " in the checkpoint but the connector configuration needs " + ShapeText(spec.shape)); + } + std::vector bytes; + const vt::DType dtype = MaterializeDitTensor(file, plan, {spec.name, spec.shape}, bytes); + int64_t numel = 1; + for (const int64_t d : spec.shape) numel *= d; + std::vector widened(static_cast(numel)); + if (dtype == vt::DType::kBF16) { + const auto* src = reinterpret_cast(bytes.data()); + for (int64_t i = 0; i < numel; ++i) widened[static_cast(i)] = Bf16ToF32(src[static_cast(i)]); + } else if (dtype == vt::DType::kF32) { + std::memcpy(widened.data(), bytes.data(), static_cast(numel) * sizeof(float)); + } else { + Fail("'" + spec.name + "' materialized as a dtype the connector bag cannot hold"); + } + out.tensors[spec.name] = std::move(widened); + } + + // The other half of the contract check: a tensor of this family that the + // configuration did NOT enumerate. Without it a config declaring FEWER layers + // than the file carries binds a valid prefix and runs — the exact shape of the + // defect the shape check above cannot see. + int64_t extra = 0; + std::string first_extra; + for (const auto& kv : plan.logical) { + if (!StartsWith(kv.first, config.prefix)) continue; + if (out.tensors.count(kv.first) != 0) continue; + if (extra == 0) first_extra = kv.first; + ++extra; + } + if (extra != 0) { + Fail("the checkpoint carries " + std::to_string(extra) + " '" + config.prefix + + "' tensors the connector configuration does not name (first: '" + first_extra + + "'). The configuration resolved " + std::to_string(config.num_layers) + + " layers, gated_attention=" + (config.apply_gated_attention ? "true" : "false") + + ", ff_bias=" + (config.ff_bias ? "true" : "false") + + ". Refusing rather than binding the prefix that matches and leaving the rest."); + } + return out; +} + std::string Ltx2ReadCheckpointModelVersion(const SafetensorsFile& file) { const std::map& meta = file.Metadata(); const auto it = meta.find("model_version"); diff --git a/src/vllm/multimodal/ltx2_video.cpp b/src/vllm/multimodal/ltx2_video.cpp index 743c19e9d..ac345a780 100644 --- a/src/vllm/multimodal/ltx2_video.cpp +++ b/src/vllm/multimodal/ltx2_video.cpp @@ -8,9 +8,11 @@ #include #include #include +#include #include #include #include +#include #include #include #include @@ -21,8 +23,10 @@ #include #include "vllm/model_executor/model_loader/safetensors_reader.h" +#include "vllm/model_executor/models/device_pool.h" // ActivePool()/DevicePool::Drain #include "vllm/model_executor/models/ltx2.h" #include "vllm/model_executor/models/ltx2_audio_vae.h" +#include "vllm/model_executor/models/ltx2_connector.h" #include "vllm/model_executor/models/ltx2_device.h" #include "vllm/model_executor/models/ltx2_loader.h" #include "vllm/model_executor/models/ltx2_pipeline.h" @@ -245,9 +249,9 @@ int64_t ExtraInt(const std::map& extras, const std::st // (minimax_h3_video.cpp): a mistyped knob that is silently dropped renders the // DEFAULT and looks like the feature not working. const char* const kKnownLoadExtras[] = { - kLtx2AudioPromptEmbedsExtra, kLtx2PipelineKindExtra, kLtx2ModelVersionExtra, - kLtx2AllowUnportedExtra, kLtx2MaxPhaseExtra, kLtx2DitConfigPathExtra, - "upsampler_path", "duration_head_path", + kLtx2AudioPromptEmbedsExtra, kLtx2PipelineKindExtra, kLtx2ModelVersionExtra, + kLtx2AllowUnportedExtra, kLtx2MaxPhaseExtra, kLtx2DitConfigPathExtra, + kLtx2PromptValidRowsExtra, "upsampler_path", "duration_head_path", }; void CheckKnownExtras(const std::map& extras) { @@ -359,9 +363,21 @@ struct Ltx2VideoEngine::Impl { // Conditioning. `has_encoder` is false for every checkpoint today; see the // header's refusal 2. bool has_encoder = false; + // What the DiT's cross-attention consumes. When the checkpoint carries the + // embeddings connector these are the CONNECTOR's output; without one they are + // the supplied prompt embeds verbatim, which is what every render before L9c + // did on every checkpoint. std::vector video_prompt_embeds, audio_prompt_embeds; int64_t prompt_tokens = 0; + // The connector's CONFIGURATION is kept; its WEIGHTS are not. They are ~8 GB + // of f32 at the shipped widths (ltx2_loader.h), the conditioning they process + // is resolved once at load, and this box reboots rather than OOM-killing — so + // they are loaded, used and dropped inside one scope below. + bool has_connector = false; + Ltx2ConnectorConfig video_connector_cfg, audio_connector_cfg; + int64_t prompt_valid_rows = 0; + std::mutex mutex; }; @@ -515,10 +531,14 @@ std::unique_ptr Ltx2VideoEngine::Load(const VideoModelParams& p "DiT is in: it carries no __metadata__ at all. Supply the config rather than " "inheriting a default that contradicts the model family."); } + // Hoisted out of the block it used to live in because the EMBEDDINGS CONNECTOR + // is configured from the SAME object: its geometry, its RoPE bounds and its + // gating are `connector_*` keys of the DiT's transformer config, and reading + // them from a second source would let the two disagree. + const nlohmann::json dit_config = + declares_config ? Ltx2ReadCheckpointConfig(dit_file) + : ReadJsonFile(kLtx2DitConfigPathExtra, config_path); { - const nlohmann::json config = - declares_config ? Ltx2ReadCheckpointConfig(dit_file) - : ReadJsonFile(kLtx2DitConfigPathExtra, config_path); const std::string source = declares_config ? std::string("the DiT checkpoint's own __metadata__[\"config\"][\"transformer\"]") @@ -527,10 +547,63 @@ std::unique_ptr Ltx2VideoEngine::Load(const VideoModelParams& p // rule lives, because the device gate drives `Ltx2StreamDitToDevice` without // this engine and owes the same check; two copies would be two rules. const Ltx2DitParams declared = Ltx2AdoptDeclaredDitParams( - config, im.dit.params, dit_options.allow_unported_modules, source); + dit_config, im.dit.params, dit_options.allow_unported_modules, source); im.dit.params = declared; } + // ── the embeddings connector (phase L9c) ────────────────────────────────── + // + // WHAT THIS CLOSES. Until L9c the conditioning this engine handed the DiT's + // cross-attention was the prompt-embeds file VERBATIM. Upstream never does + // that: `EmbeddingsProcessor` runs an 8-layer 1-D transformer over the caption + // projections before the DiT sees them (embeddings_processor.py:70-95), and + // that module SHIPS IN THE DiT FILE — 129 tensors per stream, which this + // loader used to refuse as "unported" and then step over. The bricks existed + // (`Ltx2ConnectorForward`, landed at L5, gated against upstream on five arms); + // nothing called them outside a test. This is the call. + // + // PRESENCE DECIDES, not a config flag. The connector is applied exactly when + // the checkpoint carries it, which is the same rule `DetectLtx2Video` uses on + // the DiT itself: which modules a file HOLDS is a fact, and a flag saying + // otherwise would be a second, disagreeable authority. A checkpoint carrying + // one stream's connector and not the other is refused, because rendering the + // video stream through eight transformer layers and the audio stream through + // none conditions the two halves of one clip on two different things. + const bool has_video_connector = + Ltx2CheckpointHasConnector(dit_file, Ltx2ConnectorStream::kVideo); + const bool has_audio_connector = + Ltx2CheckpointHasConnector(dit_file, Ltx2ConnectorStream::kAudio); + if (has_video_connector != has_audio_connector) { + Fail(std::string("the DiT checkpoint carries the ") + + (has_video_connector ? "video" : "audio") + + " embeddings connector but not the other stream's. Both conditioning paths run " + "through one `EmbeddingsProcessor` upstream (embeddings_processor.py:60-68 refuses " + "an audio connector without audio features and vice versa), so half a connector " + "would condition one modality on eight transformer layers and the other on none."); + } + im.has_connector = has_video_connector; + if (im.has_connector) { + im.video_connector_cfg = Ltx2ParseConnectorConfig(dit_config, Ltx2ConnectorStream::kVideo); + im.audio_connector_cfg = Ltx2ParseConnectorConfig(dit_config, Ltx2ConnectorStream::kAudio); + // The connector is dimension-PRESERVING: it consumes the caption projection's + // output and hands the DiT's cross-attention the same width. Asserted rather + // than assumed — a connector whose inner_dim disagrees with the stream it + // feeds would still forward, and the mismatch would surface as a wrong-shaped + // GEMM deep inside the DiT rather than as a load refusal. + if (im.video_connector_cfg.inner_dim() != im.dit.params.cross_attention_dim) { + Fail("the video embeddings connector is " + + std::to_string(im.video_connector_cfg.inner_dim()) + + " wide but the DiT's cross-attention takes " + + std::to_string(im.dit.params.cross_attention_dim)); + } + if (im.audio_connector_cfg.inner_dim() != im.dit.params.audio_cross_attention_dim) { + Fail("the audio embeddings connector is " + + std::to_string(im.audio_connector_cfg.inner_dim()) + + " wide but the DiT's audio cross-attention takes " + + std::to_string(im.dit.params.audio_cross_attention_dim)); + } + } + // ── the recipe ──────────────────────────────────────────────────────────── const std::string declared_version = Ltx2ReadCheckpointModelVersion(dit_file); const std::string override_version = VideoExtra(params.extras, kLtx2ModelVersionExtra); @@ -644,6 +717,74 @@ std::unique_ptr Ltx2VideoEngine::Load(const VideoModelParams& p } if (v_rows == 0) Fail("the prompt embeds are empty"); im.prompt_tokens = v_rows; + + // How much of the supplied conditioning is REAL. Everything by default; the + // extra is what a caller with a tokenizer's mask supplies. + im.prompt_valid_rows = ExtraInt(params.extras, kLtx2PromptValidRowsExtra, v_rows); + if (im.prompt_valid_rows < 0 || im.prompt_valid_rows > v_rows) { + Fail("the '" + std::string(kLtx2PromptValidRowsExtra) + "' extra is " + + std::to_string(im.prompt_valid_rows) + " but the prompt embeds hold " + + std::to_string(v_rows) + " rows"); + } + if (im.has_connector) { + const int64_t registers = im.video_connector_cfg.num_learnable_registers; + if (registers > 0 && v_rows % registers != 0) { + // Upstream asserts this (embeddings_connector.py:144) because the + // register table is TILED across the sequence rather than indexed by + // which positions were padded. + Fail("the prompt embeds hold " + std::to_string(v_rows) + + " rows, which is not a multiple of the connector's " + + std::to_string(registers) + + " learnable registers. The register table is tiled across the sequence " + "(embeddings_connector.py:144), so upstream asserts the same thing."); + } + // The additive mask `_prepare_attention_mask` would produce: 0.0 for a real + // token, -finfo(f32).max for padding. + std::vector additive(static_cast(v_rows), 0.0f); + for (int64_t s = im.prompt_valid_rows; s < v_rows; ++s) { + additive[static_cast(s)] = -std::numeric_limits::max(); + } + // The weights live and die inside this scope. See the Impl comment: 129 + // tensors per stream is ~8 GB of f32 at the shipped widths, and holding + // them for the engine's lifetime would spend that on a module that runs + // exactly once. The DiT file is still open if a per-request conditioning + // path ever needs to re-read them. + // + // THE f32 IS AN ANNOTATED ESCAPE, not an inherited default. Upstream runs + // this module at the model dtype, so f32 here is WIDER — the polarity + // AGENTS.md says a value gate cannot catch. It is taken because + // `Ltx2ConnectorForward` is L5's declared PARITY dtype and this is the + // arm its goldens cover, and it costs nothing on the production path: the + // connector runs ONCE per load over 128 rows, and its output is narrowed + // to the stream dtype on the first upload like every other activation. + // What it does cost is the transient above, which is why it is scoped. + const Ltx2VaeWeights video_connector = + Ltx2LoadConnectorWeights(dit_file, im.video_connector_cfg); + const Ltx2VaeWeights audio_connector = + Ltx2LoadConnectorWeights(dit_file, im.audio_connector_cfg); + const Ltx2ConnectorEmbeddings encoded = Ltx2ConnectorCreateEmbeddings( + im.video_connector_cfg, video_connector, im.video_prompt_embeds.data(), + im.audio_connector_cfg, audio_connector, im.audio_prompt_embeds.data(), + additive.data(), /*batch=*/1, v_rows); + // The processor returns the mask the DiT's cross-attention is supposed to + // honour (embeddings_processor.py:89). `Ltx2ModalityInput` carries no + // context mask, so a mask with a masked position would be silently dropped + // — the DiT would attend over register-free padding as if it were caption. + // With registers enabled every position is attendable and the mask is all + // ones, which is the case the shipped checkpoint is in; anything else is + // refused by name rather than ignored. + for (const float m : encoded.mask) { + if (m == 1.0f) continue; + Fail( + "the embeddings connector returned a cross-attention mask with masked " + "positions, and `Ltx2ModalityInput` carries no context mask to pass it through. " + "That happens when the connector runs with num_learnable_registers = 0, where " + "padding stays padding instead of becoming a register. Refusing rather than " + "dropping the mask, which would condition the DiT on unmasked padding."); + } + im.video_prompt_embeds = encoded.video; + im.audio_prompt_embeds = encoded.audio; + } } return engine; } @@ -970,6 +1111,40 @@ VideoResult Ltx2VideoEngine::Generate(const VideoGenParams& gen) { audio_lc = ashape.channels; audio_lf = ashape.frames; audio_lm = ashape.mel_bins; + + // PHASE CHANGE. The denoise loop just left the shared scratch pool holding + // every activation size class this phase touched, and on an UNCAPPED pool + // (`device_pool_cap_bytes == 0`, which is GB10 and Thor today) those blocks + // are never returned to the driver. What comes next allocates DIFFERENT + // classes — the next phase runs at twice the resolution, and after the last + // phase the VAE decode runs on the host — so none of them can be reused and + // all of them are headroom the next stage does not get. + // + // On GB10 that headroom is HOST memory: the unified pool is one ~119 GiB + // arena, and this class of box REBOOTS rather than OOM-killing when the + // driver runs out (`NVRM ... Out of memory [NV_ERR_NO_MEMORY]`), which it did + // twice during phase L9b's 320x192/25-frame attempt. MiniMax-H3 drains at + // exactly this boundary and for exactly this reason + // (minimax_h3_pipeline.cpp:546-563); the LTX path drained nowhere. + // + // Draining costs one free per retained block, once per phase. + if (im.on_device) { + // A MEASUREMENT LANE, in the same shape `VT_POOL_BYPASS` and + // `VT_POOL_EXACT` already take (device_pool.h): it exists so the A/B that + // establishes what the drain is worth runs on ONE binary, which is what + // AGENTS.md asks of a measurement. It is never a configuration — the drain + // is on by default and there is no supported reason to turn it off. + const char* off = std::getenv("VLLM_LTX2_POOL_DRAIN"); + if (off == nullptr || off[0] != '0') { + vt::Backend& backend = vt::GetBackend(im.device.type); + const size_t drained = ActivePool()->Drain(backend); + if (std::getenv("VT_POOL_STATS") != nullptr) { + std::fprintf(stderr, "[ltx2] phase '%s' drained %.2f GiB of denoise scratch\n", + phase.name.c_str(), + static_cast(drained) / (1024.0 * 1024.0 * 1024.0)); + } + } + } } // ── decode (distilled.py:314-315) ───────────────────────────────────────── diff --git a/tests/vllm/models/test_ltx2_loader.cpp b/tests/vllm/models/test_ltx2_loader.cpp index 7cbce1c32..509a203d1 100644 --- a/tests/vllm/models/test_ltx2_loader.cpp +++ b/tests/vllm/models/test_ltx2_loader.cpp @@ -1149,24 +1149,51 @@ TEST_CASE("ltx2 loader: the unported families are refused by name, not absorbed" } const std::string what_msg = "what: " + what; INFO(what_msg); - // All FIVE families the shipped DiT manifest names, not the three that were - // convenient: `audio_prompt_adaln_single` and `audio_embeddings_connector` go - // through the identical generic FamilyOf path, and covering them literally is - // one line each. + // THREE families now, not five. `audio_prompt_adaln_single` is here because it + // goes through the identical generic FamilyOf path as its video twin and + // covering it literally is one line. CHECK(what.find("prompt_adaln_single") != std::string::npos); CHECK(what.find("audio_prompt_adaln_single") != std::string::npos); CHECK(what.find("keyframes_abs_pos_embedding") != std::string::npos); - CHECK(what.find("video_embeddings_connector") != std::string::npos); - CHECK(what.find("audio_embeddings_connector") != std::string::npos); - // The opt-in still REPORTS every one of them; it does not make them vanish. + // THE TWO CONNECTOR FAMILIES ARE NOT UNPORTED AS OF PHASE L9c. They are + // outside the DiT's contract by design — upstream loads them into the text + // encoder's EmbeddingsProcessor (encoder_configurator.py:331-346) and so does + // this port, through `Ltx2LoadConnectorWeights`, which the video engine calls. + // Listing them here would say something untrue about the tree AND would demand + // `allow_unported_modules` from a caller whose checkpoint this port reads + // completely. Asserted as an ABSENCE, so restoring the old behaviour REDs. + CHECK(what.find("video_embeddings_connector") == std::string::npos); + CHECK(what.find("audio_embeddings_connector") == std::string::npos); + + // The opt-in still REPORTS every one of the three; it does not make them vanish. Ltx2DitLoadOptions options; options.allow_unported_modules = true; const vllm::Ltx2DitCheckpoint ck = vllm::Ltx2LoadDitFromSafetensors(file, options); - CHECK(ck.unported.size() == 5); + CHECK(ck.unported.size() == 3); + for (const std::string& family : ck.unported) { + CHECK(family != "video_embeddings_connector"); + CHECK(family != "audio_embeddings_connector"); + } CHECK(ck.checkpoint_params.use_prompt_adaln_single); CHECK_FALSE(ck.params.use_prompt_adaln_single); std::remove(path.c_str()); + + // AND THE CONNECTOR-ONLY CHECKPOINT LOADS WITH NO OPT-IN AT ALL, which is the + // half of the change a message assertion cannot see: a file whose only + // out-of-contract modules are the two connectors is now fully read by this + // port, so requiring the flag would be requiring an admission of a gap that + // has been closed. + const SyntheticDit conn_only = BuildSyntheticDit( + p, Ltx2DitQuant::kFp8, + {"video_embeddings_connector.learnable_registers", + "audio_embeddings_connector.learnable_registers"}); + const std::string conn_path = TmpPath("connector_only"); + WriteSafetensors(conn_only.entries, conn_path); + const SafetensorsFile conn_file = SafetensorsFile::Open(conn_path); + const vllm::Ltx2DitCheckpoint conn_ck = vllm::Ltx2LoadDitFromSafetensors(conn_file); + CHECK(conn_ck.unported.empty()); + std::remove(conn_path.c_str()); } TEST_CASE("ltx2 loader: the f32 widening is OPT-IN and bit-exact over bf16") { diff --git a/tests/vllm/models/test_ltx2_pipeline.cpp b/tests/vllm/models/test_ltx2_pipeline.cpp index 9baba554c..7253c482f 100644 --- a/tests/vllm/models/test_ltx2_pipeline.cpp +++ b/tests/vllm/models/test_ltx2_pipeline.cpp @@ -1727,3 +1727,193 @@ TEST_CASE("ltx2 the connector's learnable registers are stored BFLOAT16, and rou INFO("refusal = ", message); CHECK(Mentions(message, "learnable_registers")); } + +// =========================================================================== +// Section 10a — the PROCESSOR around the two connectors (phase L9c) +// =========================================================================== +// +// `EmbeddingsProcessor.create_embeddings` (embeddings_processor.py:70-95) is a +// separate module from the connector, and the two things it does that the +// connector does not are both invisible to a value golden of the connector +// itself. They are gated here by their upstream-stated PROPERTIES rather than by +// a new golden, because a golden generated through the same helper the port uses +// proves the two agree and not that either is right — this project has recorded +// that failure once already. + +namespace { + +// Two independently seeded bags at the two stream widths the shipped checkpoint +// uses, so the video and audio connectors cannot accidentally be the same module. +vllm::Ltx2ConnectorConfig ProcessorConfig(const std::string& prefix, int64_t heads, + int64_t head_dim, int64_t registers) { + vllm::Ltx2ConnectorConfig config; + config.attention_head_dim = head_dim; + config.num_attention_heads = heads; + config.num_layers = vllm_test::kLtx2ConnLayers; + config.positional_embedding_theta = vllm_test::kLtx2ConnTheta; + config.positional_embedding_max_pos = {4096}; // the shipped connector bound + config.num_learnable_registers = registers; + config.rope_type = vllm::Ltx2RopeType::kSplit; + config.double_precision_rope = true; + config.apply_gated_attention = true; + config.ff_bias = true; + config.prefix = prefix; + return config; +} + +ParamBag ProcessorBag(const vllm::Ltx2ConnectorConfig& config) { + ParamBag bag; + for (const vllm::Ltx2ConnectorTensorSpec& spec : vllm::EnumerateLtx2ConnectorTensors(config)) { + bag.Put(spec.name, spec.shape); + } + return bag; +} + +} // namespace + +TEST_CASE("ltx2 the processor is PADDING-SIDE AGNOSTIC, which is what the sort is for") { + // "Connectors expect right-padded input ([valid, pad]). Normalize layout here + // so the upstream tokenizer can keep using either side without coupling to the + // connector." — embeddings_processor.py:80-82. + // + // THE DEFECT THIS CATCHES. The register table is indexed by ABSOLUTE position + // (`s % num_registers`), not by which positions were padded, so a LEFT-padded + // batch handed straight to the connector puts registers where caption tokens + // belong. The result is finite, correctly shaped, and conditioned on the wrong + // thing — which no shape or finiteness check can see. Skipping the sort makes + // these two renders DIFFER; upstream's own contract is that they are the same. + const int64_t batch = 1, seq = 4, valid = 2; + const vllm::Ltx2ConnectorConfig vcfg = ProcessorConfig("ltx2.proc.v.", 3, 8, 2); + const vllm::Ltx2ConnectorConfig acfg = ProcessorConfig("ltx2.proc.a.", 2, 4, 2); + const ParamBag vbag = ProcessorBag(vcfg); + const ParamBag abag = ProcessorBag(acfg); + const int64_t vdim = vcfg.inner_dim(), adim = acfg.inner_dim(); + + const std::vector real_v = Make("ltx2.proc.real.v", valid * vdim, 1.0); + const std::vector real_a = Make("ltx2.proc.real.a", valid * adim, 1.0); + const std::vector junk_v = Make("ltx2.proc.junk.v", (seq - valid) * vdim, 1.0); + const std::vector junk_a = Make("ltx2.proc.junk.a", (seq - valid) * adim, 1.0); + + auto build = [&](bool pad_left, const std::vector& real, const std::vector& junk, + int64_t width) { + std::vector out; + if (pad_left) { + out.insert(out.end(), junk.begin(), junk.end()); + out.insert(out.end(), real.begin(), real.end()); + } else { + out.insert(out.end(), real.begin(), real.end()); + out.insert(out.end(), junk.begin(), junk.end()); + } + (void)width; + return out; + }; + auto mask_of = [&](bool pad_left) { + std::vector m(static_cast(seq), 0.0f); + for (int64_t s = 0; s < seq - valid; ++s) { + m[static_cast(pad_left ? s : valid + s)] = -std::numeric_limits::max(); + } + return m; + }; + + const std::vector right_v = build(false, real_v, junk_v, vdim); + const std::vector right_a = build(false, real_a, junk_a, adim); + const std::vector left_v = build(true, real_v, junk_v, vdim); + const std::vector left_a = build(true, real_a, junk_a, adim); + const std::vector right_mask = mask_of(false); + const std::vector left_mask = mask_of(true); + + const vllm::Ltx2ConnectorEmbeddings from_right = vllm::Ltx2ConnectorCreateEmbeddings( + vcfg, vbag.weights, right_v.data(), acfg, abag.weights, right_a.data(), right_mask.data(), + batch, seq); + const vllm::Ltx2ConnectorEmbeddings from_left = vllm::Ltx2ConnectorCreateEmbeddings( + vcfg, vbag.weights, left_v.data(), acfg, abag.weights, left_a.data(), left_mask.data(), + batch, seq); + + const double video_worst = + MaxAbsDiff(from_left.video, from_right.video.data(), from_right.video.size()); + const double audio_worst = + MaxAbsDiff(from_left.audio, from_right.audio.data(), from_right.audio.size()); + INFO("padding-side agnostic: video max|diff| = ", video_worst, " audio max|diff| = ", + audio_worst); + CHECK(video_worst <= kExactRoundOff); + CHECK(audio_worst <= kExactRoundOff); + + // THE CONTROL. The two inputs really are different buffers, so the equality + // above is the sort working and not two identical arrays compared to + // themselves. Without this the case passes on a port that ignores its input. + CHECK(left_v != right_v); + CHECK(left_a != right_a); + // ...and with registers on, every position is attendable (:152). + for (const float m : from_right.mask) CHECK(m == 1.0f); +} + +TEST_CASE("ltx2 the processor's binary mask mirrors a comparison that looks backwards") { + // `_to_binary_mask` is `encoded_mask < 0.000001` (embeddings_processor.py:46-48). + // An additive mask holds 0.0 for KEPT and -finfo(f32).max for PADDED, and BOTH + // are `< 0.000001` — so the mask this produces is ONE EVERYWHERE, at every + // position, for every input either reference can produce, and the video-only + // multiply that follows it is an identity. + // + // THIS CASE EXISTS BECAUSE THE INTENT-READING IS THE OPPOSITE. A port that + // reasoned "keep the unmasked ones" would write `>= 0`, get 0 at padded + // positions, zero the video encoding there, and hand the DiT a mask that + // masks. It would look more correct and it would not be a port. Written as a + // gate on the SURPRISING behaviour so that "fixing" it REDs. + // + // Confirmed on BOTH references before being pinned, because a line this odd is + // where one implementation being wrong would show: `diffusers` + // `LTX2TextConnectors.forward` writes `(video_attn_mask < 1e-6).to(torch.int64)` + // and the same video-only multiply. They agree, down to the constant. + // + // Registers are DISABLED here so the connector passes the caller's mask through + // instead of zeroing it — that is the only configuration in which the two + // readings differ at all. + const int64_t batch = 1, seq = 4, valid = 2; + const vllm::Ltx2ConnectorConfig vcfg = ProcessorConfig("ltx2.proc.nv.", 3, 8, 0); + const vllm::Ltx2ConnectorConfig acfg = ProcessorConfig("ltx2.proc.na.", 2, 4, 0); + const ParamBag vbag = ProcessorBag(vcfg); + const ParamBag abag = ProcessorBag(acfg); + const int64_t vdim = vcfg.inner_dim(), adim = acfg.inner_dim(); + + const std::vector hidden_v = Make("ltx2.proc.nv.hidden", batch * seq * vdim, 1.0); + const std::vector hidden_a = Make("ltx2.proc.na.hidden", batch * seq * adim, 1.0); + std::vector mask(static_cast(seq), 0.0f); + for (int64_t s = valid; s < seq; ++s) mask[static_cast(s)] = -std::numeric_limits::max(); + + const vllm::Ltx2ConnectorEmbeddings got = vllm::Ltx2ConnectorCreateEmbeddings( + vcfg, vbag.weights, hidden_v.data(), acfg, abag.weights, hidden_a.data(), mask.data(), + batch, seq); + + for (int64_t s = 0; s < seq; ++s) { + const bool kept = s < valid; + double video_abs = 0.0, audio_abs = 0.0; + for (int64_t i = 0; i < vdim; ++i) { + video_abs += std::fabs(static_cast(got.video[static_cast(s * vdim + i)])); + } + for (int64_t i = 0; i < adim; ++i) { + audio_abs += std::fabs(static_cast(got.audio[static_cast(s * adim + i)])); + } + INFO("position ", s, " kept = ", kept, " |video| = ", video_abs, " |audio| = ", audio_abs); + // ONE at EVERY position, padded ones included. The `>= 0` reading gives 0.0 + // here for s >= valid, which is what this pins. + CHECK(got.mask[static_cast(s)] == 1.0f); + // ...so nothing is zeroed, in either modality. Under the `>= 0` reading the + // video row would be exactly 0 at the padded positions. + CHECK(video_abs > 0.0); + CHECK(audio_abs > 0.0); + } + + // THE CONTROL. The connector really did see a mask with padded positions in + // it: with registers disabled the padded rows are NOT substituted, so they + // still derive from the caller's own features. Without this the case would + // pass on a processor that dropped the mask argument entirely. + const int64_t pad = seq - valid; + std::vector unpadded_mask(static_cast(seq), 0.0f); + const vllm::Ltx2ConnectorEmbeddings all_valid = vllm::Ltx2ConnectorCreateEmbeddings( + vcfg, vbag.weights, hidden_v.data(), acfg, abag.weights, hidden_a.data(), + unpadded_mask.data(), batch, seq); + const double masked_vs_unmasked = + MaxAbsDiff(all_valid.video, got.video.data(), got.video.size()); + INFO("padded rows = ", pad, " max|diff| vs an all-valid mask = ", masked_vs_unmasked); + CHECK(masked_vs_unmasked > 0.0); +} diff --git a/tests/vllm/multimodal/ltx2_video_fixture.h b/tests/vllm/multimodal/ltx2_video_fixture.h index d9a89ba2b..9df393fee 100644 --- a/tests/vllm/multimodal/ltx2_video_fixture.h +++ b/tests/vllm/multimodal/ltx2_video_fixture.h @@ -234,12 +234,57 @@ inline vllm::Ltx2DitParams ReducedDitParams() { return p; } +// ── the embeddings connector ─────────────────────────────────────────────── +// +// The two `*_embeddings_connector` families the shipped DiTs carry (129 tensors +// each) and that phase L9c wires into the render path. Reduced in every axis the +// shipped one is reduced in, and NOT reduced in the one that decides the +// substitution: `num_learnable_registers` tiles across the sequence, so the +// prompt row count must be a multiple of it. +struct ReducedConnectorOptions { + bool present = true; + int64_t num_layers = 2; + int64_t num_learnable_registers = 2; + bool gated = true; + bool ff_bias = true; + // Seeds the connector's own parameter stream. A second value writes a DIFFERENT + // connector into an otherwise byte-identical checkpoint, which is how a test + // proves the render actually READS these weights. + std::string tag = "a"; + // Write the audio family too. `false` is the half-a-connector checkpoint the + // engine refuses. + bool audio = true; +}; + +// The connector configuration the fixture's own tensors are written from, so the +// config the engine parses and the shapes it finds can never disagree by +// accident — only when a test makes them. +inline vllm::Ltx2ConnectorConfig ReducedConnectorConfig(const vllm::Ltx2DitParams& params, + const ReducedConnectorOptions& options, + vllm::Ltx2ConnectorStream stream) { + vllm::Ltx2ConnectorConfig c; + c.prefix = vllm::Ltx2ConnectorCheckpointPrefix(stream); + const bool video = stream == vllm::Ltx2ConnectorStream::kVideo; + c.num_attention_heads = video ? params.num_attention_heads : params.audio_num_attention_heads; + c.attention_head_dim = video ? params.attention_head_dim : params.audio_attention_head_dim; + c.num_layers = options.num_layers; + c.num_learnable_registers = options.num_learnable_registers; + c.apply_gated_attention = options.gated; + c.ff_bias = options.ff_bias; + c.rope_type = vllm::Ltx2RopeType::kSplit; + c.double_precision_rope = true; // the shipped config's frequencies_precision + c.positional_embedding_max_pos = {4096}; + return c; +} + // The DiT's `{"transformer": {...}}` object, exactly as the first-party // `Lightricks/LTX-2.5` DiT carries it in `__metadata__["config"]`. Built // separately from the file writer because two callers need it: the writer, and a // test that has to hand the SAME config to the engine through the // `dit_config_path` extra for a checkpoint that declares none. -inline nlohmann::json ReducedDitTransformerConfig(const vllm::Ltx2DitParams& params) { +inline nlohmann::json ReducedDitTransformerConfig( + const vllm::Ltx2DitParams& params, + const ReducedConnectorOptions& connector = ReducedConnectorOptions{}) { nlohmann::json transformer; transformer["_class_name"] = "AVTransformer3DModel"; // Every key `LTXModelConfigurator.from_metadata` runs `check_config_value` on @@ -290,6 +335,19 @@ inline nlohmann::json ReducedDitTransformerConfig(const vllm::Ltx2DitParams& par transformer["ff_bias"] = params.ff_bias; transformer["rope_type"] = "split"; transformer["use_middle_indices_grid"] = params.use_middle_indices_grid; + // The `connector_*` keys the two Embeddings1DConnector configurators read + // (embeddings_connector.py:194-256). `connector_positional_embedding_max_pos` + // is the one the shipped config moves OFF its class default of [1], and it + // divides every token index, so it is written at the shipped [4096]. + transformer["connector_num_attention_heads"] = params.num_attention_heads; + transformer["connector_attention_head_dim"] = params.attention_head_dim; + transformer["audio_connector_num_attention_heads"] = params.audio_num_attention_heads; + transformer["audio_connector_attention_head_dim"] = params.audio_attention_head_dim; + transformer["connector_num_layers"] = connector.num_layers; + transformer["connector_num_learnable_registers"] = connector.num_learnable_registers; + transformer["connector_apply_gated_attention"] = connector.gated; + transformer["connector_ff_bias"] = connector.ff_bias; + transformer["connector_positional_embedding_max_pos"] = std::vector{4096}; return transformer; } @@ -311,6 +369,7 @@ struct ReducedDitOptions { bool declare_config = true; bool declare_model_version = true; nlohmann::json transformer_overrides = nlohmann::json::object(); + ReducedConnectorOptions connector; }; // Write the DiT in the shipped ComfyUI + FP8 shape: every rank-2 `*.weight` @@ -343,9 +402,42 @@ inline void WriteReducedDit(const vllm::Ltx2DitParams& params, const std::string entries.push_back({full, "BF16", spec.shape, std::move(values)}); } } + // The connector families, written in the SAME FP8/BF16 split. They sit beside + // the DiT contract, not inside it — upstream loads them into the text + // encoder's EmbeddingsProcessor (encoder_configurator.py:331-346) — which is + // why they are enumerated from their own contract rather than the DiT's. + if (options.connector.present) { + std::vector streams = {vllm::Ltx2ConnectorStream::kVideo}; + if (options.connector.audio) streams.push_back(vllm::Ltx2ConnectorStream::kAudio); + for (const vllm::Ltx2ConnectorStream stream : streams) { + const vllm::Ltx2ConnectorConfig c = + ReducedConnectorConfig(params, options.connector, stream); + for (const vllm::Ltx2ConnectorTensorSpec& spec : vllm::EnumerateLtx2ConnectorTensors(c)) { + int64_t numel = 1; + for (const int64_t d : spec.shape) numel *= d; + const std::string full = prefix + spec.name; + std::vector values = + Param("ltx2.conn." + options.connector.tag + "." + spec.name, numel, 0.08); + // MEASURED from the shipped `vonkaiser` FP8 DiT header: every rank-2 + // tensor of this family is F8_E4M3 with an F32 sidecar, `learnable_registers` + // INCLUDED (`...learnable_registers` + `...learnable_registers_scale`). + // That is not the DiT's rule — there the scale-shift tables stay F32 — + // so it is stated from the file rather than inherited. + const bool quantizable = spec.shape.size() == 2; + if (quantizable) { + constexpr float kScale = 0.5F; + for (float& v : values) v /= kScale; + entries.push_back({full, "F8_E4M3", spec.shape, std::move(values)}); + entries.push_back({full + "_scale", "F32", {}, {kScale}}); + } else { + entries.push_back({full, "BF16", spec.shape, std::move(values)}); + } + } + } + } nlohmann::json metadata = nlohmann::json::object(); if (options.declare_config) { - nlohmann::json transformer = ReducedDitTransformerConfig(params); + nlohmann::json transformer = ReducedDitTransformerConfig(params, options.connector); transformer.update(options.transformer_overrides); nlohmann::json config; config["transformer"] = transformer; @@ -368,8 +460,10 @@ inline void WriteReducedDit(const vllm::Ltx2DitParams& params, const std::string // what the engine's `dit_config_path` extra reads. inline void WriteDitConfigJson(const vllm::Ltx2DitParams& params, const std::string& path, const nlohmann::json& transformer_overrides = - nlohmann::json::object()) { - nlohmann::json transformer = ReducedDitTransformerConfig(params); + nlohmann::json::object(), + const ReducedConnectorOptions& connector = + ReducedConnectorOptions{}) { + nlohmann::json transformer = ReducedDitTransformerConfig(params, connector); transformer.update(transformer_overrides); nlohmann::json config; config["transformer"] = transformer; @@ -772,7 +866,12 @@ struct Paths { std::string dit, video_vae, audio_vae, upsampler, video_embeds, audio_embeds; }; -inline Paths WriteFixture(const std::string& dir, int64_t prompt_tokens = 3) { +// `prompt_tokens` defaults to 4, not 3: the connector's register table is TILED +// across the sequence (embeddings_connector.py:144), so the row count must be a +// multiple of `ReducedConnectorOptions::num_learnable_registers` (2), and 4 makes +// the tiling repeat twice rather than once — a single repeat cannot separate a +// tiled table from an indexed one. +inline Paths WriteFixture(const std::string& dir, int64_t prompt_tokens = 4) { ::mkdir(dir.c_str(), 0755); const vllm::Ltx2DitParams dit = ReducedDitParams(); Paths p; diff --git a/tests/vllm/multimodal/test_ltx2_video.cpp b/tests/vllm/multimodal/test_ltx2_video.cpp index af5c6e268..bc4918a39 100644 --- a/tests/vllm/multimodal/test_ltx2_video.cpp +++ b/tests/vllm/multimodal/test_ltx2_video.cpp @@ -1008,3 +1008,192 @@ TEST_CASE("ltx2 video: the SHIPPED Lightricks checkpoints parse and load") { MESSAGE("shipped upsampler: " << weights.tensors.size() << " tensors"); } } + +// ─── the embeddings connector (phase L9c) ─────────────────────────────────── +// +// WHAT THESE ARE FOR. Until L9c the conditioning this engine handed the DiT was +// the prompt-embeds file VERBATIM: `Ltx2ConnectorForward` had landed at L5, had +// been gated against upstream on five arms, and was called by NOTHING but its +// own test, while the 129-tensor `*_embeddings_connector` family the shipped DiT +// carries was refused as unported and stepped over. Every case below fails on +// the pre-L9c engine, and the first two fail for the reason that matters: +// deleting the connector call leaves the render byte-identical to a render with +// DIFFERENT connector weights, because nothing read them. + +namespace { + +// Every frame byte of a render, so two renders can be compared as a whole rather +// than through a statistic that might not move. +// +// Phase 0 only, for the same reason the first render case gives: the recipe's +// second phase needs the latent spatial upsampler and running without one is a +// REFUSAL. Phase 0 already carries the conditioning through cross-attention in +// every one of its blocks, which is what these cases are about. +std::string RenderBytes(vllm::multimodal::VideoModelParams mp, const std::string& out_dir) { + mp.extras[vllm::multimodal::kLtx2MaxPhaseExtra] = "0"; + const std::unique_ptr engine = + vllm::multimodal::LoadVideoEngine(mp); + const vllm::multimodal::VideoResult result = engine->Generate(FixtureGen(out_dir)); + std::string all; + for (int64_t f = 0; f < result.frame_count; ++f) { + char name[64]; + std::snprintf(name, sizeof(name), "/frame_%06lld.ppm", static_cast(f)); + all += ReadAll(out_dir + name); + } + all += ReadAll(result.audio_path); + return all; +} + +std::string RefusalOf(const vllm::multimodal::VideoModelParams& mp) { + try { + (void)vllm::multimodal::LoadVideoEngine(mp); + } catch (const std::exception& e) { + return e.what(); + } + return std::string(); +} + +} // namespace + +TEST_CASE("ltx2 video: the render READS the checkpoint's connector weights") { + Workspace ws; + const vllm::Ltx2DitParams dit = ltx2_fixture::ReducedDitParams(); + + // A second DiT, byte-identical except that its connector's parameter stream is + // seeded differently. Same shapes, same config, same DiT weights. + ltx2_fixture::ReducedDitOptions other; + other.connector.tag = "b"; + const std::string other_dit = ws.root + "/dit_connector_b.safetensors"; + ltx2_fixture::WriteReducedDit(dit, other_dit, other); + + vllm::multimodal::VideoModelParams a = FixtureParams(ws.paths); + vllm::multimodal::VideoModelParams b = a; + b.dit_path = other_dit; + + const std::string frames_a = RenderBytes(a, ws.root + "/conn_a"); + const std::string frames_b = RenderBytes(b, ws.root + "/conn_b"); + CHECK(frames_a.size() == frames_b.size()); + // BEFORE L9c these are EQUAL: the connector weights were loaded by nobody, so + // the only thing that differed between the two files was never read. + CHECK(frames_a != frames_b); + + // ...and the same render twice is byte-identical, which is what makes the + // inequality above a statement about the connector rather than about noise. + CHECK(RenderBytes(a, ws.root + "/conn_a2") == frames_a); +} + +TEST_CASE("ltx2 video: the connector's positional bound comes from the CONFIG") { + // `connector_positional_embedding_max_pos` divides every token index + // (rope.py:132-141). LTX-2.5 declares [4096]; the class default is [1], which + // is 4096x. Nothing in a SHAPE can see the difference, so a config that is + // parsed but not USED renders confidently at the wrong RoPE — the invisible + // constant class spec section 7.0(a) names. + Workspace ws; + const vllm::Ltx2DitParams dit = ltx2_fixture::ReducedDitParams(); + ltx2_fixture::ReducedDitOptions defaulted; + defaulted.transformer_overrides["connector_positional_embedding_max_pos"] = + std::vector{1}; + const std::string other_dit = ws.root + "/dit_maxpos_1.safetensors"; + ltx2_fixture::WriteReducedDit(dit, other_dit, defaulted); + + vllm::multimodal::VideoModelParams a = FixtureParams(ws.paths); + vllm::multimodal::VideoModelParams b = a; + b.dit_path = other_dit; + CHECK(RenderBytes(a, ws.root + "/maxpos_4096") != RenderBytes(b, ws.root + "/maxpos_1")); +} + +TEST_CASE("ltx2 video: a connector config that disagrees with the FILE is refused") { + Workspace ws; + const vllm::Ltx2DitParams dit = ltx2_fixture::ReducedDitParams(); + + SUBCASE("fewer layers than the file carries binds a PREFIX, so it is refused") { + ltx2_fixture::ReducedDitOptions shrunk; + shrunk.transformer_overrides["connector_num_layers"] = 1; + const std::string path = ws.root + "/dit_conn_1layer.safetensors"; + ltx2_fixture::WriteReducedDit(dit, path, shrunk); + vllm::multimodal::VideoModelParams mp = FixtureParams(ws.paths); + mp.dit_path = path; + const std::string msg = RefusalOf(mp); + INFO(msg); + CHECK(msg.find("embeddings_connector") != std::string::npos); + CHECK(msg.find("does not name") != std::string::npos); + } + + SUBCASE("more layers than the file carries names the MISSING tensor") { + ltx2_fixture::ReducedDitOptions grown; + grown.transformer_overrides["connector_num_layers"] = 3; + const std::string path = ws.root + "/dit_conn_3layer.safetensors"; + ltx2_fixture::WriteReducedDit(dit, path, grown); + vllm::multimodal::VideoModelParams mp = FixtureParams(ws.paths); + mp.dit_path = path; + const std::string msg = RefusalOf(mp); + INFO(msg); + CHECK(msg.find("transformer_1d_blocks.2") != std::string::npos); + } + + SUBCASE("gating the config declares but the file does not carry is refused") { + ltx2_fixture::ReducedDitOptions ungated; + ungated.connector.gated = false; // writes NO to_gate_logits tensors... + ungated.transformer_overrides["connector_apply_gated_attention"] = true; // ...but claims them + const std::string path = ws.root + "/dit_conn_gate.safetensors"; + ltx2_fixture::WriteReducedDit(dit, path, ungated); + vllm::multimodal::VideoModelParams mp = FixtureParams(ws.paths); + mp.dit_path = path; + const std::string msg = RefusalOf(mp); + INFO(msg); + CHECK(msg.find("to_gate_logits") != std::string::npos); + } +} + +TEST_CASE("ltx2 video: half a connector conditions two modalities differently") { + Workspace ws; + ltx2_fixture::ReducedDitOptions video_only; + video_only.connector.audio = false; + const std::string path = ws.root + "/dit_video_connector_only.safetensors"; + ltx2_fixture::WriteReducedDit(ltx2_fixture::ReducedDitParams(), path, video_only); + vllm::multimodal::VideoModelParams mp = FixtureParams(ws.paths); + mp.dit_path = path; + const std::string msg = RefusalOf(mp); + INFO(msg); + CHECK(msg.find("embeddings connector") != std::string::npos); + CHECK(msg.find("audio") != std::string::npos); +} + +TEST_CASE("ltx2 video: prompt rows the register table cannot tile are refused") { + // `seq_len % num_learnable_registers == 0` is upstream's own assert + // (embeddings_connector.py:144), because the table is TILED across the + // sequence rather than indexed by which positions were padded. + Workspace ws; + const std::string odd = ws.root + "/three_rows.f32"; + const vllm::Ltx2DitParams dit = ltx2_fixture::ReducedDitParams(); + ltx2_fixture::WritePromptEmbeds(odd, "ltx2.embeds.video.odd", 3, dit.cross_attention_dim); + const std::string odd_audio = ws.root + "/three_rows_audio.f32"; + ltx2_fixture::WritePromptEmbeds(odd_audio, "ltx2.embeds.audio.odd", 3, + dit.audio_cross_attention_dim); + vllm::multimodal::VideoModelParams mp = FixtureParams(ws.paths); + mp.prompt_embeds_path = odd; + mp.extras[vllm::multimodal::kLtx2AudioPromptEmbedsExtra] = odd_audio; + const std::string msg = RefusalOf(mp); + INFO(msg); + CHECK(msg.find("learnable registers") != std::string::npos); +} + +TEST_CASE("ltx2 video: the valid-row count decides which positions are registers") { + // With no text tower there is no tokenizer mask, so `prompt_embeds_valid_rows` + // is what says which supplied rows are caption and which are padding — and + // padding is not inert here, it is REPLACED by the learnable register table. + // A render that ignored the extra would be identical whatever it said. + Workspace ws; + vllm::multimodal::VideoModelParams all_valid = FixtureParams(ws.paths); + vllm::multimodal::VideoModelParams half = all_valid; + half.extras[vllm::multimodal::kLtx2PromptValidRowsExtra] = "2"; + CHECK(RenderBytes(all_valid, ws.root + "/rows_4") != RenderBytes(half, ws.root + "/rows_2")); + + SUBCASE("a count past the end of the file is refused") { + vllm::multimodal::VideoModelParams over = all_valid; + over.extras[vllm::multimodal::kLtx2PromptValidRowsExtra] = "99"; + const std::string msg = RefusalOf(over); + INFO(msg); + CHECK(msg.find(vllm::multimodal::kLtx2PromptValidRowsExtra) != std::string::npos); + } +} From 708e05b56764b39f03a589c7c4ead1c0749d0ca2 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 03:23:34 +0000 Subject: [PATCH 2/2] record(ltx-2.5): the two L9C arms the GPU lock never came free for FOLLOWING_AGENTS_PROTOCOL Two bounded waits on $HOME/gpu.lock, `flock -w 2700` each, both timed out while other coordinators held the box. Waiting is normal and stealing is not, so the two arms are recorded OWED with the scripts that run them already on the box, rather than left as an absence somebody has to rediscover. The one that matters: `test_minimax_h3` and `test_capi` are UNRESOLVED, not green. The first baseline pass ran the BINARIES directly and each printed a doctest summary and then died -- `test_capi` `4 cases | 51 skipped` against a brief baseline of 55, `test_minimax_h3` `38 | 41 skipped` against 79. doctest counts the cases it never reached as "skipped", so the summary reads like a run that chose to skip them. That is the third time this campaign has recorded a green that means nothing, and it is why the second pass drives ctest and keeps the FULL output instead of grepping a bare binary. The other is the arm that would settle where L9B's ~58 GB came from at L9B's own geometry: L9B's binary, L9B's arguments. It cannot run on the L9C binary, which refuses 32 conditioning rows because that is not a multiple of the connector's 128 learnable registers. Issue: #435 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/benchmark-record.md | 23 ++++++++++++++++++++++- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 3 files changed, 24 insertions(+), 3 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 546bfce9b..60540f790 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -19552,6 +19552,27 @@ speech-shaped, or matches the mouth movements in the frames, is not something these numbers answer and is not asserted. Owed: a spectral check against the video, which is the audio half of the question the frames just answered. -Evidence: `dgx:~/work/ltx25-l9c/{hold.log,baselines.log,render-*.log,mem-*.log}`, +### Owed, and why — two arms the lock never came free for + +Two bounded waits on `$HOME/gpu.lock`, `flock -w 2700` each, both timed out +(`HOLD2_WRAPPER_DONE rc=1` at 04:36 and 05:22). The box was saturated with other +coordinators' `ctest` work for the whole 90 minutes. Waiting is normal and +stealing is not, so these are reported OWED rather than run: + +1. **The L9B repro arm** (`dgx:~/work/ltx25-l9c/dgx_repro_l9b.sh`, shipped and + syntax-checked on the box): L9B's own binary with L9B's own arguments at + 320x192/25f `--max-phase 0`, which is the one arm that separates "L9B measured + the environment" from "L9B's geometry behaves differently from ours". It cannot + run on the L9c binary, which refuses 32 conditioning rows. +2. **`test_minimax_h3` and `test_capi` through ctest** (`dgx_baselines2.sh`, also + shipped). The first baseline pass ran the BINARIES directly and both reported a + summary followed by a SIGSEGV — `test_capi` `4 cases | 51 skipped` against a + brief baseline of 55, `test_minimax_h3` `38 | 41 skipped` against 79. A summary + printed before a crash counts the unreached cases as "skipped", which is the + third instance this campaign has recorded of a run reading as a pass. **These + two are UNRESOLVED, not green**, and the instrument that would resolve them is + ctest plus the full output, not a grep of a bare binary. + +Evidence: `dgx:~/work/ltx25-l9c/{hold.log,hold2.log,hold2b.log,baselines.log,render-*.log,mem-*.log}`, renders under `dgx:~/work/ltx25-l9c/render/`, contact sheets under `dgx:~/work/ltx25-l9c/contact/`. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index be956e94d..fd4be025c 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -414,7 +414,7 @@ built on it rather than keeping the flattering one. | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | **RUNNER FOLD LANDS (ROW 7, §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (near-tie profile); FA2 MLA default-ON; SACRED green.** Server 19.0 tok/s wall; CLI 18.93 reproduced | vLLM ~21 (#111 floor; in-session re-measure ABORTED by GB10 reboot at util 0.82, §21): **~0.90×**, >= vLLM NOT met; residual = KDA host islands + grouped MoE + decode graph | | 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) | fp4 speed CLOSED. bf16-vs-quant A/B: ENCODER half MEASURED (§8.15), DiT half NOT (no bf16 render exists). Detail: benchmark-record + spec §8 | -| LTX-2.5 axes | Speed `PENDING` (vllm-omni#6066 has no native 2.5; its adapter is a black box), binding oracle too. **SIZE: 320x192/25f completes on GB10, 448x256/25f does not** | The wall is the HOST VAE decode, not the pool: the drain returns 0.11 GiB, byte-inert | +| LTX-2.5 axes | Speed `PENDING` (vllm-omni#6066 has no native 2.5), binding oracle too. **SIZE: 320x192/25f completes on GB10, 448x256 does not** | Wall is the HOST VAE decode, not the pool: drain returns 0.11 GiB, byte-inert. 2 baselines UNRESOLVED (lock) | | 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. Follow-up `H3-TASK-PARTITION-GUARD`: the task/partition mismatch now RAISES 1:1 with `_resolve_task` (spec §8.6-8.7) | | MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`) | **fl2va COHERENT; ref2va assembly bug FIXED+gated.** vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the **NVFP4 checkpoint/loader**, NOT assembly/fp4 (§8.10) | | MiniMax-H3 Thor render speed (sm_110, no FA2) | **34.6 s/step** at 864x480/124f/50 steps on Q4_K_M, **16.6x** off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 `mma.sync` (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head's K+V is 3.9 MB against 32 MB of L2) | diff --git a/docs/STATUS.md b/docs/STATUS.md index dce6457cd..e57f85614 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -106,7 +106,7 @@ token-for-token correctness against the pinned oracle. | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: was 87% of vLLM (37.55 vs 43.10, same-tool nsys)**; root cause was bf16 projections on UNIFIED/ATS host memory, and device-resident staging (byte-exact) gives 44.6, parity+ vs 43.1, default-ON | 48 layers (12 global + 36 SWA-512), 256 routed top-10 + 1 shared expert, per-head softplus attn out-gate, sigmoid `noaux_tc` router, dual per-layer RoPE, GQA 8 KV / 128 head-dim, 1M ctx. History: benchmark-record | | InternLM2 dense (fused-`wqkv` interleaved split) | Correctness-complete, speed-pending | Token-exact 16/16 (internlm2-chat-1_8b): 12/16 strict + 4/16 bf16 near-tie (max gap 0.0 nats), 0 divergent; first InternLM model; ZERO new compute kernel (reuses the Llama dense forward; the only delta is a loader-side de-interleave of the fused `wqkv`, which packs q/k/v interleaved by KV-group) | | MiniMax-H3 (`MiniMaxH3DiTModel`, video+audio DIFFUSION) | **ABI v12 ONE SURFACE; device selector uses generic `DeviceType`; DSR 32.** t2va+fl2va COHERENT; bf16 shards STREAM | ref2va ckpt fidelity §8.12; encoder A/B §8.15; GB10 re-verify residual; CPU fold 6/137 (one queue + device provenance mutation-gated) | -| LTX-2.5 (`LTX2VideoTransformer3DModel`, video+audio DIFFUSION) | **L1-L9c landed (#435).** 21.00B / 48 blocks. `VideoEngine` seam + ABI **v18**, DiT forward (CPU f32 parity, bf16 device-resident), Gemma-4 TE, both VAEs, the embeddings connector, pipeline, NVFP4/FP8 arms, `/v1/videos` | A shipped 21.00B FP8 DiT stages and RUNS device-resident on GB10, all ops `vt-native`, zero reference-tier hits. With the connector wired the frames ARE a scene at 320x192/25f; 448x256 stops in the HOST VAE decode. Speed axis `PENDING` (spec section 0: no production-config denominator); binding-oracle parity `PENDING`; the Gemma-4 TOWER still owed, so nothing encodes a prompt | +| LTX-2.5 (`LTX2VideoTransformer3DModel`, video+audio DIFFUSION) | **L1-L9c landed (#435).** 21.00B / 48 blocks. `VideoEngine` seam + ABI **v18**, DiT forward (CPU f32 parity, bf16 device-resident), Gemma-4 TE, both VAEs, the embeddings connector, pipeline, NVFP4/FP8 arms, `/v1/videos` | A shipped 21.00B FP8 DiT stages and RUNS device-resident on GB10, all ops `vt-native`, zero reference-tier hits. With the connector wired the frames ARE a scene at 320x192/25f; 448x256 stops in the HOST VAE decode. Speed axis `PENDING` (spec section 0: no production-config denominator); binding-oracle parity `PENDING`; the Gemma-4 TOWER still owed, so nothing encodes a prompt. `test_minimax_h3` and `test_capi` are UNRESOLVED on this branch, not green: two bounded GPU-lock waits timed out before the ctest re-run | | Command-R / Cohere dense (`CohereForCausalLM`) | Implemented, gate-blocked | ZERO-new-kernel port grounded in vLLM `commandr.py`: weight-only Cohere LayerNorm + GPT-J full-width RoPE + PARALLEL residual + `logit_scale` + tied embeddings, all reuse; compiles, links, self-registers. No SACRED gate yet (real checkpoints HF-gated, ungated ones tiny-random, GPU box disk-full); oracle run-verified at W0. See docs/BENCHMARKS.md | | Phi-1 / Phi-2 dense (`PhiForCausalLM`, parallel residual) | Correctness-complete, speed-pending | Token-exact 16/16 (microsoft/phi-2): 9/16 strict + 7/16 bf16 near-ties (max gap 0.25 nats), 0 forward-divergent; the OLDER Microsoft Phi arch, DISTINCT from Phi-3/Phi-4; ZERO new compute kernel (GPT-J parallel residual, LayerNorm-with-bias, biased qkv/dense, partial NeoX rope 32/80, non-gated NewGELU MLP reusing `vt::GeluTanh`, untied biased lm_head); F16 dtype-aware loader | | MiniCPM dense (`MiniCPMForCausalLM`, three scalars) | Correctness-complete, speed-pending | Token-exact 16/16 (openbmb/MiniCPM-2B-sft-bf16): 10/16 strict + 6/16 bf16 near-ties (max gap 0.0 nats), 0 forward-divergent; first OpenBMB MiniCPM model; ZERO new compute kernel (the Llama/Granite dense forward plus three scalars: scale_emb, scale_depth/sqrt(layers) residual, dim_model_base logit scaling), tied lm_head; `.bin`-only weights converted to safetensors via trusted torch |