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;