[SYCL] Fix final result saturation in mad_sat host implementation (#1025)

Signed-off-by: Joey Genfi <[email protected]>

Author: jgstarIntel

sycl/source/detail/builtins_integer.cpp

@@ -153,10 +153,11 @@ template <typename T> inline T __s_long_mad_hi(T a, T b, T c) {
 template <typename T> inline T __s_mad_sat(T a, T b, T c) {
   using UPT = typename d::make_larger<T>::type;
   UPT mul = UPT(a) * UPT(b);
+  UPT res = mul + UPT(c);
   const UPT max = d::max_v<T>();
   const UPT min = d::min_v<T>();
-  mul = std::min(std::max(mul, min), max);
-  return __s_add_sat(T(mul), c);
+  res = std::min(std::max(res, min), max);
+  return T(res);
 }
 
 template <typename T> inline T __s_long_mad_sat(T a, T b, T c) {

sycl/test/built-ins/scalar_integer.cpp

@@ -287,6 +287,28 @@ int main() {
     assert(r == 0x7FFFFFFF);
   }
 
+  // mad_sat test two
+  {
+    char r(0);
+    char exp(120);
+    {
+      cl::sycl::buffer<char, 1> buf(&r, cl::sycl::range<1>(1));
+      cl::sycl::queue q;
+      q.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
+        cgh.single_task<class kernel>([=]() {
+          signed char inputData_0(-17);
+          signed char inputData_1(-10);
+          signed char inputData_2(-50);
+          acc[0] = cl::sycl::mad_sat(inputData_0, inputData_1, inputData_2);
+        });
+      });
+    }
+    assert(r == exp); // Should return the real number of i0*i1+i2 in CPU
+                              // Only fails in vector, but passes in scalar.
+
+  }
+
   // mul_hi
   {
     s::cl_int r{ 0 };