From d1faa85d227b06efefe986082a948e1f525fd819 Mon Sep 17 00:00:00 2001
From: mixa3607 <30209772+mixa3607@users.noreply.github.com>
Date: Fri, 7 Aug 2026 23:11:53 +0500
Subject: [PATCH] ggml-cuda(?): gfx906 mmq config for Q4_K/Q5_K/Q6_K

---
 ggml/src/ggml-cuda/mmq-config-gfx906.cuh | 13 +++++++++++++
 1 file changed, 13 insertions(+)

diff --git a/ggml/src/ggml-cuda/mmq-config-gfx906.cuh b/ggml/src/ggml-cuda/mmq-config-gfx906.cuh
index c8c7702f2d42..0fdd2ab0d2a3 100644
--- a/ggml/src/ggml-cuda/mmq-config-gfx906.cuh
+++ b/ggml/src/ggml-cuda/mmq-config-gfx906.cuh
@@ -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(
@@ -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);
 }