Skip to content

[WS1][CUDA][MiniMax-H3] generation gate: 50-block generation compare and resume check (RFC #420 stack 10c/10) - #588

Draft
fusheng-ji wants to merge 239 commits into
RL-Align:test-h3from
fusheng-ji:feat/h3-generation-gate
Draft

fusheng-ji wants to merge 239 commits into
RL-Align:test-h3from
fusheng-ji:feat/h3-generation-gate

Conversation

@fusheng-ji

@fusheng-ji fusheng-ji commented Oct 25, 2026 •

Copy link
Copy Markdown
Contributor

RFC #420 stack 10c/10 (base test-h3, merge bottom-up): … → #494 → 10a (#586) → 10b (#587) → this PR (10c). This PR merges after 10b. This PR's own diff (1.5k lines). The H3 PRs #567 and #568 → #575 are stacked on the top of this layer.

Replaces #539 (closed, superseded by this stack). The split is described in 10a. The top of this stack has the same tree as #539's head 17d0e9e9 apart from nine docs/operators/h3-*.md files whose commit-hash text was updated in #483–#494 after #539 was cut (17 lines).

Evidence. Measurement outputs (logs, JSON, figures) are not part of this diff. They are on the fork's pr_assets branch and every link is fixed to commit 967000e9. Runs on the heads of this stack: B200 pending/b200-split-539, H100 pending/h100-split-539 (each with tests.json, redaction.json and SHA256SUMS). Evidence of the earlier #539 runs: pr539/.

This stack

layer PR own diff content
10a #586 3.6k block 0 replay and diagnostics
10b #587 8.8k full-model qualification
10c #588 1.5k generation gate

What

RFC #420 closeout row ws1_dual_trajectory, generation part: three-seed generation with the candidate and its comparison against the provider, built to survive preemption. Not claimed as passed (see Status).

  • Generation follows examples/models/minimax_h3/generation-contract.json: the pinned public reference image and prompt, seeds 42/43/44, 124 frames, 1344×768, 24 fps with joint 32 kHz audio. generation-compare.sbatch and h3_generation_compare.py check identical initial noise and shapes and report PSNR, SSIM, temporal RMSE and audio differences against the provider.
  • Resumable generation (h3_full_generation.py, rl_engine/validation/models/h3_generation_resume.py). Candidate generation takes about 20 hours per seed, so it runs as one resumable task per seed. Each saved step has a sidecar with the digest of the transformer's inputs and of the saved file, written atomically. On resume the saved transformer calls are replayed from disk and the pipeline's own scheduler turns them into the next latents; every replayed call's input digest is checked. Recomputing the last saved step and requiring a bitwise match is --verify-last-step (opt-in).
  • A run is tied to the git trees of the model paths, not to a commit, so a commit that changes only orchestration code can continue it; h3_generation_adopt.py moves an older commit-based run over, only if those paths are byte-identical.
  • Job scripts: generation.sbatch runs one seed per array task and refuses to start unless the numerical gates of 10b passed; generation-resume-check.sbatch checks resume; generation-compare.sbatch compares afterwards.
  • --provider-variant on h3_full_generation generates the provider's videos with one SDPA backend forced on the transformer's forward only, for the quality-bounds comparison below.

Prior art & reuse decision

No operator is written here. The reference is the diffusers pipeline; its own generation is the baseline, including variants with one SDPA backend forced.

Results

B200

Generation from clean checkouts at 4792c8a9 / 2451ae5c (same model code), 124 frames 1344×768, 49 steps. Evidence on pr_assets: h3-full-model-b200/.

Gate Result
Generation preflight (128×128, one step) complete
Three-seed generation, 124 frames 1344×768, 49 steps complete
Generation quality vs the provider, bounds fixed before generation not passed: seed 44 passes, seeds 42 and 43 do not

Generation quality per seed (bounds: PSNR ≥ 30 dB, SSIM ≥ 0.99, temporal RMSE ≤ 0.01, audio relative L2 ≤ 0.05; a failing value is bold):

seed PSNR (dB) SSIM temporal RMSE audio relative L2 passed
42 37.58 0.9917 0.00991 0.0807 no
43 32.64 0.9757 0.01056 0.0762 no
44 38.75 0.9945 0.00734 0.0456 yes

The same bounds applied to the provider with one SDPA backend forced on the transformer (same seeds, inputs and conditioning):

provider variant seeds passing PSNR (dB), seeds 42 / 43 / 44 audio relative L2, seeds 42 / 43 / 44
flash attention 0 of 3 39.08 / 23.98 / 39.33 0.054 / 0.264 / 0.056
memory-efficient attention 0 of 3 39.43 / 30.22 / 42.59 0.094 / 0.135 / 0.116
candidate 1 of 3 37.58 / 32.64 / 38.75 0.081 / 0.076 / 0.046

First denoising step replayed from its saved packed inputs (nothing has fed back yet), relative L2 of the video output against the provider's saved step (replaying the provider is bitwise equal to it):

configuration (seed 43 unless noted) video relative L2
candidate (seeds 42 / 43 / 44) 0.0285 / 0.0212 / 0.0274
provider, flash / memory-efficient attention 0.0219 / 0.0218
provider, cuBLASLt 0 (bitwise equal at this step)
candidate with attention / the six GEMMs / SwiGLU / the AdaLN projection run as the provider 0.0225 / 0.0217 / 0.0198 / 0.0193
candidate with Q/K norm, the two norms or the residuals run as the provider 0.0212 (unchanged by the swap)
candidate with all interim operators / every node run as the provider 0.0192 / 0

H100 (Hopper)

A one-step generation preflight (128×128) completed on 4 × NVIDIA H100; it checks the flow only and is not quality evidence. The three-seed generation (about 20 hours per seed) was not run on H100.

Status

  • Generation quality: not passed. Against the provider, with the bounds fixed before candidate generation, seed 44 passes and seeds 42 and 43 do not: audio relative L2 is 0.081 and 0.076 (bound 0.05), and seed 43 also misses SSIM (0.9757) and temporal RMSE (0.01056). The bounds are unchanged.
  • What the comparison says about those bounds.
    • The bounds had only been checked against a bitwise-identical repeat of the provider (PSNR 300 dB, zero error). The provider with one SDPA backend forced on the transformer, run through the same 49 steps with the same seeds, inputs and conditioning, fails them too: flash attention 0 of 3 seeds, memory-efficient attention 0 of 3, and audio relative L2 exceeds 0.05 on every seed of both (0.054 to 0.264).
    • The first denoising step shows the same: the candidate differs from the provider's step by 0.021 to 0.029 relative L2 on video, and the provider with flash or memory-efficient attention differs from its own default by 0.022. Running any one group of block nodes as the provider's leaves 0.019 to 0.023; replacing the norms or the residuals changes nothing at this step.
    • This is not a claim that the candidate is better than the provider variants: on seeds 42 and 44 its PSNR and temporal RMSE are slightly worse than both variants', on seed 43 better, and it passes one seed where they pass none.
    • The generation bounds were not changed and the candidate is not accepted on generation quality. A rule that tolerates a change of kernel, like the one for the numerical gates in 10b, would be a second rule designed after seeing a failure; it is left to the reviewers.
  • Resume checks. On the real pipeline at preflight size (128×128, 6 sampling steps), one uninterrupted run and runs preempted and resumed give byte-identical step files, initial inputs, video and audio (8 files); a resume with nothing left to compute and the opt-in verification both reproduce the uninterrupted run.
  • Cost. Candidate generation is more than 50× slower than the provider's (about 20 hours per seed, against about 1 h for the provider's).

