From 1d5faa8f16efbe5769d5741011fd43de90ab94e1 Mon Sep 17 00:00:00 2001 From: Daniel Maddern Date: Thu, 20 Aug 2026 22:08:52 +0700 Subject: [PATCH] Expand H3 runtime acceleration and validation --- Dockerfile.spark | 6 + FLASH4.md | 71 ++++ H3_LATENT_UPSCALER.md | 149 ++++++++ H3_PROMPT_GUIDE.md | 159 ++++++++ PLAN.md | 30 ++ README.md | 38 ++ ...e-864x480-141f-base12-sdpa-seed440420.json | 349 ++++++++++++++++++ compose.qwen38-vllm.yml | 46 +++ compose.spark-stack.yml | 84 +++++ compose.spark.yml | 4 +- pyproject.toml | 3 + src/h3_blackwell_runtime/attention.py | 30 +- src/h3_blackwell_runtime/latent_upscaler.py | 163 ++++++++ src/h3_blackwell_runtime/runtime.py | 43 ++- tests/test_fl2va_contracts.py | 21 +- tests/test_latent_upscaler.py | 52 +++ tools/poll_qwen38_health_spark.sh | 27 ++ tools/serve_hot_runtime.py | 16 +- tools/smoke_flash4.py | 50 +++ tools/smoke_qwen38_vllm_spark.sh | 8 + tools/spark_stack_spark.sh | 61 +++ tools/start_qwen38_vllm_spark.sh | 27 ++ tools/submit_h3_comfy_t2v.py | 91 +++++ tools/upscale_video_latent.py | 55 +++ 24 files changed, 1575 insertions(+), 8 deletions(-) create mode 100644 FLASH4.md create mode 100644 H3_LATENT_UPSCALER.md create mode 100644 H3_PROMPT_GUIDE.md create mode 100644 benchmarks/t2va-dialogue-864x480-141f-base12-sdpa-seed440420.json create mode 100644 compose.qwen38-vllm.yml create mode 100644 compose.spark-stack.yml create mode 100644 src/h3_blackwell_runtime/latent_upscaler.py create mode 100644 tests/test_latent_upscaler.py create mode 100644 tools/poll_qwen38_health_spark.sh create mode 100644 tools/smoke_flash4.py create mode 100644 tools/smoke_qwen38_vllm_spark.sh create mode 100644 tools/spark_stack_spark.sh create mode 100644 tools/start_qwen38_vllm_spark.sh create mode 100644 tools/submit_h3_comfy_t2v.py create mode 100644 tools/upscale_video_latent.py diff --git a/Dockerfile.spark b/Dockerfile.spark index 5201c69..28e9c34 100644 --- a/Dockerfile.spark +++ b/Dockerfile.spark @@ -15,6 +15,12 @@ RUN python -m pip install --no-cache-dir --no-deps comfy-kitchen==0.2.31 RUN python -m pip install --no-cache-dir "fastsafetensors>=0.1.10" +# Official CuTeDSL FlashAttention-4 beta with CUDA 13 Blackwell support. +RUN python -m pip install --no-cache-dir --pre \ + "flash-attn-4[cu13]==4.0.0b27" \ + "nvidia-cutlass-dsl[cu13]==4.6.2" \ + "quack-kernels==0.6.4" + RUN python -m pip uninstall -y pynvml \ && python -m pip install --no-cache-dir nvidia-ml-py diff --git a/FLASH4.md b/FLASH4.md new file mode 100644 index 0000000..910115e --- /dev/null +++ b/FLASH4.md @@ -0,0 +1,71 @@ +# FlashAttention-4 + +The hot runtime exposes the official FlashAttention-4 CuTeDSL Blackwell kernel +as the strict `flash4` attention backend. It does not silently fall back to +SageAttention or SDPA. + +## Versions + +- `flash-attn-4==4.0.0b27` +- `nvidia-cutlass-dsl==4.6.2` with the CUDA 13 extra +- `quack-kernels==0.6.4` +- Upstream: +- License: BSD-3-Clause + +The package is a beta release. The Spark image pins these versions rather than +tracking the latest prerelease. + +## API + +Select Flash4 per request: + +```json +{ + "prompt": "A simple cinematic scene.", + "output": "/output/flash4-example.mp4", + "width": 960, + "height": 544, + "frames": 124, + "steps": 12, + "attention": "flash4" +} +``` + +H3's main attention blocks call Flash4 in native +`[batch, sequence, heads, head_dim]` layout. The adapter requires CUDA, matching +FP16/BF16 Q/K/V tensors, matching shapes, and H3's head dimension of 128. Any +violation raises an error instead of changing kernels. + +## GB10 Validation + +Flash4 dispatches its SM120 forward implementation on the GB10, which PyTorch +reports as compute capability 12.1. + +| Validation | Result | +| --- | --- | +| BF16 `[1,257,8,128]` vs SDPA | max abs `0.001953125`, mean abs `5.04e-05` | +| BF16 `[1,20480,56,128]` vs SDPA | max abs `0.000244141`, mean abs `1.96e-08` | +| Tiny joint AV H3 smoke | passed, no fallback | +| 28 runtime contracts | passed | + +Matched base 12-step talking benchmark at `960x544x124`, seed `440410`: + +| Backend | Sampling | Relative to Flash4 | +| --- | ---: | ---: | +| Sage2 | `114.67s` | `-19.4%` | +| Flash4 | `142.23s` | baseline | +| PyTorch SDPA | `144.40s` | `+1.5%` | + +Flash4 is therefore close to SDPA and materially slower than Sage2 for this H3 +workload on GB10. Its value is as another exact-style attention and audio-quality +comparison, not as the current speed default. + +The matched Flash4 proof contains H.264 video and stereo 32 kHz AAC audio: + +```text +/home/daniel/StoryStudioAssets/H3-output/h3-native-upscale-t2v/talking-woman-base12-flash4-960x544.mp4 +``` + +The current CuTeDSL package emits a one-time `AuxData` JIT argument warning. +Compilation and inference still complete successfully, but the warning should +be rechecked when upgrading FlashAttention-4 or CUTLASS DSL. diff --git a/H3_LATENT_UPSCALER.md b/H3_LATENT_UPSCALER.md new file mode 100644 index 0000000..4160fc0 --- /dev/null +++ b/H3_LATENT_UPSCALER.md @@ -0,0 +1,149 @@ +# H3 Native Latent Upscaler + +The direct runtime can apply LBH-123-AI's learned 3D upscaler to a saved, +normalized MiniMax H3 video latent without a ComfyUI dependency. This is an +experimental spatial upscale. It preserves latent time and does not interpolate +frames. + +Upstream implementation and checkpoint: + +- +- +- Revision: `13ccf95d85d120bdbc92c05b1247a6e147bf54bf` +- FP16 SHA-256: `043e5a48e161610ef6c3ea974645220354d06fa618abca15f76d084812eb55c2` +- License: Apache-2.0 + +## Usage + +The hot service supports the complete generation, latent upscale, video/audio +decode, and mux process in one request: + +```json +{ + "prompt": "A simple cinematic scene.", + "output": "/output/h3-native-upscale/example.mp4", + "width": 960, + "height": 544, + "frames": 124, + "turbo": "4step", + "upscale": 2.0, + "mux_audio": true, + "keep_intermediates": false +} +``` + +The response reports `source_width`, `source_height`, final `width` and +`height`, and a `latent_upscaled` timing stage. H3 audio bypasses the spatial +upscaler and is decoded and muxed normally. + +Only the final muxed MP4 is retained by default. Set `keep_intermediates` to +`true` only when a separate WAV and video-only MP4 are needed for diagnostics; +`save_latent` remains a separate opt-in field. + +For offline diagnostics, save the normalized latent during generation by +setting `save_latent`, then run: + +```bash +python tools/upscale_video_latent.py \ + --latent /output/source.pt \ + --model /upscaler/minimax_h3_latent_upscaler_3d_fp16.safetensors \ + --output /output/source-h3latent2x.pt \ + --scale 2 \ + --precision float16 + +python tools/decode_video_latent.py \ + --latent /output/source-h3latent2x.pt \ + --output /output/source-h3latent2x.mp4 \ + --vae-dtype float16 \ + --vae-tile-size 256 +``` + +The upscaled state retains other source metadata and any audio latent. Its +`width` and `height` fields are updated from the resulting H3 latent dimensions. + +## GB10 Validation + +Validated on 2026-08-20 with the Posly hero image, Turbo 4-step v1.1, Sage2, +seed `440408`, and 124 frames: + +| Stage | Result | +| --- | --- | +| Source | `960x544`, latent `[1,24,37,34,60]` | +| Learned 3D upscale | `10.94s`, latent `[1,24,37,68,120]` | +| Checkpoint load | `1.36s` | +| Tiled FP16 decode command | `137.67s`, `1920x1088` H.264 | +| Source vs downscaled result | SSIM `0.924986` | + +The learned stage adds visible edge, face, and food detail. It also produces a +harsher illustrated texture, ringing/chromatic edges in some high-contrast +areas, and small facial changes. It is not fidelity-safe for protected UI, +text, QR codes, or identity-critical material without visual review. + +The upstream example follows latent upscaling with a second target-resolution +H3 sampling pass using sigmas `0.9035, 0.6316, 0.3158, 0.0`. That refinement +was deliberately excluded from this first test so the upscaler itself could be +evaluated. It adds generation cost and another opportunity to alter content. + +## Proofs + +Spark output directory: + +```text +/home/daniel/StoryStudioAssets/H3-output/posly/upscale-h3-native/ +``` + +Key files: + +- `01-posly-hero-turbo4-source-960x544.mp4` +- `01-posly-hero-turbo4-h3latent2x-1920x1088.mp4` +- `source-frame60-right-crop-lanczos.png` +- `h3latent2x-frame60-right-crop.png` + +This result is promising for cinematic footage but does not replace the proven +LTX 2.3 chain yet. Run a matched LTX 2.3 comparison and a carefully bounded H3 +refinement test before selecting a default production path. + +## Integrated T2V Validation + +The single-request path was validated with a matched 5.17-second talking-head +T2V pair at seed `440410`. Both requests sampled at `960x544`, applied the +resident 2x upscaler, decoded at `1920x1088`, and muxed native H3 audio into the +only retained artifact. + +| Variant | Sampling | Latent upscale | VAE decode | Wall time | +| --- | ---: | ---: | ---: | ---: | +| Turbo 4-step v1.1 | `56.54s` | `7.52s` | `82.74s` | `163.02s` | +| Turbo 8-step v1.0 | `113.41s` | `7.28s` | `83.69s` | `220.06s` | +| Base 12-step beta/RES | `114.67s` | `7.26s` | `83.34s` | `221.66s` | + +Both outputs contain H.264 video at 24 fps and stereo AAC audio at 32 kHz. The +4-step audio measured `-31.4 dB` mean / `-10.9 dB` peak; the 8-step audio +measured `-23.0 dB` mean / `-5.0 dB` peak; and the base 12-step audio measured +`-22.4 dB` mean / `-5.2 dB` peak. + +```text +/home/daniel/StoryStudioAssets/H3-output/h3-native-upscale-t2v/talking-woman-h3native2x-1920x1088.mp4 +/home/daniel/StoryStudioAssets/H3-output/h3-native-upscale-t2v/talking-woman-turbo8-h3native2x-1920x1088.mp4 +/home/daniel/StoryStudioAssets/H3-output/h3-native-upscale-t2v/talking-woman-base12-h3native2x-1920x1088.mp4 +``` + +### Attention Audio Isolation + +A matched base 12-step run used direct PyTorch SDPA instead of Sage2. Upscaling +was disabled, but that does not affect the audio comparison because the learned +upscaler runs after joint AV sampling and only receives the separated video +latent. + +| Attention | Sampling | Final size | Audio level | Wall time | +| --- | ---: | ---: | ---: | ---: | +| Sage2 | `114.67s` | `1920x1088` after post-sample 2x | `-22.4 dB` mean / `-5.2 dB` peak | `221.66s` | +| PyTorch SDPA | `144.40s` | `960x544`, no upscale | `-22.5 dB` mean / `-2.3 dB` peak | `171.21s` | + +The SDPA comparison file is: + +```text +/home/daniel/StoryStudioAssets/H3-output/h3-native-upscale-t2v/talking-woman-base12-sdpa-960x544.mp4 +``` + +Signal levels alone cannot establish dialogue quality or lip sync; the matched +clips require listening review to determine whether Sage2 is the cause. diff --git a/H3_PROMPT_GUIDE.md b/H3_PROMPT_GUIDE.md new file mode 100644 index 0000000..022cd5b --- /dev/null +++ b/H3_PROMPT_GUIDE.md @@ -0,0 +1,159 @@ +# H3 Prompt Guide + +Project guidance for MiniMax H3 audiovisual prompts. This document combines +MiniMax's official prompt format with observations from our direct-runtime and +ComfyUI tests. + +## Core Principle + +Treat an H3 prompt as an audiovisual timeline, not as a loose prose image +description. For T2VA and keyframe modes, use these exact fields in this order: + +```text +integrated_multimodal_description: ... + +overall_soundscape: ... + +non_diegetic_music: ... +``` + +- `integrated_multimodal_description` contains shots, visible actions, + speakers, exact dialogue, singing, and synchronized diegetic events. +- `overall_soundscape` contains ambience, Foley, and non-verbal human sounds + such as breathing, laughter, and panting. +- `non_diegetic_music` contains only music heard by the audience rather than + the characters. Use `N/A` when no score is wanted. + +T2VA begins directly with these fields. I2VA, FL2VA, and L2VA add the official +picture-alignment instruction before them. Ref2VA uses a different six-section +format and should follow the official reference guide. + +## Dialogue Syntax + +Give every speaking character a stable speaker ID such as `(S1)` or `(S2)`. +Describe the speaker, voice, action, and delivery outside the dialogue block. +Place only the language tag and exact spoken words inside ``: + +```text +The woman with a low, clear voice and measured pace (S1) says in a flat, +matter-of-fact delivery: [English] The meeting starts at three. +``` + +Rules: + +- Preserve the dialogue wording and punctuation exactly. +- Specify the language explicitly, such as `[English]`. +- Describe pitch, timbre, pace, volume, accent, and emotional restraint only + when useful. +- Keep the same speaker ID across shots. +- Use a compound ID such as `(S1,S2)` only when speakers vocalize together. +- Do not put dialogue in double quotation marks. H3 reserves double quotes for + visible on-screen text such as signs, labels, and subtitles. +- For voiceover, use `says in an off-screen voiceover` and state that the + visible character's lips remain completely closed. + +## Lip Sync And End State + +Describe speech as a physical action and specify what happens when it ends: + +```text +Her mouth movements naturally synchronize with each spoken word. Immediately +after the final word, her lips meet and her jaw ceases speaking motion. She +remains silent through the final frame. +``` + +This is especially important in short clips. Without an explicit end state, H3 +may continue mouth motion or add a non-verbal vocal reaction after the line. + +For multiple speakers, state whose lips move and whose remain still during each +line and pause. + +## Avoiding Unwanted Laughter + +Laughter is a non-verbal human sound, so control it in `overall_soundscape`, not +inside the dialogue block: + +```text +overall_soundscape: Quiet, dry indoor room tone with a faint ventilation hum. +No laughter, chuckling, giggling, sighing, gasping, audible breathing, filler +sounds, audience reaction, or other voices. +``` + +Also reinforce the visible end state in the shot description: closed lips, +stopped jaw motion, neutral expression, and silence through the final frame. + +These exclusions are project guidance, not a guaranteed negative-prompt +mechanism. H3 jointly generates video and audio, so unwanted sounds remain +probabilistic. + +Avoid positive cues that can conflict with the exclusion. In our first dialogue +test, `She smiles slightly` repeatedly became laughter across Sage2, PyTorch +SDPA, and FlashAttention-4. Since the behavior survived attention changes, the +prompt was the stronger common cause. Use observable neutral behavior instead: + +```text +She maintains a neutral, composed, closed-mouth expression. +``` + +## Short Dialogue Tests + +For a five-second diagnostic clip: + +- Use one visible speaker and one short sentence. +- Use one static shot with no cuts. +- Choose emotionally neutral wording. +- Request a flat or matter-of-fact delivery. +- Specify the post-dialogue closed-mouth state. +- Permit only a simple, positive acoustic bed such as dry room tone. +- Set `non_diegetic_music: N/A`. + +Do not use `overall_soundscape: N/A` for a speaking test. The official guide +reserves it for complete silence throughout the video. + +## Reusable Neutral T2VA Prompt + +```text +integrated_multimodal_description: [Shot 1] Live-action, cinematic, a static +medium close-up frames exactly one adult woman seated alone in an otherwise +empty, quiet cafe. Warm natural window light falls evenly across her face. She +maintains a neutral, composed, closed-mouth expression and looks steadily +toward the camera. The woman with a low, clear voice and slow, even speaking +pace (S1) physically speaks once in a flat, matter-of-fact delivery: +[English] The meeting starts at three. Her mouth movements naturally +synchronize with each spoken word. Immediately after the final word, her lips +meet and her jaw ceases speaking motion. She remains silent and maintains the +same neutral expression through the final frame. The camera remains completely +static with no cuts. + +overall_soundscape: Quiet, dry indoor room tone with a faint ventilation hum. +No laughter, chuckling, giggling, smiling vocalization, sighing, gasping, +audible breathing, filler sounds, audience reaction, or other voices. + +non_diegetic_music: N/A +``` + +## Prompt Checklist + +- Correct task mode selected. +- Exact field names and order preserved. +- Every shot has observable visual and audible events. +- Every speaker has a stable ID. +- Dialogue uses `[Language] exact words`. +- Dialogue is not enclosed in double quotes. +- Delivery and voice are outside ``. +- Lip motion and post-speech closure are explicit. +- Ambience and non-verbal sounds are in `overall_soundscape`. +- Music is isolated in `non_diegetic_music`. +- Script length fits the requested duration. +- Positive emotional cues do not conflict with audio exclusions. + +## Official Sources + +- Base prompt guide: + +- Reference prompt guide: + +- Official prompt-writing skill: + +- Official model repository: + diff --git a/PLAN.md b/PLAN.md index 35a9ec8..e6d53ff 100644 --- a/PLAN.md +++ b/PLAN.md @@ -109,6 +109,36 @@ DGX Spark caveat: - Record attention, GEMM, communication, VAE, and end-to-end timings separately. - Treat SageAttention3 as an experimental quality-gated kernel for H3. +## 2026-08-15 Handoff + +Current quality baseline: + +- `sage2` is the correct visual path for the 960x544, 124-frame, 12-step cat prompt at seed `440407`. +- Fresh Spark hot-runtime Sage2/Vortex A/B asset: `\\192.168.1.162\StoryStudioAssets\H3-output\h3-blackwell-runtime\cat-sage2-vortex-5s-440407-ab.mp4`. +- `sol_attn` is faster but visually wrong for the same prompt/seed. Keep it experimental until localized and quality-gated. +- Sol-native BSHD layout is exact versus the old Sol path, so the visual issue is likely Sol's approximation/configuration, not the layout removal. + +Tomorrow priority: + +1. Diagnose `sol_attn` quality failure. + - Save per-step latents for `sage2` and `sol_attn` with identical prompt, seed, dimensions, and sampler settings. + - Measure per-step max/mean/RMSE/cosine divergence. + - Try hybrid schedules: early steps `sage2`, later steps `sol_attn`; also try Sage2 every Nth step. + - Sweep Sol params only after the divergence point is known: `H3_SOL_TAU`, `H3_SOL_THRESH_TYPE`, `H3_SOL_INT8_QK`, `H3_SOL_INT8_PV`. +2. Optimize the quality path (`sage2`). + - Use `sage2` as visual/numerical reference. + - Attack Q/K/V HND layout and Q/K RMS+RoPE first; a fused Q/K RMS+RoPE + Sage HND layout kernel is the most plausible exact win. + - Keep Sage3 experimental; previous measurements were slower and visually/numerically different. +3. Add durable hot-runtime request logging. + - Log `generate_start`, `generate_complete`, and `generate_failed` to stdout with output path, seed, attention, dimensions, and timings. + - This prevents losing timing data when a client disconnects before reading the synchronous response. +4. Check LTX 2.5 downloads. + - Detached Spark container: `ltx25-nvfp4-download`. + - Expected files under `/home/daniel/aeon-spark-test/h3/comfy-models`: + - `diffusion_models/ltx-2.5-22b-distilled-transformer-nvfp4.safetensors` + - `text_encoders/gemma4-12b-with-proj-ltx-2.5-comfy-int8-convrot.safetensors` + - Official `Lightricks/LTX-2.5` does not list a matching NVFP4 text encoder; the int8 ConvRot Gemma4 encoder is the matching low-VRAM official text encoder. + ## 2026-08-13 VAE Debug Handoff Resolved on `vae-decode-optimization`: direct VAE temporal overlap constants now match upstream, audio decode/mux is implemented, and Comfy-equivalent FP16 video VAE is the default runtime path. The 960x544x124 cat benchmark now matches warm Comfy performance: Comfy `150.26s`, direct `149.304s` after text conditioning, direct VAE decode `25.085s`. diff --git a/README.md b/README.md index 51c4baf..a36be43 100644 --- a/README.md +++ b/README.md @@ -2,6 +2,10 @@ Direct MiniMax H3 Ref2VA runtime research project. ComfyUI is the checkpoint and correctness oracle, not the target runtime. +See [`H3_PROMPT_GUIDE.md`](H3_PROMPT_GUIDE.md) for the project's H3 audiovisual +prompt structure, dialogue syntax, lip-sync controls, soundscape rules, tested +failure modes, and reusable templates. + ## First Gate Inspect the mounted H3 NVFP4 safetensors headers before designing an importer: @@ -60,15 +64,35 @@ docker compose -f compose.spark.yml up -d h3-hot-runtime Use `GET /ready` to confirm resident model readiness. Use `POST /generate` with JSON fields like `prompt`, `output`, `width`, `height`, `frames`, `steps`, `seed`, and optional `attention`. Supported request-level attention values are reported by `/ready`; switching attention does not reload model weights. The hot image includes Sage2, forced cuDNN SDPA, and Comfy Kitchen INT8 attention. Sage2 is the default based on the 960x544x124 GB10 benchmark and the existing parity baseline. +Set `"upscale": 2.0` to apply the resident H3-native learned latent upscaler +between sampling and VAE decode in the same request. `width` and `height` remain +the low-resolution sampling canvas; the response reports both source and final +dimensions. Video is upscaled before decoding, while H3's jointly generated +audio latent follows the normal decode and mux path unchanged. Omit `upscale`, +set it to `null`, or set it to `1` to disable this stage. + +Successful requests retain only the final MP4 by default. `save_latent` is +opt-in. Set `"keep_intermediates": true` only for diagnostics that require the +separate WAV and video-only MP4; otherwise both are removed after muxing. + The Spark hot service also keeps the official FL2VA Turbo adapters resident. Set `turbo` to `"4step"` for v1.1 768p (shift 6/3) or `"8step"` for v1.0 (shift 12/3). The matching step count is selected by default and enforced when `steps` is supplied. Set `turbo` to `null` or `"none"` for the base beta/RES path. Turbo uses its separate uniform training-Euler schedule and cannot be combined with denoiser caching. See [`TURBO.md`](TURBO.md) for artifact hashes, implementation details, API examples, validation evidence, and matched GB10 performance results. +The optional H3-native learned 3D upscaler operates directly on latents saved by +the hot runtime. See [`H3_LATENT_UPSCALER.md`](H3_LATENT_UPSCALER.md) for the +pinned checkpoint, standalone command, GB10 benchmark, proof paths, and current +quality limitations. + Exact memory/lifetime options: - `attention: "kj_head_sliced"` slices attention heads and runs the slice backend from `H3_HEAD_SLICE_BACKEND` (`sage2` by default) with `H3_HEAD_SLICE_SIZE` heads per slice (`8` by default). - `attention: "cudnn_sdpa"` forces cuDNN SDPA with no fallback to another PyTorch kernel. - `attention: "ck_int8"` uses Comfy Kitchen's approximate INT8 Q/K/V attention kernel. +- `attention: "flash4"` uses the pinned official FlashAttention-4 CuTeDSL + Blackwell kernel with strict validation and no fallback. See + [`FLASH4.md`](FLASH4.md) for versions, numerical validation, limitations, and + the matched GB10 benchmark. - `attention: "sol_attn"` routes eligible H3 attention calls through the pinned ComfyUI Sol-Attn Triton kernel vendored into the Spark image. Configure with `H3_SOL_TAU` (`1.3`), `H3_SOL_MIN_TOKENS` (`4096`), `H3_SOL_THRESH_TYPE` (`diag`), `H3_SOL_INT8_QK`, `H3_SOL_INT8_PV`, `H3_SOL_FALLBACK` (`sage2`), and `H3_SOL_STRICT`. - `--mlp-chunks N` on `tools/serve_hot_runtime.py` or `tools/direct_t2v_preview.py` chunks H3 SwiGLU rows exactly to reduce peak activation memory. Default is `1` (disabled). @@ -77,3 +101,17 @@ Approximate cache options are opt-in and must be quality-gated per prompt: - `cache_mode: "easycache"` reuses cached denoised deltas while cumulative latent input change stays below `cache_threshold`. - `cache_mode: "h3_cache"` reuses cached denoised deltas when the current per-step latent input change is below `cache_threshold`. - Both modes accept `cache_start_percent`, `cache_end_percent`, and `cache_subsample_factor` in `POST /generate`; the CLI exposes equivalent `--cache-*` flags. + +## Project TODO + +- [ ] Develop an H3-specific attention backend optimized for the model's actual + GB10 tensor shapes. Validate numerical behavior and subjective output quality + against SDPA, then benchmark it with the existing two-person dialogue matrix. +- [ ] Isolate and fix the H3 startup audio artifact in the latent/AudioVAE path. + The artifact reproduced across every attention and Turbo test in the current + two-person dialogue matrix, and matched direct-runtime and Comfy SDPA outputs + have effectively identical audio. Capture affected audio latents and lossless + pre-AAC PCM, separate first-latent generation from AudioVAE boundary behavior, + and build a selective model-path fix. Do not use unconditional output trimming + or fading because valid audio can begin immediately, as demonstrated by the + nightclub music-onset test. diff --git a/benchmarks/t2va-dialogue-864x480-141f-base12-sdpa-seed440420.json b/benchmarks/t2va-dialogue-864x480-141f-base12-sdpa-seed440420.json new file mode 100644 index 0000000..b826dcf --- /dev/null +++ b/benchmarks/t2va-dialogue-864x480-141f-base12-sdpa-seed440420.json @@ -0,0 +1,349 @@ +{ + "name": "t2va-dialogue-864x480-141f-base12-sdpa-seed440420", + "measured_at": "2026-08-20", + "measurement": "Warm resident direct-runtime request including conditioning, sampling, video/audio decode, encoding, and mux", + "hardware": "NVIDIA GB10", + "runtime_image": "sha256:d842225c9c5036647ee5cab1fae2e093357cc8041bd8617e33a82c81d0bcc874", + "torch": "2.9.1+cu130", + "model": "minimax_h3_fl2va_pruned_nvfp4.safetensors", + "task": "t2va", + "resolution": [864, 480], + "frames": 141, + "fps": 24, + "duration_seconds": 5.875, + "steps": 12, + "scheduler": "beta", + "sampler": "res_multistep", + "seed": 440420, + "attention": "sdpa", + "turbo": null, + "upscale": null, + "mux_audio": true, + "keep_intermediates": false, + "vae_dtype": "float16", + "vae_tile_size": 256, + "mlp_chunks": 1, + "prompt": "integrated_multimodal_description: [Shot 1] Live-action, cinematic, a static medium two-shot frames exactly two adults seated across from each other at a small table in a quiet, otherwise empty meeting room. A composed adult man sits on the left and a composed adult woman sits on the right. Both maintain neutral, attentive expressions. The man with a low, clear baritone voice and measured speaking pace (S1) looks toward the woman and says in a calm, matter-of-fact delivery: [English] The north entrance closes at six. During S1's line, only his lips and jaw move; the woman's lips remain completely closed. Immediately after his final word, his lips meet and his jaw ceases speaking motion. After a brief silent pause, the woman with a clear alto voice and measured speaking pace (S2) looks toward the man and replies in a calm, matter-of-fact delivery: [English] Then we should leave by five thirty. During S2's line, only her lips and jaw move; the man's lips remain completely closed. Immediately after her final word, her lips meet and her jaw ceases speaking motion. Both remain silent with neutral, closed-mouth expressions through the final frame. There is no overlapping speech. The camera remains completely static with no cuts.\n\noverall_soundscape: Quiet, dry indoor room tone with a faint ventilation hum. Only S1 and S2 are audible, one at a time. No laughter, chuckling, giggling, smiling vocalization, sighing, gasping, audible breathing, filler sounds, audience reaction, narration, or other voices.\n\nnon_diegetic_music: N/A", + "timings_seconds": { + "latents_initialized": 0.018280818010680377, + "text_conditioned": 3.531974215002265, + "sampled": 125.30079188900709, + "vae_decoded": 23.84665890400356, + "pixels_cpu": 0.06607734999852255, + "raw_write": 2.75871228199685, + "video_encode": 0.39305945999512915, + "audio_decoded": 0.3400787889986532, + "audio_raw_write": 0.0008342489891219884, + "audio_encode": 0.034553833000245504, + "mux": 0.19705491400964092, + "request": 156.48807670301176, + "wall": 156.56338787100685 + }, + "throughput": { + "sampling_frames_per_second": 1.1252921699401506, + "wall_frames_per_second": 0.9005936950992043, + "sampling_realtime_factor": 21.327794364086314, + "wall_realtime_factor": 26.649087297192655 + }, + "cache": { + "mode": null, + "threshold": 0.0, + "skipped_steps": 0, + "rates": [] + }, + "output": { + "linux_path": "/home/daniel/StoryStudioAssets/H3-output/h3-baselines/dialogue-two-character-864x480-141f-base12-sdpa-seed440420.mp4", + "unc_path": "\\\\192.168.1.162\\StoryStudioAssets\\H3-output\\h3-baselines\\dialogue-two-character-864x480-141f-base12-sdpa-seed440420.mp4", + "sha256": "2eb61334978b42bcec00ec68449f3f39139d9e24ef5ba70ef2240dc8bf7116e0", + "size_bytes": 429062, + "container_bit_rate": 584254, + "duration_seconds": 5.875, + "video": { + "codec": "h264", + "profile": "High", + "pixel_format": "yuv420p", + "width": 864, + "height": 480, + "fps": 24, + "frames": 141, + "bit_rate": 439001 + }, + "audio": { + "codec": "aac", + "profile": "LC", + "sample_rate": 32000, + "channels": 2, + "channel_layout": "stereo", + "bit_rate": 136335, + "mean_volume_db": -25.8, + "max_volume_db": -4.3, + "integrated_loudness_lufs": -21.18, + "true_peak_dbtp": -4.34, + "loudness_range_lu": 4.10, + "loudness_threshold_lufs": -33.94 + } + }, + "attention_backend_matrix": { + "reference_backend": "sdpa", + "comparison_note": "PSNR measures encoded output drift from SDPA, not subjective quality. Audio is effectively identical across all backends at greater than 170 dB decoded-PCM PSNR.", + "runs": { + "sdpa": { + "sampled_seconds": 125.30079188900709, + "request_seconds": 156.48807670301176, + "wall_seconds": 156.56338787100685, + "sampling_speedup_percent_vs_sdpa": 0.0, + "output_file": "dialogue-two-character-864x480-141f-base12-sdpa-seed440420.mp4", + "sha256": "2eb61334978b42bcec00ec68449f3f39139d9e24ef5ba70ef2240dc8bf7116e0", + "size_bytes": 429062 + }, + "sage2": { + "sampled_seconds": 101.26120180900034, + "request_seconds": 132.59474182801205, + "wall_seconds": 132.66185540499282, + "sampling_speedup_percent_vs_sdpa": 19.185505, + "video_psnr_average_db_vs_sdpa": 22.4225, + "audio_psnr_db_vs_sdpa": [170.933, 170.934], + "output_file": "dialogue-two-character-864x480-141f-base12-sage2-seed440420.mp4", + "sha256": "c3e0691586c4101987f240c35a79637965af75088731108037cd084be945ba9d", + "size_bytes": 433450 + }, + "cudnn_sdpa": { + "sampled_seconds": 125.72722270998929, + "request_seconds": 157.14237468996726, + "wall_seconds": 157.2081264879962, + "sampling_speedup_percent_vs_sdpa": -0.340326, + "video_psnr_average_db_vs_sdpa": 22.093143, + "audio_psnr_db_vs_sdpa": [171.038, 171.01], + "output_file": "dialogue-two-character-864x480-141f-base12-cudnn_sdpa-seed440420.mp4", + "sha256": "cd17833c12a99cd71d1c0a9137a1dd9feba03943980d6eb8fff296cd9f9529a6", + "size_bytes": 429636 + }, + "ck_int8": { + "sampled_seconds": 104.05008868798905, + "request_seconds": 135.57199566195777, + "wall_seconds": 135.62453711099806, + "sampling_speedup_percent_vs_sdpa": 16.959752, + "video_psnr_average_db_vs_sdpa": 22.989999, + "audio_psnr_db_vs_sdpa": [170.981, 170.956], + "output_file": "dialogue-two-character-864x480-141f-base12-ck_int8-seed440420.mp4", + "sha256": "a54f9f3f2ba351d31b0b807646aea2740aae7c13c03d4eb7d44182aab3cf5618", + "size_bytes": 423190 + }, + "sage3": { + "sampled_seconds": 109.8837421490025, + "request_seconds": 142.34848491402227, + "wall_seconds": 142.41395058500348, + "sampling_speedup_percent_vs_sdpa": 12.304032, + "video_psnr_average_db_vs_sdpa": 19.63725, + "audio_psnr_db_vs_sdpa": [172.214, 172.06], + "output_file": "dialogue-two-character-864x480-141f-base12-sage3-seed440420.mp4", + "sha256": "3763ca4b79399037d997d5a37c795fdaf233aca88fdb90e122bf25450c75dddd", + "size_bytes": 596651 + }, + "sage3_mean": { + "sampled_seconds": 109.88214365398744, + "request_seconds": 141.45796999098093, + "wall_seconds": 141.51303354099218, + "sampling_speedup_percent_vs_sdpa": 12.305308, + "video_psnr_average_db_vs_sdpa": 19.63725, + "audio_psnr_db_vs_sdpa": [172.214, 172.06], + "output_file": "dialogue-two-character-864x480-141f-base12-sage3_mean-seed440420.mp4", + "sha256": "3763ca4b79399037d997d5a37c795fdaf233aca88fdb90e122bf25450c75dddd", + "size_bytes": 596651, + "identical_output_to": "sage3" + }, + "kj_sage_cuda": { + "sampled_seconds": 113.68463072601298, + "request_seconds": 145.39383504101716, + "wall_seconds": 145.45757150900317, + "sampling_speedup_percent_vs_sdpa": 9.270621, + "video_psnr_average_db_vs_sdpa": 22.83642, + "audio_psnr_db_vs_sdpa": [171.042, 171.025], + "output_file": "dialogue-two-character-864x480-141f-base12-kj_sage_cuda-seed440420.mp4", + "sha256": "d85aaca677d2c678604b0aa2aba9bbf4921934bfc2970c20d0498b1762373ab5", + "size_bytes": 427568 + }, + "kj_sage_triton": { + "sampled_seconds": 118.28390433100867, + "request_seconds": 150.37281119299587, + "wall_seconds": 150.43365036998875, + "sampling_speedup_percent_vs_sdpa": 5.600034, + "video_psnr_average_db_vs_sdpa": 22.086796, + "audio_psnr_db_vs_sdpa": [170.149, 170.084], + "output_file": "dialogue-two-character-864x480-141f-base12-kj_sage_triton-seed440420.mp4", + "sha256": "9fe63aafa51aae8ecc0940ba75fbc13a978a50091fd8f0c336c2bb04e3658a54", + "size_bytes": 430038 + }, + "kj_sage_fp8": { + "sampled_seconds": 102.6086639950081, + "request_seconds": 134.1615118059999, + "wall_seconds": 134.2236167689989, + "sampling_speedup_percent_vs_sdpa": 18.110123, + "video_psnr_average_db_vs_sdpa": 22.814833, + "audio_psnr_db_vs_sdpa": [170.719, 170.692], + "output_file": "dialogue-two-character-864x480-141f-base12-kj_sage_fp8-seed440420.mp4", + "sha256": "8c5bc2a8bf7e496c2108194479fe60e67883dea0ab50e95f3ab4bac916274816", + "size_bytes": 420189 + }, + "kj_sage_fp8pp": { + "sampled_seconds": 103.24360375599645, + "request_seconds": 134.77945204998832, + "wall_seconds": 134.84802630099875, + "sampling_speedup_percent_vs_sdpa": 17.603391, + "video_psnr_average_db_vs_sdpa": 22.296461, + "audio_psnr_db_vs_sdpa": [170.615, 170.559], + "output_file": "dialogue-two-character-864x480-141f-base12-kj_sage_fp8pp-seed440420.mp4", + "sha256": "71643c7163e5dc16d64ac0c5d567bdd846f7944471eb0a97fd54f9a1990c24e9", + "size_bytes": 433672 + }, + "kj_head_sliced": { + "sampled_seconds": 103.97128305501246, + "request_seconds": 135.67128342803335, + "wall_seconds": 135.71258915099315, + "sampling_speedup_percent_vs_sdpa": 17.022645, + "video_psnr_average_db_vs_sdpa": 21.679726, + "audio_psnr_db_vs_sdpa": [171.235, 171.237], + "output_file": "dialogue-two-character-864x480-141f-base12-kj_head_sliced-seed440420.mp4", + "sha256": "48ee4da6c0a91f18714973ca23e26f1844b302024c8111946772cbc9a4a4dd83", + "size_bytes": 421728 + }, + "sol_attn": { + "sampled_seconds": 96.45773334198748, + "request_seconds": 128.10418028896675, + "wall_seconds": 128.17140150099294, + "sampling_speedup_percent_vs_sdpa": 23.019055, + "video_psnr_average_db_vs_sdpa": 16.653003, + "audio_psnr_db_vs_sdpa": [171.239, 171.268], + "output_file": "dialogue-two-character-864x480-141f-base12-sol_attn-seed440420.mp4", + "sha256": "8b999630210692301549d11d20d9319a6af5873a60d4cd735035e2ed79c469d6", + "size_bytes": 570667, + "note": "Experimental sparse Triton dispatch; eligible calls use sol_attn and shorter calls fall back." + }, + "flash4": { + "sampled_seconds": 128.5334454880067, + "request_seconds": 160.91555801399227, + "wall_seconds": 160.96679471699463, + "sampling_speedup_percent_vs_sdpa": -2.579915, + "video_psnr_average_db_vs_sdpa": 21.238297, + "audio_psnr_db_vs_sdpa": [170.768, 170.628], + "output_file": "dialogue-two-character-864x480-141f-base12-flash4-seed440420.mp4", + "sha256": "cfb8a046cb79f8448988e7392ca748c0a08a2a034a43f6ab55a9058faa1a72ef", + "size_bytes": 449687 + } + } + }, + "turbo_sage2_matrix": { + "reference": "Base 12-step Sage2 from attention_backend_matrix", + "runs": { + "4step": { + "adapter": "Official FL2VA Turbo 4-step v1.1 768p", + "steps": 4, + "attention": "sage2", + "sampled_seconds": 51.661242722009774, + "request_seconds": 83.01122188399313, + "wall_seconds": 83.06598260099418, + "sampling_speedup_percent_vs_base_sage2": 48.982195, + "wall_speedup_percent_vs_base_sage2": 37.38518, + "output_file": "dialogue-two-character-864x480-141f-turbo4step-sage2-seed440420.mp4", + "sha256": "b02c29a8732d88c7dcf8e85462afcbe3aa83328541ce1647995e7d0d4188744a", + "size_bytes": 455325, + "container_bit_rate": 620017, + "audio": { + "integrated_loudness_lufs": -31.78, + "true_peak_dbtp": -14.55, + "first_0_25_seconds_peak_dbfs": -39.394894, + "first_0_25_seconds_rms_dbfs": -57.733654, + "first_sample_dbfs": [-54.110104, -50.553467] + } + }, + "8step": { + "adapter": "Official FL2VA Turbo 8-step v1.0", + "steps": 8, + "attention": "sage2", + "sampled_seconds": 102.70967868898879, + "request_seconds": 134.31820170898573, + "wall_seconds": 134.36813215899747, + "sampling_speedup_percent_vs_base_sage2": -1.430436, + "wall_speedup_percent_vs_base_sage2": -1.286185, + "output_file": "dialogue-two-character-864x480-141f-turbo8step-sage2-seed440420.mp4", + "sha256": "21f450848c2ded8c02054cc241c8080d55a17979022eadfd892f3ce5205833e6", + "size_bytes": 585250, + "container_bit_rate": 796936, + "audio": { + "integrated_loudness_lufs": -24.39, + "true_peak_dbtp": -7.18, + "first_0_25_seconds_peak_dbfs": -21.305752, + "first_0_25_seconds_rms_dbfs": -41.479687, + "first_sample_dbfs": [-35.859436, -39.793118] + } + } + } + }, + "comfy_sdpa_comparison": { + "comfyui_version": "0.31.0", + "attention": "PyTorch attention without --use-sage-attention", + "prompt_id": "29c45a77-efdf-4cc8-8ae9-9b8a40823b41", + "execution_seconds": 152.767, + "wall_seconds": 153.03360149999207, + "wall_difference_from_direct_percent": -2.25454138353405, + "output": { + "linux_path": "/home/daniel/StoryStudioAssets/H3-output/h3-baselines/comfy-sdpa-dialogue-two-character-864x480-141f-base12-seed440420_00001_.mp4", + "unc_path": "\\\\192.168.1.162\\StoryStudioAssets\\H3-output\\h3-baselines\\comfy-sdpa-dialogue-two-character-864x480-141f-base12-seed440420_00001_.mp4", + "sha256": "f4510d6b8ac06b8f2cb22172454c6977a214e061dcc5539ace68e1cb70abc2ee", + "size_bytes": 443723, + "container_bit_rate": 604218, + "duration_seconds": 5.875, + "video": { + "codec": "h264", + "profile": "High", + "pixel_format": "yuv420p", + "width": 864, + "height": 480, + "fps": 24, + "frames": 141, + "bit_rate": 457452 + }, + "audio": { + "codec": "aac", + "profile": "LC", + "sample_rate": 32000, + "channels": 2, + "channel_layout": "stereo", + "bit_rate": 133407, + "mean_volume_db": -25.7, + "max_volume_db": -4.3, + "integrated_loudness_lufs": -21.12, + "true_peak_dbtp": -4.25, + "loudness_range_lu": 4.10, + "loudness_threshold_lufs": -33.88 + } + }, + "first_0_25_seconds": { + "direct": { + "peak_level_db": -19.429245, + "rms_level_db": -33.601728, + "max_sample_difference": 0.022571, + "peak_count": 2, + "nan_count": 0, + "inf_count": 0 + }, + "comfy": { + "peak_level_db": -19.261140, + "rms_level_db": -33.545443, + "max_sample_difference": 0.023026, + "peak_count": 2, + "nan_count": 0, + "inf_count": 0 + }, + "decoded_pcm_psnr_db": { + "channel_1": 162.496, + "channel_2": 162.783 + }, + "conclusion": "No amplitude clipping; direct and Comfy decoded waveforms are effectively identical. Any shared audible startup artifact originates before runtime-specific audio decode and mux." + }, + "full_decoded_pcm_psnr_db": { + "channel_1": 166.581, + "channel_2": 165.385 + } + } +} diff --git a/compose.qwen38-vllm.yml b/compose.qwen38-vllm.yml new file mode 100644 index 0000000..12fc09a --- /dev/null +++ b/compose.qwen38-vllm.yml @@ -0,0 +1,46 @@ +services: + qwen38-vllm: + image: ghcr.io/aeon-7/aeon-vllm-ultimate:latest + container_name: qwen38-vllm + restart: unless-stopped + gpus: all + ipc: host + shm_size: 16g + network_mode: host + volumes: + - /home/daniel/aeon-spark-test/h3/models/Qwen3.8-27B-NVFP4:/model:ro + - /home/daniel/aeon-spark-test/h3/vllm-cache:/root/.cache/vllm + entrypoint: ["vllm"] + command: + - serve + - /model + - --served-model-name + - qwen38 + - --host + - 0.0.0.0 + - --port + - "8002" + - --tensor-parallel-size + - "1" + - --gpu-memory-utilization + - "0.45" + - --max-model-len + - "262144" + - --max-num-seqs + - "4" + - --max-num-batched-tokens + - "8192" + - --enable-chunked-prefill + - --enable-prefix-caching + - --kv-cache-dtype + - fp8_e4m3 + - --reasoning-parser + - qwen3 + - --tool-call-parser + - qwen3_xml + - --enable-auto-tool-choice + - --distributed-executor-backend + - mp + - --speculative-config + - '{"method":"mtp","num_speculative_tokens":5}' + - --trust-remote-code diff --git a/compose.spark-stack.yml b/compose.spark-stack.yml new file mode 100644 index 0000000..17049dd --- /dev/null +++ b/compose.spark-stack.yml @@ -0,0 +1,84 @@ +services: + h3-hot-runtime: + profiles: ["h3"] + image: h3-blackwell-runtime:dev + container_name: h3-hot-runtime + restart: unless-stopped + gpus: all + volumes: + - /home/daniel/aeon-spark-test/h3/comfy-models/diffusion_models:/models:ro + - /home/daniel/aeon-spark-test/h3/comfy-models/text_encoders:/text-encoders:ro + - /home/daniel/aeon-spark-test/h3/comfy-models/vae:/vae:ro + - /home/daniel/aeon-spark-test/h3/h3-blackwell-runtime/artifacts:/artifacts:ro + - /home/daniel/StoryStudioAssets/H3-output:/output + ports: + - "8001:8000" + environment: + H3_DISABLE_MMAP: "1" + H3_NVFP4_SCALE_BACKEND: "vortex" + H3_NVFP4_SCALE_VERSION: "1" + H3_SOL_QKV_LAYOUT: "native" + command: ["python", "/opt/h3-blackwell-runtime/tools/serve_hot_runtime.py", "--host", "0.0.0.0", "--port", "8000", "--attention", "sage2", "--warmup"] + + qwen38-vllm: + profiles: ["qwen"] + image: ghcr.io/aeon-7/aeon-vllm-ultimate:latest + container_name: qwen38-vllm + restart: unless-stopped + gpus: all + ipc: host + shm_size: 16g + network_mode: host + volumes: + - /home/daniel/aeon-spark-test/h3/models/Qwen3.8-27B-NVFP4:/model:ro + - /home/daniel/aeon-spark-test/h3/vllm-cache:/root/.cache/vllm + entrypoint: ["vllm"] + command: + - serve + - /model + - --served-model-name + - qwen38 + - --host + - 0.0.0.0 + - --port + - "8002" + - --tensor-parallel-size + - "1" + - --gpu-memory-utilization + - "0.45" + - --max-model-len + - "262144" + - --max-num-seqs + - "4" + - --max-num-batched-tokens + - "8192" + - --enable-chunked-prefill + - --enable-prefix-caching + - --kv-cache-dtype + - fp8_e4m3 + - --reasoning-parser + - qwen3 + - --tool-call-parser + - qwen3_xml + - --enable-auto-tool-choice + - --distributed-executor-backend + - mp + - --speculative-config + - '{"method":"mtp","num_speculative_tokens":5}' + - --trust-remote-code + + flux2-klein-comfy: + profiles: ["flux"] + image: ghcr.io/aeon-7/comfyui-aeon-spark:slim + container_name: flux2-klein-comfy + restart: unless-stopped + gpus: all + ipc: host + shm_size: 16g + volumes: + - /home/daniel/aeon-spark-test/workspace:/workspace/ComfyUI:rw + - /home/daniel/StoryStudioAssets/H3-output:/workspace/ComfyUI/output:rw + - /home/daniel/aeon-spark-test/extra_model_paths.yaml:/opt/ComfyUI/extra_model_paths.yaml:ro + - /home/daniel/comfy-data/models:/shared-models:ro + ports: + - "8192:8188" diff --git a/compose.spark.yml b/compose.spark.yml index b5aada1..f59e2ea 100644 --- a/compose.spark.yml +++ b/compose.spark.yml @@ -11,6 +11,7 @@ services: - /home/daniel/aeon-spark-test/h3/comfy-models/vae:/vae:ro - /home/daniel/aeon-spark-test/h3/h3-blackwell-runtime/artifacts:/artifacts:ro - /home/daniel/aeon-spark-test/h3/models/MiniMax-H3-Turbo-FL2VA/latest:/turbo:ro + - /home/daniel/aeon-spark-test/h3/models/Minimax-H3-Latent-Upscaler:/latent-upscaler:ro - /home/daniel/StoryStudioAssets/H3-output:/output command: ["sleep", "infinity"] h3-hot-runtime: @@ -25,6 +26,7 @@ services: - /home/daniel/aeon-spark-test/h3/comfy-models/vae:/vae:ro - /home/daniel/aeon-spark-test/h3/h3-blackwell-runtime/artifacts:/artifacts:ro - /home/daniel/aeon-spark-test/h3/models/MiniMax-H3-Turbo-FL2VA/latest:/turbo:ro + - /home/daniel/aeon-spark-test/h3/models/Minimax-H3-Latent-Upscaler:/latent-upscaler:ro - /home/daniel/StoryStudioAssets/H3-output:/output ports: - "8001:8000" @@ -33,4 +35,4 @@ services: H3_NVFP4_SCALE_BACKEND: "vortex" H3_NVFP4_SCALE_VERSION: "1" H3_SOL_QKV_LAYOUT: "native" - command: ["python", "/opt/h3-blackwell-runtime/tools/serve_hot_runtime.py", "--host", "0.0.0.0", "--port", "8000", "--attention", "sage2", "--turbo-lora", "4step=/turbo/minimax_h3_fl2v_turbo_4step_v1.1_768p_comfyui_bf16.safetensors", "--turbo-lora", "8step=/turbo/minimax_h3_fl2v_turbo_8step_v1.0_comfyui_bf16.safetensors", "--warmup"] + command: ["python", "/opt/h3-blackwell-runtime/tools/serve_hot_runtime.py", "--host", "0.0.0.0", "--port", "8000", "--attention", "sage2", "--turbo-lora", "4step=/turbo/minimax_h3_fl2v_turbo_4step_v1.1_768p_comfyui_bf16.safetensors", "--turbo-lora", "8step=/turbo/minimax_h3_fl2v_turbo_8step_v1.0_comfyui_bf16.safetensors", "--latent-upscaler", "/latent-upscaler/minimax_h3_latent_upscaler_3d_fp16.safetensors", "--warmup"] diff --git a/pyproject.toml b/pyproject.toml index b04caff..d74f797 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -11,5 +11,8 @@ dependencies = [ "transformers>=4.51,<5" ] +[project.optional-dependencies] +flash4 = ["flash-attn-4[cu13]==4.0.0b27"] + [tool.setuptools.packages.find] where = ["src"] diff --git a/src/h3_blackwell_runtime/attention.py b/src/h3_blackwell_runtime/attention.py index 07ef29b..793f630 100644 --- a/src/h3_blackwell_runtime/attention.py +++ b/src/h3_blackwell_runtime/attention.py @@ -9,8 +9,8 @@ from .checkpoint import H3Checkpoint from .nvfp4 import Nvfp4Linear -AVAILABLE_BACKENDS = ("sage2", "cudnn_sdpa", "ck_int8", "sdpa", "sage3", "sage3_mean", "kj_sage_cuda", "kj_sage_triton", "kj_sage_fp8", "kj_sage_fp8pp", "kj_head_sliced", "sol_attn") -PLANNED_BACKENDS = ("flash4", "easycache", "h3_cache", "kj_chunked_ffn") +AVAILABLE_BACKENDS = ("sage2", "cudnn_sdpa", "ck_int8", "sdpa", "flash4", "sage3", "sage3_mean", "kj_sage_cuda", "kj_sage_triton", "kj_sage_fp8", "kj_sage_fp8pp", "kj_head_sliced", "sol_attn") +PLANNED_BACKENDS = ("easycache", "h3_cache", "kj_chunked_ffn") DEFAULT_ATTENTION_BACKEND = os.getenv("H3_DEFAULT_ATTENTION", "sage2") @@ -20,7 +20,7 @@ def attention_backend_status() -> dict[str, str]: status.update({"cudnn_sdpa": "available: forced cuDNN SDPA with no backend fallback"}) status.update({"ck_int8": "available: approximate Comfy Kitchen INT8 Q/K/V attention"}) status.update({"sol_attn": "experimental: sparse Triton attention for eligible non-causal H3 attention calls; falls back below H3_SOL_MIN_TOKENS unless H3_SOL_STRICT=1"}) - status.update({"flash4": "planned: exact Blackwell kernel adapter"}) + status.update({"flash4": "available: official FlashAttention-4 CuTeDSL Blackwell kernel (strict, no fallback)"}) status.update({"easycache": "planned: approximate denoiser cache"}) status.update({"h3_cache": "planned: approximate H3-specific cache"}) status.update({"kj_chunked_ffn": "available: exact H3 MLP row chunking via H3_MLP_CHUNKS or runtime args"}) @@ -85,6 +85,23 @@ def qkv_to_bshd(qkv: torch.Tensor, heads: int, head_dim: int) -> tuple[torch.Ten ) +def run_flash4_attention_bshd(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, *, is_causal: bool) -> torch.Tensor: + """Run official FlashAttention-4 on `[batch, sequence, heads, dim]` tensors.""" + if not q.is_cuda or not k.is_cuda or not v.is_cuda: + raise ValueError("FlashAttention-4 requires CUDA tensors") + if q.dtype not in {torch.float16, torch.bfloat16} or k.dtype != q.dtype or v.dtype != q.dtype: + raise ValueError("FlashAttention-4 requires matching FP16 or BF16 Q/K/V tensors") + if q.shape != k.shape or q.shape != v.shape: + raise ValueError("FlashAttention-4 requires matching Q/K/V shapes") + if q.shape[-1] != 128: + raise ValueError(f"H3 FlashAttention-4 requires head dim 128, got {q.shape[-1]}") + + from flash_attn.cute import flash_attn_func + + result = flash_attn_func(q.contiguous(), k.contiguous(), v.contiguous(), causal=is_causal) + return result[0] if isinstance(result, tuple) else result + + def run_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, *, backend: str, is_causal: bool) -> torch.Tensor: """Run one `[batch, heads, sequence, dim]` attention operation.""" if backend == "sol_attn": @@ -103,6 +120,10 @@ def run_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, *, backend: from sageattention import sageattn return sageattn(q, k, v, is_causal=is_causal, tensor_layout="HND", smooth_k=False) + if backend == "flash4": + return run_flash4_attention_bshd( + q.transpose(1, 2), k.transpose(1, 2), v.transpose(1, 2), is_causal=is_causal, + ).transpose(1, 2) if backend == "cudnn_sdpa": from torch.nn.attention import SDPBackend, sdpa_kernel @@ -230,6 +251,9 @@ class H3SageAttention(nn.Module): q, k, v = q.contiguous(), k.contiguous(), v.contiguous() out = run_sol_attention_bshd(q, k, v, is_causal=False) return self.out_proj(out.reshape(sequence, inner).contiguous()) + if self.backend == "flash4": + out = run_flash4_attention_bshd(q, k, v, is_causal=False) + return self.out_proj(out.reshape(sequence, inner).contiguous()) q = q.transpose(1, 2).contiguous() k = k.transpose(1, 2).contiguous() diff --git a/src/h3_blackwell_runtime/latent_upscaler.py b/src/h3_blackwell_runtime/latent_upscaler.py new file mode 100644 index 0000000..08e4533 --- /dev/null +++ b/src/h3_blackwell_runtime/latent_upscaler.py @@ -0,0 +1,163 @@ +"""MiniMax H3 learned latent upscaler inference. + +Adapted from LBH-123-AI/Comfyui_Minimax_h3_latent_Upscaler (Apache-2.0). +""" + +from pathlib import Path +import re + +import torch +from safetensors.torch import load_file +from torch import nn +from torch.nn import functional as F + + +LATENTS_MEAN = ( + 0.858090341091156, -0.9606591463088989, 1.0661640167236328, -0.5090325474739075, + -0.2727581858634949, -1.3675414323806763, -0.2553254961967468, -0.26907554268836975, + -0.5376840829849243, -0.0464097298681736, 0.6657370328903198, 0.19690127670764923, + -0.5460608005523682, -0.4035342037677765, -0.23683024942874908, 0.25928452610969543, + -0.30133944749832153, 0.211341992020607, -1.1206848621368408, 0.3581933379173279, + -0.04225143790245056, 0.2604829967021942, 0.22864092886447906, 0.7056031823158264, +) +LATENTS_STD = ( + 1.2223774194717407, 1.2767263650894165, 1.6831774711608887, 1.7549455165863037, + 1.5636216402053833, 2.194143533706665, 0.9653137922286987, 1.0569885969161987, + 0.841948926448822, 0.7729952931404114, 1.8955937623977661, 0.946841835975647, + 0.7996809482574463, 0.44988900423049927, 0.7197399735450745, 0.6936293244361877, + 2.961095094680786, 2.7694199085235596, 3.0496184825897217, 2.1088054180145264, + 3.276226282119751, 3.1627357006073, 2.2816812992095947, 2.6127843856811523, +) + + +def _normalization(channels: int) -> nn.GroupNorm: + return nn.GroupNorm(32, channels) + + +class ResBlockEmb3D(nn.Module): + def __init__(self, channels: int, emb_channels: int, dropout: float = 0.0): + super().__init__() + self.in_layers = nn.Sequential( + _normalization(channels), + nn.SiLU(), + nn.Conv3d(channels, channels, 3, padding=1), + ) + self.emb_layers = nn.Sequential(nn.SiLU(), nn.Linear(emb_channels, 2 * channels)) + self.out_norm = _normalization(channels) + self.out_layers = nn.Sequential( + nn.SiLU(), + nn.Dropout(p=dropout), + nn.Conv3d(channels, channels, 3, padding=1), + ) + nn.init.zeros_(self.out_layers[-1].weight) + nn.init.zeros_(self.out_layers[-1].bias) + + def forward(self, x: torch.Tensor, emb: torch.Tensor) -> torch.Tensor: + hidden = self.in_layers(x) + scale, shift = self.emb_layers(emb).to(hidden.dtype).chunk(2, dim=1) + hidden = self.out_norm(hidden) * (1 + scale[:, :, None, None, None]) + shift[:, :, None, None, None] + return x + self.out_layers(hidden) + + +class TemporalConv(nn.Module): + def __init__(self, channels: int, kernel_size: int = 5): + super().__init__() + self.norm = _normalization(channels) + self.dwconv = nn.Conv3d( + channels, channels, kernel_size=(kernel_size, 1, 1), + padding=(kernel_size // 2, 0, 0), groups=channels, + ) + self.pwconv = nn.Conv3d(channels, channels, 1) + nn.init.zeros_(self.pwconv.weight) + nn.init.zeros_(self.pwconv.bias) + + def forward(self, x: torch.Tensor) -> torch.Tensor: + hidden = self.dwconv(F.silu(self.norm(x))) + return x + self.pwconv(hidden) + + +class H3LatentResizer3D(nn.Module): + def __init__( + self, + in_channels: int = 24, + in_blocks: int = 12, + out_blocks: int = 12, + channels: int = 512, + dropout: float = 0.1, + temporal_every: int = 2, + temporal_kernel: int = 5, + ): + super().__init__() + self.conv_in = nn.Conv3d(in_channels, channels, 3, padding=1) + embed_dim = 64 + self.embed = nn.Sequential(nn.Linear(1, embed_dim), nn.SiLU(), nn.Linear(embed_dim, embed_dim)) + self.in_blocks = self._make_blocks(in_blocks, channels, embed_dim, dropout, temporal_every, temporal_kernel) + self.out_blocks = self._make_blocks(out_blocks, channels, embed_dim, dropout, temporal_every, temporal_kernel) + self.norm_out = _normalization(channels) + self.conv_out = nn.Conv3d(channels, in_channels, 3, padding=1) + + @staticmethod + def _make_blocks(count, channels, embed_dim, dropout, temporal_every, temporal_kernel): + blocks = nn.ModuleList() + for index in range(count): + blocks.append(ResBlockEmb3D(channels, embed_dim, dropout)) + if temporal_every > 0 and index % temporal_every == 0: + blocks.append(TemporalConv(channels, temporal_kernel)) + return blocks + + def forward(self, x: torch.Tensor, *, scale: float, target_size: tuple[int, int, int]) -> torch.Tensor: + emb = self.embed(torch.tensor([[scale - 1]], dtype=x.dtype, device=x.device)).expand(x.shape[0], -1) + hidden = self.conv_in(x) + for block in self.in_blocks: + hidden = block(hidden, emb) if isinstance(block, ResBlockEmb3D) else block(hidden) + hidden = F.interpolate(hidden, size=target_size, mode="trilinear", align_corners=False) + for block in self.out_blocks: + hidden = block(hidden, emb) if isinstance(block, ResBlockEmb3D) else block(hidden) + return self.conv_out(F.silu(self.norm_out(hidden))) + + +def _checkpoint_config(state: dict[str, torch.Tensor]) -> dict: + in_ids = {int(match.group(1)) for key in state if (match := re.match(r"in_blocks\.(\d+)\.in_layers\.", key))} + out_ids = {int(match.group(1)) for key in state if (match := re.match(r"out_blocks\.(\d+)\.in_layers\.", key))} + temporal_keys = [key for key in state if key.endswith("dwconv.weight")] + conv_in = state["conv_in.weight"] + return { + "in_channels": conv_in.shape[1], + "in_blocks": len(in_ids), + "out_blocks": len(out_ids), + "channels": conv_in.shape[0], + "temporal_every": 2 if temporal_keys else 0, + "temporal_kernel": state[temporal_keys[0]].shape[2] if temporal_keys else 5, + } + + +def load_h3_latent_upscaler( + path: str | Path, + *, + device: str | torch.device = "cuda", + dtype: torch.dtype = torch.float16, +) -> H3LatentResizer3D: + state = load_file(str(path), device="cpu") + if any(key.startswith("upscaler.") for key in state): + state = {key.removeprefix("upscaler."): value for key, value in state.items() if key.startswith("upscaler.")} + model = H3LatentResizer3D(**_checkpoint_config(state)) + model.load_state_dict(state, strict=True) + return model.to(device=device, dtype=dtype).eval().requires_grad_(False) + + +@torch.inference_mode() +def upscale_h3_latent(model: H3LatentResizer3D, latent: torch.Tensor, *, scale: float = 2.0) -> torch.Tensor: + if latent.ndim != 5 or latent.shape[1] != 24: + raise ValueError("H3 latent must have shape [B,24,T,H,W]") + if not 1.0 <= scale <= 4.0: + raise ValueError("scale must be between 1.0 and 4.0") + dtype = next(model.parameters()).dtype + device = next(model.parameters()).device + source = latent.to(device=device, dtype=dtype) + if scale == 1.0: + return source + mean = torch.tensor(LATENTS_MEAN, device=device, dtype=dtype).view(1, 24, 1, 1, 1) + std = torch.tensor(LATENTS_STD, device=device, dtype=dtype).view(1, 24, 1, 1, 1) + target_size = (source.shape[2], round(source.shape[3] * scale), round(source.shape[4] * scale)) + result = model((source - mean) / std, scale=scale, target_size=target_size) + return result.mul_(std).add_(mean) diff --git a/src/h3_blackwell_runtime/runtime.py b/src/h3_blackwell_runtime/runtime.py index 8d4b4d9..c4a89aa 100644 --- a/src/h3_blackwell_runtime/runtime.py +++ b/src/h3_blackwell_runtime/runtime.py @@ -16,6 +16,7 @@ from .block import configure_mlp_chunking from .checkpoint import H3Checkpoint from .denoiser import H3PackedDenoiser from .lora import load_lora_adapter, set_active_lora +from .latent_upscaler import load_h3_latent_upscaler, upscale_h3_latent from .packing import H3PromptPacker from .qwen3vl_text import Qwen3VLPromptConditioner from .qwen3vl_vision import build_fl2va_presentation, Qwen3VL32BVision @@ -45,6 +46,7 @@ class RuntimeConfig: mlp_chunks: int = 1 mlp_chunk_threshold: int = 4096 turbo_loras: tuple[tuple[str, str], ...] = () + latent_upscaler_path: str | None = None device: str = "cuda" @@ -64,6 +66,18 @@ def _refiner_attention_backend(attention: str) -> str: return fallback if fallback in AVAILABLE_BACKENDS and fallback != "sol_attn" else "sage2" +def normalize_upscale(value) -> float | None: + if value is None or value == "" or value == "none" or value == 1 or value == 1.0: + return None + try: + scale = float(value) + except (TypeError, ValueError) as exc: + raise ValueError("upscale must be a number between 1.0 and 4.0") from exc + if not 1.0 < scale <= 4.0: + raise ValueError("upscale must be greater than 1.0 and at most 4.0") + return scale + + class H3HotRuntime: """Keep all prompt-only H3 models resident for repeated requests.""" @@ -113,6 +127,12 @@ class H3HotRuntime: "vision_tower_loaded", lambda: Qwen3VL32BVision(config.text_encoder_path, device=config.device, dtype=torch.float32), ) + self.latent_upscaler = None + if config.latent_upscaler_path is not None: + self.latent_upscaler = self._timed_load( + "latent_upscaler_loaded", + lambda: load_h3_latent_upscaler(config.latent_upscaler_path, device=config.device, dtype=torch.float16), + ) def _timed_load(self, stage: str, fn): _sync() @@ -143,6 +163,7 @@ class H3HotRuntime: "current_attention": self.attention, "available_turbos": sorted(self.turbo_loras), "current_turbo": self.turbo, + "latent_upscaler_loaded": self.latent_upscaler is not None, "vae_dtype": self.config.vae_dtype, "vae_tile_size": self.config.vae_tile_size, "mlp_chunks": self.config.mlp_chunks, @@ -224,9 +245,11 @@ class H3HotRuntime: seed: int, attention: str | None = None, turbo: str | None = None, + upscale: float | None = None, first_frame: torch.Tensor | None = None, last_frame: torch.Tensor | None = None, mux_audio: bool = True, + keep_intermediates: bool = False, ffmpeg_loglevel: str = "error", save_latent: str | Path | None = None, cache_mode: str | None = None, @@ -249,6 +272,9 @@ class H3HotRuntime: output = Path(output) output.parent.mkdir(parents=True, exist_ok=True) requested_turbo = None if turbo in {None, "", "none"} else turbo + upscale_scale = normalize_upscale(upscale) + if upscale_scale is not None and self.latent_upscaler is None: + raise ValueError("H3 latent upscaler is not loaded") if requested_turbo is not None: if requested_turbo not in self.turbo_loras: raise ValueError(f"Turbo variant {requested_turbo!r} is not loaded") @@ -301,6 +327,14 @@ class H3HotRuntime: latent, audio_latent = sampled else: latent, audio_latent = sampled, None + source_width, source_height = width, height + if upscale_scale is not None: + latent = timed( + "latent_upscaled", + lambda: upscale_h3_latent(self.latent_upscaler, latent, scale=upscale_scale), + ) + width = latent.shape[-1] * 16 + height = latent.shape[-2] * 16 if save_latent is not None: latent_path = Path(save_latent) latent_path.parent.mkdir(parents=True, exist_ok=True) @@ -351,7 +385,10 @@ class H3HotRuntime: timed("audio_encode", lambda: subprocess.run(_ffmpeg_command(ffmpeg_loglevel, "-y", "-f", "f32le", "-ar", "32000", "-ac", "2", "-i", str(audio_raw), str(audio_output)), check=True)) audio_raw.unlink() timed("mux", lambda: subprocess.run(_ffmpeg_command(ffmpeg_loglevel, "-y", "-i", str(video_output), "-i", str(audio_output), "-c:v", "copy", "-c:a", "aac", "-shortest", str(output)), check=True)) - video_output.unlink() + if not keep_intermediates: + video_output.unlink() + audio_output.unlink() + audio_output = None return { "output": str(output), @@ -359,9 +396,13 @@ class H3HotRuntime: "frames": aligned_frames, "width": width, "height": height, + "source_width": source_width, + "source_height": source_height, "seed": seed, "attention": self.attention, "turbo": self.turbo, + "upscale": upscale_scale, + "keep_intermediates": keep_intermediates, "vae_dtype": self.config.vae_dtype, "vae_tile_size": self.config.vae_tile_size, "stages": stages, diff --git a/tests/test_fl2va_contracts.py b/tests/test_fl2va_contracts.py index a8984a3..a608fef 100644 --- a/tests/test_fl2va_contracts.py +++ b/tests/test_fl2va_contracts.py @@ -63,7 +63,26 @@ class Fl2vaVAEContracts(unittest.TestCase): class AttentionBackendContracts(unittest.TestCase): def test_hot_backends_include_benchmark_candidates(self): - self.assertTrue({"sage2", "cudnn_sdpa", "ck_int8"}.issubset(AVAILABLE_BACKENDS)) + self.assertTrue({"sage2", "cudnn_sdpa", "ck_int8", "flash4"}.issubset(AVAILABLE_BACKENDS)) + + def test_flash4_dispatches_bshd_and_restores_hnd(self): + q = torch.randn(1, 2, 3, 128, dtype=torch.bfloat16) + expected_bshd = torch.randn(1, 3, 2, 128, dtype=torch.bfloat16) + flash = unittest.mock.MagicMock(return_value=(expected_bshd, torch.empty(0))) + package = SimpleNamespace(cute=SimpleNamespace(flash_attn_func=flash)) + with ( + patch.dict(sys.modules, {"flash_attn": package, "flash_attn.cute": package.cute}), + patch.object(torch.Tensor, "is_cuda", new_callable=unittest.mock.PropertyMock, return_value=True), + ): + actual = run_attention(q, q, q, backend="flash4", is_causal=False) + torch.testing.assert_close(actual, expected_bshd.transpose(1, 2)) + flash.assert_called_once() + called_q, called_k, called_v = flash.call_args.args + self.assertEqual(called_q.shape, (1, 3, 2, 128)) + self.assertTrue(called_q.is_contiguous()) + self.assertTrue(called_k.is_contiguous()) + self.assertTrue(called_v.is_contiguous()) + self.assertFalse(flash.call_args.kwargs["causal"]) def test_cudnn_backend_is_forced_without_fallback(self): q = torch.randn(1, 2, 3, 4) diff --git a/tests/test_latent_upscaler.py b/tests/test_latent_upscaler.py new file mode 100644 index 0000000..1be94ae --- /dev/null +++ b/tests/test_latent_upscaler.py @@ -0,0 +1,52 @@ +import unittest + +import torch + +from h3_blackwell_runtime.latent_upscaler import H3LatentResizer3D, _checkpoint_config, upscale_h3_latent +from h3_blackwell_runtime.runtime import normalize_upscale + + +class LatentUpscalerContracts(unittest.TestCase): + def test_request_upscale_normalization(self): + self.assertIsNone(normalize_upscale(None)) + self.assertIsNone(normalize_upscale("none")) + self.assertIsNone(normalize_upscale(1)) + self.assertEqual(normalize_upscale("2"), 2.0) + with self.assertRaisesRegex(ValueError, "greater than 1.0"): + normalize_upscale(0.5) + with self.assertRaisesRegex(ValueError, "at most 4.0"): + normalize_upscale(4.1) + with self.assertRaisesRegex(ValueError, "must be a number"): + normalize_upscale({"scale": 2}) + + def test_3d_model_preserves_time_and_scales_space(self): + model = H3LatentResizer3D(in_blocks=1, out_blocks=1, channels=32, dropout=0, temporal_every=0).eval() + latent = torch.randn(1, 24, 3, 2, 4) + output = upscale_h3_latent(model, latent, scale=2) + self.assertEqual(output.shape, (1, 24, 3, 4, 8)) + + def test_invalid_latent_shape_is_rejected(self): + model = H3LatentResizer3D(in_blocks=1, out_blocks=1, channels=32, temporal_every=0).eval() + with self.assertRaisesRegex(ValueError, r"\[B,24,T,H,W\]"): + upscale_h3_latent(model, torch.randn(1, 16, 3, 2, 4)) + + def test_scale_one_preserves_latent_exactly(self): + model = H3LatentResizer3D(in_blocks=1, out_blocks=1, channels=32, temporal_every=0).eval() + latent = torch.randn(1, 24, 3, 2, 4) + self.assertTrue(torch.equal(upscale_h3_latent(model, latent, scale=1), latent)) + + def test_checkpoint_architecture_detection_matches_module_layout(self): + model = H3LatentResizer3D(in_blocks=3, out_blocks=2, channels=32, temporal_every=2, temporal_kernel=3) + config = _checkpoint_config(model.state_dict()) + self.assertEqual(config, { + "in_channels": 24, + "in_blocks": 3, + "out_blocks": 2, + "channels": 32, + "temporal_every": 2, + "temporal_kernel": 3, + }) + + +if __name__ == "__main__": + unittest.main() diff --git a/tools/poll_qwen38_health_spark.sh b/tools/poll_qwen38_health_spark.sh new file mode 100644 index 0000000..b568d10 --- /dev/null +++ b/tools/poll_qwen38_health_spark.sh @@ -0,0 +1,27 @@ +#!/usr/bin/env bash +set -u + +for i in $(seq 0 39); do + status=$(docker ps -a --filter name=qwen38-vllm --format '{{.Status}}') + listening=$(ss -ltnp | grep -c ':8002 ' || true) + + if curl -fsS --max-time 5 http://127.0.0.1:8002/health >/tmp/qwen38-health.out 2>/tmp/qwen38-health.err; then + echo "HEALTHY poll=$i status=$status listening=$listening" + docker logs --tail 60 qwen38-vllm + exit 0 + fi + + if [[ $status != Up* ]]; then + echo "NOT_RUNNING poll=$i status=$status listening=$listening" + docker logs --tail 160 qwen38-vllm + exit 2 + fi + + echo "poll=$i status=$status listening=$listening" + docker logs --tail 10 qwen38-vllm + sleep 30 +done + +echo "TIMEOUT" +docker logs --tail 200 qwen38-vllm +exit 1 diff --git a/tools/serve_hot_runtime.py b/tools/serve_hot_runtime.py index 85d9428..adf6201 100644 --- a/tools/serve_hot_runtime.py +++ b/tools/serve_hot_runtime.py @@ -14,7 +14,7 @@ import torch from PIL import Image from h3_blackwell_runtime.attention import AVAILABLE_BACKENDS, DEFAULT_ATTENTION_BACKEND, attention_backend_status -from h3_blackwell_runtime.runtime import H3HotRuntime, RuntimeConfig, TURBO_VARIANTS +from h3_blackwell_runtime.runtime import H3HotRuntime, RuntimeConfig, TURBO_VARIANTS, normalize_upscale def _load_image(value) -> torch.Tensor | None: @@ -57,6 +57,7 @@ parser.add_argument("--vae-tile-size", type=int, default=256) parser.add_argument("--mlp-chunks", type=int, default=1) parser.add_argument("--mlp-chunk-threshold", type=int, default=4096) parser.add_argument("--turbo-lora", action="append", default=[], metavar="NAME=PATH", help="Load a resident 4step or 8step Turbo adapter.") +parser.add_argument("--latent-upscaler", help="Load the optional H3 3D latent upscaler for request-level spatial upscaling.") parser.add_argument("--warmup", action="store_true", help="Run a tiny generation before accepting traffic.") parser.add_argument("--warmup-output", type=Path, default=Path("/output/h3-blackwell-runtime/hot-runtime-warmup.mp4")) args = parser.parse_args() @@ -68,7 +69,7 @@ for value in args.turbo_lora: name, path = value.split("=", 1) turbo_loras.append((name, path)) -runtime = H3HotRuntime(RuntimeConfig(attention=args.attention, vae_dtype=args.vae_dtype, vae_tile_size=args.vae_tile_size, mlp_chunks=args.mlp_chunks, mlp_chunk_threshold=args.mlp_chunk_threshold, turbo_loras=tuple(turbo_loras))) +runtime = H3HotRuntime(RuntimeConfig(attention=args.attention, vae_dtype=args.vae_dtype, vae_tile_size=args.vae_tile_size, mlp_chunks=args.mlp_chunks, mlp_chunk_threshold=args.mlp_chunk_threshold, turbo_loras=tuple(turbo_loras), latent_upscaler_path=args.latent_upscaler)) runtime_lock = threading.Lock() warmup_result = None if args.warmup: @@ -142,6 +143,15 @@ class Handler(BaseHTTPRequestHandler): write_json(self, 400, {"error": "unsupported attention", "attention": attention, "available": list(AVAILABLE_BACKENDS)}) return mux_audio = bool(payload.get("mux_audio", True)) + keep_intermediates = bool(payload.get("keep_intermediates", False)) + try: + upscale = normalize_upscale(payload.get("upscale")) + except ValueError as exc: + write_json(self, 400, {"error": "invalid upscale", "message": str(exc)}) + return + if upscale is not None and runtime.latent_upscaler is None: + write_json(self, 400, {"error": "H3 latent upscaler is not loaded"}) + return ffmpeg_loglevel = payload.get("ffmpeg_loglevel", "error") first_frame = _load_image(payload.get("first_frame")) last_frame = _load_image(payload.get("last_frame")) @@ -166,9 +176,11 @@ class Handler(BaseHTTPRequestHandler): seed=seed, attention=attention, turbo=turbo, + upscale=upscale, first_frame=first_frame, last_frame=last_frame, mux_audio=mux_audio, + keep_intermediates=keep_intermediates, ffmpeg_loglevel=ffmpeg_loglevel, save_latent=save_latent, cache_mode=cache_mode, diff --git a/tools/smoke_flash4.py b/tools/smoke_flash4.py new file mode 100644 index 0000000..917c739 --- /dev/null +++ b/tools/smoke_flash4.py @@ -0,0 +1,50 @@ +"""Compare FlashAttention-4 against PyTorch SDPA on an H3-shaped operation.""" + +import argparse +import json + +import torch + + +parser = argparse.ArgumentParser() +parser.add_argument("--sequence", type=int, default=257) +parser.add_argument("--heads", type=int, default=8) +parser.add_argument("--head-dim", type=int, default=128) +parser.add_argument("--iterations", type=int, default=5) +args = parser.parse_args() + +from flash_attn.cute import flash_attn_func + +torch.manual_seed(440411) +shape = (1, args.sequence, args.heads, args.head_dim) +q = torch.randn(shape, device="cuda", dtype=torch.bfloat16) +k = torch.randn(shape, device="cuda", dtype=torch.bfloat16) +v = torch.randn(shape, device="cuda", dtype=torch.bfloat16) + +with torch.inference_mode(): + result = flash_attn_func(q, k, v, causal=False) + actual = result[0] if isinstance(result, tuple) else result + expected = torch.nn.functional.scaled_dot_product_attention( + q.transpose(1, 2), k.transpose(1, 2), v.transpose(1, 2), is_causal=False, + ).transpose(1, 2) + torch.cuda.synchronize() + started = torch.cuda.Event(enable_timing=True) + finished = torch.cuda.Event(enable_timing=True) + started.record() + for _ in range(args.iterations): + result = flash_attn_func(q, k, v, causal=False) + actual = result[0] if isinstance(result, tuple) else result + finished.record() + torch.cuda.synchronize() + +error = (actual.float() - expected.float()).abs() +print(json.dumps({ + "device": torch.cuda.get_device_name(), + "capability": torch.cuda.get_device_capability(), + "shape": tuple(actual.shape), + "dtype": str(actual.dtype), + "contiguous": actual.is_contiguous(), + "max_abs_error": error.max().item(), + "mean_abs_error": error.mean().item(), + "milliseconds": started.elapsed_time(finished) / args.iterations, +}, indent=2)) diff --git a/tools/smoke_qwen38_vllm_spark.sh b/tools/smoke_qwen38_vllm_spark.sh new file mode 100644 index 0000000..56d0ed8 --- /dev/null +++ b/tools/smoke_qwen38_vllm_spark.sh @@ -0,0 +1,8 @@ +#!/usr/bin/env bash +set -euo pipefail + +curl -fsS --max-time 120 http://127.0.0.1:8002/v1/chat/completions \ + -H 'Content-Type: application/json' \ + --data-binary @- <<'JSON' +{"model":"qwen38","messages":[{"role":"user","content":"Reply with exactly: ready"}],"max_tokens":64,"temperature":0,"chat_template_kwargs":{"enable_thinking":false}} +JSON diff --git a/tools/spark_stack_spark.sh b/tools/spark_stack_spark.sh new file mode 100644 index 0000000..4cac039 --- /dev/null +++ b/tools/spark_stack_spark.sh @@ -0,0 +1,61 @@ +#!/usr/bin/env bash +set -euo pipefail + +ROOT=/home/daniel/aeon-spark-test/h3 +STACK_DIR="$ROOT/spark-stack" +COMPOSE_FILE="$STACK_DIR/compose.spark-stack.yml" + +usage() { + cat <<'EOF' +Usage: + spark_stack_spark.sh start h3|qwen|flux + spark_stack_spark.sh stop + spark_stack_spark.sh status + +Ports: + h3 -> http://192.168.1.162:8001 + qwen -> http://192.168.1.162:8002 + flux -> http://192.168.1.162:8192 +EOF +} + +managed_containers=( + h3-hot-runtime + h3-blackwell-runtime-h3-hot-runtime-1 + qwen38-vllm + flux2-klein-comfy + flux-nvfp4-test +) + +stop_managed() { + for name in "${managed_containers[@]}"; do + if docker ps -a --format '{{.Names}}' | grep -qx "$name"; then + docker stop "$name" >/dev/null 2>&1 || true + docker rm "$name" >/dev/null 2>&1 || true + fi + done +} + +case "${1:-}" in + start) + target="${2:-}" + case "$target" in + h3|qwen|flux) ;; + *) usage; exit 2 ;; + esac + mkdir -p "$ROOT/vllm-cache" + stop_managed + docker compose -f "$COMPOSE_FILE" --profile "$target" up -d + ;; + stop) + stop_managed + ;; + status) + docker ps -a --format 'table {{.Names}}\t{{.Status}}\t{{.Ports}}' | grep -E 'NAMES|h3-hot-runtime|qwen38-vllm|flux2-klein-comfy|flux-nvfp4-test' || true + nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv,noheader,nounits || true + ;; + *) + usage + exit 2 + ;; +esac diff --git a/tools/start_qwen38_vllm_spark.sh b/tools/start_qwen38_vllm_spark.sh new file mode 100644 index 0000000..b1c9c01 --- /dev/null +++ b/tools/start_qwen38_vllm_spark.sh @@ -0,0 +1,27 @@ +#!/usr/bin/env bash +set -euo pipefail + +ROOT=/home/daniel/aeon-spark-test/h3 +MODEL_DIR="$ROOT/models/Qwen3.8-27B-NVFP4" +COMPOSE_FILE="$ROOT/qwen38-vllm/compose.qwen38-vllm.yml" + +if [[ ! -f "$MODEL_DIR/model.safetensors" ]]; then + echo "model.safetensors is not present yet: $MODEL_DIR/model.safetensors" >&2 + echo "Wait for qwen38-27b-nvfp4-download to finish before starting vLLM." >&2 + exit 1 +fi + +python3 - <<'PY' +import json +from pathlib import Path + +tokenizer = Path('/home/daniel/aeon-spark-test/h3/models/Qwen3.8-27B-NVFP4/tokenizer.json') +truncation = json.loads(tokenizer.read_text(encoding='utf-8')).get('truncation') +if truncation is not None: + raise SystemExit(f'tokenizer truncation must be null, got: {truncation!r}') +print('tokenizer truncation: null') +PY + +mkdir -p "$ROOT/vllm-cache" +docker compose -f "$COMPOSE_FILE" up -d +docker logs -f qwen38-vllm diff --git a/tools/submit_h3_comfy_t2v.py b/tools/submit_h3_comfy_t2v.py new file mode 100644 index 0000000..9cee354 --- /dev/null +++ b/tools/submit_h3_comfy_t2v.py @@ -0,0 +1,91 @@ +"""Submit a prompt-only MiniMax H3 T2V graph to ComfyUI.""" + +from __future__ import annotations + +import argparse +import json +import time +from urllib.request import Request, urlopen + + +def post_json(url: str, payload: dict) -> dict: + request = Request(url, data=json.dumps(payload).encode(), headers={"Content-Type": "application/json"}) + with urlopen(request, timeout=30) as response: + return json.loads(response.read().decode()) + + +def get_json(url: str) -> dict: + with urlopen(url, timeout=30) as response: + return json.loads(response.read().decode()) + + +def main() -> None: + parser = argparse.ArgumentParser() + parser.add_argument("--server", default="http://127.0.0.1:8188") + parser.add_argument("--prompt", required=True) + parser.add_argument("--filename-prefix", required=True) + parser.add_argument("--seed", type=int, default=440407) + parser.add_argument("--width", type=int, default=960) + parser.add_argument("--height", type=int, default=544) + parser.add_argument("--frames", type=int, default=124) + parser.add_argument("--steps", type=int, default=12) + parser.add_argument("--upscale", type=float) + args = parser.parse_args() + + graph = { + "1": {"class_type": "UNETLoader", "inputs": {"unet_name": "minimax_h3_fl2va_pruned_nvfp4.safetensors", "weight_dtype": "default"}}, + "3": {"class_type": "CLIPLoader", "inputs": {"clip_name": "qwen3vl_32b_minimax_h3_nvfp4_awq.safetensors", "type": "minimax"}}, + "4": {"class_type": "VAELoader", "inputs": {"vae_name": "minimax_h3_video_vae_fp16.safetensors"}}, + "5": {"class_type": "VAELoader", "inputs": {"vae_name": "minimax_h3_audio_vae_fp32.safetensors"}}, + "8": {"class_type": "MiniMaxH3ImageToVideo", "inputs": {"clip": ["3", 0], "vae": ["4", 0], "prompt": args.prompt, "width": args.width, "height": args.height, "length": args.frames}}, + "9": {"class_type": "BasicGuider", "inputs": {"model": ["1", 0], "conditioning": ["8", 0]}}, + "10": {"class_type": "RandomNoise", "inputs": {"noise_seed": args.seed}}, + "11": {"class_type": "KSamplerSelect", "inputs": {"sampler_name": "res_multistep"}}, + "12": {"class_type": "BasicScheduler", "inputs": {"model": ["1", 0], "scheduler": "beta", "steps": args.steps, "denoise": 1.0}}, + "13": {"class_type": "SamplerCustomAdvanced", "inputs": {"noise": ["10", 0], "guider": ["9", 0], "sampler": ["11", 0], "sigmas": ["12", 0], "latent_image": ["8", 1]}}, + } + decoded_latent = ["13", 0] + if args.upscale is not None: + graph.update({ + "18": {"class_type": "LTXVSeparateAVLatent", "inputs": {"av_latent": ["13", 0]}}, + "19": {"class_type": "MinimaxH3LatentUpscaler3D", "inputs": { + "latent": ["18", 0], + "model_name": "minimax_h3_latent_upscaler_3d_fp16.safetensors", + "mode": "scale by multiplier", + "mode.scale": args.upscale, + "align": 32, + "keep_proportion": True, + "device": "cuda", + "precision": "fp16", + }}, + "20": {"class_type": "LTXVConcatAVLatent", "inputs": {"video_latent": ["19", 0], "audio_latent": ["18", 1]}}, + }) + decoded_latent = ["20", 0] + graph.update({ + "14": {"class_type": "VAEDecode", "inputs": {"samples": decoded_latent, "vae": ["4", 0]}}, + "15": {"class_type": "VAEDecodeAudio", "inputs": {"samples": decoded_latent, "vae": ["5", 0]}}, + "16": {"class_type": "CreateVideo", "inputs": {"images": ["14", 0], "audio": ["15", 0], "bit_depth": 8, "fps": 24.0}}, + "17": {"class_type": "SaveVideo", "inputs": {"video": ["16", 0], "filename_prefix": args.filename_prefix, "format": "mp4", "codec": "auto"}}, + }) + + started = time.perf_counter() + response = post_json(f"{args.server}/prompt", {"prompt": graph}) + prompt_id = response["prompt_id"] + while True: + time.sleep(1) + history = get_json(f"{args.server}/history/{prompt_id}").get(prompt_id) + if history is None: + continue + status = history.get("status", {}) + if status.get("completed") or status.get("status_str") in {"success", "error"}: + break + print(json.dumps({ + "prompt_id": prompt_id, + "wall_seconds": time.perf_counter() - started, + "status": history.get("status", {}), + "outputs": history.get("outputs", {}), + }, indent=2)) + + +if __name__ == "__main__": + main() diff --git a/tools/upscale_video_latent.py b/tools/upscale_video_latent.py new file mode 100644 index 0000000..9212090 --- /dev/null +++ b/tools/upscale_video_latent.py @@ -0,0 +1,55 @@ +"""Upscale a saved direct-runtime H3 latent with the learned 3D model.""" + +import argparse +import json +import time +from pathlib import Path + +import torch + +from h3_blackwell_runtime.latent_upscaler import load_h3_latent_upscaler, upscale_h3_latent + + +parser = argparse.ArgumentParser() +parser.add_argument("--latent", type=Path, required=True) +parser.add_argument("--model", type=Path, required=True) +parser.add_argument("--output", type=Path, required=True) +parser.add_argument("--scale", type=float, default=2.0) +parser.add_argument("--precision", choices=("float16", "bfloat16", "float32"), default="float16") +args = parser.parse_args() + +dtype = getattr(torch, args.precision) +state = torch.load(args.latent, map_location="cpu", weights_only=False) +if not isinstance(state, dict) or "latent" not in state: + state = {"latent": state} +source = state["latent"] + +torch.cuda.synchronize() +started = time.perf_counter() +model = load_h3_latent_upscaler(args.model, dtype=dtype) +torch.cuda.synchronize() +loaded = time.perf_counter() +upscaled = upscale_h3_latent(model, source, scale=args.scale).cpu() +torch.cuda.synchronize() +finished = time.perf_counter() + +result = dict(state) +result["latent"] = upscaled +result["width"] = upscaled.shape[-1] * 16 +result["height"] = upscaled.shape[-2] * 16 +result["upscale"] = { + "model": args.model.name, + "scale": args.scale, + "precision": args.precision, + "source_shape": tuple(source.shape), + "output_shape": tuple(upscaled.shape), +} +args.output.parent.mkdir(parents=True, exist_ok=True) +torch.save(result, args.output) +print(json.dumps({ + "output": str(args.output), + "source_shape": tuple(source.shape), + "output_shape": tuple(upscaled.shape), + "model_load_seconds": loaded - started, + "upscale_seconds": finished - loaded, +}, indent=2))