daniel/h3-blackwell-runtime
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Expand H3 runtime acceleration and validation

Browse source
This commit is contained in:
Daniel Maddern 2026-08-20 22:08:52 +07:00
parent 837c16871d
commit 1d5faa8f16
24 changed files with 1575 additions and 8 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
Dockerfile.spark
View file

@ -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

71
FLASH4.md Normal file
View file

@ -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: <https://github.com/Dao-AILab/flash-attention>
- 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.

149
H3_LATENT_UPSCALER.md Normal file
View file

@ -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:
- <https://github.com/LBH-123-AI/Comfyui_Minimax_h3_latent_Upscaler>
- <https://huggingface.co/LBH-123-AI/Minimax_h3_latent_Upscaler>
- 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.

159
H3_PROMPT_GUIDE.md Normal file
View file

@ -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 `<d>`:
```text
The woman with a low, clear voice and measured pace (S1) says in a flat,
matter-of-fact delivery: <d>[English] The meeting starts at three.</d>
```
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:
<d>[English] The meeting starts at three.</d> 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 `<d>[Language] exact words</d>`.
- Dialogue is not enclosed in double quotes.
- Delivery and voice are outside `<d>`.
- 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:
<https://huggingface.co/MiniMaxAI/MiniMax-H3/blob/main/docs/VIDEO_PROMPT_WRITING_GUIDE_base_en.md>
- Reference prompt guide:
<https://huggingface.co/MiniMaxAI/MiniMax-H3/blob/main/docs/VIDEO_PROMPT_WRITING_GUIDE_ref_en.md>
- Official prompt-writing skill:
<https://github.com/MiniMax-AI/MiniMax-H3/tree/main/.agents/skills/h3-prompt-writing>
- Official model repository:
<https://huggingface.co/MiniMaxAI/MiniMax-H3>

30
PLAN.md
View file

@ -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`.

38
README.md
View file

@ -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.

349
benchmarks/t2va-dialogue-864x480-141f-base12-sdpa-seed440420.json Normal file
View file

@ -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: <d>[English] The north entrance closes at six.</d> 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: <d>[English] Then we should leave by five thirty.</d> 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
}
}
}

46
compose.qwen38-vllm.yml Normal file
View file

@ -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

84
compose.spark-stack.yml Normal file
View file

@ -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"

4
compose.spark.yml
View file

@ -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"]

3
pyproject.toml
View file

@ -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"]

30
src/h3_blackwell_runtime/attention.py
View file

@ -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()

163
src/h3_blackwell_runtime/latent_upscaler.py Normal file
View file

@ -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)

43
src/h3_blackwell_runtime/runtime.py
View file

@ -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,

21
tests/test_fl2va_contracts.py
View file

@ -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)

52
tests/test_latent_upscaler.py Normal file
View file

@ -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()

27
tools/poll_qwen38_health_spark.sh Normal file
View file

@ -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

16
tools/serve_hot_runtime.py
View file

@ -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,

50
tools/smoke_flash4.py Normal file
View file

@ -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))

8
tools/smoke_qwen38_vllm_spark.sh Normal file
View file

@ -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

61
tools/spark_stack_spark.sh Normal file
View file

@ -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

27
tools/start_qwen38_vllm_spark.sh Normal file
View file

@ -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

91
tools/submit_h3_comfy_t2v.py Normal file
View file

@ -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()

55
tools/upscale_video_latent.py Normal file
View file

@ -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))