feat(ltx-2.5): LTX-2.5 joint video+audio DiT, and a video seam that is no longer MiniMax-only (#435)

.agents/claims/CLAIM-LTX25-L0-L7.md

@@ -0,0 +1,5 @@
+# CLAIM-LTX25-L0-L7
+
+| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update |
+|---|---|---|---|---|---|---|---|
+| `CLAIM-LTX25-L0-L7` | `MODEL-DIFFUSION-ltx-2-5-ltx2-video-transformer-3d-model` (`SPIKE`), `ROAD-V1-LTX25` | Claude Code (opus-5), operator role | isolated worktree `/home/mudler/_git/vllm.cpp-ltx25`; GPU work on dgx.casa serialized through `flock $HOME/gpu.lock` | `row/MODEL-DIFFUSION-LTX25`, issue [#435](https://github.com/mudler/vllm.cpp/issues/435) | Owns: NEW `.agents/specs/ltx-2-5.md`; the LTX-2.5 rows in `.agents/roadmap_v1.md` and `.agents/model-matrix.md`; the LTX-2.5 rows in `docs/{STATUS,BENCHMARKS,FEATURES}.md`; NEW `src/vllm/model_executor/models/ltx2_*` and `include/vllm/model_executor/models/ltx2*`; NEW `vllm::multimodal::VideoEngine` seam plus the MiniMax-H3 move behind it (BEHAVIOUR-PRESERVING only); additive ABI v13 fields in `include/vllm.h`. EXCLUDES: any change to H3's numerics or its checkpoint handling, the v12 ABI contract, DiffVAE / `NADiffusionDecoder`, the temporal x2 upsampler, LoRA fusion, multishot, `int8-convrot`, and multi-GPU / CFG parallelism, all of which are recorded as owed in the spec rather than silently dropped | `ACTIVE` | 2026-08-13 — **MERGE AUTHORITY RECORDED**: developer-directed "keep going autonomously until you get it working, fix the bugs you find along the way and merge the PRs directly once everything is proved". Scope widened beyond L0-L7 to a working END-TO-END LTX-2.5 render: L9a NVFP4 linear scale layout, L9b real render + `--video-family`, L10 the Gemma-4 tower so a prompt works, L11 the VAE encoders for image/video/audio conditioning, and the #516 pool fix that a trustworthy render depends on. Merge is gated on PROVED, which for this campaign means: fresh review PASS, the operator's own gate rerun, and no red bought by weakening a detector. |

CMakeLists.txt

@@ -780,6 +780,36 @@ add_library(vllm STATIC
   src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp
   src/vllm/model_executor/models/minimax_h3_encoder_device.cpp
   src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp
+  # LTX-2.5 (row MODEL-DIFFUSION-LTX25, .agents/specs/ltx-2-5.md phase L2).
+  src/vllm/model_executor/models/ltx2.cpp
+  src/vllm/model_executor/models/ltx2_dit.cpp
+  # LTX-2.5 (row MODEL-DIFFUSION-LTX25, phase L8): the DEVICE-RESIDENT forward and
+  # the (kLtx2, device) resolver. The resolver is its own TU so the vt::GetOp cast
+  # links in CPU-only builds, exactly as minimax_h3_device_resolve.cpp is.
+  src/vllm/model_executor/models/ltx2_device_resolve.cpp
+  src/vllm/model_executor/models/ltx2_device.cpp
+  # LTX-2.5 (ROW MODEL-DIFFUSION-LTX25, phase L4): the Conv video VAE, the audio
+  # VAE and its vocoder. Additive files mirroring ltx_core's own structure.
+  src/vllm/model_executor/models/ltx2_audio_vae.cpp
+  src/vllm/model_executor/models/ltx2_video_vae.cpp
+  # LTX-2.5 (ROW MODEL-DIFFUSION-LTX25, phase L3): the Gemma-4 multi-layer feature
+  # aggregation, the two caption projections and the embedded asset pack. An
+  # additive file mirroring ltx_core's own structure.
+  src/vllm/model_executor/models/ltx2_text_encoder.cpp
+  # LTX-2.5 (ROW MODEL-DIFFUSION-LTX25, phase L5): the flow-matching pipeline,
+  # the latent spatial upsampler, the duration head and the embeddings
+  # connector. Additive files mirroring ltx_core's own structure.
+  src/vllm/model_executor/models/ltx2_pipeline.cpp
+  src/vllm/model_executor/models/ltx2_upsampler.cpp
+  src/vllm/model_executor/models/ltx2_duration_head.cpp
+  src/vllm/model_executor/models/ltx2_connector.cpp
+  # LTX-2.5 (ROW MODEL-DIFFUSION-LTX25, phase L11): the conditioning items that
+  # place the VAE ENCODERS' output into the denoise state.
+  src/vllm/model_executor/models/ltx2_conditioning.cpp
+  # LTX-2.5 (ROW MODEL-DIFFUSION-LTX25, phase L6): the quantized loaders — the
+  # FP8 and torchao-NVFP4 DiT arms, the torchao-NVFP4 text encoder, and the
+  # load-time device staging GB10's ATS penalty makes the default.
+  src/vllm/model_executor/models/ltx2_loader.cpp
   src/vllm/entrypoints/openai/video_api.cpp
   src/vllm/model_executor/models/kimi_kda.cpp
   src/vllm/model_executor/models/kimi_linear_registry.cpp
@@ -883,6 +913,15 @@ add_library(vllm STATIC
   # server's private /v1/videos twin. Builds artifacts + the ffmpeg argv and
   # SPAWNS NOTHING (the ratified process boundary).
   src/vllm/multimodal/minimax_h3_video.cpp
+  # The GENERALIZED video seam (vllm::multimodal::VideoEngine) + its
+  # checkpoint-detected family registry — LTX-2.5 L1, .agents/specs/ltx-2-5.md
+  # §5. Knows about no family; each family self-registers from its own TU.
+  src/vllm/multimodal/video_engine.cpp
+  # LTX-2.5 behind that seam — the SECOND family, and the driving loop that
+  # composes the L2-L6 bricks into frames + a waveform (phase L7). Additive:
+  # one file plus one REGISTER_VLLM_VIDEO_FAMILY line, which is the property
+  # the seam exists to have.
+  src/vllm/multimodal/ltx2_video.cpp
   src/vllm/v1/core/encoder_cache_manager.cpp
   src/vllm/tokenizer/unicode_data.cpp
   src/vllm/tokenizer/pretokenizer.cpp
@@ -1081,6 +1120,7 @@ add_library(vllm STATIC
   src/vt/cpu/cpu_threadpool.cpp
   src/vt/cpu/cpu_ops.cpp
   src/vt/cpu/cpu_minimax_h3.cpp
+  src/vt/cpu/cpu_ltx2.cpp
   src/vt/cpu/cpu_isa_arm.cpp
   src/vt/cpu/cpu_isa_x86.cpp
   src/vt/cpu/cpu_matmul_elem.cpp
@@ -1487,7 +1527,9 @@ if(VLLM_CPP_CUDA)
     src/vt/cuda/cuda_sample.cu
     src/vt/cuda/cuda_deepseek_v4.cu
     src/vt/cuda/cuda_laguna.cu
-    src/vt/cuda/cuda_minimax_h3.cu)
+    src/vt/cuda/cuda_minimax_h3.cu
+    src/vt/cuda/cuda_ltx2.cu
+    src/vt/cuda/cuda_attention_cross.cu)
   find_package(CUDAToolkit REQUIRED)
   # cublasLt is linked now so the Task 4 matmul lands without a build change.
   target_link_libraries(vllm PUBLIC CUDA::cudart CUDA::cublasLt)

README.md

@@ -186,7 +186,7 @@ configs, token-for-token the same output. Switching to it should be boring. Ever
 you get on top, most of it borrowed from whichever engine does it best:
 
 - **One 66 MiB binary instead of a 9.1 GiB install.** A flat, exception-free, llama.cpp-style C ABI
-  ([`include/vllm.h`](include/vllm.h), ABI v17, 35 functions) for C, C++, Go, or Rust. No Python
+  ([`include/vllm.h`](include/vllm.h), ABI v18, 36 functions) for C, C++, Go, or Rust. No Python
   interpreter in the process.
 - **GGUF as a first-class citizen.** Load the same quantized files llama.cpp uses, and on CPU
   **compute directly on the compressed blocks** (Q4_0/Q8_0/Q3_K/Q4_K/Q5_K/Q6_K) with no BF16
@@ -385,7 +385,7 @@ behind a model gallery, multi-model serving, the full OpenAI API surface, auth,
 ## Use it as a library (C API)
 
 Link `libvllm` and include [`include/vllm.h`](include/vllm.h): a flat, exception-free,
-llama.cpp-style C ABI (`VLLM_ABI_VERSION 17`, 35 exported functions) suitable for `dlopen` / FFI.
+llama.cpp-style C ABI (`VLLM_ABI_VERSION 18`, 36 exported functions) suitable for `dlopen` / FFI.
 
 ```c
 vllm_model_params mp = vllm_model_params_default();

docs/BENCHMARKS.md

@@ -435,6 +435,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), binding oracle too. **SIZE: 320x192/25f completes on GB10, 448x256 does not**; that render was REGISTER-conditioned, not prompted | Wall is the HOST VAE decode, not the pool: drain returns 0.11 GiB, byte-inert. 2 baselines UNRESOLVED (lock). A real-checkpoint PROMPTED render is OWED |
 | 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) |

docs/ENVIRONMENT.md

@@ -161,6 +161,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 |

docs/FEATURES.md

@@ -58,6 +58,7 @@ are our reading of their documented behavior, not measurements.
 | KV events (block create / evict publish) | ◐ no transport | ✅ | ☐ | ☐ |
 | Prefix-cache matching unit | ◐ resolver only | ✅ | ☐ | ☐ |
 | Compute directly on quantized blocks | ✅ | ☐ | ☐ | ✅ |
+| Scratch allocator keyed by device (two backends, one process) | ✅ since [#516](https://github.com/mudler/vllm.cpp/issues/516); a pool is bound to one backend and refuses any other, and a backend with no registered platform is refused rather than given another's residency cap | ✅ device is field 0 of the allocation handle | ✅ | ✅ |
 | Automatic memory sizing (no hand-tuned budget) | ☐ hand-typed block count | ☐ percent, hand-tuned | ☐ | ◐ |
 | Memory cap with a pre-flight error instead of an OOM | ☐ | ◐ KV pool only | ◐ | ☐ |
 
@@ -147,13 +148,20 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks.
 ### Standalone and non-registered lanes
 
 These run through dedicated forwards, not the `REGISTER_VLLM_MODEL` registry, so
-they sit outside the gated list above.
+they sit outside the gated list above. One caveat the LTX-2.5 row is too narrow
+to carry: its text tower's prompt tokenization mirrors upstream only while the
+checkpoint's tokenizer `post_processor` adds nothing. The shipped one is MEASURED
+empty, so this port's plain encode plus an explicit BOS prepend matches
+upstream's `add_special_tokens=True` today; a checkpoint with a non-empty
+`post_processor` would tokenize differently here, and `Ltx2TokenizeGemmaPrompt`
+in `ltx2_text_encoder.cpp` is the call that would have to change.
 
 | Lane | Tested checkpoint(s) | Correctness gate | Speed vs reference |
 |---|---|---|---|
 | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder FORWARD 15.90x of vLLM's whole TTFT (pin 46.02 ms), or 2.89x with opt-in `VT_WHISPER_ENC_FA2=1` (costs 3 near-tie divergences vs 0). Not a TTFT ratio. 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, VAEs+ENCODERS, conditioning, pipeline, quant loaders gated at reduced dims. Typed prompt to Gemma-4 to cross-attn, FIXTURE-gated. The 320x192/25f scene was register-conditioned; a prompted render is OWED | Family `ltx-2.5` via `ltx2-gen`. ~29 GB NVFP4/GB10, FP8 ~44 GB, +~24 GB tower. FP8, torchao and first-party NVFP4 all load. DiffVAE, LoRA, image conditioning refused AT THE ENGINE. Speed PENDING |
 | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) |
 | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on |
 | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending |
@@ -269,7 +277,7 @@ Build with `-DVLLM_CPP_VULKAN=ON`; off by default.
 | Multiple engines in one process (build, destroy, rebuild) | ✅ resident device state is owned by the weights, so a new engine never inherits a freed one's pointers | ✅ | ✅ | ✅ |
 | LoRA adapters | ☐ CPU brick only | ✅ | ✅ | ✅ |
 | Embedding / pooling endpoints | ◐ `/v1/embeddings` live (task=embed; score/rerank/classify pending) | ✅ | ✅ | ✅ |
-| OpenAI video generation `/v1/videos` (Sora shape) | ✅ `model`/`size`/`seconds` aliases + `GET /{id}/content`; `input_reference` and the `metadata` video/audio references condition the render | ◐ (vllm-omni, its own request shape) | ☐ | ☐ |
+| OpenAI video generation `/v1/videos` (Sora shape) | ✅ `model`/`size`/`seconds` aliases + `GET /{id}/content`; `input_reference` and `metadata` references condition the render; `--video-family` pins the family (default DETECT), `--video-extra K=V` carries family knobs | ◐ (vllm-omni, its own request shape) | ☐ | ☐ |
 | Flat C ABI for embedding in other languages | ✅ versioned | ☐ | ☐ | ✅ |
 
 #### C-ABI capability coverage <!-- abi-capability-table:begin -->
@@ -287,9 +295,9 @@ Build with `-DVLLM_CPP_VULKAN=ON`; off by default.
 | Custom logits processor | `vllm_logits_processor` | reachable |
 | Embeddings / pooling (task=embed) | `vllm_embed`, `vllm_embedding_result_free` (ABI v15; pooling checkpoints load via `vllm_engine_load`) | reachable |
 | Audio transcription (Parakeet ASR) | `vllm_transcribe`, `vllm_transcription_params_default`, `vllm_transcription_free` | reachable |
-| Video+audio generation (MiniMax-H3) | `vllm_video_engine_load`, `vllm_video_generate`, `vllm_video_result_free`, `vllm_video_mux_argv` | reachable |
+| Video+audio generation (MiniMax-H3, LTX-2.5) | `vllm_video_engine_load`, `vllm_video_generate`, `vllm_video_result_free`, `vllm_video_mux_argv`, `vllm_video_engine_family` (ABI v18 family registry) | reachable |
 | Explicit device selection (auto/cpu/cuda) | `device` field on `vllm_model_params` (ABI v14; 0=auto keeps the probe, explicit absent device fails loud) | reachable |
-| Run the OpenAI server (server as a thin ABI client) | `vllm_server_main` (ABI v17) | reachable |
+| Run the OpenAI server (server as a thin ABI client) | `vllm_server_main` (ABI v18) | reachable |
 | Multimodal input (image/audio/video) | none | embedder-unreachable | <!-- abi-capability-table:end -->
 
 ## Parallelism and scale-out