DPC++ Runtime
Runtime libraries for oneAPI DPC++
Public Member Functions | List of all members
cl::sycl::ext::oneapi::detail::combiner< Reducer > Class Template Reference

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

#include <sycl/ext/oneapi/reduction.hpp>

Public Member Functions

template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsPlus< _T, BinaryOperation >::value &&sycl::detail::is_geninteger< _T >::value > operator++ ()
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsPlus< _T, BinaryOperation >::value &&sycl::detail::is_geninteger< _T >::value > operator++ (int)
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsPlus< _T, BinaryOperation >::value > operator+= (const _T &Partial)
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsMultiplies< _T, BinaryOperation >::value > operator*= (const _T &Partial)
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsBitOR< _T, BinaryOperation >::value > operator|= (const _T &Partial)
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsBitXOR< _T, BinaryOperation >::value > operator^= (const _T &Partial)
 
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims==0) &&sycl::detail::IsBitAND< _T, BinaryOperation >::value > operator&= (const _T &Partial)
 
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&(IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value||IsReduOptForAtomic64Add< T, _BinaryOperation >::value) &&sycl::detail::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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value &&sycl::detail::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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value &&sycl::detail::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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t< std::is_same< typename remove_AS< _T >::type, T >::value &&IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value &&sycl::detail::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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value &&sycl::detail::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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t< BasicCheck< _T, Space, _BinaryOperation > &&IsReduOptForFastAtomicFetch< T, _BinaryOperation >::value &&sycl::detail::IsMaximum< T, _BinaryOperation >::value > atomic_combine (_T *ReduVarPtr) const
 Atomic MAX operation: *ReduVarPtr = sycl::maximum(*ReduVarPtr, MValue);. More...
 

Detailed Description

template<class Reducer>
class cl::sycl::ext::oneapi::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 153 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 = T, class _BinaryOperation = BinaryOperation>
enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && (IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value || IsReduOptForAtomic64Add<T, _BinaryOperation>::value) && sycl::detail::IsPlus<T, _BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

Atomic ADD operation: *ReduVarPtr += MValue;.

Definition at line 242 of file reduction.hpp.

◆ atomic_combine() [2/6]

template<class Reducer >
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value && sycl::detail::IsBitOR<T, _BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

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

Definition at line 253 of file reduction.hpp.

◆ atomic_combine() [3/6]

template<class Reducer >
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value && sycl::detail::IsBitXOR<T, _BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

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

Definition at line 264 of file reduction.hpp.

◆ atomic_combine() [4/6]

template<class Reducer >
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value && IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value && sycl::detail::IsBitAND<T, _BinaryOperation>::value && (Space == access::address_space::global_space || Space == access::address_space::local_space)> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

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

Definition at line 277 of file reduction.hpp.

◆ atomic_combine() [5/6]

template<class Reducer >
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value && sycl::detail::IsMinimum<T, _BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

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

Definition at line 288 of file reduction.hpp.

◆ atomic_combine() [6/6]

template<class Reducer >
template<access::address_space Space = access::address_space::global_space, typename _T = T, class _BinaryOperation = BinaryOperation>
enable_if_t<BasicCheck<_T, Space, _BinaryOperation> && IsReduOptForFastAtomicFetch<T, _BinaryOperation>::value && sycl::detail::IsMaximum<T, _BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::atomic_combine ( _T *  ReduVarPtr) const
inline

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

Definition at line 299 of file reduction.hpp.

◆ operator&=()

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsBitAND<_T, BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator&= ( const _T &  Partial)
inline

Definition at line 205 of file reduction.hpp.

◆ operator*=()

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsMultiplies<_T, BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator*= ( const _T &  Partial)
inline

Definition at line 185 of file reduction.hpp.

◆ operator++() [1/2]

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsPlus<_T, BinaryOperation>::value && sycl::detail::is_geninteger<_T>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator++ ( )
inline

Definition at line 164 of file reduction.hpp.

◆ operator++() [2/2]

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsPlus<_T, BinaryOperation>::value && sycl::detail::is_geninteger<_T>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator++ ( int  )
inline

Definition at line 172 of file reduction.hpp.

◆ operator+=()

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsPlus<_T, BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator+= ( const _T &  Partial)
inline

Definition at line 178 of file reduction.hpp.

◆ operator^=()

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsBitXOR<_T, BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator^= ( const _T &  Partial)
inline

Definition at line 198 of file reduction.hpp.

◆ operator|=()

template<class Reducer >
template<typename _T = T, int _Dims = Dims>
enable_if_t<(_Dims == 0) && sycl::detail::IsBitOR<_T, BinaryOperation>::value> cl::sycl::ext::oneapi::detail::combiner< Reducer >::operator|= ( const _T &  Partial)
inline

Definition at line 191 of file reduction.hpp.

The documentation for this class was generated from the following file: