Use CRTP to avoid redefining shorthand operators in terms of combine. More...

#include <sycl/reduction.hpp>

Public Member Functions
template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsPlus< _T, BinaryOp >::value &&is_geninteger< _T >::value, Reducer & >	operator++ ()

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsPlus< _T, BinaryOp >::value &&is_geninteger< _T >::value, Reducer & >	operator++ (int)

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsPlus< _T, BinaryOp >::value, Reducer & >	operator+= (const _T &Partial)

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsMultiplies< _T, BinaryOp >::value, Reducer & >	operator*= (const _T &Partial)

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsBitOR< _T, BinaryOp >::value, Reducer & >	operator\|= (const _T &Partial)

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsBitXOR< _T, BinaryOp >::value, Reducer & >	operator^= (const _T &Partial)

template<typename _T = Ty, int _Dims = Dims>
enable_if_t<(_Dims==0) &&IsBitAND< _T, BinaryOp >::value, Reducer & >	operator&= (const _T &Partial)

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&(IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value\|\|IsReduOptForAtomic64Op< _T, _BinaryOperation >::value) &&IsPlus< _T, _BinaryOperation >::value >	atomic_combine (_T *ReduVarPtr) const
	Atomic ADD operation: *ReduVarPtr += MValue;. More...

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value &&IsBitOR< _T, _BinaryOperation >::value >	atomic_combine (_T *ReduVarPtr) const
	Atomic BITWISE OR operation: *ReduVarPtr \|= MValue;. More...

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value &&IsBitXOR< _T, _BinaryOperation >::value >	atomic_combine (_T *ReduVarPtr) const
	Atomic BITWISE XOR operation: *ReduVarPtr ^= MValue;. More...

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< std::is_same< remove_decoration_t< _T >, _T >::value &&IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value &&IsBitAND< _T, _BinaryOperation >::value &&(Space==access::address_space::global_space\|\|Space==access::address_space::local_space)>	atomic_combine (_T *ReduVarPtr) const
	Atomic BITWISE AND operation: *ReduVarPtr &= MValue;. More...

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&(IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value\|\|IsReduOptForAtomic64Op< _T, _BinaryOperation >::value) &&IsMinimum< _T, _BinaryOperation >::value >	atomic_combine (_T *ReduVarPtr) const
	Atomic MIN operation: ReduVarPtr = sycl::minimum(ReduVarPtr, MValue);. More...

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&(IsReduOptForFastAtomicFetch< _T, _BinaryOperation >::value\|\|IsReduOptForAtomic64Op< _T, _BinaryOperation >::value) &&IsMaximum< _T, _BinaryOperation >::value >	atomic_combine (_T *ReduVarPtr) const
	Atomic MAX operation: ReduVarPtr = sycl::maximum(ReduVarPtr, MValue);. More...

Detailed Description

template<class Reducer>
class sycl::_V1::detail::combiner< Reducer >

Use CRTP to avoid redefining shorthand operators in terms of combine.

Also, for many types with known identity the operation 'atomic_combine()' is implemented here, which allows to use more efficient version of kernels using those operations, which are based on functionality provided by sycl::atomic class.

For example, it is known that 0 is identity for sycl::plus operations accepting native scalar types to which scalar 0 is convertible. Also, for int32/64 types the atomic_combine() is lowered to sycl::atomic::fetch_add().

Definition at line 213 of file reduction.hpp.

Member Function Documentation

◆ atomic_combine() [1/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && (IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value || IsReduOptForAtomic64Op<_T, _BinaryOperation>::value) && IsPlus<_T, _BinaryOperation>::value> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic ADD operation: *ReduVarPtr += MValue;.

Definition at line 299 of file reduction.hpp.

◆ atomic_combine() [2/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value && IsBitOR<_T, _BinaryOperation>::value> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic BITWISE OR operation: *ReduVarPtr |= MValue;.

Definition at line 310 of file reduction.hpp.

◆ atomic_combine() [3/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value && IsBitXOR<_T, _BinaryOperation>::value> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic BITWISE XOR operation: *ReduVarPtr ^= MValue;.

Definition at line 321 of file reduction.hpp.

◆ atomic_combine() [4/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<std::is_same<remove_decoration_t<_T>, _T>::value && IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value && IsBitAND<_T, _BinaryOperation>::value && (Space == access::address_space::global_space || Space == access::address_space::local_space)> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic BITWISE AND operation: *ReduVarPtr &= MValue;.

Definition at line 334 of file reduction.hpp.

◆ atomic_combine() [5/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && (IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value || IsReduOptForAtomic64Op<_T, _BinaryOperation>::value) && IsMinimum<_T, _BinaryOperation>::value> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic MIN operation: *ReduVarPtr = sycl::minimum(*ReduVarPtr, MValue);.

Definition at line 346 of file reduction.hpp.

◆ atomic_combine() [6/6]

template<class Reducer >

template<access::address_space Space = access::address_space::global_space, typename _T = Ty, class _BinaryOperation = BinaryOp>

enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && (IsReduOptForFastAtomicFetch<_T, _BinaryOperation>::value || IsReduOptForAtomic64Op<_T, _BinaryOperation>::value) && IsMaximum<_T, _BinaryOperation>::value> sycl::_V1::detail::combiner< Reducer >::atomic_combine ( _T * ReduVarPtr ) const

inline

Atomic MAX operation: *ReduVarPtr = sycl::maximum(*ReduVarPtr, MValue);.

Definition at line 358 of file reduction.hpp.