Renamed advance_indexing.*pp into integer_advanced_indexing.*pp

Streamlined call operator implementation for projection classes.
Added missing includes.

Author: oleksandr-pavlyk

dpctl/tensor/CMakeLists.txt

@@ -31,7 +31,7 @@ pybind11_add_module(${python_module_name} MODULE
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/copy_numpy_ndarray_into_usm_ndarray.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/copy_for_reshape.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/linear_sequences.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/advanced_indexing.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/integer_advanced_indexing.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/eye_ctor.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/full_ctor.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/triul_ctor.cpp

…or/include/kernels/advanced_indexing.hpp …de/kernels/integer_advanced_indexing.hppdpctl/tensor/libtensor/include/kernels/advanced_indexing.hpp renamed to dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp dpctl/tensor/libtensor/include/kernels/advanced_indexing.hpp renamed to dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp

@@ -26,6 +26,7 @@
 #include "utils/strided_iters.hpp"
 #include "utils/type_utils.hpp"
 #include <CL/sycl.hpp>
+#include <algorithm>
 #include <complex>
 #include <cstdint>
 #include <pybind11/pybind11.h>
@@ -52,10 +53,9 @@ class ClipIndex
 
     void operator()(py::ssize_t max_item, py::ssize_t &ind) const
     {
-        max_item = (max_item > 0) ? max_item : 1;
-        ind = (ind < 0)           ? (ind <= -max_item) ? (0) : (ind + max_item)
-                        : (ind >= max_item) ? (max_item - 1)
-                                  : ind;
+        max_item = std::max<py::ssize_t>(max_item, 1);
+        ind = std::clamp<py::ssize_t>(ind, -max_item, max_item - 1);
+        ind = (ind < 0) ? ind + max_item : ind;
         return;
     }
 };
@@ -67,10 +67,8 @@ class WrapIndex
 
     void operator()(py::ssize_t max_item, py::ssize_t &ind) const
     {
-        max_item = (max_item > 0) ? max_item : 1;
-        ind = (ind < 0)           ? max_item - (-ind % max_item)
-              : (ind >= max_item) ? ind % max_item
-                                  : ind;
+        max_item = std::max<py::ssize_t>(max_item, 1);
+        ind = ind % max_item;
         return;
     }
 };
@@ -136,9 +134,9 @@ template <typename ProjectorT, typename T, typename indT> class TakeFunctor
             dst_orthog_idx);
 
         ProjectorT proj{};
