With enable_torch_compile_vae (#1734, opt-in) inductor autotunes kernels per
process, so chunk decodes on other ranks differ from the serial rank's decode
the way two serial processes differ (GB200 @124f: max 63/255 on <0.5% of
pixels, mean ~1e-2/255; audio and chunk 0 bit-identical). Eager decoder (the
default) stays bitwise-equal to serial decode_to_pixels - measured, both
strategies, x3, 124f+345f.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two schedule fixes sized from the first GB200 tray measurements (job 2659,
serial 7.7s/21.2s at 124f/345f):
1. Every rank now decodes its round-0 chunk BEFORE the metadata broadcast.
Non-leader ranks previously blocked on the broadcast until the leader
finished chunk 0, serializing a full extra chunk-decode into round 0
(visible as 1.9x instead of ~2.6x at 7 chunks / 4 ranks). Same reorder
on the encode path.
2. The leader's per-chunk joining work (blend, denormalize, clamp, output
copies) moves to a dedicated CUDA side stream. It depends only on
already-gathered segments, but on the main stream it delayed the
leader's next-round decode and therefore every rank's next collective
(~0.1s/chunk on the critical path). Gathered storage is pinned to the
assembly stream via record_stream; the driver drains the stream in a
finally so an exception cannot leave an in-flight DMA into the output
buffer. Stream placement does not change op order or values, so the
bitwise-parity contract is untouched.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Under SP>1 the H3 video VAE decoded all temporal chunks serially on the
output rank while the other ranks idled (#1703's gate), and every rank
encoded the full reference video redundantly. Chunk decodes and clip
encodes have no cross-chunk data dependency - only the joining (overlap
blend, trim, denormalize, moment concat) is sequential - so both are
round-robined across the sequence-parallel ranks:
- fastvideo/models/vaes/minimax_h3_parallel.py: decode_to_pixels_parallel
gathers each round's decoded segments (body+halo tail, one contiguous
slice per chunk) to the SP group's first rank via NCCL gather (or
all_gather, FASTVIDEO_VAE_PARALLEL_DECODE_STRATEGY), which replays the
serial blend/trim/denormalize/copy semantics with the same VAE methods -
bitwise-equal to serial decode_to_pixels by construction. Placeholder
rounds keep collective participation uniform; the leader decodes chunk 0
first and broadcasts dtype/shape metadata so placeholders never guess the
autocast dtype. encode_pixels_parallel all-gathers per-clip moments
(latent-sized) so every rank keeps the identical full posterior,
preserving the all-ranks-hold-latents contract.
- decoding stage: output gate moves from world rank 0 to the SP group's
first rank (identical in the single-group e2e case; correct for the
trainer validation callback, which consumes each group leader's batch);
with vae_parallel_decode every rank enters the decode body so no
rank-dependent branch guards the collectives.
- latent preparation: opt-in clip-parallel reference encode on the same
seam (vae_parallel_encode).
- knobs: FastVideoArgs.vae_parallel_decode/encode (+ --vae-parallel-decode,
--vae-parallel-encode, FASTVIDEO_VAE_PARALLEL_DECODE/ENCODE env
parse-once adapters), default OFF.
- _copy_chunk_pixels factored out of _decode_to_pixels so serial and
parallel share one output-copy path (behavior unchanged).
Tests: threaded fake-group CPU suite drives the real SPMD functions
end-to-end (world sizes 2-5, both strategies, pad/blend/trim geometries,
token_drop=0, batched slicing, placeholder rounds) bit-exact vs the serial
APIs; GPU regression (torchrun world>1 gated) asserts bitwise parity under
fp16 autocast with real NCCL plus repeat-determinism x3.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Carried blocker fix discovered while benchmarking this stack (GB200 job 2594,
all four legs): with text_encoder_cpu_offload=True - the FastVideoArgs DEFAULT
and the shipped H3 example configuration - TextEncoderLoader FSDP2-fully_shards
the H3 conditioner, and #1732's @torch.inference_mode() on
MiniMaxH3Qwen3VLConditioner.encode_ids then kills the very first encode:
File torch/distributed/fsdp/_fully_shard/_fsdp_param_group.py, in
wait_for_unshard: with torch.autograd._unsafe_preserve_version_counter(t):
RuntimeError: Inference tensors do not track version counter.
FSDP2's lazy unshard runs inside the inference-mode region, so its all-gather
tensors are inference tensors, and the version-counter preservation hook
cannot read t._version. The PR's own benchmarks ran with offload disabled,
which is exactly the unexercised-default-path gap called out as finding 2 of
the pr1732 review (there for FP8; the same gap bites plain bf16 via
inference_mode).
Fix: @torch.no_grad() instead. It frees the same activation memory (the
-262 MiB claim comes from the early-return slim contract, not from
inference_mode's bookkeeping), is fully FSDP2-compatible, and as a bonus
retires review finding 7: prompt_embeds are ordinary tensors again, so any
future on-the-fly-conditioning training can backprop through them without a
clone at the stage boundary.
Verified: 1x GPU FastH3 leg boots and generates after this change (leg reruns
on GB200); the 1732 unit suites (truncation + checkpoint-fp8) still pass.
Note: fastvideo/tests/stages/test_text_encoding.py has 7 pre-existing failures
that reproduce byte-identically on plain origin/main 56d4a6074 (unrelated
generic-stage tests; not introduced by the stack, verified by A/B).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Carried blocker fix for PR #1734 on this integration stack, per pr1734.md
finding F1 (reproduced on GB200/torch 2.12): with
@torch.compile(mode="reduce-overhead") on _stitch_tiles, the stitched canvas
is a CUDA-graph static buffer, and the collect-then-torch.cat consumers -
_decode (via _decode_chunks), _encode, _encode_pixels, encode_keyframe - hold
each chunk/clip result across the next _stitch_tiles replay, which overwrites
the pooled storage. First tiled decode() with >=2 temporal chunks (any real
video) and tiled encode() of >17 frames raised:
RuntimeError: Error: accessing tensor output of CUDAGraphs that has been
overwritten by a subsequent run. ... line ..., in _stitch_tiles:
return torch.cat(result_rows, dim=-2)
Fix: .clone() the stitch output at the two eager call sites (_encode_clip
tiled return, _decode_clip tiled return) so no cudagraph-owned storage
escapes the tile driver; the clone is read before the next replay, so it is
race-free by construction. The streaming _decode_to_pixels path was already
safe (full copy-out per chunk before the next decode) and stays correct at
one extra D2D copy per chunk (~tens of microseconds vs the decode compute).
This follows the option (a) recommendation in the review; reduce-overhead is
retained on _stitch_tiles/_project_decoder_tile.
Tests:
- test_decode_clip_emits_tiled_stage_ranges updated: the tile driver now
returns a caller-owned copy (is-not + equal), pinning the ownership
contract at the mock level.
- New CUDA regression gate (the review's must-add test):
test_tiled_decode_and_encode_survive_cudagraph_buffer_reuse_on_cuda - real
tiled decode() (2 temporal chunks, 2x2 spatial tiles) and encode()/
encode_pixels() (2 clips) with _stitch_tiles unmocked, asserting bitwise
repeat-consistency plus value parity against the fully eager tile helpers
(via _torchdynamo_orig_callable). Red on the unfixed merge (cudagraph
overwrite RuntimeError on the first decode), green with this fix.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
[perf] MiniMax-H3 VAE decode optimization: _compile_conditions for the video +
audio VAE decoders (makes enable_torch_compile_vae effective for H3),
torch.compile on _stitch_tiles/_project_decoder_tile, VAE attention through the
FastVideo selector (TORCH_SDPA/FLASH_ATTN), opt-in FASTVIDEO_NVTX_PROFILE
ranges, per-worker attention-backend receipt.
Clean textual merge; semantic overlaps verified by hand:
- minimax_h3_conditioning.py: 1732's rewritten _encode_fl2va/_encode_ref2va
kept their names, 1734's nvtx_range wrap composes over them (checked).
- minimax_h3.py: 1735 fusion routing + 1734 per-block nvtx_range coexist
(fusions at attention/mlp/modulate seams, nvtx at the block loop).
- component_loader.py: 1732 quant plumbing at load_model (~L369) vs 1734
backend receipt (~L1103) - disjoint.
- envs.py: FASTVIDEO_NVTX_PROFILE (1734) + FASTVIDEO_MINIMAX_H3_FUSIONS (1735)
are distinct additions.
- minimax_h3_video.py: 1734 was authored on a base that already contains the
merged #1703 (incl. pr1703-fixes content), so the 1734-vs-1703 conflict the
reviews predicted was pre-resolved by the author's rebase.
KNOWN CARRIED DEFECT at this point in the stack: review finding F1 of
pr1734.md - reduce-overhead CUDA-graph outputs of _stitch_tiles escape into
collect-then-cat consumers (_decode/_encode/_encode_pixels), crashing tiled
eager decode/encode on CUDA. Fixed in a follow-up commit on this branch.
[perf] Opt-in MiniMax-H3 Sol-Engine Triton fusions (FASTVIDEO_MINIMAX_H3_FUSIONS):
fused rmsnorm+modulate, residual+gate+rmsnorm+modulate, per-head qknorm+partial
RoPE, and packed SwiGLU. Default-off; fused path is disclosed NON-PARITY
numerics (same-seed decoded SSIM 0.7403 vs eager per the PR body).
Head cbab605ef already carries the review's F1 blocker fix upstream
(ac98869aa: int64 row offsets in the fused qknorm+RoPE kernel - verified
present at qknorm_rope.py:43, tl.program_id(0).to(tl.int64)) plus the
engagement-test/logging hardening (cbab605ef), so no fix needs to be carried
by this stack for #1735.
Clean merge, no conflicts. Note: #1732 does not touch minimax_h3.py or envs.py
(its true merge-base is 0462e1b0e; earlier apparent overlap was #1712/#1290
noise from diffing against the wrong base), so 1735's DiT + envs.py edits had
no sibling edits to reconcile.
[perf] Optimize MiniMax-H3 text encoder memory: slim single-tensor forward
contract for the Qwen3-VL conditioner (early return at the layer-50 tap,
-262 MiB peak on top of merged #1711), TextEncoder base relaxed to
Generic[TextEncoderOutputT], and opt-in checkpoint-serialized block-FP8
(new text_encoder_quantization.py + minimax_h3_checkpoint_fp8.py).
Clean merge, no conflicts (PR base c4ad4227c == main's state for all touched
files; the tests/local_tests/minimax_h3/README.md hunk applied cleanly on top
of the #1703 dedupe).
Review status (pr1732.md): the two blockers are sm12x/FP8-only - the cutlass
m%4 crash is on the sm12x route (GB200/sm100 uses trtllm, verified clean at
m=559), and the FP8+cpu-offload gap only matters with FP8 enabled. This stack
keeps text-encoder FP8 OFF for all benches, so neither blocker is reachable.
- _can_run_minimax_h3_fusion now also requires Triton availability, so an
enabled fusion on a CUDA build without a working Triton falls back to
eager instead of hitting the strict wrappers' RuntimeError mid-forward.
- One-time logger.info of the resolved fusion set at model init (and a
warning when the set is requested without Triton), plus a
prepare_for_compile hook warning that torch.compile capture makes the
fusions inert inside compiled block forwards.
- Positive-engagement routing test: counts 1/1/2/1 fused-kernel calls in
one CUDA inference block forward and asserts a grad-enabled forward
leaves the counters unchanged, so a guard regression to always-eager
can no longer pass silently.
- Document the index-bounds contract ([0, table_rows)) on the modulation
wrappers; a device-side check would synchronize, and in-model callers
are safe by construction.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
_qknorm_partial_rope_kernel left tl.program_id(0) in int32, so
row * head_dim wrapped once the flattened input reached 2**31 elements
and the kernel read/wrote out of bounds (CUDA illegal memory access).
The PR's other two kernels (modulation.py, swiglu.py) already cast
tl.program_id(0).to(tl.int64); this one now matches, and seq_index /
table_offset inherit int64 from row.
Confirmed on GB200: (1, 8_500_000, 2, 128) bf16 (2.176e9 elements)
crashed before the cast and matches eager after it; the just-under-2**31
control shape matched all along. For H3 (56 heads x 128 head_dim) the
boundary is batch*seq >= 299_593 tokens per rank, reachable at SP=1.
Adds a GPU regression test at the over-2**31 shape that compares the
head and tail rows against eager.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
