Fixed dequant precision issues in Q4_1 and Q5_1 (#9711)

ggerganov · Oct 3, 2024 · 5639971 · 5639971
1 parent c83ad6d
commit 5639971
Showing 1 changed file with 8 additions and 8 deletions.
diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp
@@ -55,12 +55,12 @@ static __dpct_inline__ void dequantize_q4_1(const void *vx, const int64_t ib,
 #ifdef GGML_SYCL_F16
  // v = v * {d, d};
  // v = v + {m, m};
- v.s0() = (v.s0() * d) + m;
- v.s1() = (v.s1() * d) + m;
+ v.s0() = sycl::fma(v.s0(), d, m);
+ v.s1() = sycl::fma(v.s1(), d, m);
 
 #else
- v.x() = (v.x() * d) + m;
- v.y() = (v.y() * d) + m;
+ v.x() = sycl::fma(v.x(), d, m);
+ v.y() = sycl::fma(v.y(), d, m);
 #endif // GGML_SYCL_F16
 }
 
@@ -110,11 +110,11 @@ static __dpct_inline__ void dequantize_q5_1(const void *vx, const int64_t ib,
 #ifdef GGML_SYCL_F16
  // v = v * {d, d};
  // v = v + {m, m};
- v.s0() = (v.s0() * d) + m;
- v.s1() = (v.s1() * d) + m;
+ v.s0() = sycl::fma(v.s0(), d, m);
+ v.s1() = sycl::fma(v.s1(), d, m);
 #else
- v.x() = (v.x() * d) + m;
- v.y() = (v.y() * d) + m;
+ v.x() = sycl::fma(v.x(), d, m);
+ v.y() = sycl::fma(v.y(), d, m);
 #endif // GGML_SYCL_F16
 }