From 2ad5382f2977d4eaff5c9f8afb4022aa855159b9 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 15 Nov 2024 07:45:28 -0800 Subject: [PATCH 1/2] Fix build break Moved common constexpr variables out of branches. Replaced `static constexpr` with `constexpr`. Since these are defined in procedure scope, `static` is not required. Introduced typed temporary variables, so that type deduction for `sycl::min`, `sycl::max`, `sycl::clamp` can work and removed explicit use of their template parameter. Added explicit static_cast on value of `projected` variable computed as IndT type. --- .../include/utils/indexing_utils.hpp | 81 +++++++++---------- 1 file changed, 39 insertions(+), 42 deletions(-) diff --git a/dpctl/tensor/libtensor/include/utils/indexing_utils.hpp b/dpctl/tensor/libtensor/include/utils/indexing_utils.hpp index 80eb3df186..979c367a3c 100644 --- a/dpctl/tensor/libtensor/include/utils/indexing_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/indexing_utils.hpp @@ -49,41 +49,40 @@ template struct WrapIndex ssize_t operator()(ssize_t max_item, IndT ind) const { ssize_t projected; - max_item = sycl::max(max_item, 1); + constexpr ssize_t unit(1); + max_item = sycl::max(max_item, unit); + + constexpr std::uintmax_t ind_max = std::numeric_limits::max(); + constexpr std::uintmax_t ssize_max = + std::numeric_limits::max(); if constexpr (std::is_signed_v) { - static constexpr std::uintmax_t ind_max = - std::numeric_limits::max(); - static constexpr std::uintmax_t ssize_max = - std::numeric_limits::max(); - static constexpr std::intmax_t ind_min = - std::numeric_limits::min(); - static constexpr std::intmax_t ssize_min = + constexpr std::intmax_t ind_min = std::numeric_limits::min(); + constexpr std::intmax_t ssize_min = std::numeric_limits::min(); if constexpr (ind_max <= ssize_max && ind_min >= ssize_min) { - projected = sycl::clamp(static_cast(ind), - -max_item, max_item - 1); + const ssize_t ind_ = static_cast(ind); + const ssize_t lb = -max_item; + const ssize_t ub = max_item - 1; + projected = sycl::clamp(ind_, lb, ub); } else { - projected = sycl::clamp(ind, static_cast(-max_item), - static_cast(max_item - 1)); + const IndT lb = static_cast(-max_item); + const IndT ub = static_cast(max_item - 1); + projected = static_cast(sycl::clamp(ind, lb, ub)); } return (projected < 0) ? projected + max_item : projected; } else { - static constexpr std::uintmax_t ind_max = - std::numeric_limits::max(); - static constexpr std::uintmax_t ssize_max = - std::numeric_limits::max(); - if constexpr (ind_max <= ssize_max) { - projected = - sycl::min(static_cast(ind), max_item - 1); + const ssize_t ind_ = static_cast(ind); + const ssize_t ub = max_item - 1; + projected = sycl::min(ind_, ub); } else { - projected = - sycl::min(ind, static_cast(max_item - 1)); + const IndT ub = static_cast(max_item - 1); + projected = static_cast(sycl::min(ind, ub)); } return projected; } @@ -95,40 +94,38 @@ template struct ClipIndex ssize_t operator()(ssize_t max_item, IndT ind) const { ssize_t projected; - max_item = sycl::max(max_item, 1); + constexpr ssize_t unit(1); + max_item = sycl::max(max_item, unit); + constexpr std::uintmax_t ind_max = std::numeric_limits::max(); + constexpr std::uintmax_t ssize_max = + std::numeric_limits::max(); if constexpr (std::is_signed_v) { - static constexpr std::uintmax_t ind_max = - std::numeric_limits::max(); - static constexpr std::uintmax_t ssize_max = - std::numeric_limits::max(); - static constexpr std::intmax_t ind_min = - std::numeric_limits::min(); - static constexpr std::intmax_t ssize_min = + constexpr std::intmax_t ind_min = std::numeric_limits::min(); + constexpr std::intmax_t ssize_min = std::numeric_limits::min(); if constexpr (ind_max <= ssize_max && ind_min >= ssize_min) { - projected = sycl::clamp(static_cast(ind), - ssize_t(0), max_item - 1); + const ssize_t ind_ = static_cast(ind); + constexpr ssize_t lb(0); + const ssize_t ub = max_item - 1; + projected = sycl::clamp(ind_, lb, ub); } else { - projected = sycl::clamp(ind, IndT(0), - static_cast(max_item - 1)); + constexpr IndT lb(0); + const IndT ub = static_cast(max_item - 1); + projected = static_cast(sycl::clamp(ind, lb, ub)); } } else { - static constexpr std::uintmax_t ind_max = - std::numeric_limits::max(); - static constexpr std::uintmax_t ssize_max = - std::numeric_limits::max(); - if constexpr (ind_max <= ssize_max) { - projected = - sycl::min(static_cast(ind), max_item - 1); + const ssize_t ind_ = static_cast(ind); + const ssize_t ub = max_item - 1; + projected = sycl::min(ind_, ub); } else { - projected = - sycl::min(ind, static_cast(max_item - 1)); + const IndT ub = static_cast(max_item - 1); + projected = static_cast(sycl::min(ind, ub)); } } return projected; From 0493fdde830794ab2af4f60fdd10185872093e34 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 15 Nov 2024 07:49:47 -0800 Subject: [PATCH 2/2] Made `ProjectorT proj` constexpr This is possible because ProjectorT is literal type (no state and default constructor). --- .../libtensor/include/kernels/integer_advanced_indexing.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index e2a51a98e7..462323f001 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -93,7 +93,7 @@ class TakeFunctor ssize_t src_offset = orthog_offsets.get_first_offset(); ssize_t dst_offset = orthog_offsets.get_second_offset(); - const ProjectorT proj{}; + constexpr ProjectorT proj{}; for (int axis_idx = 0; axis_idx < k_; ++axis_idx) { indT *ind_data = reinterpret_cast(ind_[axis_idx]); @@ -239,7 +239,7 @@ class PutFunctor ssize_t dst_offset = orthog_offsets.get_first_offset(); ssize_t val_offset = orthog_offsets.get_second_offset(); - const ProjectorT proj{}; + constexpr ProjectorT proj{}; for (int axis_idx = 0; axis_idx < k_; ++axis_idx) { indT *ind_data = reinterpret_cast(ind_[axis_idx]);