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
ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType > Struct Template Reference

#include <dynamic_buffer.hpp>

Public Types

using type = T

Public Member Functions

__host__ __device__ constexpr DynamicBuffer (T *p_data, ElementSpaceSize element_space_size)
__host__ __device__ constexpr DynamicBuffer (T *p_data, ElementSpaceSize element_space_size, T invalid_element_value)
__host__ __device__ constexpr const T & operator[] (IndexType i) const
__host__ __device__ constexpr T & operator() (IndexType i)
template<typename X, bool DoTranspose = false, 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__ constexpr 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 __device__ constexpr AddressSpaceEnum GetAddressSpace ()
__host__ static __device__ constexpr bool IsStaticBuffer ()
__host__ static __device__ constexpr bool IsDynamicBuffer ()

Public Attributes

T * p_data_
ElementSpaceSize element_space_size_
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__ __device__ constexpr 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__ __device__ constexpr 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, bool DoTranspose = false, 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__ constexpr 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 __device__ constexpr 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 __device__ constexpr 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 __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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>
index_t ck::DynamicBuffer< BufferAddressSpace, T, ElementSpaceSize, InvalidElementUseNumericalZeroValue, coherence, IndexType >::PackedSize
staticconstexpr
Initial value:
= []() {
return 2;
else
return 1;
}()
constexpr bool is_same_v
Definition type.hpp:283
Definition data_type.hpp:187

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