Skip to content

Introduce dedicated type support matrices for in-place element-wise operations #1816

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Aug 30, 2024
Merged

Introduce dedicated type support matrices for in-place element-wise operations #1816

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
1b26f0e
Refactor in-place division to use `TypePairDefinedEntry`
ndgrigorian Aug 26, 2024
1755f66
Decouples in-place and out-of-place type support tables
ndgrigorian Aug 26, 2024
ab068f5
Fixes a comment in `_acceptance_fn_reciprocal`
ndgrigorian Aug 26, 2024
0aa5321
Remove unnecessary comments about `std::disjunction` throughout project
ndgrigorian Aug 29, 2024
2ffedad
Add `is_defined` static boolean to output type tables and uses them i…
ndgrigorian Aug 29, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 3 additions & 5 deletions dpctl/tensor/_type_utils.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -140,11 +140,9 @@ def _acceptance_fn_default_unary(arg_dtype, ret_buf_dt, res_dt, sycl_dev):


def _acceptance_fn_reciprocal(arg_dtype, buf_dt, res_dt, sycl_dev):
# if the kind of result is different from
# the kind of input, use the default data
# we use default dtype for the resulting kind.
# This guarantees alignment of reciprocal and
# divide output types.
# if the kind of result is different from the kind of input, we use the
# default floating-point dtype for the resulting kind. This guarantees
# alignment of reciprocal and divide output types.
if buf_dt.kind != arg_dtype.kind:
default_dt = _get_device_default_dtype(res_dt.kind, sycl_dev)
if res_dt == default_dt:
Expand Down
13 changes: 5 additions & 8 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -102,8 +102,7 @@ using AbsContigFunctor =

