Replace dpctl.tensor with dpnp.tensor in dpnp.tensor

Author: vlad-perevezentsev

dpnp/tensor/_device.py

@@ -40,7 +40,7 @@ class Device:
 
     This is a wrapper around :class:`dpctl.SyclQueue` with custom
     formatting. The class does not have public constructor,
-    but a class method :meth:`dpctl.tensor.Device.create_device` to construct
+    but a class method :meth:`dpnp.tensor.Device.create_device` to construct
     it from `device` keyword argument in Array-API functions.
 
     Instance can be queried for ``sycl_queue``, ``sycl_context``,

dpnp/tensor/libtensor/include/kernels/clip.hpp

@@ -29,7 +29,7 @@
 //===---------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines kernels for dpctl.tensor.clip.
+/// This file defines kernels for dpnp.tensor.clip.
 //===---------------------------------------------------------------------===//
 
 #pragma once
@@ -73,7 +73,8 @@ T clip(const T &x, const T &min, const T &max)
         return min_complex(max_complex(x, min), max);
     }
     else if constexpr (std::is_floating_point_v<T> ||
-                       std::is_same_v<T, sycl::half>) {
+                       std::is_same_v<T, sycl::half>)
+    {
         auto tmp = (std::isnan(x) || x > min) ? x : min;
         return (std::isnan(tmp) || tmp < max) ? tmp : max;
     }

dpnp/tensor/libtensor/include/kernels/dpctl_tensor_types.hpp

@@ -27,7 +27,7 @@
 //*****************************************************************************
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_impl extensions
+/// This file defines functions of dpnp.tensor._tensor_impl extensions
 //===--------------------------------------------------------------------===//
 
 #pragma once

dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp

@@ -29,7 +29,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// This file defines functions of dpnp.tensor._tensor_sorting_impl
 /// extension.
 //===----------------------------------------------------------------------===//
 