-        py::ssize_t ind_arr_idx(0);
         CIndexer_vector<py::ssize_t> ind_indxr(ind_nd_);
         for (int axis_idx = 0; axis_idx < k_; ++axis_idx) {
+            py::ssize_t ind_arr_idx(0);
             ind_indxr.get_displacement<const py::ssize_t *>(
                 static_cast<py::ssize_t>(i_along), ind_shape_and_strides_,
                 ind_shape_and_strides_ + ((axis_idx + 1) * ind_nd_),

…r/libtensor/source/advanced_indexing.cpp …sor/source/integer_advanced_indexing.cppdpctl/tensor/libtensor/source/advanced_indexing.cpp renamed to dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp dpctl/tensor/libtensor/source/advanced_indexing.cpp renamed to dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp

@@ -27,16 +27,19 @@
 #include <algorithm>
 #include <complex>
 #include <cstdint>
+#include <iostream>
 #include <pybind11/complex.h>
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 #include <utility>
 
 #include "dpctl4pybind11.hpp"
-#include "kernels/advanced_indexing.hpp"
+#include "kernels/integer_advanced_indexing.hpp"
 #include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
 
+#include "integer_advanced_indexing.hpp"
+
 #define INDEXING_MODES 2
 #define CLIP_MODE 0
 #define WRAP_MODE 1
@@ -85,8 +88,8 @@ std::vector<sycl::event> _populate_packed_shapes_strides_for_indexing(
     int arr_nd)
 {
 
-    int orthog_sh_elems = ((inp_nd - k) > 1) ? (inp_nd - k) : 1;
-    int along_sh_elems = (ind_nd > 1) ? ind_nd : 1;
+    int orthog_sh_elems = std::max<int>(inp_nd - k, 1);
+    int along_sh_elems = std::max<int>(ind_nd, 1);
 
     using usm_host_allocatorT =
         sycl::usm_allocator<py::ssize_t, sycl::usm::alloc::host>;
@@ -284,7 +287,7 @@ usm_ndarray_take(dpctl::tensor::usm_ndarray src,
                  int axis_start,
                  uint8_t mode,
                  sycl::queue exec_q,
-                 const std::vector<sycl::event> &depends = {})
+                 const std::vector<sycl::event> &depends)
 {
     int k = ind.size();
 
@@ -328,7 +331,7 @@ usm_ndarray_take(dpctl::tensor::usm_ndarray src,
     const py::ssize_t *src_shape = src.get_shape_raw();
     const py::ssize_t *dst_shape = dst.get_shape_raw();
 
-    int orthog_nd = ((src_nd - k) > 0) ? src_nd - k : 1;
+    int orthog_nd = std::max<int>(src_nd - k, 1);
 
     bool orthog_shapes_equal(true);
     size_t orthog_nelems(1);
@@ -412,7 +415,7 @@ usm_ndarray_take(dpctl::tensor::usm_ndarray src,
         }
     }
 
-    auto ind_sh_elems = (ind_nd > 0) ? ind_nd : 1;
+    int ind_sh_elems = std::max<int>(ind_nd, 1);
 
     std::vector<char *> ind_ptrs;
     ind_ptrs.reserve(k);
@@ -633,12 +636,15 @@ usm_ndarray_take(dpctl::tensor::usm_ndarray src,
                                  std::to_string(ind_type_id));
     }
 
+    std::cout << "Submitting take" << std::endl;
     sycl::event take_generic_ev =
         fn(exec_q, orthog_nelems, ind_nelems, orthog_nd, ind_nd, k,
            packed_shapes_strides, packed_axes_shapes_strides,
            packed_ind_shapes_strides, src_data, dst_data, packed_ind_ptrs,
            src_offset, dst_offset, packed_ind_offsets, all_deps);
 
+    std::cout << "Submitting take clean-up host task" << std::endl;
+
     // free packed temporaries
     auto ctx = exec_q.get_context();
     exec_q.submit([&](sycl::handler &cgh) {
@@ -666,7 +672,7 @@ usm_ndarray_put(dpctl::tensor::usm_ndarray dst,
                 int axis_start,
                 uint8_t mode,
                 sycl::queue exec_q,
-                const std::vector<sycl::event> &depends = {})
+                const std::vector<sycl::event> &depends)
 {
     int k = ind.size();
 

…r/libtensor/source/advanced_indexing.hpp …sor/source/integer_advanced_indexing.hppdpctl/tensor/libtensor/source/advanced_indexing.hpp renamed to dpctl/tensor/libtensor/source/integer_advanced_indexing.hpp dpctl/tensor/libtensor/source/advanced_indexing.hpp renamed to dpctl/tensor/libtensor/source/integer_advanced_indexing.hpp

@@ -33,13 +33,13 @@
 
 #include "dpctl4pybind11.hpp"
 
-#include "advanced_indexing.hpp"
 #include "copy_and_cast_usm_to_usm.hpp"
 #include "copy_for_reshape.hpp"
 #include "copy_numpy_ndarray_into_usm_ndarray.hpp"
 #include "device_support_queries.hpp"
 #include "eye_ctor.hpp"
 #include "full_ctor.hpp"
+#include "integer_advanced_indexing.hpp"
 #include "linear_sequences.hpp"
 #include "triul_ctor.hpp"
 #include "utils/strided_iters.hpp"