From a13695bdc7c6fb3cd71171bc6fe07615f37fd93c Mon Sep 17 00:00:00 2001 From: thelittlefireman Date: Sat, 3 Oct 2026 16:54:35 +0200 Subject: [PATCH 1/4] ggml-cuda: HIP: fix Q2_K spill on CDNA, RDNA3+ --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index a75dd046e5c3..2ad59312d0aa 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -693,6 +693,7 @@ template static __device__ __forceinline_ const int i0 = (threadIdx.y / ntx) * rows_per_warp; +#pragma unroll 1 for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; From edfa142bc10440206aafa285f3bae65c70640c5d Mon Sep 17 00:00:00 2001 From: thelittlefireman Date: Sat, 3 Oct 2026 16:56:56 +0200 Subject: [PATCH 2/4] ggml-cuda: fix Q2_K spill on AMD: GCN, RDNA1, RDNA2 and NVIDIA Pascal/Volta. --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index 2ad59312d0aa..eef000d43143 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -631,7 +631,7 @@ template static __device__ __forceinline_ for (int k01 = 0; k01 < MMQ_TILE_NE_K/2; k01 += QR2_K*VDR_Q2_K_Q8_1_MMQ) { const int k0 = k00 + k01; -#pragma unroll +#pragma unroll 2 for (int j0 = 0; j0 < J; j0 += nwarps) { const int j = j0 + threadIdx.y; @@ -654,7 +654,7 @@ template static __device__ __forceinline_ for (int k01 = MMQ_TILE_NE_K/2; k01 < MMQ_TILE_NE_K; k01 += QR2_K*VDR_Q2_K_Q8_1_MMQ) { const int k0 = k00 + k01; -#pragma unroll +#pragma unroll 2 for (int j0 = 0; j0 < J; j0 += nwarps) { const int j = j0 + threadIdx.y; From d00c5844c35d64222f11585c129148beb0f5e6b0 Mon Sep 17 00:00:00 2001 From: thelittlefireman Date: Sat, 3 Oct 2026 16:57:46 +0200 Subject: [PATCH 3/4] ggml-cuda: Optimize and reduce spill on Q2_K by remove temporary y_df loop. --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index eef000d43143..d8721109209f 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -619,14 +619,6 @@ template static __device__ __forceinline_ const int * y_qs = (const int *) y + 4; const half2 * y_ds = (const half2 *) y; - float2 y_df[J/nwarps]; -#pragma unroll - for (int j0 = 0; j0 < J; j0 += nwarps) { - const int j = j0 + threadIdx.y; - - y_df[j0/nwarps] = __half22float2(y_ds[j*MMQ_TILE_Y_K]); - } - #pragma unroll for (int k01 = 0; k01 < MMQ_TILE_NE_K/2; k01 += QR2_K*VDR_Q2_K_Q8_1_MMQ) { const int k0 = k00 + k01; @@ -642,7 +634,7 @@ template static __device__ __forceinline_ constexpr int ns = 2; sum[j0/nwarps*I/warp_size + i0/warp_size] += vec_dot_q2_K_q8_1_impl_mmq( &x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0], &y_qs[j*MMQ_TILE_Y_K + k01], - &x_dm[i*(MMQ_TILE_NE_K + 1) + k0/4], k01 < MMQ_TILE_NE_K/2 ? y_df[j0/nwarps].x : y_df[j0/nwarps].y, + &x_dm[i*(MMQ_TILE_NE_K + 1) + k0/4], __half22float2(y_ds[j*MMQ_TILE_Y_K]).x, &y_ds[j*MMQ_TILE_Y_K + (1 + k01/QI8_1)]); } } @@ -665,7 +657,7 @@ template static __device__ __forceinline_ constexpr int ns = 1; sum[j0/nwarps*I/warp_size + i0/warp_size] += vec_dot_q2_K_q8_1_impl_mmq( &x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0], &y_qs[j*MMQ_TILE_Y_K + k01], - &x_dm[i*(MMQ_TILE_NE_K + 1) + k0/4], k01 < MMQ_TILE_NE_K/2 ? y_df[j0/nwarps].x : y_df[j0/nwarps].y, + &x_dm[i*(MMQ_TILE_NE_K + 1) + k0/4], __half22float2(y_ds[j*MMQ_TILE_Y_K]).y, &y_ds[j*MMQ_TILE_Y_K + (1 + k01/QI8_1)]); } } From f2179ef59a49ec63d0f174d78715c38948c9cd3b Mon Sep 17 00:00:00 2001 From: thelittlefireman Date: Mon, 12 Oct 2026 23:12:01 +0200 Subject: [PATCH 4/4] ggml-cuda: rollback unroll --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index d8721109209f..a9fca1f86d50 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -623,7 +623,7 @@ template static __device__ __forceinline_ for (int k01 = 0; k01 < MMQ_TILE_NE_K/2; k01 += QR2_K*VDR_Q2_K_Q8_1_MMQ) { const int k0 = k00 + k01; -#pragma unroll 2 +#pragma unroll for (int j0 = 0; j0 < J; j0 += nwarps) { const int j = j0 + threadIdx.y; @@ -646,7 +646,7 @@ template static __device__ __forceinline_ for (int k01 = MMQ_TILE_NE_K/2; k01 < MMQ_TILE_NE_K; k01 += QR2_K*VDR_Q2_K_Q8_1_MMQ) { const int k0 = k00 + k01; -#pragma unroll 2 +#pragma unroll for (int j0 = 0; j0 < J; j0 += nwarps) { const int j = j0 + threadIdx.y; @@ -685,7 +685,7 @@ template static __device__ __forceinline_ const int i0 = (threadIdx.y / ntx) * rows_per_warp; -#pragma unroll 1 +#pragma unroll for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01;