From dbe732c0a67dce84fe9c48112fba88ba829d26bf Mon Sep 17 00:00:00 2001 From: xrip Date: Tue, 8 Sep 2026 17:17:17 +0300 Subject: [PATCH] wip: probe Turing nwarps for low-bpw types (experiment, do not merge) --- ggml/src/ggml-cuda/mmvq.cu | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 5bc1297ff..d5a93992a 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -565,6 +565,18 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d case GGML_TYPE_Q6_K: case GGML_TYPE_IQ4_XS: return 2; + // EXPERIMENT: low-bpw types sit at 23-41% of the f16 bandwidth + // ceiling while these fast types reach 78-95%. They currently get + // nwarps=4. Probe both directions in one build: + // IQ2_XXS, IQ3_XXS -> 2 (fewer warps, less register pressure) + // IQ2_XS, IQ3_S -> 8 (more warps, more memory in flight) + // IQ2_S, IQ1_M -> 4 unchanged, as controls + case GGML_TYPE_IQ2_XXS: + case GGML_TYPE_IQ3_XXS: + return 2; + case GGML_TYPE_IQ2_XS: + case GGML_TYPE_IQ3_S: + return 8; default: return 4; }