Skip to content

[SYCL] Use macro SYCL_REDUCTION_DETERMINISTIC for stable reduction re… #3876

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 3 commits into from
Jun 7, 2021
Merged

[SYCL] Use macro SYCL_REDUCTION_DETERMINISTIC for stable reduction re… #3876

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
8b9f826
[SYCL] Use macro SYCL_REDUCTION_DETERMINISTIC for stable reduction re…
v-klochkov Jun 3, 2021
ead7383
disable reduce() algorithm if macro SYCL_REDUCTION_DETERMINISTIC is d…
v-klochkov Jun 4, 2021
1025dd3
Merge remote-tracking branch 'intel_llvm/sycl' into public_reduction_…
v-klochkov Jun 4, 2021
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
87 changes: 61 additions & 26 deletions sycl/include/CL/sycl/ONEAPI/reduction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,11 +27,50 @@ namespace detail {

using cl::sycl::detail::bool_constant;
using cl::sycl::detail::enable_if_t;
using cl::sycl::detail::is_sgenfloat;
using cl::sycl::detail::is_sgeninteger;
using cl::sycl::detail::queue_impl;
using cl::sycl::detail::remove_AS;

// This type trait is used to detect if the atomic operation BinaryOperation
// used with operands of the type T is available for using in reduction.
// The order in which the atomic operations are performed may be arbitrary and
// thus may cause different results from run to run even on the same elements
// and on same device. The macro SYCL_REDUCTION_DETERMINISTIC prohibits using
// atomic operations for reduction and helps to produce stable results.
// SYCL_REDUCTION_DETERMINISTIC is a short term solution, which perhaps become
// deprecated eventually and is replaced by a sycl property passed to reduction.
template <typename T, class BinaryOperation>
using IsReduOptForFastAtomicFetch =
#ifdef SYCL_REDUCTION_DETERMINISTIC
bool_constant<false>;
#else
Comment on lines +43 to +45
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this sufficient to guarantee determinism? I don't think the work-group/sub-group reduction functions are guaranteed to be deterministic across all devices. Since there's only one implementation that's guaranteed to be deterministic (i.e. the manual tree reduction), why not guard all but that implementation behind ifndef SYCL_REDUCTION_DETERMINISTIC?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Right. I disabled reduce() algorithm when the macro is defined: ead7383

bool_constant<sycl::detail::is_sgeninteger<T>::value &&
sycl::detail::IsValidAtomicType<T>::value &&
(sycl::detail::IsPlus<T, BinaryOperation>::value ||
sycl::detail::IsMinimum<T, BinaryOperation>::value ||
sycl::detail::IsMaximum<T, BinaryOperation>::value ||
sycl::detail::IsBitOR<T, BinaryOperation>::value ||
sycl::detail::IsBitXOR<T, BinaryOperation>::value ||
sycl::detail::IsBitAND<T, BinaryOperation>::value)>;
#endif

// This type trait is used to detect if the group algorithm reduce() used with
// operands of the type T and the operation BinaryOperation is available
// for using in reduction.
// The macro SYCL_REDUCTION_DETERMINISTIC prohibits using the reduce() algorithm
// to produce stable results across same type devices.
template <typename T, class BinaryOperation>
using IsReduOptForFastReduce =
#ifdef SYCL_REDUCTION_DETERMINISTIC
bool_constant<false>;
#else
bool_constant<((sycl::detail::is_sgeninteger<T>::value &&
(sizeof(T) == 4 || sizeof(T) == 8)) ||
sycl::detail::is_sgenfloat<T>::value) &&
(sycl::detail::IsPlus<T, BinaryOperation>::value ||
sycl::detail::IsMinimum<T, BinaryOperation>::value ||
sycl::detail::IsMaximum<T, BinaryOperation>::value)>;
#endif

// std::tuple seems to be a) too heavy and b) not copyable to device now
// Thus sycl::detail::tuple is used instead.
// Switching from sycl::device::tuple to std::tuple can be done by re-defining
Expand All @@ -46,10 +85,6 @@ __SYCL_EXPORT size_t reduGetMaxWGSize(shared_ptr_class<queue_impl> Queue,
__SYCL_EXPORT size_t reduComputeWGSize(size_t NWorkItems, size_t MaxWGSize,
size_t &NWorkGroups);





/// Class that is used to represent objects that are passed to user's lambda
/// functions and representing users' reduction variable.
/// The generic version of the class represents those reductions of those
Expand All @@ -64,45 +99,45 @@ class reducer {
T getIdentity() const { return MIdentity; }

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value &&
sycl::detail::is_geninteger<_T>::value>
operator++() {
combine(static_cast<T>(1));
}

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value &&
sycl::detail::is_geninteger<_T>::value>
operator++(int) {
combine(static_cast<T>(1));
}

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value>
operator+=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduMultiplies<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsMultiplies<_T, BinaryOperation>::value>
operator*=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitOR<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitOR<_T, BinaryOperation>::value>
operator|=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitXOR<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitXOR<_T, BinaryOperation>::value>
operator^=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitAND<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitAND<_T, BinaryOperation>::value>
operator&=(const _T &Partial) {
combine(Partial);
}
Expand Down Expand Up @@ -150,45 +185,45 @@ class reducer<T, BinaryOperation,
}

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value &&
sycl::detail::is_geninteger<_T>::value>
operator++() {
combine(static_cast<T>(1));
}

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value &&
sycl::detail::is_geninteger<_T>::value>
operator++(int) {
combine(static_cast<T>(1));
}

template <typename _T = T>
enable_if_t<IsReduPlus<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsPlus<_T, BinaryOperation>::value>
operator+=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduMultiplies<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsMultiplies<_T, BinaryOperation>::value>
operator*=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitOR<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitOR<_T, BinaryOperation>::value>
operator|=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitXOR<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitXOR<_T, BinaryOperation>::value>
operator^=(const _T &Partial) {
combine(Partial);
}

template <typename _T = T>
enable_if_t<IsReduBitAND<_T, BinaryOperation>::value>
enable_if_t<sycl::detail::IsBitAND<_T, BinaryOperation>::value>
operator&=(const _T &Partial) {
combine(Partial);
}
Expand All @@ -197,7 +232,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduPlus<T, _BinaryOperation>::value>
sycl::detail::IsPlus<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_add(MValue);
Expand All @@ -207,7 +242,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduBitOR<T, _BinaryOperation>::value>
sycl::detail::IsBitOR<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_or(MValue);
Expand All @@ -217,7 +252,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduBitXOR<T, _BinaryOperation>::value>
sycl::detail::IsBitXOR<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_xor(MValue);
Expand All @@ -227,7 +262,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduBitAND<T, _BinaryOperation>::value>
sycl::detail::IsBitAND<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_and(MValue);
Expand All @@ -237,7 +272,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduMinimum<T, _BinaryOperation>::value>
sycl::detail::IsMinimum<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_min(MValue);
Expand All @@ -247,7 +282,7 @@ class reducer<T, BinaryOperation,
template <typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value &&
IsReduMaximum<T, _BinaryOperation>::value>
sycl::detail::IsMaximum<T, _BinaryOperation>::value>
atomic_combine(_T *ReduVarPtr) const {
atomic<T, access::address_space::global_space>(global_ptr<T>(ReduVarPtr))
.fetch_max(MValue);
Expand Down
56 changes: 17 additions & 39 deletions sycl/include/CL/sycl/known_identity.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,93 +17,71 @@ __SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace detail {

using cl::sycl::detail::is_sgeninteger;

template <typename T, class BinaryOperation>
using IsReduPlus =
using IsPlus =
bool_constant<std::is_same<BinaryOperation, ONEAPI::plus<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::plus<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduMultiplies =
bool_constant<std::is_same<BinaryOperation, std::multiplies<T>>::value ||
std::is_same<BinaryOperation, std::multiplies<void>>::value>;
using IsMultiplies = bool_constant<
std::is_same<BinaryOperation, ONEAPI::multiplies<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::multiplies<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduMinimum =
using IsMinimum =
bool_constant<std::is_same<BinaryOperation, ONEAPI::minimum<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::minimum<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduMaximum =
using IsMaximum =
bool_constant<std::is_same<BinaryOperation, ONEAPI::maximum<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::maximum<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduBitOR =
using IsBitOR =
bool_constant<std::is_same<BinaryOperation, ONEAPI::bit_or<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::bit_or<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduBitXOR =
using IsBitXOR =
bool_constant<std::is_same<BinaryOperation, ONEAPI::bit_xor<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::bit_xor<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduBitAND =
using IsBitAND =
bool_constant<std::is_same<BinaryOperation, ONEAPI::bit_and<T>>::value ||
std::is_same<BinaryOperation, ONEAPI::bit_and<void>>::value>;

template <typename T, class BinaryOperation>
using IsReduOptForFastAtomicFetch =
bool_constant<is_sgeninteger<T>::value &&
sycl::detail::IsValidAtomicType<T>::value &&
(IsReduPlus<T, BinaryOperation>::value ||
IsReduMinimum<T, BinaryOperation>::value ||
IsReduMaximum<T, BinaryOperation>::value ||
IsReduBitOR<T, BinaryOperation>::value ||
IsReduBitXOR<T, BinaryOperation>::value ||
IsReduBitAND<T, BinaryOperation>::value)>;

template <typename T, class BinaryOperation>
using IsReduOptForFastReduce =
bool_constant<((is_sgeninteger<T>::value &&
(sizeof(T) == 4 || sizeof(T) == 8)) ||
is_sgenfloat<T>::value) &&
(IsReduPlus<T, BinaryOperation>::value ||
IsReduMinimum<T, BinaryOperation>::value ||
IsReduMaximum<T, BinaryOperation>::value)>;

// Identity = 0
template <typename T, class BinaryOperation>
using IsZeroIdentityOp = bool_constant<
(is_sgeninteger<T>::value && (IsReduPlus<T, BinaryOperation>::value ||
IsReduBitOR<T, BinaryOperation>::value ||
IsReduBitXOR<T, BinaryOperation>::value)) ||
(is_sgenfloat<T>::value && IsReduPlus<T, BinaryOperation>::value)>;
(is_sgeninteger<T>::value &&
(IsPlus<T, BinaryOperation>::value || IsBitOR<T, BinaryOperation>::value ||
IsBitXOR<T, BinaryOperation>::value)) ||
(is_sgenfloat<T>::value && IsPlus<T, BinaryOperation>::value)>;

// Identity = 1
template <typename T, class BinaryOperation>
using IsOneIdentityOp =
bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
IsReduMultiplies<T, BinaryOperation>::value>;
IsMultiplies<T, BinaryOperation>::value>;

// Identity = ~0
template <typename T, class BinaryOperation>
using IsOnesIdentityOp = bool_constant<is_sgeninteger<T>::value &&
IsReduBitAND<T, BinaryOperation>::value>;
IsBitAND<T, BinaryOperation>::value>;

// Identity = <max possible value>
template <typename T, class BinaryOperation>
using IsMinimumIdentityOp =
bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
IsReduMinimum<T, BinaryOperation>::value>;
IsMinimum<T, BinaryOperation>::value>;

// Identity = <min possible value>
template <typename T, class BinaryOperation>
using IsMaximumIdentityOp =
bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
IsReduMaximum<T, BinaryOperation>::value>;
IsMaximum<T, BinaryOperation>::value>;

template <typename T, class BinaryOperation>
using IsKnownIdentityOp =
Expand Down