small change

Author: JohannesGaessler

ggml-cuda.cu

@@ -1405,7 +1405,8 @@ static __device__ __forceinline__ float vec_dot_q2_K_q8_1(
 
     const int bq8_offset = 4 * (iqs/8);
 
-    float sumf = 0;
+    float sumf_d = 0;
+    float sumf_m = 0;
 
     const float d = bq2_K->d;
     const float dmin = bq2_K->dmin;
@@ -1414,19 +1415,19 @@ static __device__ __forceinline__ float vec_dot_q2_K_q8_1(
 
     for (int i = 0; i < 4; ++i) {
         const int sc = bq2_K->scales[iqs - iqs%8 + (iqs%8) / 4 + 2*i];
-        const float dl = d    * (sc & 0xF);
-        const float ml = dmin * (sc >> 4);
 
         const int vii = (vi >> (2*i)) & 0x03030303;
 
         const block_q8_1 * bq8i = bq8_1 + bq8_offset + i;
         const float d8 = bq8i->d;
         const int qs8 = *((int*) &bq8i->qs[4*(iqs%8)]);
 
-        sumf += d8*(dl*__dp4a(vii, qs8, 0) - ml*__dp4a(0x01010101, qs8, 0));
+        sumf_d += d8 * __dp4a(vii,        qs8, 0) * (sc & 0xF);
+        sumf_m += d8 * __dp4a(0x01010101, qs8, 0) * (sc >> 4);
     }
 
-    return sumf;
+
+    return d*sumf_d - dmin*sumf_m;
 // #else
 //     return 0.0f; // only to satisfy the compiler
 // #endif // __CUDA_ARCH__ >= 600