From bffb6e4b0fbde3f7d4846004f20c063db755f5be Mon Sep 17 00:00:00 2001 From: R0CKSTAR Date: Thu, 1 Oct 2026 04:33:22 +0800 Subject: [PATCH] cuda: guard the iq4_nl dequantize row kernel against short rows (llama/29683) dequantize_block_iq4_nl writes QK_K values per block, but a row can be shorter than that (an IQ4_NL row is only guaranteed to be a multiple of QK4_NL). Threads whose 32-value sub-block starts at or past k currently read and write past the end of the row. Skip those sub-blocks; for rows that are a multiple of QK_K the check never fires. --- ggml/src/ggml-cuda/convert.cu | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 0619f4760..77fd28894 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -223,9 +223,14 @@ static __global__ void dequantize_block_iq1_m(const void * __restrict__ vx, dst_ } template -static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst_t * __restrict__ yy) { +static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst_t * __restrict__ yy, int nb32) { const int64_t i = blockIdx.x; + const int64_t ib = 8*i + threadIdx.x%8; + if (ib >= nb32) { + return; + } + dequantize_iq4_nl(vx, i, yy + i*QK_K, threadIdx.x); } @@ -352,8 +357,9 @@ static void dequantize_row_iq1_s_cuda(const void * vx, dst_t * y, const int64_t template static void dequantize_row_iq4_nl_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { + const int nb32 = k / QK4_NL; const int nb = (k + QK_K - 1) / QK_K; - dequantize_block_iq4_nl<<>>(vx, y); + dequantize_block_iq4_nl<<>>(vx, y, nb32); } template