Repository navigation
MiniMax-H3: long-sequence flash attention, fused DiT glue and a faster video VAE decode (G4 render 63.6 to 37.0 s, bit-identical) - #18
Merged
Conversation
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.
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. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: b92b2a2cbc
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
Comment on lines
+327
to
+329
| q = ggml_rope_pe_permute(ctx->ggml_ctx, q, pe, n_rot, 1.f, GGML_TYPE_F32); | ||
| k = ggml_rope_pe_permute(ctx->ggml_ctx, k, pe, n_rot, kv_scale, GGML_TYPE_F16); | ||
| auto v = ggml_rope_pe_permute(ctx->ggml_ctx, part(2), nullptr, 0, kv_scale, GGML_TYPE_F16); |
There was a problem hiding this comment.
When users build a normal recursive checkout with the documented default SD_USE_UPSTREAM_GGML=OFF, this call is compiled, but the commit leaves the ggml gitlink at f583f393..., whose headers do not contain ggml_rope_pe_permute, ggml_modulate_rows, or ggml_swiglu_scaled. Those APIs are supplied only by the new patch files, and the only repository code applying those patches is the Unsloth prebuilt workflow at .github/workflows/unsloth-sd-prebuilt.yml:227-230; regular CMake builds never apply them. Consequently, the default CPU/CUDA build from docs/build.md fails on undeclared operators. Update the pinned submodule, apply the patches in the regular build path, or feature-gate these calls independently of SD_USE_UPSTREAM_GGML.
Useful? React with 👍 / 👎.
Member
There was a problem hiding this comment.
Fixed in 5ab239d: cmake/ggml.cmake checks whether the selected ggml.h declares ggml_rope_pe_permute and defines SD_GGML_H3_FUSED_OPS only then, and minimax_h3.hpp gates the fused ops on that macro instead of SD_USE_UPSTREAM_GGML, so an unpatched default build compiles and runs the unfused graph. 703f519 documents how to apply scripts/unsloth/ggml-patches to a source checkout.
|
You have reached your Codex usage limits for code reviews. You can see your limits in the Codex usage dashboard. |
|
You have reached your Codex usage limits for code reviews. You can see your limits in the Codex usage dashboard. |
Member
|
Reached Codex review convergence at oobabooga#4. |
Member
|
No changes were needed beyond merging #17's fix down into this branch ( Tested
A 2-step HIP bisect over the patch prefixes 0001 to 0004 (0001-0002, 0001-0003, etc.) shows no slowdown from any of them: every variant runs at 64 to 71 s per step. Output:
To check that the HIP and Windows differences add no error, each GPU render was compared with a CPU-backend render of the same frames (640x384x56, music prompt, 2 steps):
test-backend-ops on sm86, sm89 and HIP all pass:
Metal reports the fused ops as unsupported and passes FLASH_ATTN_EXT 4834/4834; tiled renders are identical there.
|
Member
|
Reached Codex review convergence at oobabooga#4. |
oobabooga
added a commit
that referenced
this pull request
Oct 9, 2026
* MiniMax-H3: flash attention in the video VAE under --diffusion-fa; carry 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. * tiling: walk tile planes contiguously when splitting and merging 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. * Carry the pipelined long-sequence ggml-cuda FlashAttention patch 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. * ggml patch 0002: 128-column long-sequence FlashAttention tile on sm80 / sm89 * docs: MiniMax-H3 long-sequence flash attention switches * tiling: batched tile callback, merge tile planes on several threads 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. * MiniMax-H3: batch video VAE tiles per decoder graph, keep weights resident, 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. * ggml-patches: 0003 fused cuBLAS bias epilogues and narrow short-row RMS 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. * MiniMax-H3: value-preserving DiT graph rewrites 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. * Carry the fused DiT ops ggml patch; document the H3 graph switches 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. * MiniMax-H3: keep the folded MLP path compiling with SD_USE_UPSTREAM_GGML 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. * ggml patch 0002: long-sequence flash attention keeps the stock grid, 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). * MiniMax-H3: video VAE tile batching opt-in (SD_H3_VAE_TILE_BATCH=auto|N) 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. * MiniMax-H3: sage attention takes the fused RoPE / head-major q/k/v op 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. * MiniMax-H3 video VAE: keep the device busy between tiles 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. * MiniMax-H3: audio VAE anti-aliased activations use direct depthwise convs 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. * MiniMax-H3 video VAE: fused RoPE / head-major q/k/v for decoder attention 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. * MiniMax-H3 video VAE: fold the q/k RMS norm into the fused RoPE op 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. * MiniMax-H3 video VAE: SD_H3_VAE_TILE opt-in latent tile size 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. * Tile merge / H3 assembly / capacity reading: harden the reuse paths - 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. * MiniMax-H3 / ggml-cuda: cuDNN fused attention on Ampere and newer 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. * Tighten H3 VAE and tiling comments * Build without the fused H3 ops when ggml lacks patch 0004 * Document that the H3 ggml speedups need the carried patches * Gate the H3 VAE fused attention on ggml patch 0004 and its fused norm on 0005 * Tighten the H3 VAE, sage and tiling comments * Ship cudnn-frontend's MIT notice in the CUDA bundle * Keep the H3 VAE on its old attention outside CUDA and ROCm * Take cuDNN attention by default only where it measured faster * Install the NVRTC header the cuDNN frontend includes in the CUDA prebuilt leg --------- Co-authored-by: Daniel Han <danielhanchen@gmail.com>
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.
Stacked on #17 (base branch
perf/h3-gguf-speed); review that first.Problem
With #17, a MiniMax-H3 GGUF render on an RTX PRO 6000 still takes 63.6 s against ComfyUI's 14.6 s warm. A denoising step is 11.3 s, about 4x ComfyUI's 2.7 s, and the video VAE decode is 12.9 s. Profiling one step and one decode:
Change
Each lever has its own switch; every default-on lever is bit-identical to #17.
GGML_CUDA_FA_LONGSEQ=0GGML_CUDA_FA_LONGSEQ_NCOLS=128opts inSD_H3_GRAPH_FAST=0(finer:SD_H3_FAST_QKV,_MLP,_SEGMENTS,_VIEWS)SD_H3_VAE_KEEP_RESIDENT=0SD_H3_VAE_GRAPH_OPT=0GGML_CUDA_CUBLAS_EPILOGUE_FUSION=0GGML_CUDA_NORM_SMALL_ROWS=0SD_H3_VAE_TILE_BATCH=autoor=Nopts inscripts/unsloth/ggml-patches/0002to0004. With 0001 they apply cleanly to the pinned ggml commit, and the patched tree matches the development clone exactly.tiling.cpp: a batched tile callback, and tile planes merged on several threads.docs/minimax_h3.mdlists the switches.Results
Colab, MiniMax-H3 UD-Q3_K_XL, 960x544, 124 frames, 4 steps, seed 42, one run per arm on the same VM. ComfyUI v0.38.0 with its int8 H3 graph on the same VM. Total seconds, s/step in brackets.
speed_mode=max: #17speed_mode=max: this PRAccuracy
speed_mode=maxoutput is bit-identical to MiniMax-H3: flash attention in the video VAE, faster tile stitching, ggml-cuda speed patch #17's max output.Tests
test-backend-ops, unmasked FLASH_ATTN_EXT: 844 / 844 on G4, L4 and A100 under stock, with the new path forced on, and with the 128-column tile forced on, including new cases for partial last KV tiles and GQA.Validation on other hardware
Base #17, same seed, Studio's H3 flags.
GGML_CUDA_FA_LONGSEQ=0and_NCOLS=128), RMS_NORM, MUL_MAT (also with the epilogue off), ROPE_PE_PERMUTE, MODULATE_ROWS, SWIGLU_SCALED pass.SD_H3_GRAPH_FAST=0the DiT matches #17 again, so the DiT rewrites round differently on that driver.--backend cpu)Not covered