From 7914776a72b04602e4c5718d0014428579c6e166 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrien=20Gallou=C3=ABt?= Date: Sun, 4 Oct 2026 15:47:21 +0200 Subject: [PATCH] cuda : move blocks_per_col to where it is used (llama/29939) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Adrien Gallouët --- ggml/src/ggml-cuda/quantize.cu | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index bcc772395..935b2bf4b 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -131,8 +131,6 @@ static __global__ void quantize_mmq_nvfp4( const int64_t ne0, const int64_t ne1, const int64_t ne2, const int n_expert_used) { #if defined(BLACKWELL_MMA_AVAILABLE) - const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ; - int64_t base_idx; if constexpr (scatter) { base_idx = (int64_t) blockIdx.x * s02; // one physical row per token @@ -317,7 +315,8 @@ static __global__ void quantize_mmq_nvfp4( reinterpret_cast(yb->d4)[sub] = fp8_code; } } else { - block_fp4_mmq * yb = y + (blockIdx.y * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x); + const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ; + block_fp4_mmq * yb = y + (blockIdx.y * (blocks_per_col * ne1) + k_block * ne1 + blockIdx.x); uint32_t * yqs = reinterpret_cast(yb->qs); yqs[2 * sub + 0] = q0; yqs[2 * sub + 1] = q1;