Moved functions in strided_iters into dpctl::tensor::strides namespace

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/constructors.hpp

@@ -609,11 +609,11 @@ sycl::event tri_impl(sycl::queue exec_q,
                 py::ssize_t outer_gid = idx[0] / inner_range;
                 py::ssize_t inner_gid = idx[0] - inner_range * outer_gid;
 
-                py::ssize_t src_inner_offset, dst_inner_offset;
-                bool to_copy;
+                py::ssize_t src_inner_offset = 0, dst_inner_offset = 0;
+                bool to_copy(true);
 
                 {
-                    // py::ssize_t inner_gid = idx.get_id(0);
+                    using dpctl::tensor::strides::CIndexer_array;
                     CIndexer_array<d2, py::ssize_t> indexer_i(
                         {shape_and_strides[nd_2], shape_and_strides[nd_1]});
                     indexer_i.set(inner_gid);
@@ -634,7 +634,7 @@ sycl::event tri_impl(sycl::queue exec_q,
                 py::ssize_t src_offset = 0;
                 py::ssize_t dst_offset = 0;
                 {
-                    // py::ssize_t outer_gid = idx.get_id(1);
+                    using dpctl::tensor::strides::CIndexer_vector;
                     CIndexer_vector<py::ssize_t> outer(nd - d2);
                     outer.get_displacement(
                         outer_gid, shape_and_strides, shape_and_strides + src_s,

dpctl/tensor/libtensor/include/utils/offset_utils.hpp

@@ -141,6 +141,8 @@ struct StridedIndexer
 
     size_t operator()(size_t gid) const
     {
+        using dpctl::tensor::strides::CIndexer_vector;
+
         CIndexer_vector _ind(nd);
         py::ssize_t relative_offset(0);
         _ind.get_displacement<const py::ssize_t *, const py::ssize_t *>(
@@ -233,6 +235,8 @@ struct TwoOffsets_StridedIndexer
 
     TwoOffsets<py::ssize_t> operator()(py::ssize_t gid) const
     {
+        using dpctl::tensor::strides::CIndexer_vector;
+
         CIndexer_vector _ind(nd);
         py::ssize_t relative_first_offset(0);
         py::ssize_t relative_second_offset(0);
@@ -285,7 +289,7 @@ struct NthStrideOffset
     }
 
 private:
-    CIndexer_vector<py::ssize_t> _ind;
+    dpctl::tensor::strides::CIndexer_vector<py::ssize_t> _ind;
 
     int nd;
     py::ssize_t const *offsets;
@@ -302,11 +306,13 @@ template <int nd> struct FixedDimStridedIndexer
     }
     size_t operator()(size_t gid) const
     {
+        dpctl::tensor::strides::CIndexer_array<nd, py::ssize_t> local_indexer(
+            std::move(_ind));
+        local_indexer.set(gid);
+        auto mi = local_indexer.get();
+
         py::ssize_t relative_offset = 0;
-        CIndexer_array<nd, py::ssize_t> local_indxr(std::move(_ind));
 
-        local_indxr.set(gid);
-        auto mi = local_indxr.get();
 #pragma unroll
         for (int i = 0; i < nd; ++i) {
             relative_offset += mi[i] * strides[i];
@@ -315,7 +321,7 @@ template <int nd> struct FixedDimStridedIndexer
     }
 
 private:
-    CIndexer_array<nd, py::ssize_t> _ind;
+    dpctl::tensor::strides::CIndexer_array<nd, py::ssize_t> _ind;
 
     const std::array<py::ssize_t, nd> strides;
     py::ssize_t starting_offset;
@@ -336,26 +342,29 @@ template <int nd> struct TwoOffsets_FixedDimStridedIndexer
 
     TwoOffsets<py::ssize_t> operator()(size_t gid) const
     {
-        py::ssize_t relative_offset1 = 0;
-        py::ssize_t relative_offset2 = 0;
+        dpctl::tensor::strides::CIndexer_array<nd, py::ssize_t> local_indexer(
+            std::move(_ind));
+        local_indexer.set(gid);
+        auto mi = local_indexer.get();
 
-        CIndexer_array<nd, py::ssize_t> local_indxr(std::move(_ind));
-        local_indxr.set(gid);
-        auto mi = local_indxr.get();
+        py::ssize_t relative_offset1 = 0;
 #pragma unroll
         for (int i = 0; i < nd; ++i) {
             relative_offset1 += mi[i] * strides1[i];
         }
+
+        py::ssize_t relative_offset2 = 0;
 #pragma unroll
         for (int i = 0; i < nd; ++i) {
             relative_offset2 += mi[i] * strides2[i];
         }
+
         return TwoOffsets<py::ssize_t>(starting_offset1 + relative_offset1,
                                        starting_offset2 + relative_offset2);
     }
 
 private:
-    CIndexer_array<nd, py::ssize_t> _ind;
+    dpctl::tensor::strides::CIndexer_array<nd, py::ssize_t> _ind;
 
     const std::array<py::ssize_t, nd> strides1;
     const std::array<py::ssize_t, nd> strides2;

dpctl/tensor/libtensor/include/utils/strided_iters.hpp

@@ -32,6 +32,13 @@
 #include <tuple>
 #include <vector>
 
+namespace dpctl
+{
+namespace tensor
+{
+namespace strides
+{
+
 /* An N-dimensional array can be stored in a single
  * contiguous chunk of memory by contiguously laying
  * array elements in lexicographinc order of their
@@ -708,3 +715,7 @@ contract_iter3(vecT shape, vecT strides1, vecT strides2, vecT strides3)
     return std::make_tuple(out_shape, out_strides1, disp1, out_strides2, disp2,
                            out_strides3, disp3);
 }
+
+} // namespace strides
+} // namespace tensor
+} // namespace dpctl

dpctl/tensor/libtensor/source/simplify_iteration_space.cpp

@@ -24,6 +24,7 @@
 
 #include "simplify_iteration_space.hpp"
 #include "dpctl4pybind11.hpp"
+#include "utils/strided_iters.hpp"
 #include <pybind11/pybind11.h>
 #include <vector>
 
@@ -49,6 +50,7 @@ void simplify_iteration_space_1(int &nd,
                                 std::vector<py::ssize_t> &simplified_strides,
                                 py::ssize_t &offset)
 {
+    using dpctl::tensor::strides::simplify_iteration_stride;
     if (nd > 1) {
         // Simplify iteration space to reduce dimensionality
         // and improve access pattern
@@ -135,6 +137,7 @@ void simplify_iteration_space(int &nd,
                               py::ssize_t &src_offset,
                               py::ssize_t &dst_offset)
 {
+    using dpctl::tensor::strides::simplify_iteration_two_strides;
     if (nd > 1) {
         // Simplify iteration space to reduce dimensionality
         // and improve access pattern
@@ -280,6 +283,7 @@ void simplify_iteration_space_3(
     py::ssize_t &src2_offset,
     py::ssize_t &dst_offset)
 {
+    using dpctl::tensor::strides::simplify_iteration_three_strides;
     if (nd > 1) {
         // Simplify iteration space to reduce dimensionality
         // and improve access pattern

dpctl/tensor/libtensor/source/simplify_iteration_space.hpp

@@ -23,7 +23,6 @@
 //===--------------------------------------------------------------------===//
 
 #pragma once
-#include "utils/strided_iters.hpp"
 #include <pybind11/pybind11.h>
 #include <vector>
 

dpctl/tensor/libtensor/source/tensor_py.cpp

@@ -128,6 +128,7 @@ PYBIND11_MODULE(_tensor_impl, m)
     init_dispatch_tables();
     init_dispatch_vectors();
 
+    using dpctl::tensor::strides::contract_iter;
     m.def(
         "_contract_iter", &contract_iter<py::ssize_t, py::value_error>,
         "Simplifies iteration of array of given shape & stride. Returns "
@@ -143,6 +144,7 @@ PYBIND11_MODULE(_tensor_impl, m)
           py::arg("src"), py::arg("dst"), py::arg("sycl_queue"),
           py::arg("depends") = py::list());
 
+    using dpctl::tensor::strides::contract_iter2;
     m.def(
         "_contract_iter2", &contract_iter2<py::ssize_t, py::value_error>,
         "Simplifies iteration over elements of pair of arrays of given shape "
@@ -152,6 +154,7 @@ PYBIND11_MODULE(_tensor_impl, m)
         "as the original "
         "iterator, possibly in a different order.");
 
+    using dpctl::tensor::strides::contract_iter3;
     m.def(
         "_contract_iter3", &contract_iter3<py::ssize_t, py::value_error>,
         "Simplifies iteration over elements of 3-tuple of arrays of given "