Merge pull request #1582 from IntelPython/avoid-copy-in-functor-constructors

No copying of indexers in functor constructors

Author: oleksandr-pavlyk

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

@@ -107,8 +107,8 @@ inclusive_scan_base_step(sycl::queue &exec_q,
                          outputT *output,
                          const size_t s0,
                          const size_t s1,
-                         IndexerT indexer,
-                         TransformerT transformer,
+                         const IndexerT &indexer,
+                         const TransformerT &transformer,
                          size_t &n_groups,
                          const std::vector<sycl::event> &depends = {})
 {
@@ -234,8 +234,8 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
                                 outputT *output,
                                 const size_t s0,
                                 const size_t s1,
-                                IndexerT indexer,
-                                TransformerT transformer,
+                                const IndexerT &indexer,
+                                const TransformerT &transformer,
                                 std::vector<sycl::event> &host_tasks,
                                 const std::vector<sycl::event> &depends = {})
 {

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

@@ -56,9 +56,9 @@ struct MaskedExtractStridedFunctor
                                 char *dst_data_p,
                                 size_t orthog_iter_size,
                                 size_t masked_iter_size,
-                                OrthogIndexerT orthog_src_dst_indexer_,
-                                MaskedSrcIndexerT masked_src_indexer_,
-                                MaskedDstIndexerT masked_dst_indexer_)
+                                const OrthogIndexerT &orthog_src_dst_indexer_,
+                                const MaskedSrcIndexerT &masked_src_indexer_,
+                                const MaskedDstIndexerT &masked_dst_indexer_)
         : src_cp(src_data_p), cumsum_cp(cumsum_data_p), dst_cp(dst_data_p),
           orthog_nelems(orthog_iter_size), masked_nelems(masked_iter_size),
           orthog_src_dst_indexer(orthog_src_dst_indexer_),
@@ -106,13 +106,14 @@ struct MaskedExtractStridedFunctor
     char *dst_cp = nullptr;
     size_t orthog_nelems = 0;
     size_t masked_nelems = 0;
-    OrthogIndexerT
-        orthog_src_dst_indexer; // has nd, shape, src_strides, dst_strides for
-                                // dimensions that ARE NOT masked
-    MaskedSrcIndexerT masked_src_indexer; // has nd, shape, src_strides for
-                                          // dimensions that ARE     masked
-    MaskedDstIndexerT
-        masked_dst_indexer; // has 1, dst_strides for dimensions that ARE masked
+    // has nd, shape, src_strides, dst_strides for
+    // dimensions that ARE NOT masked
+    const OrthogIndexerT orthog_src_dst_indexer;
+    // has nd, shape, src_strides for
+    // dimensions that ARE masked
+    const MaskedSrcIndexerT masked_src_indexer;
+    // has 1, dst_strides for dimensions that ARE masked
+    const MaskedDstIndexerT masked_dst_indexer;
 };
 
 template <typename OrthogIndexerT,
@@ -127,9 +128,9 @@ struct MaskedPlaceStridedFunctor
                               const char *rhs_data_p,
                               size_t orthog_iter_size,
                               size_t masked_iter_size,
-                              OrthogIndexerT orthog_dst_rhs_indexer_,
-                              MaskedDstIndexerT masked_dst_indexer_,
-                              MaskedRhsIndexerT masked_rhs_indexer_)
+                              const OrthogIndexerT &orthog_dst_rhs_indexer_,
+                              const MaskedDstIndexerT &masked_dst_indexer_,
+                              const MaskedRhsIndexerT &masked_rhs_indexer_)
         : dst_cp(dst_data_p), cumsum_cp(cumsum_data_p), rhs_cp(rhs_data_p),
           orthog_nelems(orthog_iter_size), masked_nelems(masked_iter_size),
           orthog_dst_rhs_indexer(orthog_dst_rhs_indexer_),
@@ -177,13 +178,14 @@ struct MaskedPlaceStridedFunctor
     const char *rhs_cp = nullptr;
     size_t orthog_nelems = 0;
     size_t masked_nelems = 0;