Tests

Layer head 6fe346df; same commands as 10a.

RL_KERNEL_H3_STRICT=1 python setup.py build_ext --inplace
python -m pytest tests/models/minimax_h3 tests/backends/cuda/test_h3_full_attention.py -q
Check B200 H100
tests/models/minimax_h3 + test_h3_full_attention.py (2 GPUs) 571 passed, 4 skipped (sglang and matplotlib on the path) 573 passed, 2 skipped
full tests/ (minus tests/distributed and a ROCm import-error file) 73 failed, 3541 passed; the same 73 fail on upstream/test-h3 (measured on the heads before the re-cut; same layer code)
pre-commit, mypy clean

Skips as in 10a. tests/models/minimax_h3/test_h3_generation_resume.py (22 CPU tests, included in the counts) drives a fake denoising loop through the same resume code: a run preempted once or repeatedly equals an uninterrupted one, nothing is recomputed by default, the opt-in verification catches a drifting recompute, other identities, corrupted files and signature changes are rejected, and adopting a commit-based run is tested.

Scope / Known limitations

  • Generation quality is not passed (1 of 3 seeds), against bounds that the provider with another attention kernel does not pass either.
  • The three-seed generation was run on B200 only.
  • CUDA only; not run on AMD hardware.

…+ shared H3 harness

