Repository navigation
feat(h3): extend block validation to full-model and generation qualification - #539
Open
fusheng-ji wants to merge 218 commits into
Open
fusheng-ji wants to merge 218 commits into
fusheng-ji wants to merge 218 commits into
Conversation
…+ 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>
Brings in tools/validation/models/h3_bitwise_reuse.py (direct bitwise comparison of reuse candidates) and its evidence. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Written by tools/validation/models/h3_bitwise_reuse.py from a clean tree at 4430ef1. There is no reuse candidate beyond diffusers' MiniMaxH3AdaLayerNormOut. This op does not promise bitwise equality: its shift/scale projection uses RL-Align#485's fixed-order linear, not cuBLAS. The output differs from diffusers in 0.05-0.07% of BF16 elements at every S (1 to 32768), and every gradient differs as expected for the deterministic FP32 reductions. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in tools/validation/models/h3_bitwise_reuse.py (direct bitwise comparison of reuse candidates) and its evidence. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in tools/validation/models/h3_bitwise_reuse.py (direct bitwise comparison of reuse candidates) and its evidence. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in tools/validation/models/h3_bitwise_reuse.py (direct bitwise comparison of reuse candidates) and its evidence. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
test_backward_against_golden rounded an FP64 SiLU to BF16 for its declared cast, while the kernel and diffusers round the FP32 SiLU. Near a BF16 tie the two pick different neighbours and shift a whole dW column. The test was seeded (seed 5 scores 1.0), so it never failed, but over 300 temb draws 8 fell below 0.9999 (worst 0.999646). With the FP32-SiLU golden every dW/db element is correctly rounded for all 300 draws; d_temb is unchanged. Same fix as test_h3_adaln_modulation.py. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's FP32-SiLU golden fix (8bb2e09). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
With the FP32-SiLU golden (8bb2e09), the BF16 weight and bias gradients equal the correctly rounded golden at every element for all 500 temb draws tried (T = 2, two products per element). Replace the '> 0.999 of elements' threshold with torch.equal; the FP64 closeness check stays. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection test's exact dW/db equality check (8195820). Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…uire 0.9995 The two forward checks compared the kernel with forward_fp32, which rounds an FP64 SiLU to BF16 while the kernel and diffusers round the FP32 SiLU. Near a BF16 tie that shifts a whole output column: over 300 temb draws the old golden matched as little as 95.7% of the outputs, so the '> 0.999' checks only passed for their fixed seeds. A test-local _declared_golden now rounds F.silu(temb.float()). Exact equality is not reachable: the FP32 tensor-core accumulation over K = 2688 leaves 32-104 outputs off on every draw (ties, and a few ULPs where near-zero sums cancel). Over 300 draws per T = 1..4 the worst match is 0.999618-0.999731, so both checks now require >= 0.9995. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Brings in the projection forward-golden fix and 0.9995 threshold. Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Current qualification status — 2026-10-13
Full H3 candidate acceptance remains incomplete. Independent reference job 1106656 completed from clean source
2681cadb: FP32/BF16 runs across 50 distinct transformer blocks, both refiners and boundaries, with 638 parameter-gradient artifacts per precision. Reference generation and repeats passed for seeds 42/43/44 at 124 frames, 1344×768, 24 fps and 32 kHz audio. All 49 predicted-velocity steps matched their repeated provider controls; these are provider results, not candidate acceptance.Candidate source
98b07a3fcommits the unchanged 764 tensor bounds and a 1,635-file SHA-256 reference inventory before execution. The runner builds an independent H3 CUDA extension, verifies native identity and all consumed reference/model hashes, then gates block-0 golden tests, complete forward/backward, numerical comparison, preflight and three-seed candidate generation. Actual per-layer candidate-node calls are mandatory; fallback or missing nodes fail. First drift is ordered by real forward/backward execution. Row/interim/reference labels and the existing quality thresholds are retained.98b07a3f, one B200, exited 1 after 15m34s. Native build and all reference/checkpoint hash checks passed. Block-0 golden regression reported 8 passed, 1 failed, 0 skipped:attn.norm_q.weightcandidate gradient relative error 0.03398705939217804 exceeds the unchanged bound 0.03044327425907444 (provider error 0.016962182839382958). Full-model candidate and generation stages correctly did not execute. The original logs/JUnit and binary are retained.9317a5e1, is queued for 34 single-node substitution controls (one provider node in the candidate graph, and the reverse). These intentional hybrids are diagnostic only and cannot establish candidate acceptance. The native binary is pinned to the failed run; thresholds remain unchanged.9317a5e1, lint, unit tests and docs all actually ran and passed. GPU qualification remains incomplete.GPU jobs use partition
overflow, nametest, 512 GB host memory and a 24-hour limit. Candidate numerical/quality results and raw candidate evidence remain outstanding. Historical block-0 results below do not qualify the complete model. Seeexamples/models/minimax_h3/README.mdfor the committed candidate gates.Historical implementation and evidence
Evidence directory migration
Commit
03cd824bmoves the one-blockreport.jsonandfigure.pngtoreports/experiments/h3-one-block-b200/, updates documentation and the plotting example, and records both moves in the layout audit map.Validation: both moved files retain identical SHA-256 hashes; the report retains its original
41e4293measurement revision and environment. JSON parsing, migration mappings, documentation link targets, Python syntax, Ruff lint/format checks, andgit diff --checkpass. The existing plotting CLI successfully renders a temporary copy of the report without overwriting committed evidence. Old paths remain only in the migration map.This follow-up changes static evidence locations and references only. No GPU numerical tests or benchmarks were rerun; the historical two-GPU results and four-/eight-GPU coverage limitations below remain unchanged.
Fast-math guard (2026-10-13)
From the #483 review.
KERNEL_ALIGN_USE_FAST_MATH=1adds--use_fast_mathto the wholerl_engine._Cbuild. That replacesexpf/sinf/cosf/rsqrtfand FP32 division with approximations and enables FTZ, which breaks the H3 kernels' bitwise promises. Before the fix, a fast-math build of the top of the stack failed 33 H3 tests: thetimestep_sinusoid_h3(#483) andh3_rmsnorm(#487) bitwise tests, and the chain/block tests built on them.det_linear,adaln_row_gatherandgate_residualstill passed.nvcc 13 defines no macro for fast math, so the guard lives in the build:
build_tools/extensions.pylists everycsrc/cuda/h3source inH3_BITWISE_SOURCESand leaves them all out of fast-math builds;KERNEL_ALIGN_WITH_H3_KERNELS, so dispatch falls back to the PyTorch references;require_h3_kernelsfixture, and a build test asserts that nocsrc/cuda/h3source survives a fast-math build.On this branch: merge
4ba3a8fband9cb10d03(block dispatch-check skip). Tested on 2× B200 with pinned weights (tests/models/minimax_h3,tests/build,tests/runtime/test_dispatch.py):_CKERNEL_ALIGN_USE_FAST_MATH=1Status
test_h3_adaln_modulation.py::TestFusedModulation::test_backward_keeps_table_gradient_fp32in [WS1][CUDA][MiniMax-H3] adaln_row_gather: fused six-way AdaLN row gather (RFC #420 stack 4/9) #486 (46347118), merged here. The test's golden rounded an FP64 SiLU where the kernel rounds the FP32 SiLU, andtembwas unseeded; the kernel itself was deterministic. Full suite at27bb0295on 2× B200: 550 passed, 7 skipped in the default build, and 201 passed, 356 skipped in the fast-math build, 0 failed.No known failures at
41e4293(code) /caaf754(evidence). 11 of the 17 block nodes run on interim operators or the provider replay until their rows land (see Bindings). The block-level promises below hold with those bindings.What
RFC #420 row
ws1_one_h3_block: one real H3 block, forward and backward, with every intermediate and a first-drift report.MiniMax-H3@42ed227runs as a 17-node graph, from the AdaLN projection to the second gated residual. It runs on packed FL2VA layouts whoseposition_idsand tags are bitwise equal to the pinned diffusers pipeline'sbuild_packed_sequence; the hashes are pinned in a CPU test.first_{,isolated_,repeat_,batch_}drift.temband all 12 block parameters (repeat, error vs golden);first_grad_{repeat,batch}_drift.block_tensors). They are extracted byprepare_h3_weights.py --blockfrom the same shard, so the existing conditioning file is unchanged.Bindings
adaln_projectionadaln_projection_3modnorm1,norm2h3_rmsnormresidual_attn,residual_mlpadaln_gate_residualq/k/v_proj,o_proj,ffn_gate_up,ffn_downh3_qkv_gemm,h3_attention_o_gemm,h3_ffn_gate_up_gemm,h3_ffn_down_gemmq_norm,k_normh3_qk_rmsnorm_d128H3RMSNormCudaOpon the 128-wide heads (bitwise equal tonn.RMSNorm)rope_q,rope_kh3_mm_rope_3axis_partialattentionh3_full_attentionDeterministicAttentionOp(causal=False)swigluh3_swigluSwiGLUCudaOpWhen an owning row lands, swapping its node's
candidatecallable is the whole integration, and the report entries become that operator's block-level evidence.Prior art & reuse decision
This is a validation closeout, not a new kernel; no operator is written here.
MiniMaxH3TransformerBlockis the executable reference. It is replayed op for op as the provider (h3_provider), which pins its dtypes and op order without importing diffusers.rl_engine/validation/models/chain_gate.py). Its method is reused: per-node digests and first drift in graph order. Its graph is Qwen3's (causal, no AdaLN, Qwen shapes), so the H3 block is built on theh3_chainstage pattern of [WS1][CUDA][MiniMax-H3] timestep_sinusoid_h3: FP32 timestep features + shared H3 harness (RFC #420 stack 1/9) #483.TritonDetGemmOp.__call__reduces K as a tree with BF16 nodes; its FP32-accumulate kernel is equally deterministic and row-invariant, matches cuBLAS's accuracy, and is 9–17× faster (B200, error vs FP64):Direct bitwise comparison of reuse candidates (2026-10-13). The H3 block is composed of the WS1 rows below it, and their reuse candidates are compared byte for byte with
tools/validation/models/h3_bitwise_reuse.py(committed in #483, merged down the stack). It runs every candidate thath3_prior_art.candidates()defines on identical inputs and compares every output and leaf gradient with diffusers' path and with the rl-kernel op. Each row's report, written from a clean tree on its own branch: #483timestep_sinusoid; #484timestep_mlp; #485adaln_projection; #486adaln_row_gather; #487norm_modulate; #488gate_residual; #489final_adaln_out.Results
Generated from a clean tree at
41e4293on an otherwise idle B200 (torch 2.13.0+cu130), withtools/validation/models/h3_block_replay.py --layouts tiny,small,medium:report.json.adaln_projection/adaln_projectionadaln_projection/adaln_projectionadaln_projection/adaln_projectionnorm1/norm2,q_norm/k_norm, RoPE, and both gated residuals. The first drift is the first reduction whose tree differs from cuBLAS.adaln_proj.weight,adaln_proj.biasandtemb. diffusers'index_selectbackward accumulates in BF16 atomics; the row operators use FP32 segment sums.attn.norm_q.weightat S = 3160 (9.5e-3 vs 5.5e-3). At S = 925 the same leaf is 2× more accurate than diffusers, so the 128-element gradient's error follows BF16 rounding, not one side.h3_full_attention.Tests
501 passed, 5 skipped on two B200s, including the 9 new tests in
test_h3_one_block.py. The new tests check:Scope / Known limitations
h3_mm_rope_3axis_partiallands.tembis the provider's (the conditioning chain is [WS1][CUDA][MiniMax-H3] timestep_sinusoid_h3: FP32 timestep features + shared H3 harness (RFC #420 stack 1/9) #483–[WS1][CUDA][MiniMax-H3] adaln_row_gather: fused six-way AdaLN row gather (RFC #420 stack 4/9) #486).ws1_full_h3_transformer.Documentation repair at
9317a5e1: replace 14 report links outside the MkDocs docs tree with immutable GitHub/raw URLs at the evidence-containing commit68e472b7. The strict remote docs build passed. No evidence files or numerical results were changed.