Repository navigation
MiniMax-H3 / ggml-cuda: cuDNN fused attention on Ampere and newer (B200 max render 15.4 to 11.3 s) - #20
Merged
Conversation
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.
|
You have reached your Codex usage limits for security reviews. Please try again later. |
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. |
|
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. |
![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: 778ccadbfb
ℹ️ 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
+937
to
+939
| + for (int64_t n : { 1023, 4096 }) { | ||
| + test_cases.emplace_back(new test_flash_attn_ext(128, 128, 42, {1, 1}, n, n, false, false, 0, 0, | ||
| + GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); |
There was a problem hiding this comment.
Remove this evaluation-test hunk and the performance cases later in the patch: the repository's contribution rules explicitly prohibit committing test code and require keeping it local while reporting the verification results instead. Retaining these additions makes the commit non-compliant even though the tests validate the new backend path.
AGENTS.md reference: AGENTS.md:L120-L122
Useful? React with 👍 / 👎.
Member
There was a problem hiding this comment.
These are ggml's own test-backend-ops cases inside a ggml patch, not sd.cpp test code. The prebuilt builds only sd-cli and sd-server, so they never reach a shipped binary, and they are the coverage for the new cuDNN path. Patches 0002, 0004 and 0005 carry test-backend-ops cases the same way. Keeping them.
| + | ||
| +namespace fe = cudnn_frontend; | ||
| + | ||
| +namespace { |
There was a problem hiding this comment.
Replace this anonymous namespace with static file-local declarations or an explicit named namespace. The repository's documented C++ convention expressly prohibits introducing anonymous namespaces in new or modified code, including code carried through this patch.
AGENTS.md reference: AGENTS.md:L120-L122
Useful? React with 👍 / 👎.
Member
There was a problem hiding this comment.
This file is ggml code carried in a ggml patch, so sd.cpp's C++ style rules do not apply to it, and the namespace changes no behavior. Keeping it.
| --- /dev/null | ||
| +++ b/src/ggml-cuda/cudnn-attn-graph.h | ||
| @@ -0,0 +1,31 @@ | ||
| +#pragma once |
There was a problem hiding this comment.
Replace #pragma once with the repository's required macro-based include guards in this header and cudnn-attn.cuh; the contribution rules explicitly disallow #pragma once for newly introduced headers.
AGENTS.md reference: AGENTS.md:L120-L122
Useful? React with 👍 / 👎.
Member
There was a problem hiding this comment.
These are ggml headers. ggml uses #pragma once (every public header, and the ggml-cuda .cuh headers that have a guard at all), and sd.cpp's SD{PATH}_ guard format does not fit a file under ggml/src. Keeping it.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2db859407a
ℹ️ 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
+92
to
+96
| + # v1.26.0, the last MIT-licensed release (later ones are Apache-2.0). Only the headers are | ||
| + # used; SOURCE_SUBDIR points at a directory that does not exist so its CMake project stays out. | ||
| + FetchContent_Declare( | ||
| + ggml_cudnn_frontend | ||
| + URL https://github.com/NVIDIA/cudnn-frontend/archive/refs/tags/v1.26.0.tar.gz |
There was a problem hiding this comment.
When the CUDA prebuilt is enabled, this fetches the MIT-licensed cudnn-frontend headers and compiles their template implementation into the shipped binaries, but scripts/unsloth/package_bundle.py includes only this repository's LICENSE; the dependency's copyright and permission notice is absent from both the source artifact and release zip. Publishing these binaries therefore fails the MIT notice requirement, so copy the pinned dependency's license into the packaged notices.
Useful? React with 👍 / 👎.
Member
There was a problem hiding this comment.
Fixed in aff34f1: the CUDA leg appends cudnn-frontend's LICENSE.txt to the bundle's LICENSE, and the packaging step fails if that file is missing.
Member
|
Reached Codex review convergence at oobabooga#3. |
This was referenced Oct 9, 2026
Member
|
I pushed two changes:
Tested against #19, with cuDNN 9.27 from the
Output and correctness:
Metal is also inert (FLASH_ATTN_EXT 4840/4840, renders identical). Harness. |
This was referenced Oct 9, 2026
Member
|
Reached Codex review convergence at oobabooga#3. |
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 #19 (base branch
perf/h3-gguf-speed-v3). Review that first.Problem
After #19, a warm MiniMax-H3 render on a B200 in max mode (960x544, 124 frames, 4 steps, UD-Q3_K_XL DiT) spends most of each denoising step in attention. The DiT attention at this size is 19315 tokens x 56 heads x 128. Measured per call on the B200:
--sage-attn, int8 Q/K)With 50 of these calls per step, that is about 1.4 s of the 2.37 s step in max mode and about 2.1 s of the 7.15 s step in default mode. The video VAE decoder's attention (1797 tokens x 32 heads x 64, 3780 calls per decode) has the same gap: 157 us per call against 64 us.
Change
New ggml patch
0006-ggml-cuda-cudnn-attention.patchadds an optional cuDNN path to ggml-cuda'sFLASH_ATTN_EXTandSAGE_ATTNops. Nothing in the sd.cpp sources changes.-DGGML_CUDA_CUDNN=ON(default OFF, Linux only) compiles it againstcudnn.hand the header-only cudnn-frontend v1.26.0. v1.26.0 is the last MIT-licensed release (v1.27 moved to Apache-2.0). CMake downloads it, pinned by SHA256, unlessGGML_CUDA_CUDNN_FRONTEND_DIRpoints at a checkout.libcudnn.so.9is opened withdlopenwhen the first attention op runs. The binaries have no link-time cuDNN dependency.GGML_CUDA_CUDNN_LIBnames the library to open.FLASH_ATTN_EXTon sm89 and sm100,SAGE_ATTNon sm100. Unmasked attention without sinks, ALiBi or softcap, equal Q/K/V head sizes (a multiple of 8, up to 256), GQA and batch included, at least 2^20 attention scores.SAGE_ATTNonly in AUTO mode; an explicitly requested Sage variant keeps its kernel.GGML_CUDA_CUDNN_ATTN=0restores the previous kernels everywhere;=1uses cuDNN forFLASH_ATTN_EXTon any sm80+ GPU.GGML_CUDA_CUDNN_SAGE=0/=1does the same for--sage-attnonly.GGML_CUDA_CUDNN_ATTN_BF16=1runs BF16.GGML_CUDA_CUDNN_ATTN_LOG=1logs each call.libcudnn9-headers-cuda-129.27.0.42 package, pinned by SHA256, and installs the CUDAnvrtc-devpackage (cudnn-frontend's headers includenvrtc.h; it opens libnvrtc at run time);GGML_CUDA_CUDNN=ON;sd-cliorsd-serverends up with aNEEDEDentry for libcudnn or libnvrtc.LICENSE, since its header-only templates are compiled into the binaries.docs/minimax_h3.mddocuments the run-time requirements.F16 rather than BF16: sd.cpp already hands the attention F16 K/V, scaled by the kv scale. BF16 would narrow them further for no speed gain (same 8.2 ms).
Results
B200, exclusive card per server, base = #19 head (
ee09dfa7), cuDNN 9.27 for CUDA 12 from thenvidia-cudnn-cu12wheel. Base and this branch were interleaved, two servers each, 3 renders per server; warm = renders 2 and 3. Seconds:--sage-attn, BF16 cuBLAS)GGML_CUDA_CUDNN_ATTN_LOG=1:SAGE_ATTNcalls and 3780 video VAEFLASH_ATTN_EXTcalls take cuDNN, with no fallbacks;Per call (test-backend-ops perf, us):
Accuracy
Kernel against fp64: 22440 tokens x 4 heads x 128 with the H3 kv-scale layout, 512 query rows per head against an fp64 softmax. cuDNN F16 is nearer fp64 than the ggml kernel on 94% of output elements. A 2^-7 relative perturbation of the cuDNN output raises its error to 7.8e-3, so the check can fail. Relative L2 error:
Switches off:
GGML_CUDA_CUDNN_ATTN=0, and a build with no loadable cuDNN, produce video and audio bit-identical to the base build.Renders: this changes frames and audio, because the audio is generated by the same DiT. Six prompt/seed pairs were each rendered by both arms and by two F32-GEMM references that differ only in attention: cuDNN everywhere, or the ggml kernel. Mean LPIPS:
Tests
test-backend-ops -o FLASH_ATTN_EXTon CUDA (B200): 2968/2968 pass with cuDNN (33 cases take it) and withGGML_CUDA_CUDNN_ATTN=0. The 848 unmasked cases also pass in BF16 mode.test_sage_attnpasses (49 cases).GGML_CUDA_CUDNN=ON(the workflow's header recipe and the fetched cudnn-frontend) and with it off;NEEDEDentry.Validation on other hardware
Base #19, same seed, Studio's H3 flags; cuDNN 9.27 from the
nvidia-cudnn-cu12wheel viaGGML_CUDA_CUDNN_LIB.-DGGML_CUDA_CUDNN=ON)test-backend-ops -o FLASH_ATTN_EXT: 2968/2968 on sm86 and sm89 with and without cuDNN and with the opt-in switches.GGML_CUDA_CUDNN_ATTN=0: identical to MiniMax-H3: sage takes the fused RoPE op, faster video and audio VAE decode (B200 max render 22.6 to 15.4 s, video bit-identical) #19 on the Ada.NEEDEDcudnn or nvrtc, theLICENSEcarries the cudnn-frontend notice, zip 1240 MB. The bundle loads cuDNN on the Ada (sampling 57.6 s against 79.3 s without it) and builds no plans on the 3090. Withoutnvrtc-devthe leg failed to compile, and since it iscontinue-on-errorthe release would have shipped without the CUDA asset.Not covered
nvidia-cudnn-cu12, pointed at byGGML_CUDA_CUDNN_LIB). Studio already provides one on Linux sm80+; it is only used by default on sm89 and sm100.