RFC RL-Align#420 step 3, row `timestep_sinusoid_h3`.

- CUDA kernel: one thread per (t, k), the provider's exact FP32 op
  sequence (__fmul_rn, precise expf/sinf/cosf); bitwise equal to
  diffusers' get_timestep_embedding on CUDA. Analytic row-local backward.
  Non-1-D, empty, integer, non-finite or out-of-[0, 1] timesteps fail
  closed (RFC probe H10).
- PyTorch reference (provider replay + FP64 golden), registry, gtest.
- Shared H3 harness used by the following rows:
  * h3_manifest.json + prepare_h3_weights.py: pinned revision and sha256,
    only shard 1 of 14 is downloaded;
  * h3_provider.py: diffusers replayed op for op;
  * h3_chain.py: stage-wise chain replay, each stage declaring its
    promises (provider-bitwise, golden tolerance); used by
    scripts/h3_chain_replay.py and tests/h3/test_h3_conditioning_e2e.py;
  * h3_report.py + scripts/h3_evidence.py + scripts/plot_h3_evidence.py:
    per-op performance/accuracy JSON and figure;
  * benchmarks/benchmark_h3_conditioning.py.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Produced by scripts/h3_evidence.py at 0522865 on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…p MLP

RFC RL-Align#420 step 3, row `timestep_mlp_fp32` (stacked on timestep_sinusoid_h3).

- New CUDA module csrc/cuda/h3/det_linear.cu, contract h3-det-linear-v1:
  warp per output-column group, lane-strided 16-byte K chunks, ascending
  fmaf chains, xor butterfly, bias then fused SiLU. The launch shape never
  changes a column's summation order, so rows are batch/position
  invariant and runs are repeat-bitwise.
- Backward without cuBLAS or atomics: dinput via fixed 64-row N chunks
  folded in order (row-local), dweight/dbias as ascending-row FP32 folds.
- PyTorch reference (provider path + FP64 golden); BF16 anywhere is
  rejected (RFC probe H7: the time_embedder is a declared FP32 module).
- On the pinned weights both CUDA and cuBLAS are ~1e-6 from the FP64
  golden (100x inside the contract); the op is 2x faster for T >= 2.
- Chain stage (covered by the end-to-end test), report/figure hooks,
  registry, gtest, doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Produced by scripts/h3_evidence.py at 65ef7f6 on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…rojection

RFC RL-Align#420 step 3, row `adaln_projection_3mod` (stacked on timestep_mlp_fp32).

- SiLU at temb's FP32 precision and exactly one cast to the projection
  dtype (bitwise equal to the provider activation); a BF16 temb is
  rejected (RFC probe H7).
- 2688 -> 96768 projection on a BF16 tensor-core path in det_linear.cu,
  contract h3-det-linear-bf16-mma-v1: warp = 16 output columns x 8 input
  rows (zero-padded), K in ascending 16-wide groups, each mma.sync
  m16n8k16 from a zero accumulator added into an FP32 running sum. Same
  instruction sequence for every T <= 8, so rows are bitwise
  batch/order/tile invariant; 99.98% of outputs correctly rounded vs
  99.84% for cuBLAS; kernel time flat at ~78 us for T = 1..8. SM80+.