template <typename T> struct AbsOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::TypeMapResultEntry<T, bool>,
td_ns::TypeMapResultEntry<T, std::uint8_t>,
td_ns::TypeMapResultEntry<T, std::uint16_t>,
Expand All @@ -119,6 +118,8 @@ template <typename T> struct AbsOutputType
td_ns::TypeMapResultEntry<T, std::complex<float>, float>,
td_ns::TypeMapResultEntry<T, std::complex<double>, double>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
Expand All @@ -140,9 +141,7 @@ template <typename fnT, typename T> struct AbsContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AbsOutputType<T>::value_type,
void>)
{
if constexpr (!AbsOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -191,9 +190,7 @@ template <typename fnT, typename T> struct AbsStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AbsOutputType<T>::value_type,
void>)
{
if constexpr (!AbsOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
13 changes: 5 additions & 8 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -145,14 +145,15 @@ using AcosStridedFunctor = elementwise_common::

template <typename T> struct AcosOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::TypeMapResultEntry<T, sycl::half>,
td_ns::TypeMapResultEntry<T, float>,
td_ns::TypeMapResultEntry<T, double>,
td_ns::TypeMapResultEntry<T, std::complex<float>>,
td_ns::TypeMapResultEntry<T, std::complex<double>>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
Expand All @@ -174,9 +175,7 @@ template <typename fnT, typename T> struct AcosContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AcosOutputType<T>::value_type,
void>)
{
if constexpr (!AcosOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -222,9 +221,7 @@ template <typename fnT, typename T> struct AcosStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AcosOutputType<T>::value_type,
void>)
{
if constexpr (!AcosOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
13 changes: 5 additions & 8 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -172,14 +172,15 @@ using AcoshStridedFunctor = elementwise_common::

template <typename T> struct AcoshOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::TypeMapResultEntry<T, sycl::half>,
td_ns::TypeMapResultEntry<T, float>,
td_ns::TypeMapResultEntry<T, double>,
td_ns::TypeMapResultEntry<T, std::complex<float>>,
td_ns::TypeMapResultEntry<T, std::complex<double>>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
Expand All @@ -201,9 +202,7 @@ template <typename fnT, typename T> struct AcoshContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AcoshOutputType<T>::value_type,
void>)
{
if constexpr (!AcoshOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -249,9 +248,7 @@ template <typename fnT, typename T> struct AcoshStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AcoshOutputType<T>::value_type,
void>)
{
if constexpr (!AcoshOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
76 changes: 56 additions & 20 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,8 +132,7 @@ using AddStridedFunctor =

template <typename T1, typename T2> struct AddOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::BinaryTypeMapResultEntry<T1, bool, T2, bool, bool>,
td_ns::BinaryTypeMapResultEntry<T1,
std::uint8_t,
Expand Down Expand Up @@ -193,6 +192,8 @@ template <typename T1, typename T2> struct AddOutputType
std::complex<double>,
std::complex<double>>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename argT1,
Expand Down Expand Up @@ -223,9 +224,7 @@ template <typename fnT, typename T1, typename T2> struct AddContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
void>)
{
if constexpr (!AddOutputType<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -273,9 +272,7 @@ template <typename fnT, typename T1, typename T2> struct AddStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
void>)
{
if constexpr (!AddOutputType<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -324,12 +321,12 @@ struct AddContigMatrixContigRowBroadcastFactory
{
fnT get()
{
using resT = typename AddOutputType<T1, T2>::value_type;
if constexpr (std::is_same_v<resT, void>) {
if constexpr (!AddOutputType<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
else {
using resT = typename AddOutputType<T1, T2>::value_type;
if constexpr (dpctl::tensor::type_utils::is_complex<T1>::value ||
dpctl::tensor::type_utils::is_complex<T2>::value ||
dpctl::tensor::type_utils::is_complex<resT>::value)
Expand Down Expand Up @@ -371,12 +368,12 @@ struct AddContigRowContigMatrixBroadcastFactory
{
fnT get()
{
using resT = typename AddOutputType<T1, T2>::value_type;
if constexpr (std::is_same_v<resT, void>) {
if constexpr (!AddOutputType<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
else {
using resT = typename AddOutputType<T1, T2>::value_type;
if constexpr (dpctl::tensor::type_utils::is_complex<T1>::value ||
dpctl::tensor::type_utils::is_complex<T2>::value ||
dpctl::tensor::type_utils::is_complex<resT>::value)
Expand Down Expand Up @@ -438,6 +435,50 @@ template <typename argT,
unsigned int n_vecs>
class add_inplace_contig_kernel;

/* @brief Types supported by in-place add */
template <typename argTy, typename resTy> struct AddInplaceTypePairSupport
{
/* value if true a kernel for <argTy, resTy> must be instantiated */
static constexpr bool is_defined = std::disjunction<
td_ns::TypePairDefinedEntry<argTy, bool, resTy, bool>,
td_ns::TypePairDefinedEntry<argTy, std::int8_t, resTy, std::int8_t>,
td_ns::TypePairDefinedEntry<argTy, std::uint8_t, resTy, std::uint8_t>,
td_ns::TypePairDefinedEntry<argTy, std::int16_t, resTy, std::int16_t>,
td_ns::TypePairDefinedEntry<argTy, std::uint16_t, resTy, std::uint16_t>,
td_ns::TypePairDefinedEntry<argTy, std::int32_t, resTy, std::int32_t>,
td_ns::TypePairDefinedEntry<argTy, std::uint32_t, resTy, std::uint32_t>,
td_ns::TypePairDefinedEntry<argTy, std::int64_t, resTy, std::int64_t>,
td_ns::TypePairDefinedEntry<argTy, std::uint64_t, resTy, std::uint64_t>,
td_ns::TypePairDefinedEntry<argTy, sycl::half, resTy, sycl::half>,
td_ns::TypePairDefinedEntry<argTy, float, resTy, float>,
td_ns::TypePairDefinedEntry<argTy, double, resTy, double>,
td_ns::TypePairDefinedEntry<argTy,
std::complex<float>,
resTy,
std::complex<float>>,
td_ns::TypePairDefinedEntry<argTy,
std::complex<double>,
resTy,
std::complex<double>>,
// fall-through
td_ns::NotDefinedEntry>::is_defined;
};

template <typename fnT, typename argT, typename resT>
struct AddInplaceTypeMapFactory
{
/*! @brief get typeid for output type of x += y */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
if constexpr (AddInplaceTypePairSupport<argT, resT>::is_defined) {
return td_ns::GetTypeid<resT>{}.get();
}
else {
return td_ns::GetTypeid<void>{}.get();
}
}
};

template <typename argTy, typename resTy>
sycl::event
add_inplace_contig_impl(sycl::queue &exec_q,
Expand All @@ -457,9 +498,7 @@ template <typename fnT, typename T1, typename T2> struct AddInplaceContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
void>)
{
if constexpr (!AddInplaceTypePairSupport<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -497,9 +536,7 @@ struct AddInplaceStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
void>)
{
if constexpr (!AddInplaceTypePairSupport<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -544,8 +581,7 @@ struct AddInplaceRowMatrixBroadcastFactory
{
fnT get()
{
using resT = typename AddOutputType<T1, T2>::value_type;
if constexpr (!std::is_same_v<resT, T2>) {
if constexpr (!AddInplaceTypePairSupport<T1, T2>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
13 changes: 5 additions & 8 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,11 +91,12 @@ using AngleStridedFunctor = elementwise_common::

template <typename T> struct AngleOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::TypeMapResultEntry<T, std::complex<float>, float>,
td_ns::TypeMapResultEntry<T, std::complex<double>, double>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
Expand All @@ -117,9 +118,7 @@ template <typename fnT, typename T> struct AngleContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AngleOutputType<T>::value_type,
void>)
{
if constexpr (!AngleOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -165,9 +164,7 @@ template <typename fnT, typename T> struct AngleStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AngleOutputType<T>::value_type,
void>)
{
if constexpr (!AngleOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
13 changes: 5 additions & 8 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -165,14 +165,15 @@ using AsinStridedFunctor = elementwise_common::

template <typename T> struct AsinOutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
using value_type = typename std::disjunction<
td_ns::TypeMapResultEntry<T, sycl::half>,
td_ns::TypeMapResultEntry<T, float>,
td_ns::TypeMapResultEntry<T, double>,
td_ns::TypeMapResultEntry<T, std::complex<float>>,
td_ns::TypeMapResultEntry<T, std::complex<double>>,
td_ns::DefaultResultEntry<void>>::result_type;

static constexpr bool is_defined = !std::is_same_v<value_type, void>;
};

template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
Expand All @@ -194,9 +195,7 @@ template <typename fnT, typename T> struct AsinContigFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AsinOutputType<T>::value_type,
void>)
{
if constexpr (!AsinOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down Expand Up @@ -242,9 +241,7 @@ template <typename fnT, typename T> struct AsinStridedFactory
{
fnT get()
{
if constexpr (std::is_same_v<typename AsinOutputType<T>::value_type,
void>)
{
if constexpr (!AsinOutputType<T>::is_defined) {
fnT fn = nullptr;
return fn;
}
Expand Down
Loading
Loading