DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence > Struct Template Reference

DynamicBuffer&lt; BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence &gt; Struct Template Reference#

Composable Kernel: ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence > Struct Template Reference

#include <dynamic_buffer.hpp>

Public Types

using type = T
 

Public Member Functions

__host__ constexpr __device__ DynamicBuffer (T *p_data, ElementSpaceSize element_space_size)
 
__host__ constexpr __device__ DynamicBuffer (T *p_data, ElementSpaceSize element_space_size, T invalid_element_value)
 
__host__ constexpr __device__ const T & operator[] (index_t i) const
 
__host__ constexpr __device__ T & operator() (index_t i)
 
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value||!is_native_type< X >(), bool >::type = false>
__host__ constexpr __device__ auto Get (index_t i, bool is_valid_element) const
 
template<InMemoryDataOperationEnum Op, typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void Update (index_t i, bool is_valid_element, const X &x)
 
template<typename DstBuffer , index_t NumElemsPerThread>
__host__ __device__ void DirectCopyToLds (DstBuffer &dst_buf, index_t src_offset, index_t dst_offset, bool is_valid_element) const
 
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value||!is_native_type< X >(), bool >::type = false>
__host__ __device__ void Set (index_t i, bool is_valid_element, const X &x)
 
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void AtomicAdd (index_t i, bool is_valid_element, const X &x)
 
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void AtomicMax (index_t i, bool is_valid_element, const X &x)
 

Static Public Member Functions

__host__ static constexpr __device__ AddressSpaceEnum GetAddressSpace ()
 
__host__ static constexpr __device__ bool IsStaticBuffer ()
 
__host__ static constexpr __device__ bool IsDynamicBuffer ()
 

Public Attributes

T * p_data_
 
ElementSpaceSize element_space_size_
 
T invalid_element_value_ = T{0}
 

Static Public Attributes

static constexpr index_t PackedSize
 

Member Typedef Documentation

◆ type

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
using ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::type = T

Constructor & Destructor Documentation

◆ DynamicBuffer() [1/2]

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ constexpr __device__ ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::DynamicBuffer ( T *  p_data,
ElementSpaceSize  element_space_size 
)
inlineconstexpr

◆ DynamicBuffer() [2/2]

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ constexpr __device__ ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::DynamicBuffer ( T *  p_data,
ElementSpaceSize  element_space_size,
T  invalid_element_value 
)
inlineconstexpr

Member Function Documentation

◆ AtomicAdd()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::AtomicAdd ( index_t  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ AtomicMax()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::AtomicMax ( index_t  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ DirectCopyToLds()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<typename DstBuffer , index_t NumElemsPerThread>
__host__ __device__ void ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::DirectCopyToLds ( DstBuffer &  dst_buf,
index_t  src_offset,
index_t  dst_offset,
bool  is_valid_element 
) const
inline

◆ Get()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value||!is_native_type< X >(), bool >::type = false>
__host__ constexpr __device__ auto ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::Get ( index_t  i,
bool  is_valid_element 
) const
inlineconstexpr

◆ GetAddressSpace()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ static constexpr __device__ AddressSpaceEnum ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::GetAddressSpace ( )
inlinestaticconstexpr

◆ IsDynamicBuffer()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ static constexpr __device__ bool ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::IsDynamicBuffer ( )
inlinestaticconstexpr

◆ IsStaticBuffer()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ static constexpr __device__ bool ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::IsStaticBuffer ( )
inlinestaticconstexpr

◆ operator()()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ constexpr __device__ T& ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::operator() ( index_t  i)
inlineconstexpr

◆ operator[]()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
__host__ constexpr __device__ const T& ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::operator[] ( index_t  i) const
inlineconstexpr

◆ Set()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value||!is_native_type< X >(), bool >::type = false>
__host__ __device__ void ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::Set ( index_t  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ Update()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
template<InMemoryDataOperationEnum Op, typename X , typename enable_if< is_same< typename scalar_type< remove_cvref_t< X >>::type, typename scalar_type< remove_cvref_t< T >>::type >::value, bool >::type = false>
__host__ __device__ void ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::Update ( index_t  i,
bool  is_valid_element,
const X &  x 
)
inline

Member Data Documentation

◆ element_space_size_

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
ElementSpaceSize ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::element_space_size_

◆ invalid_element_value_

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
T ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::invalid_element_value_ = T{0}

◆ p_data_

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
T* ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::p_data_

◆ PackedSize

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
constexpr index_t ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence >::PackedSize
staticconstexpr
Initial value:
= []() {
if constexpr(is_same_v<remove_cvref_t<T>, pk_i4_t>)
return 2;
else
return 1;
}()
ck::is_same_v
constexpr bool is_same_v
Definition: type.hpp:283
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/utility/dynamic_buffer.hpp