assembled_value< T,... > Class Template Reference

Class for value assembly using warp-based partial summation to reduce conflict occurences in downstream atomic addition. More...

#include <cs_math_hip.h>

Collaboration diagram for assembled_value< T,... >:

Public Types
using	inner_type = T

Public Member Functions
	assembled_value () noexcept=default
__device__	assembled_value (T val) noexcept
__device__ void	add (const assembled_value &restrict other) restrict noexcept
__device__ void	atomic_add (const assembled_value &restrict other) restrict noexcept
__device__ assembled_value	exchange (const assembled_value &restrict other) restrict noexcept
__device__ assembled_value	atomic_exchange (const assembled_value &restrict other) restrict noexcept
__device__ assembled_value	shuffle (cs_hip_mask_t mask, unsigned laneid) const noexcept
__device__ cs_hip_mask_t	conflict_mask (cs_hip_mask_t mask) const noexcept
__device__ bool	reduce_add (cs_hip_mask_t mask, cs_hip_mask_t peers) noexcept
__device__ void	conflict_free_add (cs_hip_mask_t mask, assembled_value other) noexcept
__device__ inner_type &	operator* () noexcept
__device__ inner_type const &	operator* () const noexcept
__device__ inner_type *	operator-> () noexcept
__device__ inner_type const *	operator-> () const noexcept
__device__ inner_type &	get () noexcept
__device__ inner_type const &	get () const noexcept

Static Public Member Functions
static __device__ assembled_value &	ref (inner_type &r) noexcept
static __device__ assembled_value const &	ref (inner_type const &r) noexcept

Detailed Description

template<class T, size_t...>
class assembled_value< T,... >

Class for value assembly using warp-based partial summation to reduce conflict occurences in downstream atomic addition.

Member Typedef Documentation

inner_type

using inner_type = T

Constructor & Destructor Documentation

assembled_value() [1/2]

assembled_value ( )

defaultnoexcept

assembled_value() [2/2]

__device__ assembled_value ( T  val )

inlinenoexcept

Member Function Documentation

add()

__device__ void add ( const assembled_value< T,... > &restrict  other )

inlinenoexcept

atomic_add()

__device__ void atomic_add ( const assembled_value< T,... > &restrict  other )

inlinenoexcept

atomic_exchange()

__device__ assembled_value atomic_exchange ( const assembled_value< T,... > &restrict  other )

inlinenoexcept

conflict_free_add()

__device__ void conflict_free_add ( cs_hip_mask_t  mask, assembled_value< T,... >  other  )

inlinenoexcept

conflict_mask()

__device__ cs_hip_mask_t conflict_mask ( cs_hip_mask_t  mask ) const

inlinenoexcept

exchange()

__device__ assembled_value exchange ( const assembled_value< T,... > &restrict  other )

inlinenoexcept

get() [1/2]

__device__ inner_type const & get ( ) const

inlinenoexcept

get() [2/2]

__device__ inner_type & get ( )

inlinenoexcept

operator*() [1/2]

__device__ inner_type const & operator* ( ) const

inlinenoexcept

operator*() [2/2]

__device__ inner_type & operator* ( )

inlinenoexcept

operator->() [1/2]

__device__ inner_type const * operator-> ( ) const

inlinenoexcept

operator->() [2/2]

__device__ inner_type * operator-> ( )

inlinenoexcept

reduce_add()

__device__ bool reduce_add ( cs_hip_mask_t  mask, cs_hip_mask_t  peers  )

inlinenoexcept

ref() [1/2]

static __device__ assembled_value & ref ( inner_type &  r )

inlinestaticnoexcept

ref() [2/2]

static __device__ assembled_value const & ref ( inner_type const &  r )

inlinestaticnoexcept

shuffle()

__device__ assembled_value shuffle ( cs_hip_mask_t  mask, unsigned  laneid  ) const

inlinenoexcept

---

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

• src/base/cs_math_hip.h