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

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

Composable Kernel: ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType > 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, IndexType > 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[] (IndexType i) const
 
__host__ constexpr __device__ T & operator() (IndexType 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 (IndexType 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 (IndexType i, bool is_valid_element, const X &x)
 
template<typename DstBuffer , index_t NumElemsPerThread>
__host__ __device__ void DirectCopyToLds (DstBuffer &dst_buf, IndexType src_offset, IndexType 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 (IndexType 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 (IndexType 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 (IndexType 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, typename IndexType = index_t>
using ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::type = T

Constructor & Destructor Documentation

◆ DynamicBuffer() [1/2]

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
__host__ constexpr __device__ ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::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, typename IndexType = index_t>
__host__ constexpr __device__ ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::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, typename IndexType = index_t>
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, IndexType >::AtomicAdd ( IndexType  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ AtomicMax()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
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, IndexType >::AtomicMax ( IndexType  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ DirectCopyToLds()

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

◆ Get()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
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, IndexType >::Get ( IndexType  i,
bool  is_valid_element 
) const
inlineconstexpr

◆ GetAddressSpace()

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

◆ IsDynamicBuffer()

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

◆ IsStaticBuffer()

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

◆ operator()()

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

◆ operator[]()

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

◆ Set()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
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, IndexType >::Set ( IndexType  i,
bool  is_valid_element,
const X &  x 
)
inline

◆ Update()

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
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, IndexType >::Update ( IndexType  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, typename IndexType = index_t>
ElementSpaceSize ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::element_space_size_

◆ invalid_element_value_

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

◆ p_data_

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

◆ PackedSize

template<AddressSpaceEnum BufferAddressSpace, typename T , typename ElementSpaceSize , bool InvalidElementUseNumericalZeroValue, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence, typename IndexType = index_t>
constexpr index_t ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::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/develop/include/ck/utility/dynamic_buffer.hpp