modification to handle a special case

Author: vtavana

dpctl/tensor/_elementwise_funcs.py

@@ -131,9 +131,8 @@
         Default: "K".
 Returns:
     usm_narray:
-        An array containing the element-wise inverse cosine, in radians
-        and in the closed interval `[-pi/2, pi/2]`. The data type
-        of the returned array is determined by the Type Promotion Rules.
+        An array containing the element-wise ceiling of input array.
+        The returned array has the same data type as `x`.
 """
 
 ceil = UnaryElementwiseFunc(
@@ -312,9 +311,8 @@
         Default: "K".
 Returns:
     usm_narray:
-        An array containing the element-wise floor, in radians
-        and in the closed interval `[-pi/2, pi/2]`. The data type
-        of the returned array is determined by the Type Promotion Rules.
+        An array containing the element-wise floor of input array.
+        The returned array has the same data type as `x`.
 """
 
 floor = UnaryElementwiseFunc(
@@ -972,9 +970,8 @@
         Default: "K".
 Returns:
     usm_narray:
-        An array containing the element-wise inverse cosine, in radians
-        and in the closed interval `[-pi/2, pi/2]`. The data type
-        of the returned array is determined by the Type Promotion Rules.
+        An array containing the element-wise truncated value of input array.
+        The returned array has the same data type as `x`.
 """
 
 trunc = UnaryElementwiseFunc(

dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp

@@ -69,7 +69,10 @@ template <typename argT, typename resT> struct CeilFunctor
             return in;
         }
         else {
-            return sycl::ceil(in);
+            if (in == 0 && std::signbit(in)) {
+                return in;
+            }
+            return std::ceil(in);
         }
     }
 };

dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor.hpp

@@ -69,7 +69,10 @@ template <typename argT, typename resT> struct FloorFunctor
             return in;
         }
         else {
-            return sycl::floor(in);
+            if (in == 0 && std::signbit(in)) {
+                return in;
+            }
+            return std::floor(in);
         }
     }
 };

dpctl/tests/elementwise/test_floor_ceil_trunc.py

@@ -18,7 +18,11 @@
 
 import numpy as np
 import pytest
-from numpy.testing import assert_allclose, assert_raises_regex
+from numpy.testing import (
+    assert_allclose,
+    assert_array_equal,
+    assert_raises_regex,
+)
 
 import dpctl
 import dpctl.tensor as dpt
@@ -196,12 +200,14 @@ def test_floor_ceil_trunc_special_cases(np_call, dpt_call, dtype):
     q = get_queue_or_skip()
     skip_if_dtype_not_supported(dtype, q)
 
-    x = [np.nan, np.inf, -np.inf]
+    x = [np.nan, np.inf, -np.inf, +0.0, -0.0]
 
     xf = np.array(x, dtype=dtype)
     yf = dpt.asarray(xf, dtype=dtype, sycl_queue=q)
 
     Y_np = np_call(xf)
+    Y = dpt.asnumpy(dpt_call(yf))
 
     tol = 8 * dpt.finfo(dtype).resolution
-    assert_allclose(dpt.asnumpy(dpt_call(yf)), Y_np, atol=tol, rtol=tol)
+    assert_allclose(Y, Y_np, atol=tol, rtol=tol)
+    assert_array_equal(np.signbit(Y), np.signbit(Y_np))