Commit 8fc037672 for llama.cpp
commit 8fc0376721e9662d498a5706d7ee8e7574ce99f4
Author: mmalpe-nv <mmalpe@nvidia.com>
Date: Tue Nov 3 15:54:42 2026 +0530
cuda: potential fix for Q8_1 correction sum (#31247)
* cuda: potential fix for Q8_1 correction sum
* Update ggml/src/ggml-cuda/quantize.cu
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
* addressed review comment to change the quantize_mmq_q8_1 flow as well
---------
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu
index ea089a8c6..727764cdd 100644
--- a/ggml/src/ggml-cuda/quantize.cu
+++ b/ggml/src/ggml-cuda/quantize.cu
@@ -83,21 +83,21 @@ static __global__ void quantize_q8_1(
ggml_cuda_pdl_sync();
const float xi = i0 < ne00 ? x[i03*s03 + i02*s02 + i01*s01 + i00] : 0.0f;
float amax = fabsf(xi);
- float sum = xi;
amax = warp_reduce_max<QK8_1>(amax);
- sum = warp_reduce_sum<QK8_1>(sum);
const float d = amax / 127.0f;
const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
+ const int sum = warp_reduce_sum<QK8_1>((int) q);
+
y[ib].qs[iqs] = q;
if (iqs > 0) {
return;
}
- y[ib].ds = make_half2(d, sum);
+ y[ib].ds = make_half2(d, d * sum);
}
__device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) {
@@ -500,17 +500,6 @@ static __global__ void quantize_mmq_q8_1(
amax = fmaxf(amax, __shfl_xor_sync(0xFFFFFFFF, amax, offset, WARP_SIZE));
}
- float sum;
- if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4) {
- sum = xi.x + xi.y + xi.z + xi.w;
-
- // Calculate sums across vals_per_sum/4 threads.
-#pragma unroll
- for (int offset = vals_per_sum/8; offset > 0; offset >>= 1) {
- sum += __shfl_xor_sync(0xFFFFFFFF, sum, offset, WARP_SIZE);
- }
- }
-
const float d_inv = 127.0f / amax;
char4 q;
q.x = roundf(xi.x*d_inv);
@@ -519,6 +508,19 @@ static __global__ void quantize_mmq_q8_1(
q.w = roundf(xi.w*d_inv);
const float d = 1.0f / d_inv;
+ float sum = 0.0f;
+ if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4) {
+ int sum_q = int(q.x) + int(q.y) + int(q.z) + int(q.w);
+
+ // Calculate sums across vals_per_sum/4 threads.
+#pragma unroll
+ for (int offset = vals_per_sum/8; offset > 0; offset >>= 1) {
+ sum_q += __shfl_xor_sync(0xFFFFFFFF, sum_q, offset, WARP_SIZE);
+ }
+
+ sum = d * float(sum_q);
+ }
+
// write the block once (normal) or to each of the token's compact rows (scatter)
const int nwrite = scatter ? n_expert_used : 1;
#pragma unroll