Skip to content

feat(h3): extend block validation to full-model and generation qualification - #539

Open
fusheng-ji wants to merge 218 commits into
RL-Align:test-h3from
fusheng-ji:feat/h3-one-block
Open

fusheng-ji wants to merge 218 commits into
RL-Align:test-h3from
fusheng-ji:feat/h3-one-block

Conversation

@fusheng-ji

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

Copy link
Copy Markdown

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 98b07a3f commits 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.

  • Failed: candidate job 1107493, clean source 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.weight candidate 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.
  • Diagnostic completed; localization continues: job 1107566 completed in 25 seconds, reproducing the block q-norm gradient failure. With identical inputs/upstream gradients, isolated Q/K norm candidate and provider errors were equal (Q 0.0016656962, K 0.0029305728). This does not prove the full-chain error is fixed. New one-B200 job 1107836, clean source 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.
  • CPU validation: changed-file pre-commit passed. At latest repair commit 9317a5e1, lint, unit tests and docs all actually ran and passed. GPU qualification remains incomplete.
  • History retained: reference job 1105544 failed generation because torchvision/processor was missing after completing its reference backward stages. The processor repair and matching torchvision overlay enabled successful retry 1106656. Earlier local expanded CPU results (149 passed, 358 hardware skips, three subprocess-startup timeouts) remain historical, not an all-pass claim.

GPU jobs use partition overflow, name test, 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. See examples/models/minimax_h3/README.md for the committed candidate gates.

Historical implementation and evidence

Draft. Opened before the claim on #420; it will be marked ready once the row is assigned. Stack: #483 → … → #494 → this PR, so it merges last. This PR's own diff contains the code and evidence commits plus the evidence-directory migration.

Evidence directory migration

Commit 03cd824b moves the one-block report.json and figure.png to reports/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 41e4293 measurement revision and environment. JSON parsing, migration mappings, documentation link targets, Python syntax, Ruff lint/format checks, and git diff --check pass. 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=1 adds --use_fast_math to the whole rl_engine._C build. That replaces expf/sinf/cosf/rsqrtf and 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: the timestep_sinusoid_h3 (#483) and h3_rmsnorm (#487) bitwise tests, and the chain/block tests built on them. det_linear, adaln_row_gather and gate_residual still passed.

nvcc 13 defines no macro for fast math, so the guard lives in the build:

  • build_tools/extensions.py lists every csrc/cuda/h3 source in H3_BITWISE_SOURCES and leaves them all out of fast-math builds;
  • the bindings are gated on KERNEL_ALIGN_WITH_H3_KERNELS, so dispatch falls back to the PyTorch references;
  • tests that assert CUDA dispatch skip through a shared require_h3_kernels fixture, and a build test asserts that no csrc/cuda/h3 source survives a fast-math build.

On this branch: merge 4ba3a8fb and 9cb10d03 (block dispatch-check skip). Tested on 2× B200 with pinned weights (tests/models/minimax_h3, tests/build, tests/runtime/test_dispatch.py):

Build H3 kernels in _C passed / skipped / failed
default yes 545 / 7 / 0
KERNEL_ALIGN_USE_FAST_MATH=1 no (reference fallback) 196 / 356 / 0

Status

  • Fixed: the flaky test_h3_adaln_modulation.py::TestFusedModulation::test_backward_keeps_table_gradient_fp32 in [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, and temb was unseeded; the kernel itself was deterministic. Full suite at 27bb0295 on 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.

  • The block. Block 0 of MiniMax-H3@42ed227 runs as a 17-node graph, from the AdaLN projection to the second gated residual. It runs on packed FL2VA layouts whose position_ids and tags are bitwise equal to the pinned diffusers pipeline's build_packed_sequence; the hashes are pinned in a CPU test.
  • Three executions per node. Each node runs as:
    • the diffusers replay (op for op, no diffusers import);
    • the RL-Kernel binding, both chained and isolated (fed diffusers' inputs);
    • an FP32 golden.
  • Ownership. Each node names the RFC row that owns its arithmetic, and whether it is bound to that row's operator, an interim deterministic operator, or the provider replay.
  • Forward report, per node:
    • repeat and batch-row bitwise checks;
    • chained and isolated agreement with diffusers (including the first differing packed row);
    • candidate and diffusers error against the golden;
    • the node's sha256;
    • first_{,isolated_,repeat_,batch_}drift.
  • Backward report:
    • gradients of the block input, temb and all 12 block parameters (repeat, error vs golden);
    • the gradient reaching every node (repeat, batch row), with first_grad_{repeat,batch}_drift.
  • Weights. Block-0 attention/FFN weights are pinned in a new manifest section (block_tensors). They are extracted by prepare_h3_weights.py --block from the same shard, so the existing conditioning file is unchanged.

Bindings

Node Owning row Bound to
adaln_projection adaln_projection_3mod row (#485), via the registry
norm1, norm2 h3_rmsnorm row (#487), via the registry
residual_attn, residual_mlp adaln_gate_residual row (#488), via the registry
q/k/v_proj, o_proj, ffn_gate_up, ffn_down h3_qkv_gemm, h3_attention_o_gemm, h3_ffn_gate_up_gemm, h3_ffn_down_gemm interim: Triton FP32-accumulate GEMM, one BF16 rounding
q_norm, k_norm h3_qk_rmsnorm_d128 interim: #487's plain H3RMSNormCudaOp on the 128-wide heads (bitwise equal to nn.RMSNorm)
rope_q, rope_k h3_mm_rope_3axis_partial reference: the provider replay, since no MM-RoPE kernel exists yet
attention h3_full_attention interim: DeterministicAttentionOp(causal=False)
swiglu h3_swiglu interim: SwiGLUCudaOp

When an owning row lands, swapping its node's candidate callable 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.

  • diffusers. MiniMaxH3TransformerBlock is the executable reference. It is replayed op for op as the provider (h3_provider), which pins its dtypes and op order without importing diffusers.
  • [ws1]: WS1 Full Qwen3-8B Dense Train-Inference Closeout #315 Qwen3-8B chain gate (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 the h3_chain stage pattern of [WS1][CUDA][MiniMax-H3] timestep_sinusoid_h3: FP32 timestep features + shared H3 harness (RFC #420 stack 1/9) #483.
  • Interim operators are reused from RL-Kernel. The one decision taken here is the GEMM entry point. 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):
Shape Tree path FP32 K loop → BF16 cuBLAS Tree time FP32 K loop time
264 × 5376 → 7168 7.0e-3 2.4e-3 2.4e-3 10.4 ms 1.2 ms
925 × 14336 → 5376 7.5e-3 2.6e-3 2.6e-3 47 ms 2.7 ms

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 that h3_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: #483 timestep_sinusoid; #484 timestep_mlp; #485 adaln_projection; #486 adaln_row_gather; #487 norm_modulate; #488 gate_residual; #489 final_adaln_out.

Results

Generated from a clean tree at 41e4293 on an otherwise idle B200 (torch 2.13.0+cu130), with tools/validation/models/h3_block_replay.py --layouts tiny,small,medium: report.json.

one-block evidence

S First drift (chained / isolated) Repeat or batch drift, any node or gradient Output err: RL-Kernel / diffusers Fwd ms Fwd+bwd ms
264 adaln_projection / adaln_projection none 5.21e-3 / 5.21e-3 7.4 vs 0.61 22.7 vs 2.8
925 adaln_projection / adaln_projection none 7.87e-3 / 7.87e-3 30.1 vs 1.12 81.7 vs 5.0
3160 adaln_projection / adaln_projection none 8.97e-3 / 8.97e-3 225 vs 3.8 577 vs 11.5
  • Isolated bitwise promises. Every node that promises it is bitwise equal to diffusers on diffusers' inputs at every size: norm1/norm2, q_norm/k_norm, RoPE, and both gated residuals. The first drift is the first reduction whose tree differs from cuBLAS.
  • Forward accuracy. Per-node error vs the FP32 golden is 0.93–1.20× diffusers'. The largest is the interim attention at S = 3160.
  • Gradient accuracy:
    • The candidate is up to 7× more accurate than diffusers on adaln_proj.weight, adaln_proj.bias and temb. diffusers' index_select backward accumulates in BF16 atomics; the row operators use FP32 segment sums.
    • The candidate is within 1.16× on the other leaves, except attn.norm_q.weight at 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.
  • Control H-C3, reported, not promised. Permuting the packed rows with their positions and indices changes 79 / 450 / 2199 output rows (max 8). Attention sums keys in packed order; settling this belongs to h3_full_attention.
  • Speed. The 12–60× forward cost is the interim GEMMs and attention (FP32 scores materialised), not the row operators.

Tests

RL_KERNEL_H3_WEIGHTS=<dir> CUDA_VISIBLE_DEVICES=0,1 python -m pytest tests/models/minimax_h3 -q

501 passed, 5 skipped on two B200s, including the 9 new tests in test_h3_one_block.py. The new tests check:

  • CPU:
    • the layouts are pinned to diffusers;
    • the graph is well formed;
    • MM-RoPE rotates 96 of 128 channels.
  • GPU forward (264, 925):
    • every node repeats and is batch-row bitwise;
    • promised nodes are isolated-bitwise;
    • node and output errors are at most 1.25 × diffusers' + 1e-3.
  • GPU backward (264):
    • every leaf repeats;
    • dX is batch-row bitwise;
    • there is no gradient drift;
    • leaf errors are at most 1.5 × diffusers' + 5e-3.

Scope / Known limitations

Documentation repair at 9317a5e1: replace 14 report links outside the MkDocs docs tree with immutable GitHub/raw URLs at the evidence-containing commit 68e472b7. The strict remote docs build passed. No evidence files or numerical results were changed.

…+ 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

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

Labels

MiniMax-H3 new-model platform: cuda Specific optimizations or bugs in NVIDIA graphics cards (such as FlashInfer, TMA optimizations) WS1

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants