Added tests for in-place remainder and pow

Fixed in-place remainder for devices that do not support 64-bit floating point data types

Author: ndgrigorian

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

@@ -316,10 +316,11 @@ template <typename fnT, typename T1, typename T2> struct RemainderStridedFactory
 template <typename argT, typename resT> struct RemainderInplaceFunctor
 {
 
-    using supports_sg_loadstore = std::negation<
-        std::disjunction<tu_ns::is_complex<argT>, tu_ns::is_complex<resT>>>;
-    using supports_vec = std::negation<
-        std::disjunction<tu_ns::is_complex<argT>, tu_ns::is_complex<resT>>>;
+    using supports_sg_loadstore = std::true_type;
+    using supports_vec = std::true_type;
+
+    // functor is only well-defined when argT and resT are the same
+    static_assert(std::is_same_v<argT, resT>);
 
     void operator()(resT &res, const argT &in)
     {
@@ -331,7 +332,7 @@ template <typename argT, typename resT> struct RemainderInplaceFunctor
             if constexpr (std::is_signed_v<argT> || std::is_signed_v<resT>) {
                 auto tmp = res;
                 res %= in;
-                if (res != 0 && l_xor(tmp < 0, in < 0)) {
+                if (res != resT(0) && l_xor(tmp < 0, in < 0)) {
                     res += in;
                 }
             }
@@ -347,7 +348,7 @@ template <typename argT, typename resT> struct RemainderInplaceFunctor
                 }
             }
             else {
-                res = std::copysign(0, in);
+                res = sycl::copysign(resT(0), in);
             }
         }
     }
@@ -384,7 +385,7 @@ template <typename argT, typename resT> struct RemainderInplaceFunctor
                     }
                 }
                 else {
-                    res[i] = std::copysign(0, in[i]);
+                    res[i] = sycl::copysign(resT(0), in[i]);
                 }
             }
         }