- Six (3T, 5376) views in diffusers' view(-1, 6H).chunk(6) layout, row
  t*3+m; a channel-tagging test pins chunk/modality placement (probe H4).
- Backward: FP32 d_temb through the identity-VJP cast, ascending-row
  dW/db folds, no cuBLAS or atomics. The FP64 golden applies the cast
  straight-through so its own gradient is not rounded to BF16.
- Chain stage (end-to-end test), report/figure hooks, registry, gtest, doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Produced by scripts/h3_evidence.py at d06e120 on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
RFC RL-Align#420 step 3, row `adaln_row_gather` (stacked on adaln_projection_3mod).

- Forward: one launch gathers all six (S, H) modulation tensors by
  timestep_index * 3 + token_tag with 16-byte copies; bitwise equal to
  diffusers' six index_select calls. Out-of-range tags or timesteps fail
  closed (RFC probes H2/H3), never clamped.
- Backward: deterministic FP32 segmented sum (stable sort by row, fixed
  256-position tiles, ordered fold, one cast; no atomics) instead of
  index_select's atomic BF16 scatter-add. The PyTorch reference gets an
  equally deterministic backward.
- H3AdaLNModulationCudaOp: projection + gather as one autograd node,
  bitwise-equal forward, FP32 table gradient into the projection backward.
- Chain replay gains the gather stage and a backward mode (separate,
  fused and diffusers chains vs an FP64 golden); the end-to-end test now
  covers the whole chain forward and backward.
- Report/figure hooks, registry, gtest (reduction class: the VJP is a
  segmented sum), doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…y and figure

Produced by scripts/h3_evidence.py and scripts/h3_chain_replay.py --backward at fa551c2 on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
RFC RL-Align#420 step 4, row `h3_rmsnorm` (stacked on adaln_row_gather).

- CUDA kernel replaying PyTorch's own RMSNorm reduction order
  (vectorized_layer_norm_kernel: (32, 4) block per row, 4-element
  vectors, shuffle-down and cross-warp trees, rsqrtf, w * (rstd * x)):
  bitwise equal to nn.RMSNorm on all four pinned H3 norm weights and
  batch/position invariant.
- Fused AdaLN modulation n * (1 + scale[i]) + shift[i] with the rows
  gathered in-kernel from strided table views and each step rounded where
  the eager expression rounds: bitwise equal to diffusers for the block
  (adaln_indices) and norm_out (timestep_indices), without materialising
  the (S, H) gathered tensors; 2.5x faster forward at S = 131072.
- Deterministic backward: row-local dx, dweight over fixed 256-row tiles,
  dshift/dscale as sorted segment sums (diffusers: non-deterministic BF16
  atomics, ~15x larger error).
- Manifest pins the block/refiner-final/norm_out norm weights and
  norm_out.linear (all in shard 1); chain gains a per-mode history and the
  block norm1 + modulation stage (covered by the end-to-end test).
- Registry, gtest, report/figure hooks, doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Produced by scripts/h3_evidence.py at f36aa60 on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…ernel gate gather

RFC RL-Align#420 step 6, row `adaln_gate_residual` (stacked on h3_rmsnorm).

- residual + gate[adaln_index] * sublayer_output in one pass, the gate row
  read from the AdaLN table view in the kernel; the two roundings sit where
  the eager expression rounds (__fmul_rn/__fadd_rn, no FMA contraction), so
  the forward is bitwise equal to diffusers in bf16/fp16/fp32. 1.7x faster
  forward at S = 131072.
- Backward: d_residual = grad; d_sublayer bitwise equal to the eager VJP
  (exact product, one rounding); d_gate a deterministic FP32 segment sum
  (diffusers: non-deterministic BF16 atomics, 13-49x larger error, varying
  run to run). The PyTorch reference gets an equally deterministic gate
  gradient.
