Leverage dpctl.tensor.copy() implementation

.github/workflows/conda-package.yml

@@ -15,6 +15,7 @@ env:
       test_arraycreation.py
       test_dot.py
       test_dparray.py
+      test_copy.py
       test_fft.py
       test_linalg.py
       test_logic.py

dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

@@ -1326,8 +1326,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap)
             {                                                                  \
                 constexpr size_t lws = 64;                                     \
                 constexpr unsigned int vec_sz = 8;                             \
-                constexpr sycl::access::address_space global_space =           \
-                    sycl::access::address_space::global_space;                 \
                                                                                \
                 auto gws_range = sycl::range<1>(                               \
                     ((result_size + lws * vec_sz - 1) / (lws * vec_sz)) *      \

dpnp/backend/kernels/dpnp_krnl_logic.cpp

@@ -521,8 +521,6 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef,
         else {                                                                 \
             constexpr size_t lws = 64;                                         \
             constexpr unsigned int vec_sz = 8;                                 \
-            constexpr sycl::access::address_space global_space =               \
-                sycl::access::address_space::global_space;                     \
                                                                                \
             auto gws_range = sycl::range<1>(                                   \
                 ((result_size + lws * vec_sz - 1) / (lws * vec_sz)) * lws);    \

dpnp/dpnp_algo/dpnp_algo_mathematical.pxi

@@ -352,7 +352,7 @@ cpdef tuple dpnp_modf(utils.dpnp_descriptor x1):
 
 
 cpdef utils.dpnp_descriptor dpnp_nancumprod(utils.dpnp_descriptor x1):
-    cur_x1 = dpnp_copy(x1).get_pyobj()
+    cur_x1 = x1.get_pyobj().copy()
 
     cur_x1_flatiter = cur_x1.flat
 
@@ -365,7 +365,7 @@ cpdef utils.dpnp_descriptor dpnp_nancumprod(utils.dpnp_descriptor x1):
 
 
 cpdef utils.dpnp_descriptor dpnp_nancumsum(utils.dpnp_descriptor x1):
-    cur_x1 = dpnp_copy(x1).get_pyobj()
+    cur_x1 = x1.get_pyobj().copy()
 
     cur_x1_flatiter = cur_x1.flat
 

dpnp/dpnp_array.py

@@ -168,7 +168,15 @@ def __complex__(self):
         return self._array_obj.__complex__()
 
     # '__contains__',
-    # '__copy__',
+
+    def __copy__(self):
+        """
+        Used if :func:`copy.copy` is called on an array. Returns a copy of the array.
+
+        Equivalent to ``a.copy(order="K")``.
+        """
+        return self.copy(order="K")
+
     # '__deepcopy__',
     # '__delattr__',
     # '__delitem__',
@@ -651,7 +659,47 @@ def conjugate(self):
         else:
             return dpnp.conjugate(self)
 
-    # 'copy',
+    def copy(self, order="C"):
+        """
+        Return a copy of the array.
+
+        Returns
+        -------
+        out : dpnp.ndarray
+            A copy of the array.
+
+        See also
+        --------
+        :obj:`dpnp.copy` : Similar function with different default behavior
+        :obj:`dpnp.copyto` : Copies values from one array to another.
+
+        Notes
+        -----
+        This function is the preferred method for creating an array copy. The
+        function :func:`dpnp.copy` is similar, but it defaults to using order 'K'.
+
+        Examples
+        --------
+        >>> import dpnp as np
+        >>> x = np.array([[1, 2, 3], [4, 5, 6]], order='F')
+        >>> y = x.copy()
+        >>> x.fill(0)
+
+        >>> x
+        array([[0, 0, 0],
+               [0, 0, 0]])
+
+        >>> y
+        array([[1, 2, 3],
+               [4, 5, 6]])
+
+        >>> y.flags['C_CONTIGUOUS']
+        True
+
+        """
+
+        return dpnp.copy(self, order=order)
+
     # 'ctypes',
     # 'cumprod',
 

dpnp/dpnp_container.py

@@ -43,6 +43,7 @@
 __all__ = [
     "arange",
     "asarray",
+    "copy",
     "empty",
     "eye",
     "full",
@@ -135,6 +136,15 @@ def asarray(
     return dpnp_array(array_obj.shape, buffer=array_obj, order=order)
 
 
+def copy(x1, /, *, order="K"):
+    """Creates `dpnp_array` as a copy of given instance of input array."""
+    if order is None:
+        order = "K"
+
+    array_obj = dpt.copy(dpnp.get_usm_ndarray(x1), order=order)
+    return dpnp_array(array_obj.shape, buffer=array_obj, order="K")
+
+
 def empty(
     shape,
     *,
@@ -264,9 +274,7 @@ def meshgrid(*xi, indexing="xy"):
     """Creates list of `dpnp_array` coordinate matrices from vectors."""
     if len(xi) == 0:
         return []
-    arrays = tuple(
-        x.get_array() if isinstance(x, dpnp_array) else x for x in xi
-    )
+    arrays = tuple(dpnp.get_usm_ndarray(x) for x in xi)
     arrays_obj = dpt.meshgrid(*arrays, indexing=indexing)
     return [
         dpnp_array._create_from_usm_ndarray(array_obj)
@@ -304,17 +312,13 @@ def ones(
 
 def tril(x1, /, *, k=0):
     """Creates `dpnp_array` as lower triangular part of an input array."""
-    array_obj = dpt.tril(
-        x1.get_array() if isinstance(x1, dpnp_array) else x1, k
-    )
+    array_obj = dpt.tril(dpnp.get_usm_ndarray(x1), k)
     return dpnp_array(array_obj.shape, buffer=array_obj, order="K")
 
 
 def triu(x1, /, *, k=0):
     """Creates `dpnp_array` as upper triangular part of an input array."""
-    array_obj = dpt.triu(
-        x1.get_array() if isinstance(x1, dpnp_array) else x1, k
-    )
+    array_obj = dpt.triu(dpnp.get_usm_ndarray(x1), k)
     return dpnp_array(array_obj.shape, buffer=array_obj, order="K")