From e6b6ae55f41e338871283606912856a771b01741 Mon Sep 17 00:00:00 2001 From: Henri Vasserman Date: Fri, 11 Aug 2023 09:30:28 +0300 Subject: [PATCH] Undo mess --------- Co-authored-by: ardfork <134447697+ardfork@users.noreply.github.com> --- ggml-cuda.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml-cuda.cu b/ggml-cuda.cu index a89cbab7f..c1b179c6b 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -1663,8 +1663,8 @@ template static __device__ __forceinline__ float vec_dot_q8_1_q8_1_imp #else const float2 dm8f = __half22float2(dm8); const float2 ds8f = __half22float2(ds8); - const float d8d8 = __low2float(dm8) * __low2float(ds8); - const float m8s8 = __high2float(dm8) * __high2float(ds8); + const float d8d8 = dm8f.x * ds8f.x; + const float m8s8 = dm8f.y * ds8f.y; #endif // GGML_CUDA_F16 // scale second part of sum by QI8_1/ vdr to compensate for multiple threads adding it