Skip to content

mmq config for Q4_K/Q5_K/Q6_K #3

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
mixa3607 wants to merge 2 commits into
base: master
Choose a base branch
from
Open

mmq config for Q4_K/Q5_K/Q6_K #3

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d1faa85
ggml-cuda(?): gfx906 mmq config for Q4_K/Q5_K/Q6_K
mixa3607 Aug 7, 2026
40418ad
Merge branch 'mxxm-t:master' into gfx906-mmq-overrides
mixa3607 Aug 27, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 13 additions & 0 deletions ggml/src/ggml-cuda/mmq-config-gfx906.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,10 @@
//
// MXFP4: 8 warps as well - rdna2's table with nthreads overridden, so tile
// widths and layout stay in sync with upstream.
//
// Q4_K/Q5_K/Q6_K: measured on MI50. I=64 halves the accumulator array and the X
// LDS tile, cutting VGPR/scratch and improving Q5_K/Q6_K throughput ~15-35%.
// Q4_K stays at I=128 (I=64 makes it ~53% slower). J=64 only.
static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_gfx906(ggml_type type, int J, bool fallback) {
if (type == GGML_TYPE_Q8_0 && J >= 8 && J <= 128 && (J % 8) == 0) {
return ggml_cuda_mmq_config(
Expand All @@ -24,5 +28,14 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
return ggml_cuda_mmq_config(
rdna2.type, 512, rdna2.occupancy, rdna2.I, rdna2.J, rdna2.sram_layout, rdna2.K_vram, rdna2.stream_k, rdna2.fallback);
}
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);

CASE(GGML_TYPE_Q5_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);

CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);

return ggml_cuda_mmq_get_config_rdna2(type, J, fallback);
}