-    OrthogIndexerT
-        orthog_dst_rhs_indexer; // has nd, shape, dst_strides, rhs_strides for
-                                // dimensions that ARE NOT masked
-    MaskedDstIndexerT masked_dst_indexer; // has nd, shape, dst_strides for
-                                          // dimensions that ARE     masked
-    MaskedRhsIndexerT
-        masked_rhs_indexer; // has 1, rhs_strides for dimensions that ARE masked
+    // has nd, shape, dst_strides, rhs_strides for
+    // dimensions that ARE NOT masked
+    const OrthogIndexerT orthog_dst_rhs_indexer;
+    // has nd, shape, dst_strides for
+    // dimensions that ARE masked
+    const MaskedDstIndexerT masked_dst_indexer;
+    // has 1, rhs_strides for dimensions that ARE masked
+    const MaskedRhsIndexerT masked_rhs_indexer;
 };
 
 // ======= Masked extraction ================================
@@ -226,12 +228,12 @@ sycl::event masked_extract_all_slices_strided_impl(
     //  using StridedIndexer;
     //  using TwoZeroOffsets_Indexer;
 
-    TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
+    constexpr TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
 
     /* StridedIndexer(int _nd, ssize_t _offset, ssize_t const
      * *_packed_shape_strides) */
-    StridedIndexer masked_src_indexer(nd, 0, packed_src_shape_strides);
-    Strided1DIndexer masked_dst_indexer(0, dst_size, dst_stride);
+    const StridedIndexer masked_src_indexer(nd, 0, packed_src_shape_strides);
+    const Strided1DIndexer masked_dst_indexer(0, dst_size, dst_stride);
 
     sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
@@ -283,17 +285,16 @@ sycl::event masked_extract_some_slices_strided_impl(
     const char *cumsum_p,
     char *dst_p,
     int orthog_nd,
-    const ssize_t
-        *packed_ortho_src_dst_shape_strides, // [ortho_shape, ortho_src_strides,
-                                             // ortho_dst_strides], length
-                                             // 3*ortho_nd
+    // [ortho_shape, ortho_src_strides, // ortho_dst_strides],
+    // length 3*ortho_nd
+    const ssize_t *packed_ortho_src_dst_shape_strides,
     ssize_t ortho_src_offset,
     ssize_t ortho_dst_offset,
     int masked_nd,
-    const ssize_t *packed_masked_src_shape_strides, // [masked_src_shape,
-                                                    // masked_src_strides],
-                                                    // length 2*masked_nd
-    ssize_t masked_dst_size,                        // mask_dst is 1D
+    // [masked_src_shape, masked_src_strides],
+    // length 2*masked_nd, mask_dst is 1D
+    const ssize_t *packed_masked_src_shape_strides,
+    ssize_t masked_dst_size,
     ssize_t masked_dst_stride,
     const std::vector<sycl::event> &depends = {})
 {
@@ -302,13 +303,14 @@ sycl::event masked_extract_some_slices_strided_impl(
     //  using StridedIndexer;
     //  using TwoOffsets_StridedIndexer;
 
-    TwoOffsets_StridedIndexer orthog_src_dst_indexer{
+    const TwoOffsets_StridedIndexer orthog_src_dst_indexer{
         orthog_nd, ortho_src_offset, ortho_dst_offset,
         packed_ortho_src_dst_shape_strides};
 
-    StridedIndexer masked_src_indexer{masked_nd, 0,
-                                      packed_masked_src_shape_strides};
-    Strided1DIndexer masked_dst_indexer{0, masked_dst_size, masked_dst_stride};
+    const StridedIndexer masked_src_indexer{masked_nd, 0,
+                                            packed_masked_src_shape_strides};
+    const Strided1DIndexer masked_dst_indexer{0, masked_dst_size,
+                                              masked_dst_stride};
 
     sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
@@ -403,12 +405,12 @@ sycl::event masked_place_all_slices_strided_impl(
     ssize_t rhs_stride,
     const std::vector<sycl::event> &depends = {})
 {
-    TwoZeroOffsets_Indexer orthog_dst_rhs_indexer{};
+    constexpr TwoZeroOffsets_Indexer orthog_dst_rhs_indexer{};
 
     /* StridedIndexer(int _nd, ssize_t _offset, ssize_t const
      * *_packed_shape_strides) */
-    StridedIndexer masked_dst_indexer(nd, 0, packed_dst_shape_strides);
-    Strided1DCyclicIndexer masked_rhs_indexer(0, rhs_size, rhs_stride);
+    const StridedIndexer masked_dst_indexer(nd, 0, packed_dst_shape_strides);
+    const Strided1DCyclicIndexer masked_rhs_indexer(0, rhs_size, rhs_stride);
 
     sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
@@ -460,30 +462,29 @@ sycl::event masked_place_some_slices_strided_impl(
     const char *cumsum_p,
     const char *rhs_p,
     int orthog_nd,
-    const ssize_t
-        *packed_ortho_dst_rhs_shape_strides, // [ortho_shape, ortho_dst_strides,
-                                             // ortho_rhs_strides], length
-                                             // 3*ortho_nd
+    // [ortho_shape, ortho_dst_strides, ortho_rhs_strides],
+    // length 3*ortho_nd
+    const ssize_t *packed_ortho_dst_rhs_shape_strides,
     ssize_t ortho_dst_offset,
     ssize_t ortho_rhs_offset,
     int masked_nd,
-    const ssize_t *packed_masked_dst_shape_strides, // [masked_dst_shape,
-                                                    // masked_dst_strides],
-                                                    // length 2*masked_nd
-    ssize_t masked_rhs_size,                        // mask_dst is 1D
+    // [masked_dst_shape, masked_dst_strides],
+    // length 2*masked_nd, mask_dst is 1D
+    const ssize_t *packed_masked_dst_shape_strides,
+    ssize_t masked_rhs_size,
     ssize_t masked_rhs_stride,
     const std::vector<sycl::event> &depends = {})
 {
-    TwoOffsets_StridedIndexer orthog_dst_rhs_indexer{
+    const TwoOffsets_StridedIndexer orthog_dst_rhs_indexer{
         orthog_nd, ortho_dst_offset, ortho_rhs_offset,
         packed_ortho_dst_rhs_shape_strides};
 
     /* StridedIndexer(int _nd, ssize_t _offset, ssize_t const
      * *_packed_shape_strides) */
-    StridedIndexer masked_dst_indexer{masked_nd, 0,
-                                      packed_masked_dst_shape_strides};
-    Strided1DCyclicIndexer masked_rhs_indexer{0, masked_rhs_size,
-                                              masked_rhs_stride};
+    const StridedIndexer masked_dst_indexer{masked_nd, 0,
+                                            packed_masked_dst_shape_strides};
+    const Strided1DCyclicIndexer masked_rhs_indexer{0, masked_rhs_size,
+                                                    masked_rhs_stride};
 
     sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);

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

@@ -235,14 +235,14 @@ template <typename T, typename IndexerT> class ClipStridedFunctor
     const T *min_p = nullptr;
     const T *max_p = nullptr;
     T *dst_p = nullptr;
-    IndexerT indexer;
+    const IndexerT indexer;
 
 public:
     ClipStridedFunctor(const T *x_p_,
                        const T *min_p_,
                        const T *max_p_,
                        T *dst_p_,
-                       IndexerT indexer_)
+                       const IndexerT &indexer_)
         : x_p(x_p_), min_p(min_p_), max_p(max_p_), dst_p(dst_p_),
           indexer(indexer_)
     {
@@ -298,7 +298,7 @@ sycl::event clip_strided_impl(sycl::queue &q,
     sycl::event clip_ev = q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
 
-        FourOffsets_StridedIndexer indexer{
+        const FourOffsets_StridedIndexer indexer{
             nd, x_offset, min_offset, max_offset, dst_offset, shape_strides};
 
         cgh.parallel_for<clip_strided_kernel<T, FourOffsets_StridedIndexer>>(

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

@@ -78,10 +78,10 @@ class GenericCopyFunctor
 private:
     const srcT *src_ = nullptr;
     dstT *dst_ = nullptr;
-    IndexerT indexer_;
+    const IndexerT indexer_;
 
 public:
-    GenericCopyFunctor(const srcT *src_p, dstT *dst_p, IndexerT indexer)
+    GenericCopyFunctor(const srcT *src_p, dstT *dst_p, const IndexerT &indexer)
         : src_(src_p), dst_(dst_p), indexer_(indexer)
     {
     }
@@ -169,8 +169,8 @@ copy_and_cast_generic_impl(sycl::queue &q,
         cgh.depends_on(depends);
         cgh.depends_on(additional_depends);
 
-        TwoOffsets_StridedIndexer indexer{nd, src_offset, dst_offset,
-                                          shape_and_strides};
+        const TwoOffsets_StridedIndexer indexer{nd, src_offset, dst_offset,
+                                                shape_and_strides};
         const srcTy *src_tp = reinterpret_cast<const srcTy *>(src_p);
         dstTy *dst_tp = reinterpret_cast<dstTy *>(dst_p);
 
@@ -460,8 +460,8 @@ copy_and_cast_nd_specialized_impl(sycl::queue &q,
 
     sycl::event copy_and_cast_ev = q.submit([&](sycl::handler &cgh) {
         using IndexerT = TwoOffsets_FixedDimStridedIndexer<nd>;
-        IndexerT indexer{shape, src_strides, dst_strides, src_offset,
-                         dst_offset};
+        const IndexerT indexer{shape, src_strides, dst_strides, src_offset,
+                               dst_offset};
         const srcTy *src_tp = reinterpret_cast<const srcTy *>(src_p);
         dstTy *dst_tp = reinterpret_cast<dstTy *>(dst_p);
 
@@ -515,12 +515,12 @@ class GenericCopyFromHostFunctor
 private:
     const AccessorT src_acc_;
     dstTy *dst_ = nullptr;
-    IndexerT indexer_;
+    const IndexerT indexer_;
 
 public:
     GenericCopyFromHostFunctor(const AccessorT &src_acc,
                                dstTy *dst_p,
-                               IndexerT indexer)
+                               const IndexerT &indexer)
         : src_acc_(src_acc), dst_(dst_p), indexer_(indexer)
     {
     }
@@ -618,7 +618,7 @@ void copy_and_cast_from_host_impl(
 
         sycl::accessor npy_acc(npy_buf, cgh, sycl::read_only);
 
-        TwoOffsets_StridedIndexer indexer{
+        const TwoOffsets_StridedIndexer indexer{
             nd, src_offset - src_min_nelem_offset, dst_offset,
             const_cast<const ssize_t *>(shape_and_strides)};
 
@@ -666,14 +666,14 @@ class GenericCopyForReshapeFunctor
 private:
     const Ty *src_p = nullptr;
     Ty *dst_p = nullptr;
-    SrcIndexerT src_indexer_;
-    DstIndexerT dst_indexer_;
+    const SrcIndexerT src_indexer_;
+    const DstIndexerT dst_indexer_;
 
 public:
     GenericCopyForReshapeFunctor(const char *src_ptr,
                                  char *dst_ptr,
-                                 SrcIndexerT src_indexer,
-                                 DstIndexerT dst_indexer)
+                                 const SrcIndexerT &src_indexer,
+                                 const DstIndexerT &dst_indexer)
         : src_p(reinterpret_cast<const Ty *>(src_ptr)),
           dst_p(reinterpret_cast<Ty *>(dst_ptr)), src_indexer_(src_indexer),
           dst_indexer_(dst_indexer)
@@ -747,8 +747,8 @@ copy_for_reshape_generic_impl(sycl::queue &q,
         const ssize_t *dst_shape_and_strides = const_cast<const ssize_t *>(
             packed_shapes_and_strides + (2 * src_nd));
 
-        StridedIndexer src_indexer{src_nd, 0, src_shape_and_strides};
-        StridedIndexer dst_indexer{dst_nd, 0, dst_shape_and_strides};
+        const StridedIndexer src_indexer{src_nd, 0, src_shape_and_strides};
+        const StridedIndexer dst_indexer{dst_nd, 0, dst_shape_and_strides};
 
         using KernelName =
             copy_for_reshape_generic_kernel<Ty, StridedIndexer, StridedIndexer>;
@@ -864,14 +864,14 @@ class StridedCopyForRollFunctor
 private:
     const Ty *src_p = nullptr;
     Ty *dst_p = nullptr;
-    SrcIndexerT src_indexer_;
-    DstIndexerT dst_indexer_;
+    const SrcIndexerT src_indexer_;
+    const DstIndexerT dst_indexer_;
 
 public:
     StridedCopyForRollFunctor(const Ty *src_ptr,
                               Ty *dst_ptr,
-                              SrcIndexerT src_indexer,
-                              DstIndexerT dst_indexer)
+                              const SrcIndexerT &src_indexer,
+                              const DstIndexerT &dst_indexer)
         : src_p(src_ptr), dst_p(dst_ptr), src_indexer_(src_indexer),
           dst_indexer_(dst_indexer)
     {
@@ -946,14 +946,15 @@ sycl::event copy_for_roll_strided_impl(sycl::queue &q,
         //   USM array of size 3 * nd
         //   [ common_shape; src_strides; dst_strides ]
 
-        StridedIndexer src_indexer{nd, src_offset, packed_shapes_and_strides};
-        LeftRolled1DTransformer left_roll_transformer{shift, nelems};
+        const StridedIndexer src_indexer{nd, src_offset,
+                                         packed_shapes_and_strides};
+        const LeftRolled1DTransformer left_roll_transformer{shift, nelems};
 
         using CompositeIndexerT =
             CompositionIndexer<StridedIndexer, LeftRolled1DTransformer>;
 
-        CompositeIndexerT rolled_src_indexer(src_indexer,
-                                             left_roll_transformer);
+        const CompositeIndexerT rolled_src_indexer(src_indexer,
+                                                   left_roll_transformer);
 
         UnpackedStridedIndexer dst_indexer{nd, dst_offset,
                                            packed_shapes_and_strides,
@@ -1024,12 +1025,12 @@ sycl::event copy_for_roll_contig_impl(sycl::queue &q,
     sycl::event copy_for_roll_ev = q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
 
-        NoOpIndexer src_indexer{};
-        LeftRolled1DTransformer roller{shift, nelems};
+        constexpr NoOpIndexer src_indexer{};
+        const LeftRolled1DTransformer roller{shift, nelems};
 
-        CompositionIndexer<NoOpIndexer, LeftRolled1DTransformer>
+        const CompositionIndexer<NoOpIndexer, LeftRolled1DTransformer>
             left_rolled_src_indexer{src_indexer, roller};
-        NoOpIndexer dst_indexer{};
+        constexpr NoOpIndexer dst_indexer{};
 
         using KernelName = copy_for_roll_contig_kernel<Ty>;
 
@@ -1119,11 +1120,11 @@ sycl::event copy_for_roll_ndshift_strided_impl(
         const ssize_t *shifts_ptr =
             packed_shapes_and_strides_and_shifts + 3 * nd;
 
-        RolledNDIndexer src_indexer{nd, shape_ptr, src_strides_ptr, shifts_ptr,
-                                    src_offset};
+        const RolledNDIndexer src_indexer{nd, shape_ptr, src_strides_ptr,
+                                          shifts_ptr, src_offset};
 
-        UnpackedStridedIndexer dst_indexer{nd, dst_offset, shape_ptr,
-                                           dst_strides_ptr};
+        const UnpackedStridedIndexer dst_indexer{nd, dst_offset, shape_ptr,
+                                                 dst_strides_ptr};
 
         using KernelName = copy_for_roll_strided_kernel<Ty, RolledNDIndexer,
                                                         UnpackedStridedIndexer>;