Repository navigation
ggml-cuda: Q8_0 weights take dequant-to-BF16 + cuBLAS by default; k-quants stay on MMQ - #39
Merged
Merged
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 Summary: ✅ CompletedThis 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.
💡 Codex Reviewstable-diffusion.cpp/include/stable-diffusion.h Lines 556 to 558 in ef55f7d  Preserve the existing C API signatures
Existing clients linked against the shared library still call AGENTS.md reference: AGENTS.md:L74-L76 stable-diffusion.cpp/examples/server/routes_openai.cpp Lines 160 to 161 in ef55f7d Restore uploaded-image initialization for OpenAI edits
For ℹ️ 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 Etherll/unsloth
that referenced
this pull request
Nov 8, 2026
…F, not its companions (1.57x on A100) (unslothai#17098) ## Problem On CUDA, Studio sets `GGML_CUDA_QUANT_CUBLAS_MIN_BATCH` for the sd.cpp child when the model is Q8_0, so its large matmuls run as BF16 dequant + cuBLAS instead of int8 MMQ (unslothai#14983: faster on A100 / Hopper / datacenter Blackwell, and closer to the unquantized render). Since unslothai#15054 the check takes the union of the quantized types of the diffusion GGUF and every companion GGUF, so a Q8_0 DiT next to a k-quant encoder gets no env and the whole denoiser stays on MMQ. That pairing is common: the Qwen-Image families (text-to-image, edit, layered) load Qwen2.5-VL Q4_K_M as their encoder, `UNSLOTH_SD_CPP_TEXT_ENCODER_QUANT=q4_k_m` picks Qwen3-4B Q4_K_M, and a cache-only load falls back to an older cached Q4_K_M encoder. A100 40 GB, Studio's pinned prebuilt (`master-982-0518b34-uaba3a1c`, Linux CUDA 13 bundle), Z-Image Turbo Q8_0 DiT + Qwen3-4B Q4_K_M encoder, 1024 px, 8 steps, cfg 1, seed 42, Studio's native argv (`--diffusion-fa --diffusion-conv-direct`), one warm-up per arm then 3 interleaved reps. Each arm gets the env its own tree computes from the same GGUF files (`main`: none; this branch: `1024`): | | sampling s, median [min, max] | encoder s | VAE s | peak VRAM | |---|---|---|---|---| | `main` (MMQ) | 8.02 [8.01, 8.06] | 0.71-0.73 | 0.62 | 13148 MiB | | this branch (BF16 cuBLAS) | 5.10 [5.10, 5.30] | 0.71 | 0.62 | 13148 MiB | 1.57x faster sampling. Each arm renders the same pixels in every rep; the two arms differ, as expected when the denoiser changes matmul path. ## Why the companion rule can go unslothai#15054 added the union because the env is process-wide: with it set, an encoder's quantized matmuls with at least that many rows also take BF16 cuBLAS. In practice: - A text prompt is far below 1024 rows, so the encoder keeps MMQ. The encoder time above is unchanged (0.71 s both arms). On an earlier A100 run of the same Z-Image Q8_0 + Q4_K_M pair (fork master build), the env render was pixel-identical to the fork's Q8_0-only default from unslothai/stable-diffusion.cpp#39, which never moves a k-quant. So the encoder took the same path either way. - Where an encoder does reach the threshold (a large reference image through Qwen2.5-VL in an edit), those k-quant matmuls take BF16 cuBLAS. That was faster on these cards in unslothai#12527 (MiniMax-H3, k-quant DiT: 2.3x on A100 and B200) and closer to an F32 reference there. The cost is a BF16 copy of one weight at a time as scratch. ## Change - The load path now decides from `gguf_quantized_weight_types(files.diffusion_model)` instead of the union, and `model_files_quantized_weight_types` (its only caller) is removed. - Unchanged: - the per-GPU table: sm80 / sm90 / sm100 1024, sm120 4096, Ada / consumer Ampere / Orin / older MMQ; - every card sd.cpp may pick must agree; - a mixed, k-quant, 4-bit, unquantized or unreadable diffusion GGUF keeps MMQ; - a user-exported value always wins, `0` included. - Once the pin carries the fork's Q8_0-only default (unslothai/stable-diffusion.cpp#39, patch 0016), Studio can stop setting the env and leave a user value as the only override. The comment at `QUANT_CUBLAS_ENV` says so. ## Cells Each tree's own `sd_cpp_args` and load glue on real GGUFs (Z-Image Turbo Q8_0 / Q4_K_M DiT x Qwen3-4B Q4_K_M / Q8_0 / BF16 / safetensors encoder x sm75 / 80 / 86 / 89 / 90 / 100 / 120): 56 cells, 52 identical, 4 moved. The 4 moved cells are all Q8_0 DiT + Q4_K_M encoder on sm80 / sm90 / sm100 (none -> 1024) and sm120 (none -> 4096). L4 (sm89), consumer Ampere and T4 do not move. ## Tests `tests/test_sd_cpp_backend.py`: `test_a_q8_0_diffusion_gguf_with_a_k_quant_companion_keeps_mmq` becomes `test_a_q8_0_diffusion_gguf_takes_the_route_whatever_its_companions_hold`: a k-quant encoder, a Q8_0 encoder with a Q5_K VAE, an unreadable encoder, and a Qwen2.5-VL k-quant encoder with its projector all get `1024`, and only the diffusion GGUF is read. On `main` all 4 cases fail; the existing k-quant / mixed / unknown diffusion GGUF cases still keep MMQ. `tests/test_sd_cpp_args.py` checks the tensor-table reader and the route on real tiny GGUFs. `test_sd_cpp_args.py` + `test_sd_cpp_backend.py`: 291 passed. ## Not covered - MiniMax-H3 keeps its own setting (`h3_quant_cublas_env`, 1024 on every sm80+ card, every type). It was measured faster with the route on L4 too (unslothai#12527: 54.6 -> 44.3 s/step), so the image table is not applied to it. - Timed on A100 only. H100 / B200 / RTX PRO 6000 take the same env value they already got for an all-Q8_0 load (unslothai#14983), now also with a k-quant encoder. - The Qwen2.5-VL edit case (encoder rows above the threshold) was not timed here.
Member
Author
|
@codex review |
1 similar comment
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". |
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 master (ggml
135ce1e0, merged in on Nov 7; the patch applies unchanged). 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
Measured: B200 (accuracy and timing), A100, H100, RTX PRO 6000 (Q8_0 timing).
Measured (Nov 7, base ba6ac30 vs this head ef55f7d, CUDA 13.0, 512x512, 3 interleaved repeats, sampling median):
GGML_CUDA_QUANT_CUBLAS_MIN_BATCH=1024)The forced L4 arm is why sm89 keeps MMQ: on Ada the BF16 path is 1.19x slower on Z-Image.
Measured (Nov 7, Colab A100-SXM4-40GB, sm80): Z-Image-Turbo Q8_0 DiT with a Qwen3-4B Q4_K_M text encoder, 1024x1024, 8 steps, 3 interleaved repeats. Master without the env (what a mixed-type load gets today, since the env would also move the Q4_K encoder): 7.74 s sampling [7.87, 7.74, 7.72]; this PR's default: 4.84 s [4.86, 4.81, 4.84], 1.60x, PSNR 30.5 vs master and pixel-identical to master with
GGML_CUDA_QUANT_CUBLAS_MIN_BATCH=1024. The Q4_K encoder stays on MMQ, because the default is per tensor and Q8_0 only.Not measured: consumer RTX 30 / 40 / 50, Jetson Orin, B300 (sm103), GB10 (sm121). sm86 / sm87 / sm89 keep MMQ by rule; sm103 follows sm100 and sm121 follows sm120 by rule.
No change: ROCm, Vulkan, Metal, CPU.
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.