- Chain stage residual + gate_msa * (stand-in attention output), covered
  by the end-to-end test; registry, gtest, report/figure hooks, doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Produced by scripts/h3_evidence.py at a26b41a on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…norm/modulation

RFC RL-Align#420 step 7, row `final_adaln_out` (stacked on adaln_gate_residual).

- norm_out as one autograd node: FP32 SiLU -> one BF16 cast -> the
  deterministic tensor-core GEMV on norm_out.linear (shift first), then
  the h3_rmsnorm kernel indexed by timestep_indices. Forward equal to
  diffusers on 99.96% of elements (1-ULP projection ties); rows bitwise
  batch/position invariant; 2.4x faster forward at S = 131072.
- Backward keeps the table gradient in FP32 into the projection backward:
  d_temb/dW/db 13-57x closer to FP64 than diffusers (whose error varies
  run to run), all gradients repeat-bitwise.
- Golden rounds only at the declared boundaries (SiLU cast, BF16 table),
  straight-through for the gradient; gtest inputs at realistic scale.
- Chain stage: the gated residual stream -> norm_out with the MLP stage's
  temb (end-to-end test now runs timestep -> ... -> norm_out); registry,
  report/figure hooks, doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
… and figure

Produced by scripts/h3_evidence.py and scripts/h3_chain_replay.py --backward at 38d575f on a clean tree; rendered with scripts/plot_h3_evidence.py and embedded in the operator doc.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>

# Conflicts:
#	csrc/ops.cpp
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>

# Conflicts:
#	scripts/plot_h3_evidence.py
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Validate pinned tensor contents at preparation and load time. Execute chain prerequisites for selected stages and require checkpoint weights for weighted evidence. Reject invalid CUDA tile and backward inputs, align the CUDA wrapper guards, skip unsupported devices, and rotate comparative timing with raw samples.

Validation: 143 H3 tests passed on B200 after a fresh CUDA build; illegal FP32 tiles fail compilation; isolated AdaLN replay and pinned-weight evidence generation passed; Black, isort, Ruff, and Flake8 passed.
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Check timestep and modality bounds on the device before computing gather addresses, including unchecked calls and overflowing int64 indices. Reject non-prefix replay stage selections and partial backward chains before loading weights.

Validation: rebuilt the CUDA extension on B200; tests/h3: 175 passed, including 12 isolated native bounds probes and 16 CLI regressions.
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
CodeRabbit skips automatic reviews targeting non-default branches by default. Allow the test-h3 base branch so pushes to the H3 stack receive incremental reviews.

Validation: configuration validated against the official CodeRabbit JSON schema.
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Include the RFC 420 target branch in CodeRabbit automatic reviews so pushes to the H3 stack are reviewed instead of silently skipped.

Validation: configuration parses and validates against the published CodeRabbit v2 schema; the additional branch pattern matches only test-h3.
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
The raw results (JSON, figures, audit files) now live on the fork's pr_assets branch;
the documents link to them at a fixed commit.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
The history of this branch was rewritten on 2026-10-23 (evidence directories removed from every
commit, sign-offs added). Commit hashes quoted in the reports and documents now name the
rewritten commits; the evidence archives still record the pre-rewrite hashes.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
… and diagnostics (RFC RL-Align#420, RL-Align#539 split 1/3)

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…e rule and reference manifest (RFC RL-Align#420, RL-Align#539 split 2/3)

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…and resume check (RFC RL-Align#420, RL-Align#539 split 3/3)

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@coderabbitai

coderabbitai Bot commented Oct 25, 2026

Copy link
Copy Markdown

Important

Draft PR not reviewed

Draft PRs are not automatically reviewed by default.

  • Trigger a manual review

To automatically review draft PRs, update your CodeRabbit configuration:

reviews:
  auto_review:
    drafts: true
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Comment @coderabbitai help to get the list of available commands.

This was referenced Oct 25, 2026
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…fication

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…on-gate

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant