[SYCL] Fix for cl::sycl::sub_sat() issue on host

Signed-off-by: Sergey Dmitriev <serguei.n.dmitriev@intel.com>
intel · Sep 23, 2019 · 7865dfc · 7865dfc
1 parent 5635959
commit 7865dfc
Show file tree

Hide file tree

Showing 3 changed files with 61 additions and 33 deletions.
diff --git a/sycl/source/detail/builtins_integer.cpp b/sycl/source/detail/builtins_integer.cpp
@@ -167,11 +167,12 @@ template <typename T> inline T __u_sub_sat(T x, T y) {
 }
 
 template <typename T> inline T __s_sub_sat(T x, T y) {
-  if (y > 0)
-    return (y < (x - d::min_v<T>()) ? x - y : d::min_v<T>());
-  if (y < 0)
-    return (y > (x - d::max_v<T>()) ? x - y : d::max_v<T>());
-  return x;
+  using UT = typename std::make_unsigned<T>::type;
+  T result = UT(x) - UT(y);
+  // Saturate result if (+) - (-) = (-) or (-) - (+) = (+).
+  if (((x < 0) ^ (y < 0)) && ((x < 0) ^ (result < 0)))
+    result = result < 0 ? d::max_v<T>() : d::min_v<T>();
+  return result;
 }
 
 template <typename T1, typename T2>

diff --git a/sycl/test/built-ins/scalar_integer.cpp b/sycl/test/built-ins/scalar_integer.cpp
@@ -337,18 +337,29 @@ int main() {
   }
   // sub_sat
   {
-    s::cl_int r{ 0 };
-    {
-      s::buffer<s::cl_int, 1> BufR(&r, s::range<1>(1));
-      s::queue myQueue;
-      myQueue.submit([&](s::handler &cgh) {
-        auto AccR = BufR.get_access<s::access::mode::write>(cgh);
-        cgh.single_task<class sub_satSI1SI1>([=]() {
-          AccR[0] = s::sub_sat(s::cl_int{ 10 }, s::cl_int(0x80000000));
-        }); // 10 - (-2^31(minimum value)) = saturates on Maximum value
-      });
-    }
-    assert(r == 0x7FFFFFFF);
+    auto TestSubSat = [](s::cl_int x, s::cl_int y) {
+      s::cl_int r{ 0 };
+      {
+        s::buffer<s::cl_int, 1> BufR(&r, s::range<1>(1));
+        s::queue myQueue;
+        myQueue.submit([&](s::handler &cgh) {
+          auto AccR = BufR.get_access<s::access::mode::write>(cgh);
+          cgh.single_task<class sub_satSI1SI1>([=]() {
+            AccR[0] = s::sub_sat(x, y);
+          });
+        });
+      }
+      return r;
+    };
+    // 10 - (-2^31(minimum value)) = saturates on Maximum value
+    s::cl_int r1 = TestSubSat(10, 0x80000000);
+    assert(r1 == 0x7FFFFFFF);
+    s::cl_int r2 = TestSubSat(0x7FFFFFFF, 0xFFFFFFFF);
+    assert(r2 == 0x7FFFFFFF);
+    s::cl_int r3 = TestSubSat(0x80000000, 0x00000001);
+    assert(r3 == 0x80000000);
+    s::cl_int r4 = TestSubSat(10499, 30678);
+    assert(r4 == -20179);
   }
 
   // upsample - 1

diff --git a/sycl/test/built-ins/vector_integer.cpp b/sycl/test/built-ins/vector_integer.cpp
@@ -481,22 +481,38 @@ int main() {
 
   // sub_sat
   {
-    s::cl_int2 r{ 0 };
-    {
-      s::buffer<s::cl_int2, 1> BufR(&r, s::range<1>(1));
-      s::queue myQueue;
-      myQueue.submit([&](s::handler &cgh) {
-        auto AccR = BufR.get_access<s::access::mode::write>(cgh);
-        cgh.single_task<class sub_satSI2SI2>([=]() {
-          AccR[0] = s::sub_sat(s::cl_int2{ 10, 10 },
-                               s::cl_int2{ int(0x80000000), int(0x80000000) });
-        });
-      });
-    }
-    s::cl_int r1 = r.x();
-    s::cl_int r2 = r.y();
-    assert(r1 == 0x7FFFFFFF);
-    assert(r2 == 0x7FFFFFFF);
+    auto TestSubSat = [](s::cl_int2 x, s::cl_int2 y) {
+      s::cl_int2 r{ 0 };
+      {
+        s::buffer<s::cl_int2, 1> BufR(&r, s::range<1>(1));
+        s::queue myQueue;
+        myQueue.submit([&](s::handler &cgh) {
+          auto AccR = BufR.get_access<s::access::mode::write>(cgh);
+          cgh.single_task<class sub_satSI2SI2>([=]() {
+            AccR[0] = s::sub_sat(x, y);
+          });
+        });
+      }
+      return r;
+    };
+    s::cl_int2 r1 = TestSubSat(s::cl_int2{ 10, 10 },
+                               s::cl_int2{ 0x80000000, 0x80000000 });
+    s::cl_int r1x = r1.x();
+    s::cl_int r1y = r1.y();
+    assert(r1x == 0x7FFFFFFF);
+    assert(r1y == 0x7FFFFFFF);
+    s::cl_int2 r2 = TestSubSat(s::cl_int2{ 0x7FFFFFFF, 0x80000000 },
+                               s::cl_int2{ 0xFFFFFFFF, 0x00000001 });
+    s::cl_int r2x = r2.x();
+    s::cl_int r2y = r2.y();
+    assert(r2x == 0x7FFFFFFF);
+    assert(r2y == 0x80000000);
+    s::cl_int2 r3 = TestSubSat(s::cl_int2{ 10499, 30678 },
+                               s::cl_int2{ 30678, 10499 });
+    s::cl_int r3x = r3.x();
+    s::cl_int r3y = r3.y();
+    assert(r3x == -20179);
+    assert(r3y ==  20179);
   }
 
   // upsample - 1