diff --git a/sycl/source/detail/builtins_integer.cpp b/sycl/source/detail/builtins_integer.cpp index 49f83becfb6dd..8b3ad1efe54c6 100644 --- a/sycl/source/detail/builtins_integer.cpp +++ b/sycl/source/detail/builtins_integer.cpp @@ -153,10 +153,11 @@ template inline T __s_long_mad_hi(T a, T b, T c) { template inline T __s_mad_sat(T a, T b, T c) { using UPT = typename d::make_larger::type; UPT mul = UPT(a) * UPT(b); + UPT res = mul + UPT(c); const UPT max = d::max_v(); const UPT min = d::min_v(); - 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 inline T __s_long_mad_sat(T a, T b, T c) { diff --git a/sycl/test/built-ins/scalar_integer.cpp b/sycl/test/built-ins/scalar_integer.cpp index 61d8e88542694..528f4fb18aa07 100644 --- a/sycl/test/built-ins/scalar_integer.cpp +++ b/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 buf(&r, cl::sycl::range<1>(1)); + cl::sycl::queue q; + q.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { + 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 };