diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 8fb435b355..cec323a038 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -187,22 +187,10 @@ class dpctl_capi return api; } - py::object default_sycl_queue_pyobj() - { - return *default_sycl_queue_; - } - py::object default_usm_memory_pyobj() - { - return *default_usm_memory_; - } - py::object default_usm_ndarray_pyobj() - { - return *default_usm_ndarray_; - } - py::object as_usm_memory_pyobj() - { - return *as_usm_memory_; - } + py::object default_sycl_queue_pyobj() { return *default_sycl_queue_; } + py::object default_usm_memory_pyobj() { return *default_usm_memory_; } + py::object default_usm_ndarray_pyobj() { return *default_usm_ndarray_; } + py::object as_usm_memory_pyobj() { return *as_usm_memory_; } private: struct Deleter diff --git a/dpctl/tensor/libtensor/include/kernels/accumulators.hpp b/dpctl/tensor/libtensor/include/kernels/accumulators.hpp index 18e0e1bc8a..47465d0be4 100644 --- a/dpctl/tensor/libtensor/include/kernels/accumulators.hpp +++ b/dpctl/tensor/libtensor/include/kernels/accumulators.hpp @@ -47,10 +47,7 @@ namespace accumulators using namespace dpctl::tensor::offset_utils; -template T ceiling_quotient(T n, T m) -{ - return (n + m - 1) / m; -} +template T ceiling_quotient(T n, T m) { return (n + m - 1) / m; } template struct NonZeroIndicator { @@ -70,10 +67,7 @@ template struct NoOpTransformer { constexpr NoOpTransformer() {} - T operator()(const T &val) const - { - return val; - } + T operator()(const T &val) const { return val; } }; template struct CastTransformer @@ -108,20 +102,11 @@ template class stack_t } ~stack_t(){}; - T *get_src_ptr() const - { - return src_; - } + T *get_src_ptr() const { return src_; } - size_t get_size() const - { - return size_; - } + size_t get_size() const { return size_; } - T *get_local_scans_ptr() const - { - return local_scans_; - } + T *get_local_scans_ptr() const { return local_scans_; } }; template class stack_strided_t @@ -140,25 +125,13 @@ template class stack_strided_t } ~stack_strided_t(){}; - T *get_src_ptr() const - { - return src_; - } + T *get_src_ptr() const { return src_; } - size_t get_size() const - { - return size_; - } + size_t get_size() const { return size_; } - T *get_local_scans_ptr() const - { - return local_scans_; - } + T *get_local_scans_ptr() const { return local_scans_; } - size_t get_local_stride() const - { - return local_stride_; - } + size_t get_local_stride() const { return local_stride_; } }; } // end of anonymous namespace @@ -283,7 +256,8 @@ inclusive_scan_base_step(sycl::queue &exec_q, outputT wg_iscan_val; if constexpr (can_use_inclusive_scan_over_group::value) { + outputT>::value) + { wg_iscan_val = sycl::inclusive_scan_over_group( it.get_group(), local_iscan.back(), scan_op, identity); } @@ -305,7 +279,8 @@ inclusive_scan_base_step(sycl::queue &exec_q, } for (nwiT m_wi = 0; (m_wi < n_wi) && (i + m_wi < acc_nelems); - ++m_wi) { + ++m_wi) + { output[out_iter_offset + out_indexer(i + m_wi)] = local_iscan[m_wi]; } diff --git a/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp index 14cb49e06e..de55854768 100644 --- a/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp @@ -573,28 +573,27 @@ sycl::event non_zero_indexes_impl(sycl::queue &exec_q, sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { cgh.depends_on(depends); cgh.parallel_for>( - sycl::range<1>(iter_size), [=](sycl::id<1> idx) - { - auto i = idx[0]; - - auto cs_curr_val = cumsum_data[i] - 1; - auto cs_prev_val = (i > 0) ? cumsum_data[i - 1] : indT1(0); - bool cond = (cs_curr_val == cs_prev_val); - - ssize_t i_ = static_cast(i); - for (int dim = nd; --dim > 0;) { - auto sd = mask_shape[dim]; - ssize_t q = i_ / sd; - ssize_t r = (i_ - q * sd); + sycl::range<1>(iter_size), [=](sycl::id<1> idx) { + auto i = idx[0]; + + auto cs_curr_val = cumsum_data[i] - 1; + auto cs_prev_val = (i > 0) ? cumsum_data[i - 1] : indT1(0); + bool cond = (cs_curr_val == cs_prev_val); + + ssize_t i_ = static_cast(i); + for (int dim = nd; --dim > 0;) { + auto sd = mask_shape[dim]; + ssize_t q = i_ / sd; + ssize_t r = (i_ - q * sd); + if (cond) { + indexes_data[cs_curr_val + dim * nz_elems] = + static_cast(r); + } + i_ = q; + } if (cond) { - indexes_data[cs_curr_val + dim * nz_elems] = - static_cast(r); + indexes_data[cs_curr_val] = static_cast(i_); } - i_ = q; - } - if (cond) { - indexes_data[cs_curr_val] = static_cast(i_); - } }); }); diff --git a/dpctl/tensor/libtensor/include/kernels/clip.hpp b/dpctl/tensor/libtensor/include/kernels/clip.hpp index 72fcd2c35a..7b422c1281 100644 --- a/dpctl/tensor/libtensor/include/kernels/clip.hpp +++ b/dpctl/tensor/libtensor/include/kernels/clip.hpp @@ -60,7 +60,8 @@ template T clip(const T &x, const T &min, const T &max) return min_complex(max_complex(x, min), max); } else if constexpr (std::is_floating_point_v || - std::is_same_v) { + std::is_same_v) + { auto tmp = (std::isnan(x) || x > min) ? x : min; return (std::isnan(tmp) || tmp < max) ? tmp : max; } @@ -121,7 +122,8 @@ class ClipContigFunctor sg.get_group_id()[0] * max_sgSize); if (base + n_vecs * vec_sz * sgSize < nelems && - sgSize == max_sgSize) { + sgSize == max_sgSize) + { sycl::vec x_vec; sycl::vec min_vec; sycl::vec max_vec; @@ -155,7 +157,8 @@ class ClipContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems; - k += sgSize) { + k += sgSize) + { dst_p[k] = clip(x_p[k], min_p[k], max_p[k]); } } diff --git a/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp b/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp index 04969f740a..46f49eb01a 100644 --- a/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp +++ b/dpctl/tensor/libtensor/include/kernels/copy_and_cast.hpp @@ -248,7 +248,8 @@ class ContigCopyFunctor sg.get_group_id()[0] * max_sgSize); if (base + n_vecs * vec_sz * sgSize < nelems && - sgSize == max_sgSize) { + sgSize == max_sgSize) + { sycl::vec src_vec; sycl::vec dst_vec; @@ -273,7 +274,8 @@ class ContigCopyFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems; - k += sgSize) { + k += sgSize) + { dst_p[k] = fn(src_p[k]); } } @@ -803,10 +805,7 @@ template struct CompositionIndexer { CompositionIndexer(IndexerT f, TransformerT t) : f_(f), t_(t) {} - auto operator()(size_t gid) const - { - return f_(t_(gid)); - } + auto operator()(size_t gid) const { return f_(t_(gid)); } private: IndexerT f_; @@ -827,10 +826,7 @@ struct RolledNDIndexer { } - ssize_t operator()(size_t gid) const - { - return compute_offset(gid); - } + ssize_t operator()(size_t gid) const { return compute_offset(gid); } private: int nd_ = -1; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp index 5ae983803a..5de9024b6f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp @@ -141,7 +141,8 @@ template struct AbsContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -191,7 +192,8 @@ template struct AbsStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp index 5b8204f8b1..a9bf000a20 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp @@ -175,7 +175,8 @@ template struct AcosContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -222,7 +223,8 @@ template struct AcosStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp index 6ecfe02765..94c9c5e56e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp @@ -202,7 +202,8 @@ template struct AcoshContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -249,7 +250,8 @@ template struct AcoshStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp index bbd7045f04..e77068b5e1 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp @@ -94,7 +94,8 @@ template struct AddFunctor { auto tmp = in1 + in2; if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { @@ -223,7 +224,8 @@ template struct AddContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -272,7 +274,8 @@ template struct AddStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -398,10 +401,7 @@ template struct AddInplaceFunctor using supports_vec = std::negation< std::disjunction, tu_ns::is_complex>>; - void operator()(resT &res, const argT &in) - { - res += in; - } + void operator()(resT &res, const argT &in) { res += in; } template void operator()(sycl::vec &res, @@ -458,7 +458,8 @@ template struct AddInplaceContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -497,7 +498,8 @@ struct AddInplaceStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp index e34159d8c1..9622a7a207 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp @@ -118,7 +118,8 @@ template struct AngleContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -165,7 +166,8 @@ template struct AngleStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp index e98693cdfc..3a0d6efecf 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp @@ -195,7 +195,8 @@ template struct AsinContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -242,7 +243,8 @@ template struct AsinStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp index 9850ab6a3c..d64f3f0233 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp @@ -178,7 +178,8 @@ template struct AsinhContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -225,7 +226,8 @@ template struct AsinhStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp index 0e78c01903..64cd3c316f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp @@ -185,7 +185,8 @@ template struct AtanContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -232,7 +233,8 @@ template struct AtanStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp index c9130619a5..5002a18b19 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp @@ -179,7 +179,8 @@ template struct AtanhContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -226,7 +227,8 @@ template struct AtanhStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_left_shift.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_left_shift.hpp index 39075d91b8..18a87e5287 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_left_shift.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_left_shift.hpp @@ -278,10 +278,7 @@ template struct BitwiseLeftShiftInplaceFunctor using supports_sg_loadstore = typename std::true_type; using supports_vec = typename std::true_type; - void operator()(resT &res, const argT &in) const - { - impl(res, in); - } + void operator()(resT &res, const argT &in) const { impl(res, in); } template void operator()(sycl::vec &res, diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_right_shift.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_right_shift.hpp index 21410e0267..2fbee2e49d 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_right_shift.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/bitwise_right_shift.hpp @@ -280,10 +280,7 @@ template struct BitwiseRightShiftInplaceFunctor using supports_sg_loadstore = typename std::true_type; using supports_vec = typename std::true_type; - void operator()(resT &res, const argT &in) const - { - impl(res, in); - } + void operator()(resT &res, const argT &in) const { impl(res, in); } template void operator()(sycl::vec &res, diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cbrt.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cbrt.hpp index 37adb46723..083b7bee9d 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cbrt.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cbrt.hpp @@ -60,10 +60,7 @@ template struct CbrtFunctor // do both argTy and resTy support sugroup store/load operation using supports_sg_loadstore = typename std::true_type; - resT operator()(const argT &in) const - { - return sycl::cbrt(in); - } + resT operator()(const argT &in) const { return sycl::cbrt(in); } }; template struct CbrtContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -160,7 +158,8 @@ template struct CbrtStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp index 2fce16e84f..3cb90df632 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp @@ -131,7 +131,8 @@ template struct CeilContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -178,7 +179,8 @@ template struct CeilStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp index c6f3e5b617..b57a4eeb6c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp @@ -84,7 +84,8 @@ struct UnaryContigFunctor (ndit.get_group(0) * ndit.get_local_range(0) + sg.get_group_id()[0] * sgSize); if (base + n_vecs * vec_sz * sgSize < nelems_ && - max_sgSize == sgSize) { + max_sgSize == sgSize) + { sycl::vec res_vec(const_val); #pragma unroll for (std::uint8_t it = 0; it < n_vecs * vec_sz; it += vec_sz) { @@ -99,7 +100,8 @@ struct UnaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { out[k] = const_val; } } @@ -115,7 +117,8 @@ struct UnaryContigFunctor (ndit.get_group(0) * ndit.get_local_range(0) + sg.get_group_id()[0] * max_sgSize); if (base + n_vecs * vec_sz * sgSize < nelems_ && - sgSize == max_sgSize) { + sgSize == max_sgSize) + { sycl::vec x; #pragma unroll @@ -136,7 +139,8 @@ struct UnaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { // scalar call out[k] = op(in[k]); } @@ -156,7 +160,8 @@ struct UnaryContigFunctor sg.get_group_id()[0] * maxsgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (maxsgSize == sgSize)) { + (maxsgSize == sgSize)) + { sycl::vec arg_vec; #pragma unroll @@ -180,7 +185,8 @@ struct UnaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { out[k] = op(in[k]); } } @@ -198,7 +204,8 @@ struct UnaryContigFunctor sg.get_group_id()[0] * maxsgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (maxsgSize == sgSize)) { + (maxsgSize == sgSize)) + { sycl::vec arg_vec; sycl::vec res_vec; @@ -223,7 +230,8 @@ struct UnaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { out[k] = op(in[k]); } } @@ -410,7 +418,8 @@ struct BinaryContigFunctor sg.get_group_id()[0] * sgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (sgSize == maxsgSize)) { + (sgSize == maxsgSize)) + { sycl::vec arg1_vec; sycl::vec arg2_vec; sycl::vec res_vec; @@ -437,7 +446,8 @@ struct BinaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { out[k] = op(in1[k], in2[k]); } } @@ -454,7 +464,8 @@ struct BinaryContigFunctor sg.get_group_id()[0] * sgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (sgSize == maxsgSize)) { + (sgSize == maxsgSize)) + { sycl::vec arg1_vec; sycl::vec arg2_vec; sycl::vec res_vec; @@ -485,7 +496,8 @@ struct BinaryContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { out[k] = op(in1[k], in2[k]); } } @@ -599,7 +611,8 @@ struct BinaryContigMatrixContigRowBroadcastingFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < n_elems; - k += sgSize) { + k += sgSize) + { res[k] = op(mat[k], padded_vec[k % n1]); } } @@ -663,7 +676,8 @@ struct BinaryContigRowContigMatrixBroadcastingFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < n_elems; - k += sgSize) { + k += sgSize) + { res[k] = op(padded_vec[k % n1], mat[k]); } } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp index d9ad2a7482..7bf2be10c9 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp @@ -84,7 +84,8 @@ struct BinaryInplaceContigFunctor sg.get_group_id()[0] * sgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (sgSize == maxsgSize)) { + (sgSize == maxsgSize)) + { sycl::vec arg_vec; sycl::vec res_vec; @@ -107,7 +108,8 @@ struct BinaryInplaceContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { op(lhs[k], rhs[k]); } } @@ -124,7 +126,8 @@ struct BinaryInplaceContigFunctor sg.get_group_id()[0] * sgSize); if ((base + n_vecs * vec_sz * sgSize < nelems_) && - (sgSize == maxsgSize)) { + (sgSize == maxsgSize)) + { sycl::vec arg_vec; sycl::vec res_vec; @@ -148,7 +151,8 @@ struct BinaryInplaceContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems_; - k += sgSize) { + k += sgSize) + { op(lhs[k], rhs[k]); } } @@ -248,7 +252,8 @@ struct BinaryInplaceRowMatrixBroadcastingFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < n_elems; - k += sgSize) { + k += sgSize) + { op(mat[k], padded_vec[k % n1]); } } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp index e0123e8d91..5348173856 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp @@ -138,7 +138,8 @@ template struct ConjContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -185,7 +186,8 @@ template struct ConjStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/copysign.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/copysign.hpp index d2ce413943..d226422494 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/copysign.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/copysign.hpp @@ -66,7 +66,8 @@ template struct CopysignFunctor { auto tmp = sycl::copysign(in1, in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp index 308f6b0d44..14b2345788 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp @@ -211,7 +211,8 @@ template struct CosContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -257,7 +258,8 @@ template struct CosStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp index 96de1639ca..866bd2731d 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp @@ -200,7 +200,8 @@ template struct CoshContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -247,7 +248,8 @@ template struct CoshStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp index b3a1627a9a..5086a89cec 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp @@ -77,12 +77,14 @@ template struct EqualFunctor std::is_signed_v != std::is_signed_v) { if constexpr (std::is_signed_v && - !std::is_signed_v) { + !std::is_signed_v) + { return (in1 < 0) ? false : (static_cast(in1) == in2); } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? false : (in1 == static_cast(in2)); } @@ -101,7 +103,8 @@ template struct EqualFunctor { auto tmp = (in1 == in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp index cd451dda5a..38abed80cb 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp @@ -169,7 +169,8 @@ template struct ExpContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -215,7 +216,8 @@ template struct ExpStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp index dbe32b47d1..9d244a0375 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp @@ -171,7 +171,8 @@ template struct Exp2ContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -218,7 +219,8 @@ template struct Exp2StridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp index 1002ad1f1e..2ab077ab76 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp @@ -184,7 +184,8 @@ template struct Expm1ContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -232,7 +233,8 @@ template struct Expm1StridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor.hpp index a709787a65..90e6941bdd 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor.hpp @@ -131,7 +131,8 @@ template struct FloorContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -178,7 +179,8 @@ template struct FloorStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor_divide.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor_divide.hpp index 231e18f4da..ce89b0778f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor_divide.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor_divide.hpp @@ -120,10 +120,7 @@ struct FloorDivideFunctor } private: - bool l_xor(bool b1, bool b2) const - { - return (b1 != b2); - } + bool l_xor(bool b1, bool b2) const { return (b1 != b2); } }; template struct FloorDivideInplaceFunctor } private: - bool l_xor(bool b1, bool b2) const - { - return (b1 != b2); - } + bool l_xor(bool b1, bool b2) const { return (b1 != b2); } }; template struct GreaterFunctor std::is_signed_v != std::is_signed_v) { if constexpr (std::is_signed_v && - !std::is_signed_v) { + !std::is_signed_v) + { return (in1 < 0) ? false : (static_cast(in1) > in2); } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? true : (in1 > static_cast(in2)); } @@ -102,7 +104,8 @@ template struct GreaterFunctor auto tmp = (in1 > in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/greater_equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/greater_equal.hpp index 8569eb0216..b429a3d00c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/greater_equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/greater_equal.hpp @@ -77,12 +77,14 @@ struct GreaterEqualFunctor std::is_signed_v != std::is_signed_v) { if constexpr (std::is_signed_v && - !std::is_signed_v) { + !std::is_signed_v) + { return (in1 < 0) ? false : (static_cast(in1) >= in2); } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? true : (in1 >= static_cast(in2)); } @@ -103,7 +105,8 @@ struct GreaterEqualFunctor auto tmp = (in1 >= in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp index 53817373a0..52498f76d5 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp @@ -68,7 +68,8 @@ template struct HypotFunctor { auto res = sycl::hypot(in1, in2); if constexpr (std::is_same_v) { + typename decltype(res)::element_type>) + { return res; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/imag.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/imag.hpp index e422be02c2..6b937b3071 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/imag.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/imag.hpp @@ -134,7 +134,8 @@ template struct ImagContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -181,7 +182,8 @@ template struct ImagStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less.hpp index 294a78ba2f..88d0e6e19f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less.hpp @@ -75,12 +75,14 @@ template struct LessFunctor std::is_signed_v != std::is_signed_v) { if constexpr (std::is_signed_v && - !std::is_signed_v) { + !std::is_signed_v) + { return (in1 < 0) ? true : (static_cast(in1) < in2); } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? false : (in1 < static_cast(in2)); } @@ -100,7 +102,8 @@ template struct LessFunctor auto tmp = (in1 < in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { @@ -213,7 +216,8 @@ template struct LessContigFactory fnT get() { if constexpr (std::is_same_v< - typename LessOutputType::value_type, void>) { + typename LessOutputType::value_type, void>) + { fnT fn = nullptr; return fn; } @@ -264,7 +268,8 @@ template struct LessStridedFactory fnT get() { if constexpr (std::is_same_v< - typename LessOutputType::value_type, void>) { + typename LessOutputType::value_type, void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less_equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less_equal.hpp index 7b18a0b045..97400aa475 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less_equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/less_equal.hpp @@ -75,12 +75,14 @@ template struct LessEqualFunctor std::is_signed_v != std::is_signed_v) { if constexpr (std::is_signed_v && - !std::is_signed_v) { + !std::is_signed_v) + { return (in1 < 0) ? true : (static_cast(in1) <= in2); } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? false : (in1 <= static_cast(in2)); } @@ -101,7 +103,8 @@ template struct LessEqualFunctor auto tmp = (in1 <= in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp index 5fddfd9b64..6a4fd4e34e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp @@ -126,7 +126,8 @@ template struct LogContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -173,7 +174,8 @@ template struct LogStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp index d062da6c6c..af2ad072c5 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp @@ -145,7 +145,8 @@ template struct Log10ContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -193,7 +194,8 @@ template struct Log10StridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp index 9cc0927cca..fe08dc805f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp @@ -150,7 +150,8 @@ template struct Log1pContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -198,7 +199,8 @@ template struct Log1pStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp index dac784c7a4..d4ea0aca47 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp @@ -146,7 +146,8 @@ template struct Log2ContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -194,7 +195,8 @@ template struct Log2StridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp index eb621cdbf0..403a8a2799 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp @@ -77,7 +77,8 @@ struct LogicalAndFunctor auto tmp = (in1 && in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp index 8d1dec5751..4706c7936c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp @@ -76,7 +76,8 @@ template struct LogicalOrFunctor auto tmp = (in1 || in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp index ba4540fa44..a444bc5395 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp @@ -80,7 +80,8 @@ struct LogicalXorFunctor auto tmp = (tmp1 != tmp2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp index b98ca2ef16..147f62f53e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp @@ -82,7 +82,8 @@ template struct MultiplyFunctor { auto tmp = in1 * in2; if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { @@ -395,10 +396,7 @@ template struct MultiplyInplaceFunctor using supports_vec = std::negation< std::disjunction, tu_ns::is_complex>>; - void operator()(resT &res, const argT &in) - { - res *= in; - } + void operator()(resT &res, const argT &in) { res *= in; } template void operator()(sycl::vec &res, diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/negative.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/negative.hpp index 31bba228cd..d52c2f33ee 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/negative.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/negative.hpp @@ -60,10 +60,7 @@ template struct NegativeFunctor using supports_sg_loadstore = typename std::negation< std::disjunction, is_complex>>; - resT operator()(const argT &x) const - { - return -x; - } + resT operator()(const argT &x) const { return -x; } }; template struct NegativeContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -173,7 +171,8 @@ template struct NegativeStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/nextafter.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/nextafter.hpp index 7e96de4a6a..d5137f0c6d 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/nextafter.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/nextafter.hpp @@ -66,7 +66,8 @@ template struct NextafterFunctor { auto res = sycl::nextafter(in1, in2); if constexpr (std::is_same_v) { + typename decltype(res)::element_type>) + { return res; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/not_equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/not_equal.hpp index c31a05b266..437ceb2da8 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/not_equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/not_equal.hpp @@ -69,7 +69,8 @@ template struct NotEqualFunctor } else { if constexpr (!std::is_signed_v && - std::is_signed_v) { + std::is_signed_v) + { return (in2 < 0) ? true : (in1 != static_cast(in2)); } } @@ -86,7 +87,8 @@ template struct NotEqualFunctor { auto tmp = (in1 != in2); if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/positive.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/positive.hpp index e77b823f51..92eaf3c0d2 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/positive.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/positive.hpp @@ -61,10 +61,7 @@ template struct PositiveFunctor using supports_sg_loadstore = typename std::negation< std::disjunction, is_complex>>; - resT operator()(const argT &x) const - { - return x; - } + resT operator()(const argT &x) const { return x; } template sycl::vec operator()(const sycl::vec &in) const @@ -135,7 +132,8 @@ template struct PositiveContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -188,7 +186,8 @@ template struct PositiveStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp index 7b2fe38af5..a21b2d4318 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp @@ -264,7 +264,8 @@ template struct PowContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -314,7 +315,8 @@ template struct PowStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -464,7 +466,8 @@ template struct PowInplaceContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -503,7 +506,8 @@ struct PowInplaceStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/proj.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/proj.hpp index 3e220b2b13..92a8535309 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/proj.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/proj.hpp @@ -135,7 +135,8 @@ template struct ProjContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -188,7 +189,8 @@ template struct ProjStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/real.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/real.hpp index b5197b8fde..8949a79955 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/real.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/real.hpp @@ -134,7 +134,8 @@ template struct RealContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -181,7 +182,8 @@ template struct RealStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/remainder.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/remainder.hpp index 15c0d6d070..585d1c6d7f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/remainder.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/remainder.hpp @@ -102,7 +102,8 @@ template struct RemainderFunctor else { rem[i] = in1[i] % in2[i]; if constexpr (std::is_signed_v || - std::is_signed_v) { + std::is_signed_v) + { if (rem[i] != 0 && l_xor(in1[i] < 0, in2[i] < 0)) { rem[i] += in2[i]; } @@ -137,10 +138,7 @@ template struct RemainderFunctor } private: - bool l_xor(bool b1, bool b2) const - { - return (b1 != b2); - } + bool l_xor(bool b1, bool b2) const { return (b1 != b2); } }; template struct RemainderInplaceFunctor else { auto rem = res[i] % in[i]; if constexpr (std::is_signed_v || - std::is_signed_v) { + std::is_signed_v) + { if (rem != 0 && l_xor(res[i] < 0, in[i] < 0)) { rem += in[i]; } @@ -394,10 +393,7 @@ template struct RemainderInplaceFunctor } private: - bool l_xor(bool b1, bool b2) const - { - return (b1 != b2); - } + bool l_xor(bool b1, bool b2) const { return (b1 != b2); } }; template struct RoundContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -189,7 +190,8 @@ template struct RoundStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/rsqrt.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/rsqrt.hpp index 614f067323..541b036931 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/rsqrt.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/rsqrt.hpp @@ -63,10 +63,7 @@ template struct RsqrtFunctor // do both argTy and resTy support sugroup store/load operation using supports_sg_loadstore = typename std::true_type; - resT operator()(const argT &in) const - { - return sycl::rsqrt(in); - } + resT operator()(const argT &in) const { return sycl::rsqrt(in); } }; template struct RsqrtContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -163,7 +161,8 @@ template struct RsqrtStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp index db24a2cf14..554398ae56 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp @@ -154,7 +154,8 @@ template struct SignContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -206,7 +207,8 @@ template struct SignStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/signbit.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/signbit.hpp index 345073d8a9..ee4a97d9b6 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/signbit.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/signbit.hpp @@ -58,10 +58,7 @@ template struct SignbitFunctor using supports_vec = std::true_type; using supports_sg_loadstore = std::true_type; - resT operator()(const argT &in) const - { - return std::signbit(in); - } + resT operator()(const argT &in) const { return std::signbit(in); } template sycl::vec operator()(const sycl::vec &in) const @@ -122,7 +119,8 @@ template struct SignbitContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -170,7 +168,8 @@ template struct SignbitStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp index d8425cd5a0..ba46affae6 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp @@ -233,7 +233,8 @@ template struct SinContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -279,7 +280,8 @@ template struct SinStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp index c5ed20ad1b..01a87d923f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp @@ -202,7 +202,8 @@ template struct SinhContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -249,7 +250,8 @@ template struct SinhStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp index 032666618b..6b63f74fe7 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp @@ -128,7 +128,8 @@ template struct SqrtContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -175,7 +176,8 @@ template struct SqrtStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp index eb6ddb1783..72f3bda389 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp @@ -153,7 +153,8 @@ template struct SquareContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -200,7 +201,8 @@ template struct SquareStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/subtract.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/subtract.hpp index 7206f09bbf..4a8cfb50a7 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/subtract.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/subtract.hpp @@ -69,7 +69,8 @@ template struct SubtractFunctor { auto tmp = in1 - in2; if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { @@ -396,10 +397,7 @@ template struct SubtractInplaceFunctor using supports_vec = std::negation< std::disjunction, tu_ns::is_complex>>; - void operator()(resT &res, const argT &in) - { - res -= in; - } + void operator()(resT &res, const argT &in) { res -= in; } template void operator()(sycl::vec &res, diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp index 03bb5baba1..c745bc011b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp @@ -177,7 +177,8 @@ template struct TanContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -223,7 +224,8 @@ template struct TanStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp index db275bc5ea..2e23e46f4a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp @@ -171,7 +171,8 @@ template struct TanhContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -218,7 +219,8 @@ template struct TanhStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp index fd2b4a2ddd..396e5c995e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp @@ -96,7 +96,8 @@ struct TrueDivideFunctor { auto tmp = in1 / in2; if constexpr (std::is_same_v) { + typename decltype(tmp)::element_type>) + { return tmp; } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/trunc.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/trunc.hpp index a91545c64d..5740bc0ef2 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/trunc.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/trunc.hpp @@ -128,7 +128,8 @@ template struct TruncContigFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } @@ -175,7 +176,8 @@ template struct TruncStridedFactory fnT get() { if constexpr (std::is_same_v::value_type, - void>) { + void>) + { fnT fn = nullptr; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index ac4cf5c3c0..485fbef513 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -350,7 +350,8 @@ template struct TakeWrapFactory fnT get() { if constexpr (std::is_integral::value && - !std::is_same::value) { + !std::is_same::value) + { fnT fn = take_impl; return fn; } @@ -366,7 +367,8 @@ template struct TakeClipFactory fnT get() { if constexpr (std::is_integral::value && - !std::is_same::value) { + !std::is_same::value) + { fnT fn = take_impl; return fn; } @@ -382,7 +384,8 @@ template struct PutWrapFactory fnT get() { if constexpr (std::is_integral::value && - !std::is_same::value) { + !std::is_same::value) + { fnT fn = put_impl; return fn; } @@ -398,7 +401,8 @@ template struct PutClipFactory fnT get() { if constexpr (std::is_integral::value && - !std::is_same::value) { + !std::is_same::value) + { fnT fn = put_impl; return fn; } diff --git a/dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp b/dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp index 11eb710c30..15f7f4c291 100644 --- a/dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp +++ b/dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp @@ -1074,7 +1074,8 @@ sycl::event dot_product_tree_impl(sycl::queue &exec_q, sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -1326,7 +1327,8 @@ dot_product_contig_tree_impl(sycl::queue &exec_q, sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); diff --git a/dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp b/dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp index b92ed0be40..03f4780b43 100644 --- a/dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp +++ b/dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp @@ -1037,7 +1037,8 @@ class GemmBatchFunctorThreadNM_vecm slmB_t vec{}; #pragma unroll for (std::uint32_t lane_id = 0; lane_id < m_vec_size; - ++lane_id) { + ++lane_id) + { const size_t g_j1 = g_j + lane_id; vec[lane_id] = (g_j1 < m && g_s < k) ? static_cast( @@ -1105,7 +1106,8 @@ class GemmBatchFunctorThreadNM_vecm j + (local_j + pr_j * wg_delta_m) * m_vec_size; #pragma unroll for (std::uint32_t lane_id = 0; lane_id < m_vec_size; - ++lane_id) { + ++lane_id) + { const size_t out_flat_id = out_i * c_st0 + (out_j + lane_id) * c_st1; if (out_j + lane_id < m) { @@ -1139,18 +1141,12 @@ struct GemmBatchFunctorThreadNM_vecm_HyperParameters { } - constexpr std::uint32_t get_wi_delta_n() const - { - return wi_delta_n; - } + constexpr std::uint32_t get_wi_delta_n() const { return wi_delta_n; } constexpr std::uint32_t get_wi_delta_m_vecs() const { return wi_delta_m_vecs; } - constexpr std::uint32_t get_m_vec_size() const - { - return m_vec_size; - } + constexpr std::uint32_t get_m_vec_size() const { return m_vec_size; } }; template diff --git a/dpctl/tensor/libtensor/include/kernels/reductions.hpp b/dpctl/tensor/libtensor/include/kernels/reductions.hpp index 992fe3b388..ed8f7dbb23 100644 --- a/dpctl/tensor/libtensor/include/kernels/reductions.hpp +++ b/dpctl/tensor/libtensor/include/kernels/reductions.hpp @@ -383,7 +383,8 @@ struct CustomReductionOverGroupWithAtomicFunctor res_ref.fetch_min(red_val_over_wg); } else if constexpr (su_ns::IsSyclLogicalAnd::value) { + ReductionOp>::value) + { res_ref.fetch_and(red_val_over_wg); } else if constexpr (su_ns::IsSyclLogicalOr::value) @@ -1290,7 +1291,8 @@ sycl::event reduction_over_group_temps_strided_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -1540,7 +1542,8 @@ sycl::event reduction_axis1_over_group_temps_contig_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -1783,7 +1786,8 @@ sycl::event reduction_axis0_over_group_temps_contig_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -2098,7 +2102,8 @@ struct SearchReduction } } else if constexpr (su_ns::IsMaximum::value) { + ReductionOp>::value) + { if (val > local_red_val) { local_red_val = val; if constexpr (!First) { @@ -2290,7 +2295,8 @@ struct CustomSearchReduction } } else if constexpr (su_ns::IsMaximum::value) { + ReductionOp>::value) + { using dpctl::tensor::type_utils::is_complex; if constexpr (is_complex::value) { using dpctl::tensor::math_utils::greater_complex; @@ -2707,7 +2713,8 @@ sycl::event search_over_group_temps_strided_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -3000,7 +3007,8 @@ sycl::event search_axis1_over_group_temps_contig_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); @@ -3286,7 +3294,8 @@ sycl::event search_axis0_over_group_temps_contig_impl( sycl::event dependent_ev = first_reduction_ev; while (remaining_reduction_nelems > - preferred_reductions_per_wi * max_wg) { + preferred_reductions_per_wi * max_wg) + { size_t reduction_groups_ = (remaining_reduction_nelems + preferred_reductions_per_wi * wg - 1) / (preferred_reductions_per_wi * wg); diff --git a/dpctl/tensor/libtensor/include/kernels/sorting/sort.hpp b/dpctl/tensor/libtensor/include/kernels/sorting/sort.hpp index cd5ba4b025..b26638ff75 100644 --- a/dpctl/tensor/libtensor/include/kernels/sorting/sort.hpp +++ b/dpctl/tensor/libtensor/include/kernels/sorting/sort.hpp @@ -125,7 +125,8 @@ void merge_impl(const std::size_t offset, if (r_search_bound_2 == l_search_bound_2) { const std::size_t shift_2 = l_search_bound_2 - start_2; for (std::size_t idx = local_start_1 + 1; idx < local_end_1 - 1; - ++idx) { + ++idx) + { const auto intermediate_item_1 = in_acc[idx]; const std::size_t shift_1 = idx - start_1; out_acc[start_out + shift_1 + shift_2] = @@ -134,7 +135,8 @@ void merge_impl(const std::size_t offset, } else { for (std::size_t idx = local_start_1 + 1; idx < local_end_1 - 1; - ++idx) { + ++idx) + { const auto intermediate_item_1 = in_acc[idx]; // we shouldn't seek in whole 2nd sequence. Just for the // part where the 1st sequence should be @@ -286,10 +288,7 @@ struct GetValueType> template struct GetReadOnlyAccess { - Iter operator()(Iter it, sycl::handler &) - { - return it; - } + Iter operator()(Iter it, sycl::handler &) { return it; } }; template @@ -305,10 +304,7 @@ struct GetReadOnlyAccess> template struct GetWriteDiscardAccess { - Iter operator()(Iter it, sycl::handler &) - { - return it; - } + Iter operator()(Iter it, sycl::handler &) { return it; } }; template @@ -324,10 +320,7 @@ struct GetWriteDiscardAccess> template struct GetReadWriteAccess { - Iter operator()(Iter it, sycl::handler &) - { - return it; - } + Iter operator()(Iter it, sycl::handler &) { return it; } }; template diff --git a/dpctl/tensor/libtensor/include/kernels/sorting/sort_detail.hpp b/dpctl/tensor/libtensor/include/kernels/sorting/sort_detail.hpp index 57f47fc756..b286f04dfe 100644 --- a/dpctl/tensor/libtensor/include/kernels/sorting/sort_detail.hpp +++ b/dpctl/tensor/libtensor/include/kernels/sorting/sort_detail.hpp @@ -38,10 +38,7 @@ namespace kernels namespace sort_detail { -template T quotient_ceil(T n, T m) -{ - return (n + m - 1) / m; -} +template T quotient_ceil(T n, T m) { return (n + m - 1) / m; } template std::size_t lower_bound_impl(const Acc acc, diff --git a/dpctl/tensor/libtensor/include/kernels/where.hpp b/dpctl/tensor/libtensor/include/kernels/where.hpp index d6465b2f59..b356c256c3 100644 --- a/dpctl/tensor/libtensor/include/kernels/where.hpp +++ b/dpctl/tensor/libtensor/include/kernels/where.hpp @@ -108,7 +108,8 @@ class WhereContigFunctor sg.get_group_id()[0] * max_sgSize); if (base + n_vecs * vec_sz * sgSize < nelems && - sgSize == max_sgSize) { + sgSize == max_sgSize) + { sycl::vec dst_vec; sycl::vec x1_vec; sycl::vec x2_vec; @@ -142,7 +143,8 @@ class WhereContigFunctor } else { for (size_t k = base + sg.get_local_id()[0]; k < nelems; - k += sgSize) { + k += sgSize) + { dst_p[k] = cond_p[k] ? x1_p[k] : x2_p[k]; } } diff --git a/dpctl/tensor/libtensor/include/utils/offset_utils.hpp b/dpctl/tensor/libtensor/include/utils/offset_utils.hpp index 3906522817..1ad89c4fac 100644 --- a/dpctl/tensor/libtensor/include/utils/offset_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/offset_utils.hpp @@ -126,10 +126,7 @@ device_allocate_and_pack(sycl::queue &q, struct NoOpIndexer { constexpr NoOpIndexer() {} - constexpr size_t operator()(size_t gid) const - { - return gid; - } + constexpr size_t operator()(size_t gid) const { return gid; } }; using dpctl::tensor::ssize_t; @@ -145,10 +142,7 @@ struct StridedIndexer { } - ssize_t operator()(ssize_t gid) const - { - return compute_offset(gid); - } + ssize_t operator()(ssize_t gid) const { return compute_offset(gid); } ssize_t operator()(size_t gid) const { @@ -186,10 +180,7 @@ struct UnpackedStridedIndexer { } - ssize_t operator()(ssize_t gid) const - { - return compute_offset(gid); - } + ssize_t operator()(ssize_t gid) const { return compute_offset(gid); } ssize_t operator()(size_t gid) const { @@ -263,14 +254,8 @@ template struct TwoOffsets { } - constexpr displacementT get_first_offset() const - { - return first_offset; - } - constexpr displacementT get_second_offset() const - { - return second_offset; - } + constexpr displacementT get_first_offset() const { return first_offset; } + constexpr displacementT get_second_offset() const { return second_offset; } private: displacementT first_offset = 0; @@ -368,18 +353,9 @@ template struct ThreeOffsets { } - constexpr displacementT get_first_offset() const - { - return first_offset; - } - constexpr displacementT get_second_offset() const - { - return second_offset; - } - constexpr displacementT get_third_offset() const - { - return third_offset; - } + constexpr displacementT get_first_offset() const { return first_offset; } + constexpr displacementT get_second_offset() const { return second_offset; } + constexpr displacementT get_third_offset() const { return third_offset; } private: displacementT first_offset = 0; @@ -497,22 +473,10 @@ template struct FourOffsets { } - constexpr displacementT get_first_offset() const - { - return first_offset; - } - constexpr displacementT get_second_offset() const - { - return second_offset; - } - constexpr displacementT get_third_offset() const - { - return third_offset; - } - constexpr displacementT get_fourth_offset() const - { - return fourth_offset; - } + constexpr displacementT get_first_offset() const { return first_offset; } + constexpr displacementT get_second_offset() const { return second_offset; } + constexpr displacementT get_third_offset() const { return third_offset; } + constexpr displacementT get_fourth_offset() const { return fourth_offset; } private: displacementT first_offset = 0; diff --git a/dpctl/tensor/libtensor/include/utils/strided_iters.hpp b/dpctl/tensor/libtensor/include/utils/strided_iters.hpp index e12bbebbd5..880278dd48 100644 --- a/dpctl/tensor/libtensor/include/utils/strided_iters.hpp +++ b/dpctl/tensor/libtensor/include/utils/strided_iters.hpp @@ -305,14 +305,8 @@ template class CIndexer_array elem_count = s; } - indT size() const - { - return elem_count; - } - indT rank() const - { - return ndim; - } + indT size() const { return elem_count; } + indT rank() const { return ndim; } void set(const indT i) { @@ -332,10 +326,7 @@ template class CIndexer_array multi_index[0] = i_; } - const index_t &get() const - { - return multi_index; - } + const index_t &get() const { return multi_index; } }; /* @@ -723,7 +714,8 @@ std::tuple contract_iter3(const vecT &shape, { const size_t dim = shape.size(); if (dim != strides1.size() || dim != strides2.size() || - dim != strides3.size()) { + dim != strides3.size()) + { throw Error("Shape and strides must be of equal size."); } vecT out_shape = shape; @@ -857,7 +849,8 @@ int simplify_iteration_four_strides(const int nd, StridesTy jump4 = strides4_w[i] - (shape_w[i + 1] - 1) * str4; if (jump1 == str1 && jump2 == str2 && jump3 == str3 && - jump4 == str4) { + jump4 == str4) + { changed = true; shape_w[i] *= shape_w[i + 1]; for (int j = i; j < nd_; ++j) { diff --git a/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp b/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp index 75ab1d9341..3ad465db6a 100644 --- a/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp @@ -193,7 +193,8 @@ template struct Maximum return max_complex(x, y); } else if constexpr (std::is_floating_point_v || - std::is_same_v) { + std::is_same_v) + { return (std::isnan(x) || x > y) ? x : y; } else if constexpr (std::is_same_v) { @@ -216,7 +217,8 @@ template struct Minimum return min_complex(x, y); } else if constexpr (std::is_floating_point_v || - std::is_same_v) { + std::is_same_v) + { return (std::isnan(x) || x < y) ? x : y; } else if constexpr (std::is_same_v) { @@ -351,10 +353,7 @@ struct GetIdentity::value>> template struct Hypot { - T operator()(const T &x, const T &y) const - { - return sycl::hypot(x, y); - } + T operator()(const T &x, const T &y) const { return sycl::hypot(x, y); } }; template diff --git a/dpctl/tensor/libtensor/include/utils/type_dispatch_building.hpp b/dpctl/tensor/libtensor/include/utils/type_dispatch_building.hpp index b133f29a97..a9f3a0c876 100644 --- a/dpctl/tensor/libtensor/include/utils/type_dispatch_building.hpp +++ b/dpctl/tensor/libtensor/include/utils/type_dispatch_building.hpp @@ -251,10 +251,7 @@ template struct NullPtrVector NullPtrVector() : val(nullptr) {} - const_reference operator[](int) const - { - return val; - } + const_reference operator[](int) const { return val; } private: value_type val; @@ -268,10 +265,7 @@ template struct NullPtrTable NullPtrTable() : val() {} - const_reference operator[](int) const - { - return val; - } + const_reference operator[](int) const { return val; } private: value_type val; diff --git a/dpctl/tensor/libtensor/source/boolean_advanced_indexing.cpp b/dpctl/tensor/libtensor/source/boolean_advanced_indexing.cpp index 582fc050fe..a1cb94a008 100644 --- a/dpctl/tensor/libtensor/source/boolean_advanced_indexing.cpp +++ b/dpctl/tensor/libtensor/source/boolean_advanced_indexing.cpp @@ -170,7 +170,8 @@ py_extract(const dpctl::tensor::usm_ndarray &src, // masked_dst_nelems is number of set elements in the mask, or last element // in cumsum if (!same_ortho_dims || - (masked_src_nelems != static_cast(cumsum_sz))) { + (masked_src_nelems != static_cast(cumsum_sz))) + { throw py::value_error("Inconsistent array dimensions"); } @@ -492,7 +493,8 @@ py_place(const dpctl::tensor::usm_ndarray &dst, } if (!same_ortho_dims || - (masked_dst_nelems != static_cast(cumsum_sz))) { + (masked_dst_nelems != static_cast(cumsum_sz))) + { throw py::value_error("Inconsistent array dimensions"); } diff --git a/dpctl/tensor/libtensor/source/copy_for_roll.cpp b/dpctl/tensor/libtensor/source/copy_for_roll.cpp index e9f5376a97..da477150c0 100644 --- a/dpctl/tensor/libtensor/source/copy_for_roll.cpp +++ b/dpctl/tensor/libtensor/source/copy_for_roll.cpp @@ -189,7 +189,8 @@ copy_usm_ndarray_for_roll_1d(const dpctl::tensor::usm_ndarray &src, src_offset, dst_offset); if (nd == 1 && simplified_src_strides[0] == 1 && - simplified_dst_strides[0] == 1) { + simplified_dst_strides[0] == 1) + { auto fn = copy_for_roll_contig_dispatch_vector[type_id]; if (fn != nullptr) { diff --git a/dpctl/tensor/libtensor/source/device_support_queries.cpp b/dpctl/tensor/libtensor/source/device_support_queries.cpp index b2a48b2e97..2d212ca0fc 100644 --- a/dpctl/tensor/libtensor/source/device_support_queries.cpp +++ b/dpctl/tensor/libtensor/source/device_support_queries.cpp @@ -103,15 +103,9 @@ std::string _default_device_complex_type(const sycl::device &d) } } -std::string _default_device_bool_type(const sycl::device &) -{ - return "b1"; -} +std::string _default_device_bool_type(const sycl::device &) { return "b1"; } -std::string _default_device_index_type(const sycl::device &) -{ - return "i8"; -} +std::string _default_device_index_type(const sycl::device &) { return "i8"; } sycl::device _extract_device(const py::object &arg) { diff --git a/dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp b/dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp index 9c823b8fc4..48cc9f5c47 100644 --- a/dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp +++ b/dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp @@ -181,7 +181,8 @@ py_unary_ufunc(const dpctl::tensor::usm_ndarray &src, src_offset, dst_offset); if (nd == 1 && simplified_src_strides[0] == 1 && - simplified_dst_strides[0] == 1) { + simplified_dst_strides[0] == 1) + { // Special case of contiguous data auto contig_fn = contig_dispatch_vector[src_typeid]; diff --git a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp index 7e8766bd61..56db97eab7 100644 --- a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp +++ b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp @@ -382,7 +382,8 @@ usm_ndarray_take(const dpctl::tensor::usm_ndarray &src, } if (!(ind_type_id == - array_types.typenum_to_lookup_id(ind_.get_typenum()))) { + array_types.typenum_to_lookup_id(ind_.get_typenum()))) + { throw py::type_error( "Indices array data types are not all the same."); } @@ -689,7 +690,8 @@ usm_ndarray_put(const dpctl::tensor::usm_ndarray &dst, } if (!(ind_type_id == - array_types.typenum_to_lookup_id(ind_.get_typenum()))) { + array_types.typenum_to_lookup_id(ind_.get_typenum()))) + { throw py::type_error( "Indices array data types are not all the same."); } diff --git a/dpctl/tensor/libtensor/source/linalg_functions/dot.cpp b/dpctl/tensor/libtensor/source/linalg_functions/dot.cpp index bc6157e085..00683a1315 100644 --- a/dpctl/tensor/libtensor/source/linalg_functions/dot.cpp +++ b/dpctl/tensor/libtensor/source/linalg_functions/dot.cpp @@ -239,7 +239,8 @@ py_dot(const dpctl::tensor::usm_ndarray &x1, } size_t x1_outer_nelems(1); for (int i = batch_dims; same_shapes && (i < (batch_dims + x1_outer_dims)); - ++i) { + ++i) + { same_shapes = same_shapes && (x1_shape_ptr[i] == dst_shape_ptr[i]); x1_outer_nelems *= x1_shape_ptr[i]; } @@ -419,7 +420,8 @@ py_dot(const dpctl::tensor::usm_ndarray &x1, bool reduce_all_elems = false; if (simplified_inner_x1_strides[0] == 1 && - simplified_inner_x2_strides[0] == 1) { + simplified_inner_x2_strides[0] == 1) + { reduce_all_elems = (simplified_batch_shape[0] == 1); dot_product_c_contig = (simplified_batch_dst_strides[0] == 1) && diff --git a/dpctl/tensor/libtensor/source/reductions/argmax.cpp b/dpctl/tensor/libtensor/source/reductions/argmax.cpp index 3331423ddc..e441a36139 100644 --- a/dpctl/tensor/libtensor/source/reductions/argmax.cpp +++ b/dpctl/tensor/libtensor/source/reductions/argmax.cpp @@ -126,7 +126,8 @@ struct ArgmaxOverAxisTempsStridedFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::maximum; // op for indices @@ -160,7 +161,8 @@ struct ArgmaxOverAxis1TempsContigFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::maximum; // op for indices @@ -194,7 +196,8 @@ struct ArgmaxOverAxis0TempsContigFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::maximum; // op for indices diff --git a/dpctl/tensor/libtensor/source/reductions/argmin.cpp b/dpctl/tensor/libtensor/source/reductions/argmin.cpp index 582a96247c..4892893cc5 100644 --- a/dpctl/tensor/libtensor/source/reductions/argmin.cpp +++ b/dpctl/tensor/libtensor/source/reductions/argmin.cpp @@ -126,7 +126,8 @@ struct ArgminOverAxisTempsStridedFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::minimum; // op for indices @@ -160,7 +161,8 @@ struct ArgminOverAxis1TempsContigFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::minimum; // op for indices @@ -194,7 +196,8 @@ struct ArgminOverAxis0TempsContigFactory dstTy>::is_defined) { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { // op for values using ReductionOpT = sycl::minimum; // op for indices diff --git a/dpctl/tensor/libtensor/source/reductions/max.cpp b/dpctl/tensor/libtensor/source/reductions/max.cpp index cfd6daf06e..0e402403a4 100644 --- a/dpctl/tensor/libtensor/source/reductions/max.cpp +++ b/dpctl/tensor/libtensor/source/reductions/max.cpp @@ -183,9 +183,11 @@ struct MaxOverAxisTempsStridedFactory fnT get() const { if constexpr (TypePairSupportDataForMaxReductionTemps< - srcTy, dstTy>::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::maximum; return dpctl::tensor::kernels:: reduction_over_group_temps_strided_impl::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::maximum; return dpctl::tensor::kernels:: reduction_axis1_over_group_temps_contig_impl::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::maximum; return dpctl::tensor::kernels:: reduction_axis0_over_group_temps_contig_impl::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::minimum; return dpctl::tensor::kernels:: reduction_over_group_temps_strided_impl::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::minimum; return dpctl::tensor::kernels:: reduction_axis1_over_group_temps_contig_impl::is_defined) { + srcTy, dstTy>::is_defined) + { if constexpr (std::is_integral_v && - !std::is_same_v) { + !std::is_same_v) + { using ReductionOpT = sycl::minimum; return dpctl::tensor::kernels:: reduction_axis0_over_group_temps_contig_impl struct ArithmeticAtomicSupportFactory template struct MinMaxAtomicSupportFactory { - fnT get() - { - return check_atomic_support; - } + fnT get() { return check_atomic_support; } }; template diff --git a/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp b/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp index f9c61db5a8..649487dd12 100644 --- a/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp +++ b/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp @@ -976,7 +976,8 @@ std::pair py_search_over_axis( reduction_nelems); } else if (static_cast(compact_reduction_src_strides[0]) == - iter_nelems) { + iter_nelems) + { mat_reduce_over_axis0 = (simplified_iteration_dst_strides[0] == 1) && (simplified_iteration_src_strides[0] == 1); @@ -1257,7 +1258,8 @@ py_boolean_reduction(const dpctl::tensor::usm_ndarray &src, red_nelems); } else if (static_cast(simplified_red_src_strides[0]) == - iter_nelems) { + iter_nelems) + { mat_reduce_over_axis0 = (simplified_iter_dst_strides[0] == 1) && (simplified_iter_src_strides[0] == 1); } diff --git a/dpctl/tensor/libtensor/source/reductions/sum.cpp b/dpctl/tensor/libtensor/source/reductions/sum.cpp index 6184f400fa..bb24da9287 100644 --- a/dpctl/tensor/libtensor/source/reductions/sum.cpp +++ b/dpctl/tensor/libtensor/source/reductions/sum.cpp @@ -257,7 +257,8 @@ struct SumOverAxisTempsStridedFactory fnT get() const { if constexpr (TypePairSupportDataForSumReductionTemps< - srcTy, dstTy>::is_defined) { + srcTy, dstTy>::is_defined) + { using ReductionOpT = sycl::plus; return dpctl::tensor::kernels:: reduction_over_group_temps_strided_impl::is_defined) { + srcTy, dstTy>::is_defined) + { using ReductionOpT = sycl::plus; return dpctl::tensor::kernels:: reduction_axis1_over_group_temps_contig_impl::is_defined) { + srcTy, dstTy>::is_defined) + { using ReductionOpT = sycl::plus; return dpctl::tensor::kernels:: reduction_axis0_over_group_temps_contig_impl src_nd && src_nd > 0) || - (axis > 0 && src_nd == 0)) { + (axis > 0 && src_nd == 0)) + { throw py::value_error("Specified axis is invalid."); } @@ -517,7 +518,8 @@ py_repeat_by_scalar(const dpctl::tensor::usm_ndarray &src, { int src_nd = src.get_ndim(); if (axis < 0 || (axis + 1 > src_nd && src_nd > 0) || - (axis > 0 && src_nd == 0)) { + (axis > 0 && src_nd == 0)) + { throw py::value_error("Specified axis is invalid."); } diff --git a/dpctl/tensor/libtensor/source/simplify_iteration_space.cpp b/dpctl/tensor/libtensor/source/simplify_iteration_space.cpp index 1230cf1d0a..1371770937 100644 --- a/dpctl/tensor/libtensor/source/simplify_iteration_space.cpp +++ b/dpctl/tensor/libtensor/source/simplify_iteration_space.cpp @@ -233,7 +233,8 @@ void simplify_iteration_space_3( simplified_dst_strides.reserve(nd); if ((src1_strides[0] < 0) && (src2_strides[0] < 0) && - (dst_strides[0] < 0)) { + (dst_strides[0] < 0)) + { simplified_src1_strides.push_back(-src1_strides[0]); simplified_src2_strides.push_back(-src2_strides[0]); simplified_dst_strides.push_back(-dst_strides[0]); diff --git a/dpctl/tensor/libtensor/source/sorting/rich_comparisons.hpp b/dpctl/tensor/libtensor/source/sorting/rich_comparisons.hpp index 8e94a271ae..e770f6582b 100644 --- a/dpctl/tensor/libtensor/source/sorting/rich_comparisons.hpp +++ b/dpctl/tensor/libtensor/source/sorting/rich_comparisons.hpp @@ -99,9 +99,9 @@ template struct ExtendedComplexFPGreater }; template -inline constexpr bool is_fp_v = (std::is_same_v || - std::is_same_v || - std::is_same_v); +inline constexpr bool is_fp_v = + (std::is_same_v || std::is_same_v || + std::is_same_v); } // end of anonymous namespace diff --git a/dpctl/utils/src/sequential_order_keeper.hpp b/dpctl/utils/src/sequential_order_keeper.hpp index 0acc4a8b87..db58b99510 100644 --- a/dpctl/utils/src/sequential_order_keeper.hpp +++ b/dpctl/utils/src/sequential_order_keeper.hpp @@ -76,10 +76,7 @@ class SequentialOrder return *this; } - size_t get_num_submitted_events() const - { - return submitted_events.size(); - } + size_t get_num_submitted_events() const { return submitted_events.size(); } const std::vector &get_host_task_events() { @@ -93,10 +90,7 @@ class SequentialOrder } */ - size_t get_num_host_task_events() const - { - return host_task_events.size(); - } + size_t get_num_host_task_events() const { return host_task_events.size(); } const std::vector &get_submitted_events() { diff --git a/examples/cython/usm_memory/src/sycl_blackscholes.hpp b/examples/cython/usm_memory/src/sycl_blackscholes.hpp index 376b0f54c5..267549255e 100644 --- a/examples/cython/usm_memory/src/sycl_blackscholes.hpp +++ b/examples/cython/usm_memory/src/sycl_blackscholes.hpp @@ -57,54 +57,53 @@ void cpp_blackscholes(sycl::queue &q, size_t n_opts, T *params, T *callput) data_t half = one / two; cgh.parallel_for>( - sycl::range<1>(n_opts), [=](sycl::id<1> idx) - { - const size_t i = n_params * idx[0]; - const data_t opt_price = params[i + PRICE]; - const data_t opt_strike = params[i + STRIKE]; - const data_t opt_maturity = params[i + MATURITY]; - const data_t opt_rate = params[i + RATE]; - const data_t opt_volatility = params[i + VOLATILITY]; - data_t a, b, c, y, z, e, d1, d1c, d2, d2c, w1, w2; - data_t mr = -opt_rate, - sig_sig_two = two * opt_volatility * opt_volatility; - - a = sycl::log(opt_price / opt_strike); - b = opt_maturity * mr; - z = opt_maturity * sig_sig_two; - - c = quarter * z; - e = sycl::exp(b); - y = sycl::rsqrt(z); - - a = b - a; - w1 = (a - c) * y; - w2 = (a + c) * y; - - if (w1 < zero) { - d1 = sycl::erfc(w1) * half; - d1c = one - d1; - } - else { - d1c = sycl::erfc(-w1) * half; - d1 = one - d1c; - } - if (w2 < zero) { - d2 = sycl::erfc(w2) * half; - d2c = one - d2; - } - else { - d2c = sycl::erfc(-w2) * half; - d2 = one - d2c; - } - - e *= opt_strike; - data_t call_price = opt_price * d1 - e * d2; - data_t put_price = e * d2c - opt_price * d1c; - - const size_t callput_i = n_prices * idx[0]; - callput[callput_i + CALL] = call_price; - callput[callput_i + PUT] = put_price; + sycl::range<1>(n_opts), [=](sycl::id<1> idx) { + const size_t i = n_params * idx[0]; + const data_t opt_price = params[i + PRICE]; + const data_t opt_strike = params[i + STRIKE]; + const data_t opt_maturity = params[i + MATURITY]; + const data_t opt_rate = params[i + RATE]; + const data_t opt_volatility = params[i + VOLATILITY]; + data_t a, b, c, y, z, e, d1, d1c, d2, d2c, w1, w2; + data_t mr = -opt_rate, + sig_sig_two = two * opt_volatility * opt_volatility; + + a = sycl::log(opt_price / opt_strike); + b = opt_maturity * mr; + z = opt_maturity * sig_sig_two; + + c = quarter * z; + e = sycl::exp(b); + y = sycl::rsqrt(z); + + a = b - a; + w1 = (a - c) * y; + w2 = (a + c) * y; + + if (w1 < zero) { + d1 = sycl::erfc(w1) * half; + d1c = one - d1; + } + else { + d1c = sycl::erfc(-w1) * half; + d1 = one - d1c; + } + if (w2 < zero) { + d2 = sycl::erfc(w2) * half; + d2c = one - d2; + } + else { + d2c = sycl::erfc(-w2) * half; + d2 = one - d2c; + } + + e *= opt_strike; + data_t call_price = opt_price * d1 - e * d2; + data_t put_price = e * d2c - opt_price * d1c; + + const size_t callput_i = n_prices * idx[0]; + callput[callput_i + CALL] = call_price; + callput[callput_i + PUT] = put_price; }); }); diff --git a/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp b/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp index 2625cd4ba4..84f3990f89 100644 --- a/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp +++ b/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp @@ -52,27 +52,12 @@ struct DMatrix DMatrix(const DMatrix &) = default; DMatrix(DMatrix &&) = default; - size_t get_n() const - { - return n_; - } - size_t get_m() const - { - return m_; - } - vec_t &get_vector() - { - return vec_; - } - sycl::queue get_queue() const - { - return q_; - } + size_t get_n() const { return n_; } + size_t get_m() const { return m_; } + vec_t &get_vector() { return vec_; } + sycl::queue get_queue() const { return q_; } - double get_element(size_t i, size_t j) - { - return vec_.at(i * m_ + j); - } + double get_element(size_t i, size_t j) { return vec_.at(i * m_ + j); } private: size_t n_; diff --git a/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp b/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp index 1ebf5e7bff..1e203fe04d 100644 --- a/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp +++ b/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp @@ -48,7 +48,8 @@ py_gemv(sycl::queue &q, const std::vector &depends = {}) { if (matrix.get_ndim() != 2 || vector.get_ndim() != 1 || - result.get_ndim() != 1) { + result.get_ndim() != 1) + { throw std::runtime_error( "Inconsistent dimensions, expecting matrix and a vector"); } diff --git a/libsyclinterface/source/dpctl_service.cpp b/libsyclinterface/source/dpctl_service.cpp index 4129edf959..f74cdaa02a 100644 --- a/libsyclinterface/source/dpctl_service.cpp +++ b/libsyclinterface/source/dpctl_service.cpp @@ -65,15 +65,12 @@ void DPCTLService_InitLogger(const char *app_name, const char *log_dir) } } -void DPCTLService_ShutdownLogger(void) -{ - google::ShutdownGoogleLogging(); -} +void DPCTLService_ShutdownLogger(void) { google::ShutdownGoogleLogging(); } #else void DPCTLService_InitLogger([[maybe_unused]] const char *app_name, - [[maybe_unused]] const char *log_dir){}; + [[maybe_unused]] const char *log_dir) {}; -void DPCTLService_ShutdownLogger(void){}; +void DPCTLService_ShutdownLogger(void) {}; #endif diff --git a/libsyclinterface/source/dpctl_utils.cpp b/libsyclinterface/source/dpctl_utils.cpp index 84ab5d47db..5b6d1e8356 100644 --- a/libsyclinterface/source/dpctl_utils.cpp +++ b/libsyclinterface/source/dpctl_utils.cpp @@ -25,12 +25,6 @@ #include "dpctl_utils.h" -void DPCTLCString_Delete(__dpctl_take const char *str) -{ - delete[] str; -} +void DPCTLCString_Delete(__dpctl_take const char *str) { delete[] str; } -void DPCTLSize_t_Array_Delete(__dpctl_take size_t *arr) -{ - delete[] arr; -} +void DPCTLSize_t_Array_Delete(__dpctl_take size_t *arr) { delete[] arr; } diff --git a/libsyclinterface/tests/test_sycl_device_interface.cpp b/libsyclinterface/tests/test_sycl_device_interface.cpp index 3569dd7ad8..fa2cf5e01c 100644 --- a/libsyclinterface/tests/test_sycl_device_interface.cpp +++ b/libsyclinterface/tests/test_sycl_device_interface.cpp @@ -57,10 +57,7 @@ struct TestDPCTLSyclDeviceInterface } } - ~TestDPCTLSyclDeviceInterface() - { - DPCTLDevice_Delete(DRef); - } + ~TestDPCTLSyclDeviceInterface() { DPCTLDevice_Delete(DRef); } }; TEST_P(TestDPCTLSyclDeviceInterface, ChkCopy) diff --git a/libsyclinterface/tests/test_sycl_device_subdevices.cpp b/libsyclinterface/tests/test_sycl_device_subdevices.cpp index 011dd57119..8c1d9fad6f 100644 --- a/libsyclinterface/tests/test_sycl_device_subdevices.cpp +++ b/libsyclinterface/tests/test_sycl_device_subdevices.cpp @@ -64,10 +64,7 @@ struct TestDPCTLSyclDeviceInterface } } - ~TestDPCTLSyclDeviceInterface() - { - DPCTLDevice_Delete(DRef); - } + ~TestDPCTLSyclDeviceInterface() { DPCTLDevice_Delete(DRef); } }; TEST_P(TestDPCTLSyclDeviceInterface, ChkCreateSubDevicesEqually) diff --git a/libsyclinterface/tests/test_sycl_event_interface.cpp b/libsyclinterface/tests/test_sycl_event_interface.cpp index b92473236d..b891b899c0 100644 --- a/libsyclinterface/tests/test_sycl_event_interface.cpp +++ b/libsyclinterface/tests/test_sycl_event_interface.cpp @@ -65,10 +65,7 @@ struct TestDPCTLSyclEventInterface : public ::testing::Test EXPECT_NO_FATAL_FAILURE(ERef = DPCTLEvent_Create()); } - void SetUp() - { - ASSERT_TRUE(ERef); - } + void SetUp() { ASSERT_TRUE(ERef); } ~TestDPCTLSyclEventInterface() { diff --git a/libsyclinterface/tests/test_sycl_platform_interface.cpp b/libsyclinterface/tests/test_sycl_platform_interface.cpp index 9b62184297..609ba79103 100644 --- a/libsyclinterface/tests/test_sycl_platform_interface.cpp +++ b/libsyclinterface/tests/test_sycl_platform_interface.cpp @@ -192,10 +192,7 @@ struct TestDPCTLSyclDefaultPlatform : public ::testing::Test EXPECT_NO_FATAL_FAILURE(PRef = DPCTLPlatform_Create()); } - void SetUp() - { - ASSERT_TRUE(PRef); - } + void SetUp() { ASSERT_TRUE(PRef); } ~TestDPCTLSyclDefaultPlatform() { @@ -283,10 +280,7 @@ TEST_P(TestDPCTLSyclPlatformInterface, ChkAreEqNullArg) ASSERT_TRUE(DPCTLPlatform_Hash(Null_PRef) == 0); } -TEST_F(TestDPCTLSyclDefaultPlatform, ChkGetName) -{ - check_platform_name(PRef); -} +TEST_F(TestDPCTLSyclDefaultPlatform, ChkGetName) { check_platform_name(PRef); } TEST_F(TestDPCTLSyclDefaultPlatform, ChkGetVendor) { diff --git a/libsyclinterface/tests/test_sycl_queue_interface.cpp b/libsyclinterface/tests/test_sycl_queue_interface.cpp index 75e083e471..2ff4810c1b 100644 --- a/libsyclinterface/tests/test_sycl_queue_interface.cpp +++ b/libsyclinterface/tests/test_sycl_queue_interface.cpp @@ -42,10 +42,7 @@ using namespace dpctl::syclinterface; namespace { -void error_handler_fn(int /*err*/) -{ - return; -} +void error_handler_fn(int /*err*/) { return; } struct TestDPCTLQueueMemberFunctions : public ::testing::TestWithParam<