@@ -340,7 +340,8 @@ sycl::event
                 sycl::min(seg_start + elems_per_segment, n);
             if (is_ascending) {
                 for (std::size_t val_id = seg_start + lid; val_id < seg_end;
-                     val_id += wg_size) {
+                     val_id += wg_size)
+                {
                     // get the bucket for the bit-ordered input value,
                     // applying the offset and mask for radix bits
                     const auto val =
@@ -355,7 +356,8 @@ sycl::event
             }
             else {
                 for (std::size_t val_id = seg_start + lid; val_id < seg_end;
-                     val_id += wg_size) {
+                     val_id += wg_size)
+                {
                     // get the bucket for the bit-ordered input value,
                     // applying the offset and mask for radix bits
                     const auto val =
@@ -691,7 +693,8 @@ void copy_func_for_radix_sort(const std::size_t n_segments,
     // find offsets for the same values within a segment and fill the resulting
     // buffer
     for (std::size_t val_id = seg_start + lid; val_id < seg_end;
-         val_id += sg_size) {
+         val_id += sg_size)
+    {
         output_ptr[val_id] = std::move(input_ptr[val_id]);
     }
 
@@ -822,7 +825,8 @@ sycl::event
             // resulting buffer
             if (is_ascending) {
                 for (std::size_t val_id = seg_start + lid; val_id < seg_end;
-                     val_id += sg_size) {
+                     val_id += sg_size)
+                {
                     ValueT in_val = std::move(b_input_ptr[val_id]);
 
                     // get the bucket for the bit-ordered input value, applying
@@ -851,7 +855,8 @@ sycl::event
             }
             else {
                 for (std::size_t val_id = seg_start + lid; val_id < seg_end;
-                     val_id += sg_size) {
+                     val_id += sg_size)
+                {
                     ValueT in_val = std::move(b_input_ptr[val_id]);
 
                     // get the bucket for the bit-ordered input value, applying
@@ -1164,10 +1169,7 @@ struct subgroup_radix_sort
             return sycl::local_accessor<KeyT>(buf_size, cgh);
         }
 
-        std::size_t get_iter_stride() const
-        {
-            return std::size_t{0};
-        }
+        std::size_t get_iter_stride() const { return std::size_t{0}; }
     };
 
     template <typename KeyT>
@@ -1185,10 +1187,7 @@ struct subgroup_radix_sort
         {
             return sycl::accessor(buf, cgh, sycl::read_write, sycl::no_init);
         }
-        std::size_t get_iter_stride() const
-        {
-            return iter_stride;
-        }
+        std::size_t get_iter_stride() const { return iter_stride; }
     };
 
     static_assert(wg_size <= 1024);
@@ -1377,7 +1376,8 @@ struct subgroup_radix_sort
                                 if (is_ascending) {
 #pragma unroll
                                     for (std::uint16_t i = 0; i < block_size;
-                                         ++i) {
+                                         ++i)
+                                    {
                                         const std::uint16_t id =
                                             wi * block_size + i;
                                         static constexpr std::uint16_t
@@ -1413,7 +1413,8 @@ struct subgroup_radix_sort
                                 else {
 #pragma unroll
                                     for (std::uint16_t i = 0; i < block_size;
-                                         ++i) {
+                                         ++i)
+                                    {
                                         const std::uint16_t id =
                                             wi * block_size + i;
                                         static constexpr std::uint16_t
@@ -1740,7 +1741,8 @@ sycl::event parallel_radix_sort_impl(sycl::queue &exec_q,
                                                depends);
 
         for (std::uint32_t radix_iter = 1; radix_iter < radix_iters;
-             ++radix_iter) {
+             ++radix_iter)
+        {
             if (radix_iter % 2 == 0) {
                 sort_ev = parallel_radix_sort_iteration_step<
                     radix_bits,
@@ -1798,10 +1800,7 @@ struct IndexedProj
     {
     }
 
-    auto operator()(IndexT i) const
-    {
-        return value_projector(ptr[i]);
-    }
+    auto operator()(IndexT i) const { return value_projector(ptr[i]); }
 
 private:
     const ValueT *ptr;

dpnp/tensor/libtensor/include/kernels/sorting/sort_impl_fn_ptr_t.hpp

@@ -29,7 +29,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// This file defines functions of dpnp.tensor._tensor_sorting_impl
 /// extension.
 //===----------------------------------------------------------------------===//
 

dpnp/tensor/libtensor/include/kernels/where.hpp

@@ -29,7 +29,7 @@
 //===---------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines kernels for dpctl.tensor.where.
+/// This file defines kernels for dpnp.tensor.where.
 //===---------------------------------------------------------------------===//
 
 #pragma once

dpnp/tensor/libtensor/include/utils/rich_comparisons.hpp

@@ -27,7 +27,7 @@
 //*****************************************************************************
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// This file defines functions of dpnp.tensor._tensor_sorting_impl
 /// extension.
 //===---------------------------------------------------------------------===//
 
@@ -112,9 +112,9 @@ struct ExtendedComplexFPGreater
 };
 
 template <typename T>
-inline constexpr bool is_fp_v = (std::is_same_v<T, sycl::half> ||
-                                 std::is_same_v<T, float> ||
-                                 std::is_same_v<T, double>);
+inline constexpr bool is_fp_v =
+    (std::is_same_v<T, sycl::half> || std::is_same_v<T, float> ||
+     std::is_same_v<T, double>);
 
 } // namespace detail
 

dpnp/tensor/libtensor/source/accumulators.cpp

@@ -29,7 +29,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_impl extensions
+/// This file defines functions of dpnp.tensor._tensor_impl extensions
 //===----------------------------------------------------------------------===//
 
 #include <cstddef>

dpnp/tensor/libtensor/source/accumulators.hpp

@@ -29,7 +29,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_impl extensions
+/// This file defines functions of dpnp.tensor._tensor_impl extensions
 //===----------------------------------------------------------------------===//
 
 #pragma once

dpnp/tensor/libtensor/source/accumulators/accumulate_over_axis.hpp

@@ -29,7 +29,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines functions of dpctl.tensor._tensor_accumulation_impl
+/// This file defines functions of dpnp.tensor._tensor_accumulation_impl
 //  extensions
 //===----------------------------------------------------------------------===//