Repository navigation
Merge the MiniMax-H3 speed stack (#18, #19, #20) into master - #23
Merged
Merged
Conversation
…rry ggml-cuda speed patch The H3 video VAE decoder is a 36-layer ViT run per 16x16 latent tile. With only --diffusion-fa (what Studio passes) it used mul_mat + scale + softmax attention, materialising a 32 x L x L f32 score matrix per tile and layer. --diffusion-fa now also covers it; SD_H3_VAE_FLASH_ATTN=0 restores the old path. The text encoder is unchanged (that is what --fa would also switch, which moves the conditioning). scripts/unsloth/ggml-patches/ carries two ggml-cuda commits on top of the pinned submodule (BF16 cuBLAS for large-batch quantized matmul behind GGML_CUDA_QUANT_CUBLAS_MIN_BATCH, a cached device integrated flag, a vectorized row copy); the prebuilt workflow applies them after the submodule checkout.
process_tiles_2d copied and blended each pixel across all planes in the inner loop, striding a whole output plane (960x544 floats for H3) per element. A video tile has frames x channels planes, so every element was a cache and TLB miss: on MiniMax-H3 960x544x124 (105 tiles) this was about 10 s of single-threaded host time per decode. Loop planes outermost and precompute the per-column and per-row blend weights; each element still computes old + new * wy * wx in the same order, so the output is bitwise identical.
scripts/unsloth/ggml-patches/0002 pipelines the unmasked, non-GQA mma FlashAttention kernels (all full KV tiles via cp.async, only the partial last tile synchronously). The MiniMax-H3 DiT (56 heads x 128, 19108 tokens) and the H3 video VAE (32 heads x 64) attention take this path under --diffusion-fa. GGML_CUDA_FA_LONGSEQ=0 restores the stock kernels.
process_tiles_2d_batched walks the same tile plan and merge order as process_tiles_2d but hands up to batch_size(remaining) consecutive tiles to one callback, which returns them stacked as contiguous blocks; the merge reads each tile in place instead of from a copy. process_tiles_2d is now a one-tile wrapper around it. The merge splits the output planes across threads. Each element is still written by one thread with the same arithmetic, so outputs are bitwise identical.
…ident, leaner rope and SwiGLU The video VAE decoder (a 36-layer ViT) ran one graph per 16x16 latent tile: 105 graph builds, allocations and output readbacks for 960x544x124, and every temporal chunk ended with runner_end(), which evicted the 5.5 GB of weights and reuploaded them for the next chunk. - Several tiles now go through one decoder graph. The first tile runs alone to measure its compute buffer and the batch is sized from free device memory, capped at 4 (SD_H3_VAE_TILE_BATCH_MAX); SD_H3_VAE_TILE_BATCH=N forces N and 1 restores one graph per tile. Projections run as one 2D matmul over all tiles' tokens and each tile keeps its own attention call: a broadcast batched cuBLAS GEMM and a batched flash-attention launch both changed the summation order relative to the single-tile decode. With this layout the decoded frames match the per-tile decode exactly. - Weights stay resident across temporal chunks and runner_end() runs once after the last chunk (SD_H3_VAE_KEEP_RESIDENT=0 restores the per-chunk release). - Rotary embedding reads the two rotated halves as permuted views of the projection layout and multiplies them by precomputed coefficient tables, with the same separate mul/mul/add per element as before, instead of per-half permutes, repeats, a tail permute and concat. The feed-forward uses ggml_swiglu instead of silu + mul. SD_H3_VAE_GRAPH_OPT=0 restores the previous graph. - Chunks are blended in place with plain loops instead of index() lookups, and the decoded chunks are concatenated once at the end instead of growing the result chunk by chunk.
…MS norm blocks Carried on top of the pinned ggml submodule like 0001 (the prebuilt workflow applies scripts/unsloth/ggml-patches/*.patch in order). In the H3 video VAE decoder every linear layer is an F16 cuBLAS matmul followed by a bias add and then either a scale + residual add or SWIGLU; these now run inside the F16 -> F32 conversion of the GEMM output instead of as separate passes over F32 activations, bitwise identical (GGML_CUDA_CUBLAS_EPILOGUE_FUSION=0 disables). The per-head q/k RMS norm (64 columns) now launches 64-thread blocks instead of 256, also bitwise identical (GGML_CUDA_NORM_SMALL_ROWS=0 restores 256). Applies cleanly with or without 0002.
The DiT spent about 1.9 s of an 8.6 s B200 step (960x544x124, Q3) outside matmul and attention: chunk copies of q/k/v and of the MLP gate, partial RoPE as slice + permute copies + repeat + mul/add + concat, K/V scale then cast, per-segment adaLN modulation and gated residuals as slice copies + mul/add + concat, and separate scale passes around every MLP Linear. Each lever builds the same arithmetic with the same rounding points, so the output is bit-identical: - q/k/v are read in place from the fused projection; one ggml_rope_pe_permute per tensor applies the partial RoPE, writes the head-major layout flash attention reads and, for K/V, the kv scale and F16 cast (SD_H3_FAST_QKV=0 restores the old graph). - The MLP gate is one swiglu on two views of fc1's output instead of two chunk copies + silu + mul (SD_H3_FAST_MLP=0). - adaLN modulation and gated residuals use one in-place ggml_modulate_rows / ggml_gated_add_rows per sequence segment, and the MLP Linear scales are folded into them and into ggml_swiglu_scaled (SD_H3_FAST_SEGMENTS=0). - Modulation vectors and segment slices are views, not copies (SD_H3_FAST_VIEWS=0). SD_H3_GRAPH_FAST=0 turns all of them off. Backends without the fused ops (and upstream ggml builds) keep the original graph.
scripts/unsloth/ggml-patches/0004 adds ggml_rope_pe_permute, ggml_modulate_rows / ggml_gated_add_rows and ggml_swiglu_scaled (CPU + CUDA, test-backend-ops cases) on top of 0001 and 0002. The MiniMax-H3 DiT uses them when the backend supports them and keeps the unfused graph otherwise.
forward_folded is only reachable when the fused ops exist; return from inside the guarded branch so the upstream-ggml build does not reference an undeclared result.
…wide tile opt-in The pipelined variant now sizes its stream-k grid from the stock kernel so the work split, and with it the result, is identical on every GPU (on L4 the two kernels' occupancy differs). The 128-column tile is opt-in only (GGML_CUDA_FA_LONGSEQ_NCOLS=128).
The batched decoder graph runs each projection as one matmul over all tiles' tokens. On an RTX PRO 6000 Blackwell (sm120) cuBLAS picks a different kernel for that larger M and the frames come out at 53 dB PSNR against the per-tile decode (identical on B200), and the batched graph measured no faster there (VAE decode 7.6 s batched vs 7.5 s per tile, the other VAE levers on). Decode one tile per graph by default; SD_H3_VAE_TILE_BATCH=auto or =N opts in.
With --sage-attn, forward_head_major returned nullptr, so every block fell back to the unfused chunk / slice / rope / concat / scale / cast chain (about 0.75 s per 960x544x124 step on B200). ggml_rope_pe_permute already writes [head_dim, tokens, heads, batch], the layout ggml_sage_attn reads: emit F32 Q, F32 K * kv_scale and F16 V * kv_scale and call sage directly. Same arithmetic and rounding points as the unfused sage branch; the rendered video and audio are bit-identical. SD_H3_FAST_SAGE_QKV=0 restores the old chain.
The 960x544x124 decode left the B200 idle for about 2 s between its 105 tile graphs: the host blended each tile, zero-filled and assembled the frames, rebuilt the rotary tables and queried free device memory four times per tile while the GPU waited. - tiling: the blend of a batch into the output runs on a worker thread while the next batch is split and computed; merges stay serial and in tile order with the same arithmetic (SD_TILE_ASYNC_MERGE=0 merges inline). - H3 decode: each temporal chunk is trimmed, cross-faded and copied straight into the trimmed output on a worker thread while the next chunk decodes, instead of concatenating and slicing at the end (SD_H3_VAE_ASYNC_ASSEMBLY=0 restores it). - rotary tables are rebuilt only when the tile shape changes. - DeviceMemoryRequest::reuse_device_query: while a runner repeats identical computes, a capacity check that allocates nothing new (no pending bytes, all parameters resident) reuses the owner's earlier free-memory reading; runner_end() drops it. The H3 decode turns it on for its duration (SD_H3_VAE_REUSE_MEMQUERY=0 disables). Warm decode on B200 6.8-7.0 s -> 5.06 s; frames and audio bit-identical.
…onvs The up/down-sampling filters of every Activation1D went through ggml_conv_1d / ggml_conv_1d_dw: an F16 im2col plus a single-column matmul per call. Route them through CONV_2D_DW (H=1, F32 per-channel kernel) instead, for the MiniMax-H3 audio VAE only. Falls back to the im2col graph when the backend lacks the op. SD_H3_AUDIO_DIRECT_DW=0 restores the previous graph.
…tion The decoder attention chunk-copied q/k/v out of the fused projection, applied the partial RoPE as table mul/add/concat chains and then permuted and cast K/V for flash attention. With one tile per graph, normalise q/k in place on the projection and let ggml_rope_pe_permute (already used by the DiT) write the F32 Q and F16 K/V head-major tensors ggml_ext_attention_prepared reads. Same products, sums and casts as the table path, so the decoded frames are bit-identical. SD_H3_VAE_FUSED_QKV=0 restores the old graph; batched tiles, sage and non-flash decodes keep it.
Carry ggml patch 0005 (ggml_rope_pe_permute_rms) and use it for the decoder's q/k: the per-head norm was 7560 launches of a 64-thread kernel, about 0.45 s per 960x544x124 decode on B200. The fused op normalises each head row with the same block reduction as the standalone norm, so the frames stay bit-identical. Backends without it (HIP, Vulkan, GGML_CUDA_NORM_SMALL_ROWS=0) keep the separate norm; SD_H3_VAE_FUSED_QK_NORM=0 forces that.
Fewer, larger tiles cut the per-tile fixed cost (20x20 decodes 960x544x124 in 3.07 s vs 3.77 s at 16x16 on B200) but move the seams: 36.2 dB PSNR / 0.09 LPIPS vs the 16x16 decode, slightly closer to an untiled decode than 16x16 is. Default stays 16.
- The saved device free-memory reading is now taken only by a check that allocates nothing, so it is never the reading from before this owner's compute buffer was allocated, and any check that does not fit drops the owner's readings so the reclaim / evict retries see fresh ones. - If a worker thread cannot be started, the tile merge and the chunk assembly run inline instead of throwing.
New ggml patch 0006 adds an optional cuDNN path to ggml-cuda's FLASH_ATTN_EXT and SAGE_ATTN (AUTO mode): GGML_CUDA_CUDNN=ON compiles it against cudnn.h and the header-only cudnn-frontend v1.26.0 (MIT, downloaded by CMake and pinned by checksum); libcudnn.so.9 is opened at run time, so the binaries keep no cuDNN dependency and run the existing kernels wherever cuDNN cannot be loaded, on Turing and older, for small attention calls and for shapes cuDNN declines. GGML_CUDA_CUDNN_ATTN=0 restores the previous kernels everywhere and GGML_CUDA_CUDNN_SAGE=0 keeps the sage kernel for --sage-attn. On a B200 the H3 DiT attention at 960x544x124 (19315 tokens x 56 heads) takes about 8 ms per call against 28 ms for the sage kernel and 43 ms for the ggml flash attention kernel, and is closer to an fp64 reference than either. The Linux CUDA prebuilt fetches the pinned cuDNN 9.27 headers package, builds with GGML_CUDA_CUDNN=ON and fails if either binary ends up linked to libcudnn. docs/minimax_h3.md describes the run-time requirements.
…/h3-gguf-speed-v3
…/h3-gguf-speed-v3
…/h3-gguf-speed-v3
…/h3-gguf-speed-v3
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
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
2 participants
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.
#18, #19 and #20 were stacked and merged into their parent branches (
perf/h3-gguf-speed,-v2,-v3) instead of master, so master currently has #17 and #21 but none of #18 to #20.This PR is a single merge of #20's head (ae9429f, which contains #17 to #20) into current master. The only conflict was
src/runtime/tiling.h, where #18 and #21 each add a declaration at the end of the file; both are kept.Checked:
ff941a8..ae9429f): same files, same changes, nothing else.Its first parent is current master, so it merges cleanly. "Create a merge commit" keeps the #18 to #20 history; a squash gives the same tree.