Repository navigation
ggml-cuda: Q8_0 weights take dequant-to-BF16 + cuBLAS by default; k-quants stay on MMQ - #39
Open
danielhanchen wants to merge 1 commit into
Open
danielhanchen wants to merge 1 commit into
danielhanchen wants to merge 1 commit into
Conversation
…100, Hopper and datacenter Blackwell k-quants and sub-8-bit types keep MMQ; GGML_CUDA_QUANT_CUBLAS_MIN_BATCH still overrides for every type.
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. |
Member
Author
|
@codex review |
|
Codex Review: Didn't find any major issues. Nice work! Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
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". |
danielhanchen
added a commit
to danielhanchen/unsloth-staging-2
that referenced
this pull request
Oct 23, 2026
…ts on MMQ (unslothai#14983) ## Problem When an image GGUF runs on the stable-diffusion.cpp CUDA build (the native engine picked for the load or forced with `UNSLOTH_DIFFUSION_ENGINE=sd_cpp`, and the families only sd.cpp can load, such as MOSS-Image GGUFs), every quantized DiT matmul runs on int8 MMQ, which also quantizes the activations to 8 bits. For a Q8_0 checkpoint that 8-bit activation step is the larger error. CUDA hosts otherwise use the diffusers path, which this PR does not touch. The pinned sd.cpp build already has the alternative (`GGML_CUDA_QUANT_CUBLAS_MIN_BATCH`: dequantize the weight to BF16, cuBLAS GEMM with FP32 accumulate), but Studio only turns it on for MiniMax-H3 video. ## Fix A CUDA sd.cpp load of a **Q8_0 diffusion GGUF** now gets `GGML_CUDA_QUANT_CUBLAS_MIN_BATCH` in its child env when every card it may run on is one where the route measured faster: - A100 (8.0), Hopper (9.x), datacenter Blackwell (10.x): `1024` - RTX PRO 6000 / RTX 50 (12.x): `4096` - Ada / L4, consumer Ampere, Orin and anything else: unset (MMQ) The env knob applies to every quantized type, so Studio reads the diffusion GGUF's tensor table (memory-mapped, no tensor data is read) and sets it only when every quantized tensor is Q8_0. K-quant and 4-bit GGUFs (Q4_K_M, Q5_K_M, Q6_K, Q4_0, ...), mixed files, safetensors and unreadable files keep MMQ. The capability comes from torch for the selected card (every visible card when unpinned; mixed cards that disagree get nothing). A user-exported value always wins, and nothing changes off CUDA. The same value reaches both the resident sd-server and the one-shot sd-cli (`child_env`). The fork change `cuda-q8-bf16-cublas` makes the same Q8_0-only choice the build's default. Once the pin carries it, this env is redundant but harmless (same value, same types). ## Accuracy Against an unquantized render of the same model (BF16 weights in the denoiser, same `sd-cli` build), B200, seeds 1-6 at 512 px plus seed 1 at 1024 px, LPIPS (AlexNet) mean over the 7 cells: | Q8_0 model | MMQ (main) | BF16 route (this PR) | Cells closer | |---|---|---|---| | Z-Image Turbo | 0.0148 | 0.0052 | 6/7 | | FLUX.2 klein 4B | 0.0057 | 0.0047 | 5/7 | | FLUX.3 Image Klein | 0.0019 | 0.0012 | 6/7 | Each arm renders identically run to run. Q4_K_M GGUFs stay on MMQ: routing them too was a coin flip against the same references (FLUX.2 klein 0.0736 -> 0.0738), because the 4-bit weights dominate both arms. ## Speed Clean VMs (Colab), published `master-982-0518b34-ud331c66` bundle, `sd-cli`, Q8_0 DiT + Qwen3-4B Q8_0 encoder, seed 42, cfg 1, MMQ (env unset), MMQ again (A/A), and `GGML_CUDA_QUANT_CUBLAS_MIN_BATCH=1024` (the kernels this branch selects for Q8_0 by default). One VM per GPU, every arm on that VM, 4 rounds interleaved, first dropped, median sampling seconds [min, max]: | GPU | Model, size | MMQ | MMQ again (A/A) | BF16 cuBLAS | Speedup | |---|---|---|---|---|---| | A100 40 GB (sm80) | Z-Image Turbo, 1024, 8 steps | 8.01 [8.01, 8.01] | 8.02 [8.01, 8.02] | 5.09 [5.09, 5.10] | 1.57x | | A100 40 GB (sm80) | Z-Image Turbo, 512 | 3.54 [3.54, 3.54] | 3.54 [3.54, 3.55] | 2.90 [2.76, 2.92] | 1.22x | | A100 40 GB (sm80) | FLUX.2 klein 4B, 1024, 4 steps | 3.08 [3.08, 3.10] | 3.09 [3.08, 3.10] | 2.20 [2.19, 2.20] | 1.40x | | A100 40 GB (sm80) | FLUX.2 klein 4B, 512 | 1.74 [1.74, 1.74] | 1.74 [1.74, 1.74] | 1.49 [1.49, 1.50] | 1.17x | | H100 80 GB (sm90) | Z-Image Turbo, 1024 | 3.75 [3.60, 3.75] | 3.60 [3.60, 3.62] | 2.46 [2.45, 2.46] | 1.46-1.52x | | H100 80 GB (sm90) | Z-Image Turbo, 512 | 1.66 [1.65, 1.68] | 1.67 [1.65, 1.67] | 1.42 [1.42, 1.58] | 1.17x | | H100 80 GB (sm90) | FLUX.2 klein 4B, 1024 | 1.58 [1.57, 1.58] | 1.58 [1.57, 1.58] | 1.23 [1.23, 1.23] | 1.28x | | H100 80 GB (sm90) | FLUX.2 klein 4B, 512 | 0.96 [0.96, 0.96] | 0.96 [0.96, 0.96] | 0.86 [0.86, 0.86] | 1.12x | | RTX PRO 6000 (sm120) | Z-Image Turbo, 1024 | 2.46 [2.45, 2.46] | 2.62 [2.45, 2.62] | 2.48 [2.48, 2.48] | even | | RTX PRO 6000 (sm120) | Z-Image Turbo, 512 | 1.01 [1.00, 1.16] | 1.00 [1.00, 1.01] | 1.08 [1.07, 1.08] | 0.93x (default keeps MMQ here) | | RTX PRO 6000 (sm120) | FLUX.2 klein 4B, 1024 | 1.11 [0.97, 1.12] | 1.12 [0.97, 1.12] | 1.10 [0.95, 1.10] | even | | RTX PRO 6000 (sm120) | FLUX.2 klein 4B, 512 | 0.51 [0.50, 0.65] | 0.65 [0.65, 0.65] | 0.67 [0.67, 0.67] | 0.76-1.0x (default keeps MMQ here) | On sm120 the default threshold is 4096 rows, so a 512 px image (1024 image tokens) stays on MMQ and only the 1024 px rows, where the route is even with MMQ, change path. Every arm rendered the same pixels in every round. L4 could not be allocated on Colab during this run; Ada keeps MMQ by rule, so this branch changes nothing there. B200 (sm100, this host, exclusive card, resident `sd-server`, Qwen3-4B BF16 encoder, 1024 px, this branch vs its base): Z-Image Turbo Q8_0 sampling 9.85 [7.85, 14.26] (MMQ), 8.16 [5.51, 9.79] (MMQ again), 5.03 [3.94, 8.61] (this branch), about 1.6x. The host's load average was 900-1200, so these spread 2-5x; the Colab table above is the clean measurement. ## Tests `tests/test_sd_cpp_args.py`: the capability table, mixed-card disagreement, the user override, Q8_0-only gating (k-quant, mixed, 4-bit, unquantized and unknown files get nothing), and the tensor-table reader on real tiny GGUFs. `tests/test_sd_cpp_backend.py`: a CUDA server load of a Q8_0 GGUF gets the env (and records it in `child_env` for one-shot sd-cli), a k-quant / mixed / unknown GGUF does not and the check reads the diffusion model, not the encoder; Ada / unknown / CPU loads get nothing; a user value is left alone. `pytest tests/test_sd_cpp_args.py tests/test_sd_cpp_backend.py tests/test_sd_cpp_engine.py tests/test_h3_native_server.py`: 378 passed. On this host's GGUFs the reader returns `{Q8_0}` for the Z-Image Turbo and FLUX.2 klein Q8_0 files and `{Q4_K, Q5_K, Q6_K}` / `{Q4_K, Q5_K}` for their Q4_K_M files. ## Not covered Q8_0 timing covers A100, H100, RTX PRO 6000 and B200; L4 could not be allocated (Ada keeps MMQ by rule). Consumer RTX 30 / 40 / 50 cards, Jetson, B300 and GB10 were not timed (the rule keeps MMQ on sm86 / sm87 / sm89). MiniMax-H3 keeps its own setting (`1024` from sm80, every type). Windows CUDA uses the same env path but was not run. ## Related Pairs with unslothai/stable-diffusion.cpp#39: the Studio side sets the env only for all-Q8_0 GGUFs on today's pin; once the pin carries the fork patch, that env is redundant but harmless. --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
danielhanchen
added a commit
to danielhanchen/unsloth-staging-2
that referenced
this pull request
Oct 24, 2026
…ess loads is Q8_0 (unslothai#15054) ## Summary Follow-up to unslothai#14983. Studio sets `GGML_CUDA_QUANT_CUBLAS_MIN_BATCH` for the sd.cpp child only when the model is all Q8_0, so large Q8_0 matmuls run as BF16 dequant + cuBLAS while k-quants stay on MMQ. The check read only the diffusion GGUF, but the variable is process-wide: - The pinned fork patch (`scripts/unsloth/ggml-patches/0001-ggml-cuda-h3-speed.patch`) reads it once per process and applies it to every quantized `src0` that has a BF16 converter. Q4_K, Q5_K and Q6_K all have one. - Nothing scopes it per tensor or per model. unslothai/stable-diffusion.cpp#39 adds a Q8_0-only default when the variable is unset, but a set value still applies to every quantized type. So a Q8_0 denoiser next to a k-quant text encoder (for example the default Qwen-Image-2.1 encoder) also moved the encoder's matmuls with 1024+ rows (4096 on sm120) off MMQ. That breaks the "k-quants stay on MMQ" rule and adds full-size BF16 dequant scratch for those weights. ## Change - `sd_cpp_args.model_files_quantized_weight_types(files)` unions the quantized tensor types of the diffusion model and every companion that sd.cpp loads: VAE, audio VAE, clip_l, clip_g, t5xxl, llm, llm_vision, qwen2vl. - Companions that are not GGUF (safetensors) load dense and add nothing. - An unreadable companion GGUF returns None, and unknown keeps MMQ, the same as an unreadable diffusion GGUF today. - The sd.cpp load path passes this union instead of the diffusion-only types. The existing `quant_cublas_env` gate is unchanged: the override is still set only when every quantized type is Q8_0, every card sd.cpp may pick agrees, and the user has not set the variable. ## Tests - `test_a_q8_0_diffusion_gguf_with_a_k_quant_companion_keeps_mmq` (backend load path) covers: - a Q8_0 diffusion GGUF with a Q4_K/Q6_K text encoder: no override - a Q8_0 TE with a Q5_K VAE: no override - an unreadable TE GGUF: no override - a Q8_0 TE with a dense GGUF VAE: override kept at 1024 - `test_gguf_quantized_weight_types_reads_the_tensor_table` now writes real Q8_0, Q4_K and F16 GGUFs and checks the union and the resulting env. On main, all five new cases fail: the k-quant TE cases get `{'GGML_CUDA_QUANT_CUBLAS_MIN_BATCH': '1024'}`, and the helper does not exist. With this change, `tests/test_sd_cpp_args.py` and `tests/test_sd_cpp_backend.py` pass (268 passed). ## Notes - This rule is conservative. A k-quant TE pinned to CPU with `--clip-on-cpu` never reaches ggml-cuda, but it still turns the route off. A narrower rule would need the final placement flags, which the fallback ladder can change after the env is chosen. - LoRA files (`--lora-model-dir`) and the upscaler are not inspected. Their behaviour is the same as before this change.
This branch has not been deployed
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
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.
Based on
sync-upstream-a77d362(#35,344a3b66). One new ggml patch, no source change.Problem
On NVIDIA every quantized DiT / UNet matmul goes to MMQ: the activations are quantized to q8_1 and multiplied with int8
mma.synctiles. For a Q8_0 checkpoint that makes the activations the larger error: the weights are 8-bit, and so are the activations. A diffusion step multiplies thousands of rows per call (4096 image tokens for a 1024 px Z-Image step), where a BF16 tensor-core GEMM from cuBLAS is also faster on datacenter cards.Patch 0001 already has the route (
GGML_CUDA_QUANT_CUBLAS_MIN_BATCH: dequantize the weight to BF16, cuBLAS GEMM with FP32 accumulate), but it defaults to off.Fix
ggml-patches/0016gives the route a default for Q8_0 weights only:K-quants (Q2_K to Q6_K), Q4_0 / Q4_1 / Q5_0 / Q5_1, IQ and other sub-8-bit types keep MMQ. For those the 4-bit weights dominate the error and neither path is reliably closer to the unquantized model (measured below), so there is no accuracy reason to switch them. Q8_0 is the only 8-bit block type GGUF weights use: Q8_1 is the activation format and Q8_K is internal to the k-quant dot products. Matmuls below the row threshold, F16 / BF16 weights, AMD and every non-CUDA backend are untouched. A set
GGML_CUDA_QUANT_CUBLAS_MIN_BATCHstill applies to every quantized type, as before (0= MMQ everywhere), so MiniMax-H3's explicit setting keeps its behaviour.Accuracy
Reference: an unquantized GGUF of the same denoiser (every tensor the Q8_0 GGUF quantizes is the original diffusers weight in BF16, every other tensor copied from the Q8_0 GGUF; the name mapping is checked against the dequantized Q8_0 tensors, cosine 1.0000), rendered by the same
sd-clibuild. FLUX.3 Image Klein uses unsloth's published BF16 GGUF. B200,sd-cli, sd.cpp defaults, Qwen3-4B BF16 encoder (FLUX.3: Qwen3-VL-4B BF16), one prompt, seeds 1-6 at 512 px plus seed 1 at 1024 px, Z-Image 8 steps, FLUX.2 klein 4 steps, FLUX.3 16 steps, cfg 1. Base = #35 with its 13 ggml patches (MMQ), this branch = the same tree plus 0016. LPIPS (AlexNet) against the reference, mean over the 7 cells:Per cell (MMQ / this branch):
Every arm renders identical pixels run to run (A/A floor 0). Closer on average on all three families, and at 1024 px on all three. Four of the 21 cells are further: Z-Image s6 (+0.0005), FLUX.2 klein s5 (+0.0007) and s6 (+0.0056), FLUX.3 s3 (+0.0001).
MMQ kept for everything else: the Q4_K_M GGUFs of the same three models (Q4_K / Q5_K / Q6_K tensors) render pixel-identical to the base on all 21 cells. The unquantized reference renders identically on both builds, and a fresh reference render matches the earlier session's pixel for pixel.
For comparison, routing Q4_K_M as well (the earlier all-types variant) was a coin flip against the same references: Z-Image 0.0905 -> 0.0889, FLUX.3 0.0325 -> 0.0311, FLUX.2 klein 0.0736 -> 0.0738. The 4-bit weights dominate both arms there, which is why this branch leaves them on MMQ.
Speed
Clean VMs (Colab), published
master-982-0518b34-ud331c66bundle,sd-cli, Q8_0 DiT + Qwen3-4B Q8_0 encoder, seed 42, cfg 1, MMQ (env unset), MMQ again (A/A), andGGML_CUDA_QUANT_CUBLAS_MIN_BATCH=1024(the kernels this branch selects for Q8_0 by default). One VM per GPU, every arm on that VM, 4 rounds interleaved, first dropped, median sampling seconds [min, max]:On sm120 the default threshold is 4096 rows, so a 512 px image (1024 image tokens) stays on MMQ and only the 1024 px rows, where the route is even with MMQ, change path. Every arm rendered the same pixels in every round. L4 could not be allocated on Colab during this run; Ada keeps MMQ by rule, so this branch changes nothing there.
B200 (sm100, this host, exclusive card, resident
sd-server, Qwen3-4B BF16 encoder, 1024 px, this branch vs its base): Z-Image Turbo Q8_0 sampling 9.85 [7.85, 14.26] (MMQ), 8.16 [5.51, 9.79] (MMQ again), 5.03 [3.94, 8.61] (this branch), about 1.6x. The host's load average was 900-1200, so these spread 2-5x; the Colab table above is the clean measurement.Reproduce
scripts/sdcpp_speed/make_ref_gguf.py <Q8_0 gguf> <diffusers transformer> <zimage|klein> <out.gguf>.scripts/sdcpp_speed/acc_q8_run.sh <family> <outdir>. Scores:scripts/sdcpp_speed/acc_q8_score.py. Timing:scripts/sdcpp_speed/q8_speed.sh.Platforms
Related
Pairs with unslothai/unsloth#14983: the Studio side sets the env only for all-Q8_0 GGUFs on today's pin; once the pin carries the fork patch, that env is redundant but harmless.