Resolve COPY_INSTEAD_OF_MOVE issues in libtensor

Author: oleksandr-pavlyk

dpctl/apis/include/dpctl4pybind11.hpp

@@ -30,6 +30,7 @@
 #include <complex>
 #include <memory>
 #include <pybind11/pybind11.h>
+#include <utility>
 #include <vector>
 
 namespace py = pybind11;
@@ -369,19 +370,19 @@ class dpctl_capi
         sycl::queue q_{};
         PySyclQueueObject *py_q_tmp =
             SyclQueue_Make(reinterpret_cast<DPCTLSyclQueueRef>(&q_));
-        py::object py_sycl_queue = py::reinterpret_steal<py::object>(
+        const py::object &py_sycl_queue = py::reinterpret_steal<py::object>(
             reinterpret_cast<PyObject *>(py_q_tmp));
 
         default_sycl_queue_ = std::shared_ptr<py::object>(
             new py::object(py_sycl_queue), Deleter{});
 
         py::module_ mod_memory = py::module_::import("dpctl.memory");
-        py::object py_as_usm_memory = mod_memory.attr("as_usm_memory");
+        const py::object &py_as_usm_memory = mod_memory.attr("as_usm_memory");
         as_usm_memory_ = std::shared_ptr<py::object>(
             new py::object{py_as_usm_memory}, Deleter{});
 
         auto mem_kl = mod_memory.attr("MemoryUSMHost");
-        py::object py_default_usm_memory =
+        const py::object &py_default_usm_memory =
             mem_kl(1, py::arg("queue") = py_sycl_queue);
         default_usm_memory_ = std::shared_ptr<py::object>(
             new py::object{py_default_usm_memory}, Deleter{});
@@ -390,7 +391,7 @@ class dpctl_capi
             py::module_::import("dpctl.tensor._usmarray");
         auto tensor_kl = mod_usmarray.attr("usm_ndarray");
 
-        py::object py_default_usm_ndarray =
+        const py::object &py_default_usm_ndarray =
             tensor_kl(py::tuple(), py::arg("dtype") = py::str("u1"),
                       py::arg("buffer") = py_default_usm_memory);
 
@@ -1032,7 +1033,7 @@ namespace utils
 {
 
 template <std::size_t num>
-sycl::event keep_args_alive(sycl::queue q,
+sycl::event keep_args_alive(sycl::queue &q,
                             const py::object (&py_objs)[num],
                             const std::vector<sycl::event> &depends = {})
 {
@@ -1043,7 +1044,7 @@ sycl::event keep_args_alive(sycl::queue q,
             shp_arr[i] = std::make_shared<py::handle>(py_objs[i]);
             shp_arr[i]->inc_ref();
         }
-        cgh.host_task([=]() {
+        cgh.host_task([shp_arr = std::move(shp_arr)]() {
             py::gil_scoped_acquire acquire;
 
             for (std::size_t i = 0; i < num; ++i) {

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

@@ -116,19 +116,20 @@ sycl::event inclusive_scan_rec(sycl::queue &exec_q,
 {
     size_t n_groups = ceiling_quotient(n_elems, n_wi * wg_size);
 
-    sycl::event inc_scan_phase1_ev = exec_q.submit([&](sycl::handler &cgh) {
-        cgh.depends_on(depends);
+    const sycl::event &inc_scan_phase1_ev =
+        exec_q.submit([&](sycl::handler &cgh) {
+            cgh.depends_on(depends);
 
-        using slmT = sycl::local_accessor<size_t, 1>;
+            using slmT = sycl::local_accessor<size_t, 1>;
 
-        auto lws = sycl::range<1>(wg_size);
-        auto gws = sycl::range<1>(n_groups * wg_size);
+            auto lws = sycl::range<1>(wg_size);
+            auto gws = sycl::range<1>(n_groups * wg_size);
 
-        slmT slm_iscan_tmp(lws, cgh);
+            slmT slm_iscan_tmp(lws, cgh);
 
         cgh.parallel_for<class inclusive_scan_rec_local_scan_krn<
             inputT, outputT, n_wi, IndexerT, decltype(transformer)>>(
-            sycl::nd_range<1>(gws, lws), [=](sycl::nd_item<1> it)
+            sycl::nd_range<1>(gws, lws), [=, slm_iscan_tmp = std::move(slm_iscan_tmp)](sycl::nd_item<1> it)
         {
             auto chunk_gid = it.get_global_id(0);
             auto lid = it.get_local_id(0);
@@ -172,7 +173,7 @@ sycl::event inclusive_scan_rec(sycl::queue &exec_q,
                 output[i + m_wi] = local_isum[m_wi];
             }
         });
-    });
+        });
 
     sycl::event out_event = inc_scan_phase1_ev;
     if (n_groups > 1) {
@@ -203,11 +204,11 @@ sycl::event inclusive_scan_rec(sycl::queue &exec_q,
 
         sycl::event e4 = exec_q.submit([&](sycl::handler &cgh) {
             cgh.depends_on(e3);
-            auto ctx = exec_q.get_context();
+            const auto &ctx = exec_q.get_context();
             cgh.host_task([ctx, temp]() { sycl::free(temp, ctx); });
         });
 
-        out_event = e4;
+        out_event = std::move(e4);
     }
 
     return out_event;
@@ -235,7 +236,7 @@ size_t accumulate_contig_impl(sycl::queue &q,
     NoOpIndexer flat_indexer{};
     transformerT non_zero_indicator{};
 
-    sycl::event comp_ev =
+    const sycl::event &comp_ev =
         inclusive_scan_rec<maskT, cumsumT, n_wi, decltype(flat_indexer),
                            decltype(non_zero_indicator)>(
             q, n_elems, wg_size, mask_data_ptr, cumsum_data_ptr, 0, 1,
@@ -321,7 +322,7 @@ size_t accumulate_strided_impl(sycl::queue &q,
     StridedIndexer strided_indexer{nd, 0, shape_strides};
     transformerT non_zero_indicator{};
 
-    sycl::event comp_ev =
+    const sycl::event &comp_ev =
         inclusive_scan_rec<maskT, cumsumT, n_wi, decltype(strided_indexer),
                            decltype(non_zero_indicator)>(
             q, n_elems, wg_size, mask_data_ptr, cumsum_data_ptr, 0, 1,

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

@@ -483,12 +483,12 @@ template <typename AccessorT,
 class GenericCopyFromHostFunctor
 {
 private:
-    AccessorT src_acc_;
+    const AccessorT src_acc_;
     dstTy *dst_ = nullptr;
     IndexerT indexer_;
 
 public:
-    GenericCopyFromHostFunctor(AccessorT src_acc,
+    GenericCopyFromHostFunctor(const AccessorT &src_acc,
                                dstTy *dst_p,
                                IndexerT indexer)
         : src_acc_(src_acc), dst_(dst_p), indexer_(indexer)

dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp

@@ -862,7 +862,7 @@ sycl::event binary_contig_matrix_contig_row_broadcast_impl(
 
     sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(comp_ev);
-        sycl::context ctx = exec_q.get_context();
+        const sycl::context &ctx = exec_q.get_context();
         cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
     });
     host_tasks.push_back(tmp_cleanup_ev);

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

@@ -28,6 +28,7 @@
 #include <cstddef>
 #include <cstdint>
 #include <type_traits>
+#include <utility>
 #include <vector>
 
 #include "pybind11/pybind11.h"
@@ -760,7 +761,8 @@ sycl::event sum_reduction_over_group_temps_strided_impl(
                 partially_reduced_tmp + reduction_groups * iter_nelems;
         }
 
-        sycl::event first_reduction_ev = exec_q.submit([&](sycl::handler &cgh) {
+        const sycl::event &first_reduction_ev = exec_q.submit([&](sycl::handler
+                                                                      &cgh) {
             cgh.depends_on(depends);
 
             using InputIndexerT = dpctl::tensor::offset_utils::StridedIndexer;
@@ -858,7 +860,7 @@ sycl::event sum_reduction_over_group_temps_strided_impl(
 
             remaining_reduction_nelems = reduction_groups_;
             std::swap(temp_arg, temp2_arg);
-            dependent_ev = partial_reduction_ev;
+            dependent_ev = std::move(partial_reduction_ev);
         }
 
         // final reduction to res
@@ -915,7 +917,7 @@ sycl::event sum_reduction_over_group_temps_strided_impl(
         sycl::event cleanup_host_task_event =
             exec_q.submit([&](sycl::handler &cgh) {
                 cgh.depends_on(final_reduction_ev);
-                sycl::context ctx = exec_q.get_context();
+                const sycl::context &ctx = exec_q.get_context();
 
                 cgh.host_task([ctx, partially_reduced_tmp] {
                     sycl::free(partially_reduced_tmp, ctx);

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

@@ -98,7 +98,7 @@ device_allocate_and_pack(sycl::queue q,
 
     usm_host_allocatorT usm_host_allocator(q);
     shT empty{0, usm_host_allocator};
-    shT packed_shape_strides = detail::concat(empty, vs...);
+    shT packed_shape_strides = detail::concat(std::move(empty), vs...);
 
     auto packed_shape_strides_owner =
         std::make_shared<shT>(std::move(packed_shape_strides));

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

@@ -541,7 +541,7 @@ int simplify_iteration_two_strides(const int nd,
 }
 
 template <typename T, class Error, typename vecT = std::vector<T>>
-std::tuple<vecT, vecT, T> contract_iter(vecT shape, vecT strides)
+std::tuple<vecT, vecT, T> contract_iter(const vecT &shape, const vecT &strides)
 {
     const size_t dim = shape.size();
     if (dim != strides.size()) {
@@ -560,7 +560,7 @@ std::tuple<vecT, vecT, T> contract_iter(vecT shape, vecT strides)
 
 template <typename T, class Error, typename vecT = std::vector<T>>
 std::tuple<vecT, vecT, T, vecT, T>
-contract_iter2(vecT shape, vecT strides1, vecT strides2)
+contract_iter2(const vecT &shape, const vecT &strides1, const vecT &strides2)
 {
     const size_t dim = shape.size();
     if (dim != strides1.size() || dim != strides2.size()) {
@@ -714,8 +714,10 @@ int simplify_iteration_three_strides(const int nd,
 }
 
 template <typename T, class Error, typename vecT = std::vector<T>>
-std::tuple<vecT, vecT, T, vecT, T, vecT, T>
-contract_iter3(vecT shape, vecT strides1, vecT strides2, vecT strides3)
+std::tuple<vecT, vecT, T, vecT, T, vecT, T> contract_iter3(const vecT &shape,
+                                                           const vecT &strides1,
+                                                           const vecT &strides2,
+                                                           const vecT &strides3)
 {
     const size_t dim = shape.size();
     if (dim != strides1.size() || dim != strides2.size() ||
@@ -899,11 +901,11 @@ int simplify_iteration_four_strides(const int nd,
 
 template <typename T, class Error, typename vecT = std::vector<T>>
 std::tuple<vecT, vecT, T, vecT, T, vecT, T, vecT, T>
-contract_iter4(vecT shape,
-               vecT strides1,
-               vecT strides2,
-               vecT strides3,
-               vecT strides4)
+contract_iter4(const vecT &shape,
+               const vecT &strides1,
+               const vecT &strides2,
+               const vecT &strides3,
+               const vecT &strides4)
 {
     const size_t dim = shape.size();
     if (dim != strides1.size() || dim != strides2.size() ||

dpctl/tensor/libtensor/source/boolean_advanced_indexing.cpp

@@ -264,7 +264,7 @@ py_extract(const dpctl::tensor::usm_ndarray &src,
         sycl::event cleanup_tmp_allocations_ev =
             exec_q.submit([&](sycl::handler &cgh) {
                 cgh.depends_on(extract_ev);
-                auto ctx = exec_q.get_context();
+                const auto &ctx = exec_q.get_context();
                 cgh.host_task([ctx, packed_src_shape_strides] {
                     sycl::free(packed_src_shape_strides, ctx);
                 });
@@ -366,7 +366,7 @@ py_extract(const dpctl::tensor::usm_ndarray &src,
         sycl::event cleanup_tmp_allocations_ev =
             exec_q.submit([&](sycl::handler &cgh) {
                 cgh.depends_on(extract_ev);
-                auto ctx = exec_q.get_context();
+                const auto &ctx = exec_q.get_context();
                 cgh.host_task([ctx, packed_shapes_strides] {
                     sycl::free(packed_shapes_strides, ctx);
                 });
@@ -693,7 +693,7 @@ py_place(const dpctl::tensor::usm_ndarray &dst,
         sycl::event cleanup_tmp_allocations_ev =
             exec_q.submit([&](sycl::handler &cgh) {
                 cgh.depends_on(place_ev);
-                auto ctx = exec_q.get_context();
+                const auto &ctx = exec_q.get_context();
                 cgh.host_task([ctx, packed_shapes_strides] {
                     sycl::free(packed_shapes_strides, ctx);
                 });
@@ -838,7 +838,7 @@ py_nonzero(const dpctl::tensor::usm_ndarray
 
     sycl::event temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(non_zero_indexes_ev);
-        auto ctx = exec_q.get_context();
+        const auto &ctx = exec_q.get_context();
         cgh.host_task([ctx, src_shape_device_ptr] {
             sycl::free(src_shape_device_ptr, ctx);
         });

dpctl/tensor/libtensor/source/boolean_reductions.hpp

@@ -292,7 +292,7 @@ py_boolean_reduction(const dpctl::tensor::usm_ndarray &src,
 
     sycl::event temp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(red_ev);
-        auto ctx = exec_q.get_context();
+        const auto &ctx = exec_q.get_context();
         cgh.host_task([ctx, packed_shapes_and_strides] {
             sycl::free(packed_shapes_and_strides, ctx);
         });

dpctl/tensor/libtensor/source/copy_and_cast_usm_to_usm.cpp

@@ -250,15 +250,15 @@ copy_usm_ndarray_into_usm_ndarray(const dpctl::tensor::usm_ndarray &src,
     if (shape_strides == nullptr) {
         throw std::runtime_error("Unable to allocate device memory");
     }
-    sycl::event copy_shape_ev = std::get<2>(ptr_size_event_tuple);
+    const sycl::event &copy_shape_ev = std::get<2>(ptr_size_event_tuple);
 
-    sycl::event copy_and_cast_generic_ev = copy_and_cast_fn(
+    const sycl::event &copy_and_cast_generic_ev = copy_and_cast_fn(
         exec_q, src_nelems, nd, shape_strides, src_data, src_offset, dst_data,
         dst_offset, depends, {copy_shape_ev});
 
     // async free of shape_strides temporary
-    auto ctx = exec_q.get_context();
-    auto temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
+    const auto &ctx = exec_q.get_context();
+    const auto &temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(copy_and_cast_generic_ev);
         cgh.host_task(
             [ctx, shape_strides]() { sycl::free(shape_strides, ctx); });

dpctl/tensor/libtensor/source/copy_for_reshape.cpp

@@ -158,7 +158,7 @@ copy_usm_ndarray_for_reshape(const dpctl::tensor::usm_ndarray &src,
 
     auto temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(copy_for_reshape_event);
-        auto ctx = exec_q.get_context();
+        const auto &ctx = exec_q.get_context();
         cgh.host_task(
             [shape_strides, ctx]() { sycl::free(shape_strides, ctx); });
     });

dpctl/tensor/libtensor/source/copy_for_roll.cpp

@@ -239,7 +239,7 @@ copy_usm_ndarray_for_roll_1d(const dpctl::tensor::usm_ndarray &src,
 
     auto temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(copy_for_roll_event);
-        auto ctx = exec_q.get_context();
+        const auto &ctx = exec_q.get_context();
         cgh.host_task(
             [shape_strides, ctx]() { sycl::free(shape_strides, ctx); });
     });
@@ -379,7 +379,7 @@ copy_usm_ndarray_for_roll_nd(const dpctl::tensor::usm_ndarray &src,
 
     auto temporaries_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(copy_for_roll_event);
-        auto ctx = exec_q.get_context();
+        const auto &ctx = exec_q.get_context();
         cgh.host_task([shape_strides_shifts, ctx]() {
             sycl::free(shape_strides_shifts, ctx);
         });

dpctl/tensor/libtensor/source/copy_numpy_ndarray_into_usm_ndarray.cpp

@@ -236,7 +236,7 @@ void copy_numpy_ndarray_into_usm_ndarray(
     if (shape_strides == nullptr) {
         throw std::runtime_error("Unable to allocate device memory");
     }
-    sycl::event copy_shape_ev = std::get<2>(ptr_size_event_tuple);
+    const sycl::event &copy_shape_ev = std::get<2>(ptr_size_event_tuple);
 
     // Get implementation function pointer
     auto copy_and_cast_from_host_blocking_fn =

dpctl/tensor/libtensor/source/elementwise_functions.hpp

@@ -227,7 +227,7 @@ py_unary_ufunc(const dpctl::tensor::usm_ndarray &src,
         q, host_tasks, simplified_shape, simplified_src_strides,
         simplified_dst_strides);
     py::ssize_t *shape_strides = std::get<0>(ptr_size_event_triple_);
-    sycl::event copy_shape_ev = std::get<2>(ptr_size_event_triple_);
+    const sycl::event &copy_shape_ev = std::get<2>(ptr_size_event_triple_);
 
     if (shape_strides == nullptr) {
         throw std::runtime_error("Device memory allocation failed");
@@ -533,7 +533,7 @@ std::pair<sycl::event, sycl::event> py_binary_ufunc(
         simplified_src2_strides, simplified_dst_strides);
 
     py::ssize_t *shape_strides = std::get<0>(ptr_sz_event_triple_);
-    sycl::event copy_shape_ev = std::get<2>(ptr_sz_event_triple_);
+    const sycl::event &copy_shape_ev = std::get<2>(ptr_sz_event_triple_);
 
     if (shape_strides == nullptr) {
         throw std::runtime_error("Unabled to allocate device memory");
@@ -799,7 +799,7 @@ py_binary_inplace_ufunc(const dpctl::tensor::usm_ndarray &lhs,
         simplified_lhs_strides);
 
     py::ssize_t *shape_strides = std::get<0>(ptr_sz_event_triple_);
-    sycl::event copy_shape_ev = std::get<2>(ptr_sz_event_triple_);
+    const sycl::event &copy_shape_ev = std::get<2>(ptr_sz_event_triple_);
 
     if (shape_strides == nullptr) {
         throw std::runtime_error("Unabled to allocate device memory");