diff --git a/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp index 4384a36e4c..84e72490cd 100644 --- a/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp @@ -116,15 +116,15 @@ struct MaskedExtractStridedFunctor const dataT *src = nullptr; const indT *cumsum = nullptr; dataT *dst = nullptr; - const std::size_t masked_nelems = 0; + std::size_t masked_nelems = 0; // has nd, shape, src_strides, dst_strides for // dimensions that ARE NOT masked - const OrthogIndexerT orthog_src_dst_indexer; + OrthogIndexerT orthog_src_dst_indexer; // has nd, shape, src_strides for // dimensions that ARE masked - const MaskedSrcIndexerT masked_src_indexer; + MaskedSrcIndexerT masked_src_indexer; // has 1, dst_strides for dimensions that ARE masked - const MaskedDstIndexerT masked_dst_indexer; + MaskedDstIndexerT masked_dst_indexer; LocalAccessorT lacc; }; @@ -198,15 +198,15 @@ struct MaskedPlaceStridedFunctor dataT *dst = nullptr; const indT *cumsum = nullptr; const dataT *rhs = nullptr; - const std::size_t masked_nelems = 0; + std::size_t masked_nelems = 0; // has nd, shape, dst_strides, rhs_strides for // dimensions that ARE NOT masked - const OrthogIndexerT orthog_dst_rhs_indexer; + OrthogIndexerT orthog_dst_rhs_indexer; // has nd, shape, dst_strides for // dimensions that ARE masked - const MaskedDstIndexerT masked_dst_indexer; + MaskedDstIndexerT masked_dst_indexer; // has 1, rhs_strides for dimensions that ARE masked - const MaskedRhsIndexerT masked_rhs_indexer; + MaskedRhsIndexerT masked_rhs_indexer; LocalAccessorT lacc; }; diff --git a/dpctl/tensor/libtensor/include/kernels/clip.hpp b/dpctl/tensor/libtensor/include/kernels/clip.hpp index 5cc7ebc533..7374774188 100644 --- a/dpctl/tensor/libtensor/include/kernels/clip.hpp +++ b/dpctl/tensor/libtensor/include/kernels/clip.hpp @@ -245,7 +245,7 @@ template class ClipStridedFunctor const T *min_p = nullptr; const T *max_p = nullptr; T *dst_p = nullptr; - const IndexerT indexer; + IndexerT indexer; public: ClipStridedFunctor(const T *x_p_, diff --git a/dpctl/tensor/libtensor/include/kernels/constructors.hpp b/dpctl/tensor/libtensor/include/kernels/constructors.hpp index 723218f661..5491726d8b 100644 --- a/dpctl/tensor/libtensor/include/kernels/constructors.hpp +++ b/dpctl/tensor/libtensor/include/kernels/constructors.hpp @@ -260,8 +260,8 @@ template class FullStridedFunctor { private: Ty *p = nullptr; - const Ty fill_v; - const IndexerT indexer; + Ty fill_v; + IndexerT indexer; public: FullStridedFunctor(Ty *p_, const Ty &fill_v_, const IndexerT &indexer_) diff --git a/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp b/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp index 3b977679e8..1db06a914a 100644 --- a/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp +++ b/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp @@ -87,7 +87,7 @@ class GenericCopyFunctor private: const srcT *src_ = nullptr; dstT *dst_ = nullptr; - const IndexerT indexer_; + IndexerT indexer_; public: GenericCopyFunctor(const srcT *src_p, dstT *dst_p, const IndexerT &indexer) @@ -219,7 +219,7 @@ template > template struct GetReadOnlyAccess { - Iter operator()(Iter it, sycl::handler &) { return it; } + Iter operator()(const Iter &it, sycl::handler &) { return it; } }; template struct GetReadOnlyAccess> { - auto operator()(sycl::buffer buf, + auto operator()(const sycl::buffer &buf, sycl::handler &cgh) { sycl::accessor acc(buf, cgh, sycl::read_only); @@ -295,7 +295,7 @@ template struct GetWriteDiscardAccess template struct GetWriteDiscardAccess> { - auto operator()(sycl::buffer buf, + auto operator()(sycl::buffer &buf, sycl::handler &cgh) { sycl::accessor acc(buf, cgh, sycl::write_only, sycl::no_init); @@ -305,13 +305,13 @@ struct GetWriteDiscardAccess> template struct GetReadWriteAccess { - Iter operator()(Iter it, sycl::handler &) { return it; } + Iter operator()(Iter &it, sycl::handler &) { return it; } }; template struct GetReadWriteAccess> { - auto operator()(sycl::buffer buf, + auto operator()(sycl::buffer &buf, sycl::handler &cgh) { sycl::accessor acc(buf, cgh, sycl::read_write); diff --git a/dpctl/tensor/libtensor/include/kernels/sorting/radix_sort.hpp b/dpctl/tensor/libtensor/include/kernels/sorting/radix_sort.hpp index e7528a26f4..8af2bc5923 100644 --- a/dpctl/tensor/libtensor/include/kernels/sorting/radix_sort.hpp +++ b/dpctl/tensor/libtensor/include/kernels/sorting/radix_sort.hpp @@ -526,11 +526,13 @@ struct peer_prefix_helper sycl::access::address_space::local_space>; using TempStorageT = sycl::local_accessor; - const sycl::sub_group sgroup; - const std::uint32_t lid; - const std::uint32_t item_mask; - const AtomicT atomic_peer_mask; +private: + sycl::sub_group sgroup; + std::uint32_t lid; + std::uint32_t item_mask; + AtomicT atomic_peer_mask; +public: peer_prefix_helper(sycl::nd_item<1> ndit, TempStorageT lacc) : sgroup(ndit.get_sub_group()), lid(ndit.get_local_linear_id()), item_mask(n_ls_bits_set(lid)), atomic_peer_mask(lacc[0]) @@ -572,9 +574,11 @@ struct peer_prefix_helper using ItemType = sycl::nd_item<1>; using SubGroupType = sycl::sub_group; - const SubGroupType sgroup; - const std::uint32_t sg_size; +private: + SubGroupType sgroup; + std::uint32_t sg_size; +public: peer_prefix_helper(sycl::nd_item<1> ndit, TempStorageT) : sgroup(ndit.get_sub_group()), sg_size(sgroup.get_local_range()[0]) { @@ -606,6 +610,10 @@ template struct peer_prefix_helper { private: + sycl::sub_group sgroup; + std::uint32_t lid; + sycl::ext::oneapi::sub_group_mask item_sg_mask; + sycl::ext::oneapi::sub_group_mask mask_builder(std::uint32_t mask, std::uint32_t sg_size) { @@ -616,10 +624,6 @@ struct peer_prefix_helper public: using TempStorageT = empty_storage; - const sycl::sub_group sgroup; - const std::uint32_t lid; - const sycl::ext::oneapi::sub_group_mask item_sg_mask; - peer_prefix_helper(sycl::nd_item<1> ndit, TempStorageT) : sgroup(ndit.get_sub_group()), lid(ndit.get_local_linear_id()), item_sg_mask( @@ -1136,7 +1140,7 @@ struct subgroup_radix_sort template class TempBuf { - const std::size_t buf_size; + std::size_t buf_size; public: TempBuf(std::size_t, std::size_t n) : buf_size(n) {} @@ -1151,7 +1155,7 @@ struct subgroup_radix_sort template class TempBuf { sycl::buffer buf; - const std::size_t iter_stride; + std::size_t iter_stride; public: TempBuf(std::size_t n_iters, std::size_t n) diff --git a/dpctl/tensor/libtensor/include/kernels/sorting/searchsorted.hpp b/dpctl/tensor/libtensor/include/kernels/sorting/searchsorted.hpp index 29a4c49c25..ddf61471c6 100644 --- a/dpctl/tensor/libtensor/include/kernels/sorting/searchsorted.hpp +++ b/dpctl/tensor/libtensor/include/kernels/sorting/searchsorted.hpp @@ -57,10 +57,10 @@ struct SearchSortedFunctor const argTy *hay_tp; const argTy *needles_tp; indTy *positions_tp; - const std::size_t hay_nelems; - const HayIndexerT hay_indexer; - const NeedlesIndexerT needles_indexer; - const PositionsIndexerT positions_indexer; + std::size_t hay_nelems; + HayIndexerT hay_indexer; + NeedlesIndexerT needles_indexer; + PositionsIndexerT positions_indexer; public: SearchSortedFunctor(const argTy *hay_, diff --git a/dpctl/tensor/libtensor/include/kernels/where.hpp b/dpctl/tensor/libtensor/include/kernels/where.hpp index 80e0262638..7c103e6ce4 100644 --- a/dpctl/tensor/libtensor/include/kernels/where.hpp +++ b/dpctl/tensor/libtensor/include/kernels/where.hpp @@ -232,7 +232,7 @@ class WhereStridedFunctor const T *x2_p = nullptr; T *dst_p = nullptr; const condT *cond_p = nullptr; - const IndexerT indexer; + IndexerT indexer; public: WhereStridedFunctor(const condT *cond_p_,