TNL::Backend::Uninitialized< T, maxWordSize > Struct Template Reference

Storage wrapper for type-punning through an array of device words. More...

#include <TNL/Backend/Functions.h>

Public Types
using	DeviceWord

Public Member Functions
__device__ T &	get ()
	Returns a reference to the stored value via reinterpret_cast.
__device__ T	load () const
	Copies the value out of the storage using memcpy.
__device__	operator T& ()
__device__ Uninitialized &	operator= (const T &other)
__device__ void	store (const T &value)
	Copies a value into the storage using memcpy.

Public Attributes
DeviceWord	storage [WORDS]

Static Public Attributes
static constexpr int	WORDS = TNL::roundUpDivision( sizeof( T ), WORD_SIZE )

Detailed Description

template<typename T, std::size_t maxWordSize = 8>
struct TNL::Backend::Uninitialized< T, maxWordSize >

Storage wrapper for type-punning through an array of device words.

The storage is an array of device words whose size is selected as the largest word that evenly divides sizeof(T), does not exceed alignof(T), and does not exceed maxWordSize.

The maxWordSize parameter controls the upper bound on the word size:

• Default (8): uses the largest efficient word for __shared__ memory access (up to 8 bytes for standard integer types).
• 4: restricts to uint32_t words, which is the native word size for CUDA/HIP __shfl_*_sync intrinsics.

Type-punning access

Two access patterns are provided:

• get() returns a T& via reinterpret_cast. This is safe when all accesses to the storage go through the same type (e.g. __shared__ variables in reducing kernels that are always accessed as T).
• load() / store() use memcpy to copy between T and the raw storage array. These must be used when the storage is also accessed through DeviceWord (e.g. in warp shuffle wrappers that shuffle individual words). The reinterpret_cast-based get() would be a strict aliasing violation in that case, which Clang (HIP) exploits to optimize away the connection between the two access paths — the shuffle then silently returns the thread's own value instead of the other lane's. HIP's own __shfl_xor for double uses the same memcpy-based pattern to avoid this.

Member Typedef Documentation

DeviceWord

template<typename T, std::size_t maxWordSize = 8>

using TNL::Backend::Uninitialized< T, maxWordSize >::DeviceWord

Initial value:

std::conditional_t<
     WORD_SIZE == 8,
     std::uint64_t,
     std::conditional_t< WORD_SIZE == 4, std::uint32_t, std::conditional_t< WORD_SIZE == 2, std::uint16_t, std::uint8_t > > >

Member Function Documentation

get()

template<typename T, std::size_t maxWordSize = 8>

__device__ T & TNL::Backend::Uninitialized< T, maxWordSize >::get ( )

inline

Returns a reference to the stored value via reinterpret_cast.

Safe only when all accesses to the storage go through the same type T. Use load() / store() instead when the storage is also accessed through the DeviceWord array (e.g. in warp shuffle wrappers).

load()

template<typename T, std::size_t maxWordSize = 8>

__device__ T TNL::Backend::Uninitialized< T, maxWordSize >::load ( ) const

inline

Copies the value out of the storage using memcpy.

Use this instead of get() when the storage is also accessed through the DeviceWord array to avoid strict aliasing violations.

store()

template<typename T, std::size_t maxWordSize = 8>

__device__ void TNL::Backend::Uninitialized< T, maxWordSize >::store ( const T & value )

inline

Copies a value into the storage using memcpy.

Use this instead of get() when the storage is also accessed through the DeviceWord array to avoid strict aliasing violations.

---

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

• src/TNL/Backend/Functions.h