Repository navigation
Conversation
|
Thank you, @drowzeys, the stored-bf16 H3 blocks provide the CUDA component to carry forward, with #947's contribution retained in the history. We checked 96f0aa2. cuda.zig:84 forces device zero despite accepting host pointers; cuda_lib.zig:60-79 casts signed TP values before validating them. Geometry also grows unchecked byte products before its region check. We can repair context sharing with #946, validate the C ABI and budget shapes before allocation, with failure controls. We will qualify single-device blocks first and TP separately, then finish native conditioning, projections, sampling and decoders with the shared H3 job/Metal design. Please retain identical-input per-stage fixtures and per-rank memory/failure cases; the current block library still leaves those surrounding pipeline stages to its host. Posted by Ash's AI Assistant Synthena |
…for a host pipeline A video model's denoising transformer: no lane round, and int8 products and attention on mma.sync m16n8k32. It is a library (zig build tf-h3) beside the engine for a host that owns the rest of the pipeline and shares its CUDA context; the lane core and the server are untouched. docs/recipes/minimax-h3-cuda.md holds the receipt.
A host that merges a LoRA into the checkpoint's weights hands the merged projection over and the library quantizes it as at load; the int8 copy saved beside the checkpoint is left as it was.
A rank runs the rows of its own routing tiles; every block it hands the other ranks its tiles' int8 keys, values, scales and pooled means and takes theirs, then attends as on one GPU. The kernels and their arithmetic are unchanged row for row, so a pass on N ranks gives the bits of a pass on one. - h3_xchg.c: the exchange over libibverbs. One registered region a rank, one reliable queue pair a peer and rail; a rank RDMA-writes its share into the same offsets of every peer's region, the last write carrying a sequence number. A rail is one RDMA device; two rails split each span. - The region is pinned host memory; the buffers stay on the device and h3_spans packs a rank's tiles into the region and unpacks the others' out of it. Two halves of the region by block parity: a rank one block ahead never writes what another is unpacking. - Inputs.tile0/tile1/cuts: a rank's tile range and every rank's. Routing, top-k and attention run for the rank's own query tiles; kernels take their tile range by where their buffers start. - tf_h3_open on a .tf-int8 path opens the saved int8 copy alone (a rank holds no checkpoint). - tf_h3_forward_host, tf_h3_host_alloc / _free: inputs in host memory, for a rank without a framework. - tf_h3_tp_init / _close / _info / _connect / _stats. The connection records travel over the host's own channel. - The sequence's buffers move to cuda_geometry.zig and the world to cuda_tp.zig, keeping each file under 600 lines.
…e recipe's receipt contract_h3.zig registers h3_gemm, h3_gemm_n256, h3_gemm_swiglu and h3_gemm_wide (exact int32 sums, fp32 scales, bf16 out) and h3_spans (a move) for gates 1 to 4, with their launches and per-SM records for gate 5. The H3 kernel image is built once for the library and the contract's runner, and kernel-contract-fatbins installs it.
…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.
Nine row, routing and attention kernels against their float64 formulas (ported from the Metal family's references), the top-k against a CPU oracle under its stated NaN and tie rule, and h3_spans as before; 11 launch records.
…e a host's process A library that runs kernels for a host framework (PyTorch, say) has to work in the host's context and on its stream. Context.shared retains the primary context of the device the calling thread's current context is on (device 0's when none is current) and refuses a context that is not a primary one; Stream.borrowed wraps a stream the caller owns, null being the legacy default stream. Nothing that exists changes; cuCtxGetDevice joins the driver table. tf-cuda-test shared-context checks both on a GPU.
…n fp32 forward, and the 1/2/3-Spark receipt
96f0aa2 to
e1c121e
Compare
|
Rebased on main |
…for the H3 library's context
Model.load always retained device 0's primary context, though the library runs inside a host process that may have made another device current. It now takes Context.shared (ashhart#946): the current context's device, and device 0 only when the host has none yet; a current context that is not its device's primary context is refused. The legacy default stream is taken as Stream.borrowed.
… them tf_h3_tp_init, tf_h3_tp_info and tf_h3_tp_connect @intcast their c_int world, rank and peer to u32 before anything looked at them, so a negative value trapped (or wrapped in a release build). cuda_bounds.zig checks the signed values first: 2 to 8 ranks, 0 <= rank < world, and a peer that is another rank of the model's world, refusing with error.BadWorld. Its host tests run in test-cuda-host.
… or allocated geometry.prepare multiplied the host's u32 row and tile counts into byte sizes unchecked, and only compared the exchange region after some buffers were already made. cuda_bounds.bytes() now computes every buffer's bytes and the region's with std.math.mul/add and refuses an overflow, or a row or slot count past the kernels' int, with error.GeometryTooLarge (more prefix tiles than tiles is error.BadGeometry). prepare budgets all of it, region included, before it releases the previous geometry or allocates. Host tests: a real sequence's sizes and each overflow.
|
Thank you. All three points are fixed on the branch; the head is now
All of these are host tests, wired into Checks on a DGX Spark (CPU only, no GPU work):
Kernels and launches are unchanged, so I didn't rerun the kernel contract. I haven't rerun the library on a GPU since these commits. The identical-input per-stage fixtures and the per-rank memory and failure cases are all kept. |
|
Thanks @drowzeys for the current-context, signed-value and checked-geometry repairs, including validation before replacing old geometry. The named refusal tests are useful host evidence; the changed library has not had its device rerun. Our landing split is the BF16 single-device block library first, with shared runtime ownership and the later partial-load cleanup repairs retained. TP stays a separate follow-up with per-rank memory, exchange and failure gates. Please keep same-input stage references and allocation/refusal controls attached. This component still relies on external encoders and decoders, so it does not establish complete native clip serving. Posted by Ash's AI Assistant Synthena |
beneath it #958's transformer blocks, kernels and contracts as a C library for a host pipeline (zig build tf-h3, libtf_h3.so) on one device; its sequence-parallel exchange over RoCE waits for a follow-up of its own. With our repairs: checkpoint and projection bounds before any allocation, the host's primary context through Context.shared, complete blocks only, partial loads keeping their destinations and parent objects until their copies are proven complete, and one recovery walker for each failed-load registry. The H3 kernels join the CUDA kernel contract. The host steps (test-h3-ownership, -metadata, -arithmetic, -blocks and -boundaries, compile-h3-library) sit beside the Metal family's in zig/build/h3_host.zig, and the tests are named for what they check.
|
Thank you, @drowzeys, for H3/FastH3's CUDA transformer blocks. 1.1.0 includes #958 ( Posted by Ash's AI Assistant Synthena |
… with ashhart#947 beneath it ashhart#958's transformer blocks, kernels and contracts as a C library for a host pipeline (zig build tf-h3, libtf_h3.so) on one device; its sequence-parallel exchange over RoCE waits for a follow-up of its own. With our repairs: checkpoint and projection bounds before any allocation, the host's primary context through Context.shared, complete blocks only, partial loads keeping their destinations and parent objects until their copies are proven complete, and one recovery walker for each failed-load registry. The H3 kernels join the CUDA kernel contract. The host steps (test-h3-ownership, -metadata, -arithmetic, -blocks and -boundaries, compile-h3-library) sit beside the Metal family's in zig/build/h3_host.zig, and the tests are named for what they check.
|
Closing: this was merged into main as 070f09a ('Merge #958') for 1.1.0, and main's own validation has since superseded the later review commits. If you'd like the RoCE part of this branch considered, please open it as a new PR on current main. #1209 tracks RoCE beyond two ranks. Thank you @drowzeys. - Synthema, Ash's AI Assistant. |
Thank you, Ash, for the review of #585 and #947, and for the runtime this builds on.
H3's 50 blocks now run on TensorFold's CUDA runtime at the checkpoint's own precision, and the same pass splits across DGX Sparks with the bits of one GPU.
Linear(zig/src/cuda/cublaslt.zig: fp32 compute, no TF32).This answers #585's objection ("the int8 projections and attention are also a precision trade we don't make anywhere in the engine"). It is the CUDA half of #949, the Metal bf16 version: the same routing in fp32, the same kernel set, and its float64 references ported. The two read as one family on both backends.
The receipts are for the same machines and clip as #947: a 5 s clip at 864x480, 124 frames, 15,433 rows and 8 passes, on DGX Spark (GB10). The reference is ComfyUI's own bf16 blocks.
main, 517 here.zig build testpasses.Stacked on #947; the commits of this change alone are drowzeys/TensorFold@upstream/h3-cuda-parallel...upstream/h3-cuda-bf16. If you would rather take one change, the two squash cleanly into a bf16-only family. The int8 variant stays on the fork.
Why bf16 only, and not bf16 beside int8
.tf-int8copy are gone: 987 lines out, 497 in, in the family and its kernels.What changed
zig/src/families/h3/:cuda_weights.zigcopies the checkpoint's bf16 block tensors to the device as stored (42 GB).tf_h3_set_weightcopies a host-merged bf16 projection in.cuda_gemm.zig(new) holds the products on the runtime'sLinear. Each (outputs, inputs) shape keeps the algorithm cuBLASLt picks at 16,384 rows, so a rank's rows sum in the same order as on one GPU. That is what keeps ranks bit-identical.cuda.zigholds the forward. The dense pass skips the routing gate's product.cuda_tp.zighandles the exchange: bf16 keys and values, plus fp32 key and value means when routing.zig/kernels/cuda/h3.cu,h3_attention.cuh:mma.sync m16n8k16: bf16 q, k and v, fp32 scores and online softmax. The weights are rounded to bf16 for the value product, as flash attention does. Q, K and V take 48 KiB, so two blocks fit on an SM, and V's load overlaps the scores.zig/tests/cuda/contract_h3*.zig: nine kernels against their float64 formulas (the Metal family's references, ported), and the top-k against a CPU oracle under a stated rule: a NaN score ranks as −inf, and ties go to the lower tile.h3_spansis as before.registry.zigholds 11 records.docs/recipes/minimax-h3-cuda.md: the arithmetic and the receipt.The library's C API is unchanged except that
tf_h3_opentakes the checkpoint alone; there is no saved copy. The lane core and the server are untouched.Receipts
Environment:
96f0aa2e, on MiniMax H3 / FastH3 video and audio generation on CUDA, one GPU or several Sparks over RoCE #947'sb570b8f0, onmain46a08c18.a4b5a045(0.39.0).FastVideo/FastVideo-FastH3-Comfyfastvideo_fasth3_8step_v2_pruned_bf16.safetensors.For comparison, #947's int8 on 2 and 3 GB10s: 4.06–4.14 s and 3.02–3.06 s a pass. bf16 doubles the bytes exchanged, about 1.0 s of a 3-rank pass against int8's 0.6 s.
Where one GPU's pass goes, profiled on the library alone:
Bit-identity:
Parity. All rows are measured on the first pass's rows with dense attention. The case was dumped from ComfyUI and replayed through each library on the same machine. Errors are relative L2. Worst row and worst column are the relative errors of the worst single row and channel.
The fp32 forward is an independent torch implementation of the blocks: fp32 weights and activations, no TF32, SDPA's memory-efficient fp32 kernel. It shares no code with ComfyUI or this library. Against it, this library's aggregate and worst-column errors are below ComfyUI's own. Its worst row is slightly above (0.341 against 0.328).
Memory, the library alone:
tf_h3_openon the checkpoint, then three passes over the 15,433-row geometry, sampled every 0.5 s.int8 needed 23,052 MiB. The process maps the checkpoint while it copies the tensors, so its resident pages peak at 36–38 GB during the load. It ran beside ComfyUI with its own checkpoint mapped and its text encoder resident, on a 128 GB Spark.
Clips (seed 7 fox, frames at 2.5 s and 4 s; and a spoken line, seed 11). Left to right: this library, #947 int8, ComfyUI's bf16 blocks. All three are clean and coherent, and differ only as sparse routing does. Whisper (small, English) heard "The last train leaves at nine tonight." from all three speech clips.
Commands:
The clip timings came from the Studio's ComfyUI node, as in #947. Parity used a case dump of the first pass, replayed through each library's
tf_h3_forward_host. The ranks ran the Studio'srank.py, and the memory figures came from itsbench.py. These are app-side harnesses and are not part of this change.Kernel contract
The H3 kernels are
h3_norm,h3_gate_add,h3_gate_mix,h3_swiglu,h3_heads,h3_pool,h3_tile_scores,h3_coarse,h3_attention,h3_topkandh3_spans.mids, and inputs are one-signed where a sum could cancel.h3_tile_scores. The softmax, norm and pool kernels state why their output hides the sums' error, asmain's softmax and RMS norm do.Checks
main46a08c1python3 tools/zig/lean_check.pyzig build test(fresh cache)zig build test-kernel-contractzig fmt --check, changed Zig filesGaps and next steps
Linear, unchanged. Its one change of use is a fixed algorithm per shape. This change checks them end to end against the fp32 forward, not class by class in FP64. I can add a per-class FP64 check of q/k/v, out, fc1 and fc2 on the engine's own rows if you want it in this PR.Licenses
The code is Apache-2.0. MiniMax H3 / FastH3 weights are under the MiniMax H3 Community License, which limits territories. No weights are included.
The Studio that drives this library is #941; the CUDA context hook it can use is #946.
🤖 Generated with Claude Code
https://claude.ai/code/session_01PfjmufXm4L7qsVNFyRw9ZK