@@ -444,8 +445,10 @@ struct RemainderInplaceContigFactory
 {
     fnT get()
     {
-        if constexpr (std::is_same_v<typename RemainderOutputType<T1, T2>::value_type,
-                                     void>) {
+        if constexpr (std::is_same_v<
+                          typename RemainderOutputType<T1, T2>::value_type,
+                          void>)
+        {
             fnT fn = nullptr;
             return fn;
         }
@@ -484,8 +487,10 @@ struct RemainderInplaceStridedFactory
 {
     fnT get()
     {
-        if constexpr (std::is_same_v<typename RemainderOutputType<T1, T2>::value_type,
-                                     void>) {
+        if constexpr (std::is_same_v<
+                          typename RemainderOutputType<T1, T2>::value_type,
+                          void>)
+        {
             fnT fn = nullptr;
             return fn;
         }

dpctl/tests/elementwise/test_pow.py

@@ -21,6 +21,7 @@
 
 import dpctl
 import dpctl.tensor as dpt
+from dpctl.tensor._type_utils import _can_cast
 from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported
 
 from .utils import _all_dtypes, _compare_dtypes, _usm_types
@@ -152,3 +153,60 @@ def test_pow_python_scalar(arr_dt):
         assert isinstance(R, dpt.usm_ndarray)
         R = dpt.pow(sc, X)
         assert isinstance(R, dpt.usm_ndarray)
+
+
+@pytest.mark.parametrize("dtype", _all_dtypes[1:])
+def test_pow_inplace_python_scalar(dtype):
+    q = get_queue_or_skip()
+    skip_if_dtype_not_supported(dtype, q)
+    X = dpt.ones((10, 10), dtype=dtype, sycl_queue=q)
+    dt_kind = X.dtype.kind
+    if dt_kind in "ui":
+        X **= int(1)
+    elif dt_kind == "f":
+        X **= float(1)
+    elif dt_kind == "c":
+        X **= complex(1)
+
+
+@pytest.mark.parametrize("op1_dtype", _all_dtypes[1:])
+@pytest.mark.parametrize("op2_dtype", _all_dtypes[1:])
+def test_pow_inplace_dtype_matrix(op1_dtype, op2_dtype):
+    q = get_queue_or_skip()
+    skip_if_dtype_not_supported(op1_dtype, q)
+    skip_if_dtype_not_supported(op2_dtype, q)
+
+    sz = 127
+    ar1 = dpt.ones(sz, dtype=op1_dtype)
+    ar2 = dpt.ones_like(ar1, dtype=op2_dtype)
+
+    dev = q.sycl_device
+    _fp16 = dev.has_aspect_fp16
+    _fp64 = dev.has_aspect_fp64
+    if _can_cast(ar2.dtype, ar1.dtype, _fp16, _fp64):
+        ar1 **= ar2
+        assert (
+            dpt.asnumpy(ar1) == np.full(ar1.shape, 1, dtype=ar1.dtype)
+        ).all()
+
+        ar3 = dpt.ones(sz, dtype=op1_dtype)
+        ar4 = dpt.ones(2 * sz, dtype=op2_dtype)
+
+        ar3[::-1] *= ar4[::2]
+        assert (
+            dpt.asnumpy(ar3) == np.full(ar3.shape, 1, dtype=ar3.dtype)
+        ).all()
+
+    else:
+        with pytest.raises(TypeError):
+            ar1 **= ar2
+
+
+def test_pow_inplace_basic():
+    get_queue_or_skip()
+
+    x = dpt.arange(10, dtype="i4")
+    expected = dpt.square(x)
+    x **= 2
+
+    assert dpt.all(x == expected)

dpctl/tests/elementwise/test_remainder.py

@@ -21,6 +21,7 @@
 
 import dpctl
 import dpctl.tensor as dpt
+from dpctl.tensor._type_utils import _can_cast
 from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported
 
 from .utils import _compare_dtypes, _no_complex_dtypes, _usm_types
@@ -206,3 +207,54 @@ def test_remainder_python_scalar(arr_dt):
         assert isinstance(R, dpt.usm_ndarray)
         R = dpt.remainder(sc, X)
         assert isinstance(R, dpt.usm_ndarray)
+
+
+@pytest.mark.parametrize("dtype", _no_complex_dtypes[1:])
+def test_remainder_inplace_python_scalar(dtype):
+    q = get_queue_or_skip()
+    skip_if_dtype_not_supported(dtype, q)
+    X = dpt.ones((10, 10), dtype=dtype, sycl_queue=q)
+    dt_kind = X.dtype.kind
+    if dt_kind in "ui":
+        X %= int(1)
+    elif dt_kind == "f":
+        X %= float(1)
+
+
+@pytest.mark.parametrize("op1_dtype", _no_complex_dtypes[1:])
+@pytest.mark.parametrize("op2_dtype", _no_complex_dtypes[1:])
+def test_remainder_inplace_dtype_matrix(op1_dtype, op2_dtype):
+    q = get_queue_or_skip()
+    skip_if_dtype_not_supported(op1_dtype, q)
+    skip_if_dtype_not_supported(op2_dtype, q)
+
+    sz = 127
+    ar1 = dpt.ones(sz, dtype=op1_dtype)
+    ar2 = dpt.ones_like(ar1, dtype=op2_dtype)
+
+    dev = q.sycl_device
+    _fp16 = dev.has_aspect_fp16
+    _fp64 = dev.has_aspect_fp64
+    if _can_cast(ar2.dtype, ar1.dtype, _fp16, _fp64):
+        ar1 %= ar2
+        assert dpt.all(ar1 == dpt.zeros(ar1.shape, dtype=ar1.dtype))
+
+        ar3 = dpt.ones(sz, dtype=op1_dtype)
+        ar4 = dpt.ones(2 * sz, dtype=op2_dtype)
+
+        ar3[::-1] %= ar4[::2]
+        assert dpt.all(ar3 == dpt.zeros(ar3.shape, dtype=ar3.dtype))
+
+    else:
+        with pytest.raises(TypeError):
+            ar1 %= ar2
+
+
+def test_remainder_inplace_basic():
+    get_queue_or_skip()
+
+    x = dpt.arange(10, dtype="i4")
+    expected = x & 1
+    x %= 2
+
+    assert dpt.all(x == expected)