coherence_default > Struct Template Reference

coherence_default > Struct Template Reference#

Composable Kernel: ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | List of all members
ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default > Struct Template Reference

#include <buffer_view.hpp>

Public Types

using type = T

Public Member Functions

CK_TILE_HOST_DEVICE constexpr buffer_view ()
CK_TILE_HOST_DEVICE constexpr buffer_view (T *__restrict__ p_data, BufferSizeType buffer_size)
CK_TILE_HOST_DEVICE constexpr buffer_view (T *__restrict__ p_data, BufferSizeType buffer_size, T invalid_element_value)
CK_TILE_HOST_DEVICE void init_raw ()
CK_TILE_DEVICE constexpr const T & operator[] (index_t i) const
CK_TILE_DEVICE constexpr T & operator() (index_t i)
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE constexpr auto get (index_t i, index_t, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
template<memory_operation_enum Op, typename X, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE void update (index_t i, index_t linear_offset, bool is_valid_element, const X &x)
template<typename X, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE void set (index_t i, index_t linear_offset, bool is_valid_element, const X &x)

Static Public Member Functions

static CK_TILE_DEVICE constexpr address_space_enum get_address_space ()
static CK_TILE_DEVICE constexpr bool is_static_buffer ()
static CK_TILE_DEVICE constexpr bool is_dynamic_buffer ()

Public Attributes

T * p_data_ = nullptr
BufferSizeType buffer_size_
remove_cvref_t< T > invalid_element_value_ = T{0}

Member Typedef Documentation

◆ type

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
using ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::type = T

Constructor & Destructor Documentation

◆ buffer_view() [1/3]

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_HOST_DEVICE constexpr ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::buffer_view ( )
inlineconstexpr

◆ buffer_view() [2/3]

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_HOST_DEVICE constexpr ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::buffer_view ( T *__restrict__ p_data,
BufferSizeType buffer_size )
inlineconstexpr

◆ buffer_view() [3/3]

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_HOST_DEVICE constexpr ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::buffer_view ( T *__restrict__ p_data,
BufferSizeType buffer_size,
T invalid_element_value )
inlineconstexpr

Member Function Documentation

◆ get()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE constexpr auto ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::get ( index_t i,
index_t ,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ get_address_space()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_DEVICE constexpr address_space_enum ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::get_address_space ( )
inlinestaticconstexpr

◆ init_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_HOST_DEVICE void ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::init_raw ( )
inline

◆ is_dynamic_buffer()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_DEVICE constexpr bool ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::is_dynamic_buffer ( )
inlinestaticconstexpr

◆ is_static_buffer()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_DEVICE constexpr bool ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::is_static_buffer ( )
inlinestaticconstexpr

◆ operator()()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_DEVICE constexpr T & ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::operator() ( index_t i)
inlineconstexpr

◆ operator[]()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
CK_TILE_DEVICE constexpr const T & ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::operator[] ( index_t i) const
inlineconstexpr

◆ set()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
template<typename X, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE void ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::set ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ update()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
template<memory_operation_enum Op, typename X, typename std::enable_if< std::is_same< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< T > >::scalar_type >::value, bool >::type = false>
CK_TILE_DEVICE void ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::update ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

Member Data Documentation

◆ buffer_size_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
BufferSizeType ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::buffer_size_

◆ invalid_element_value_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
remove_cvref_t<T> ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::invalid_element_value_ = T{0}

◆ p_data_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue>
T* ck_tile::buffer_view< address_space_enum::vgpr, T, BufferSizeType, InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum::coherence_default >::p_data_ = nullptr
The documentation for this struct was generated from the following file: