Repository navigation
Conversation
…de the engine A video and audio model's denoising transformer on the M5 tensor units, in the checkpoint's precision: bf16 weights read in place, products on the prefill kernels' tfp::gemm_nax (MLX's GEMM, fp32 sums), FastH3's tile-routed attention in bf16 with fp32 scores and an online softmax. zig build tf-h3-dit runs every denoising forward for a case file. The lane core, the server and the other families are untouched. Thirteen kernels are registered with the kernel contract and pass gates 1-4; their launches and the two GEMM entry points are in registry.zig. The recipe holds the receipt.
458be74 to
a428b8b
Compare
…Linear, a bf16 routed attention The blocks now run at the checkpoint's own precision, and int8 is gone from the family. - Weights: the checkpoint's bf16 tensors copied to the device as stored; no quantized copy. - Products: zig/src/cuda/cublaslt.zig's Linear, bf16 in and out, fp32 compute, no TF32. Each shape keeps the algorithm chosen at 16,384 rows, so a rank's rows take the same sums as on one GPU. - Attention: FastH3's tile-routed attention on mma.sync m16n8k16, fp32 scores and online softmax. - Routing: tile means, scores, top-k and the pooled branch in fp32, as the Metal family (ashhart#949) does. - Rows: norm, gated adds, SwiGLU and head layout round where torch's bf16 ops round in the reference. - Ranks trade bf16 keys and values and fp32 key and value means.
|
Rebased on main |
…Linear, a bf16 routed attention The blocks now run at the checkpoint's own precision, and int8 is gone from the family. - Weights: the checkpoint's bf16 tensors copied to the device as stored; no quantized copy. - Products: zig/src/cuda/cublaslt.zig's Linear, bf16 in and out, fp32 compute, no TF32. Each shape keeps the algorithm chosen at 16,384 rows, so a rank's rows take the same sums as on one GPU. - Attention: FastH3's tile-routed attention on mma.sync m16n8k16, fp32 scores and online softmax. - Routing: tile means, scores, top-k and the pooled branch in fp32, as the Metal family (#949) does. - Rows: norm, gated adds, SwiGLU and head layout round where torch's bf16 ops round in the reference. - Ranks trade bf16 keys and values and fp32 key and value means.
… beneath it The H3 and FastH3 transformer on Metal as a test tool beside the release's, clips that start from an image, any step count with its schedule and AdaLN tables, the released MiniMax-H3 weights with re-voicing of a finished clip, and Ref2VA's reference pictures, videos and audio. Against this branch: build.zig runs the H3 schedule's host tests beside the XPU admission tests.
|
Thank you, @drowzeys, for H3/FastH3 on Metal. The selected work from #949 is included through #961 ( Posted by Ash's AI Assistant Synthena |
|
The bf16 H3 Metal component shipped in 1.1.0 as Posted by Ash's AI Assistant Synthena |
…hart#950, ashhart#959 and ashhart#960 beneath it The H3 and FastH3 transformer on Metal as a test tool beside the release's, clips that start from an image, any step count with its schedule and AdaLN tables, the released MiniMax-H3 weights with re-voicing of a finished clip, and Ref2VA's reference pictures, videos and audio. Against this branch: build.zig runs the H3 schedule's host tests beside the XPU admission tests.
…Linear, a bf16 routed attention The blocks now run at the checkpoint's own precision, and int8 is gone from the family. - Weights: the checkpoint's bf16 tensors copied to the device as stored; no quantized copy. - Products: zig/src/cuda/cublaslt.zig's Linear, bf16 in and out, fp32 compute, no TF32. Each shape keeps the algorithm chosen at 16,384 rows, so a rank's rows take the same sums as on one GPU. - Attention: FastH3's tile-routed attention on mma.sync m16n8k16, fp32 scores and online softmax. - Routing: tile means, scores, top-k and the pooled branch in fp32, as the Metal family (ashhart#949) does. - Rows: norm, gated adds, SwiGLU and head layout round where torch's bf16 ops round in the reference. - Ranks trade bf16 keys and values and fp32 key and value means.
Thank you, Ash, for the invitation to send this for review, and for the Metal runtime it builds on.
This is the first video and audio generation model on TensorFold's Zig + Metal runtime. MiniMax H3 / FastH3 is a 21B-parameter joint video and audio diffusion transformer: 50 blocks denoise one packed sequence of text, audio and video rows. Its whole denoising forward now runs on the M5's tensor units through TensorFold's runtime, with no Python and no MLX in the forward, in the checkpoint's own precision. It is the case for TensorFold beyond text: the same runtime, kernel contract and lean code carry a model whose step is a denoising pass over 15,000 rows instead of a token.
On a Mac Studio (M5 Ultra, 256 GB), 864x480, 124 frames (5 s), 15,435 packed rows, FastH3 8-Step V2:
tfp::gemm_nax, unchanged.main.Precision is the checkpoint's own:
The routed attention rounds its softmax weights to bf16 before the value product, as MLX's attention does. Nothing is quantized. An earlier int8 variant (closed as #586) stays on the fork and is not part of this change.
What comes next, for the gate we agree on:
upstream/h3-firstframe.tensorfold generatepath. Today this is a test tool beside the engine, with no lane round. How a denoiser should sit beside or inside the lane core is a design question I would like to settle with you.I did not open an issue first, since this came by invitation. Happy to move the gate discussion into one if you prefer.
What it adds
These are new files, plus the minimal wiring: 13 lines in
build.zig, 1 import inzig/build/kernel_contract.zig, 17 lines in the Metal contract runner, 7 inmetal_limits.zig, and 15 launch records inregistry.zig.zig/src/families/h3/:h3.zigloads the model and the case file, and holds the forward around the blocks.block.zigis one block: modulated norm, q/k/v and the routing gate, head layout, routing, attention, the pooled branch, the out projection and the SwiGLU MLP.tensors.zigmaps the safetensors shards whole into Metal buffers, finding tensors by name.zig/kernels/metal/h3/:h3_rows.metal: the rows in and out, the norms, the gated adds and the SwiGLU.h3_routing.metal: the head layout with q/k norm and rotary, tile means, tile scores, top-k and the pooled branch.h3_attention.metal: tile-routed attention on the tensor units, with fp32 scores and an online softmax.gemm_naxentry points as they are.zig/tests/h3_dit.zig(zig build tf-h3-dit): every forward for a case file, the first forward against the reference velocities, a per-stage profile, and the rows as raw f32.zig/tests/contract/metal_h3*.zig: the 13 kernels' float64 references and launches, plus the top-k against a CPU oracle.docs/recipes/minimax-h3-metal.md: the arithmetic, the case file and the receipt.How it fits the engine
Receipts
Mac Studio, Apple M5 Ultra, 256 GB, macOS 27.0.1, Zig 0.17.0, Xcode Metal toolchain.
upstream/h3-metalonmain46a08c18.FastVideo/FastVideo-FastH3-8-Step-V2, 14 shards.Speed, alternated on the same machine:
The GPU's own time a forward is within 0.01 s of wall time from step 2 on.
Per stage, from one profiled forward (each stage in its own command buffer, so they sum to more than a forward):
Parity: the first forward's velocities against the reference:
At step 1, every video tile is noise. FastH3's top-k among them is decided by differences far below bf16, so any two correct implementations differ there. Video against the reference sits at the engine's own self-noise floor at both steps. The audio rows, which attend densely, show the arithmetic agreement directly. Per row (step 5): median relative error 0.047, 99th percentile 0.138.
Memory (
/usr/bin/time -l, the 8-step run):Determinism: two runs, byte-identical final video and audio rows (
cmp).Kernel contract
zig build tf-kernel-contract -Dcpu=apple_m1 -Doptimize=ReleaseFast, thenzig-out/bin/tf-kernel-contract --offline DIR(with M1/M2 compiles) on the M5 Ultra:The two known findings are
main'saffine_mmq4 entries.The registered kernels:
h3_rows_in,h3_rows_out,h3_norm,h3_final_norm,h3_gate_add,h3_gate_mix,h3_swiglu,h3_heads,h3_pool,h3_tile_scores,h3_coarse,h3_attention, andh3_topk.mids. Inputs are of one sign where a sum could cancel, so no rounding hides behind a small result.h3_topkis checked byte for byte against a CPU oracle under a stated rule: a NaN score ranks as −inf, and ties go to the lower tile. Gate 1 found that a NaN score let a row keep one tile too many and write past its row. The kernel now ranks NaN as −inf, and the test keeps that failure covered.h3_rows_in,h3_rows_out,h3_tile_scores). The softmax and norm kernels state why their bf16 output hides the sums' error, asmain's softmax and RMS norm do.--offline.h3_attentionis a tensor-op kernel at 256 threads, and the GEMMs reserve 256.Checks
main46a08c1python3 tools/zig/lean_check.pyzig build test -Dcpu=apple_m1zig fmt --checkon changed Zigzig build tf-h3-diton the M5 UltraThe 10 skips are the same on both: JPEG pixel cases without libjpeg-turbo, and others the suite skips here.
Not tested
tools/zig/h3_case.py, branchh3-1.0.4). I can attach the receipt's case file (186 MB) wherever you prefer.Licenses
The code is Apache-2.0, as TensorFold is. MiniMax H3 and FastH3 weights are under the MiniMax H3 Community License, which limits territories. No weights are included.
The CUDA side of the same model is #947, and the Studio that drives both is #941.
🤖 Generated with Claude Code
https://claude.ai/code/session_01PfjmufXm4L7qsVNFyRw9ZK