global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence > Struct Template Reference

global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence > Struct Template Reference#

Composable Kernel: ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence > 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 linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
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 transpose_get (index_t i, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, 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_raw (remove_cvref_t< X > &dst, index_t v_offset, index_t i_offset, bool is_valid_element, bool_constant< pre_nop >={}) const
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 async_get (CK_TILE_LDS_ADDR remove_cvref_t< T > *smem, index_t i, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
template<typename X, bool pre_nop = false, 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 async_get_raw (remove_cvref_t< T > *smem, index_t i, index_t linear_offset, bool, bool_constant< pre_nop >={}) const
template<memory_operation_enum Op, 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 void update (index_t i, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
template<memory_operation_enum Op, typename X, bool oob_conditional_check = true, bool pre_nop = false, 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_raw (index_t i, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})
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 void set (index_t i, index_t linear_offset, bool is_valid_element, const X &x)
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 void set_raw (index_t i, index_t linear_offset, bool is_valid_element, const X &x)
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 void atomic_add (index_t i, index_t linear_offset, bool is_valid_element, const X &x)
template<typename X, bool oob_conditional_check = true, bool pre_nop = 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 void atomic_add_raw (index_t i, index_t linear_offset, bool is_valid_element, const X &x)
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 void atomic_max (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_
int32x4_t cached_buf_res_
remove_cvref_t< T > invalid_element_value_ = T{0}

Static Public Attributes

static constexpr index_t PackedSize = ck_tile::numeric_traits<remove_cvref_t<T>>::PackedSize

Member Typedef Documentation

◆ type

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

Constructor & Destructor Documentation

◆ buffer_view() [1/3]

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

◆ buffer_view() [2/3]

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

◆ buffer_view() [3/3]

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

Member Function Documentation

◆ async_get()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::async_get ( CK_TILE_LDS_ADDR remove_cvref_t< T > * smem,
index_t i,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ async_get_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
template<typename X, bool pre_nop = false, 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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::async_get_raw ( remove_cvref_t< T > * smem,
index_t i,
index_t linear_offset,
bool ,
bool_constant< pre_nop > = {} ) const
inlineconstexpr

◆ atomic_add()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::atomic_add ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ atomic_add_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
template<typename X, bool oob_conditional_check = true, bool pre_nop = 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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::atomic_add_raw ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ atomic_max()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::atomic_max ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ get()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::get ( index_t i,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ get_address_space()

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

◆ get_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, 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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::get_raw ( remove_cvref_t< X > & dst,
index_t v_offset,
index_t i_offset,
bool is_valid_element,
bool_constant< pre_nop > = {} ) const
inlineconstexpr

◆ init_raw()

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

◆ is_dynamic_buffer()

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

◆ is_static_buffer()

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

◆ operator()()

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

◆ operator[]()

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

◆ set()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::set ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ set_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::set_raw ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x )
inline

◆ transpose_get()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::transpose_get ( index_t i,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ update()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
template<memory_operation_enum Op, 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 void ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::update ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {} )
inline

◆ update_raw()

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
template<memory_operation_enum Op, typename X, bool oob_conditional_check = true, bool pre_nop = false, 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::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::update_raw ( index_t i,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} )
inline

Member Data Documentation

◆ buffer_size_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
BufferSizeType ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::buffer_size_

◆ cached_buf_res_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
int32x4_t ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::cached_buf_res_

◆ invalid_element_value_

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

◆ p_data_

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
T* ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::p_data_ = nullptr

◆ PackedSize

template<typename T, typename BufferSizeType, bool InvalidElementUseNumericalZeroValue, amd_buffer_coherence_enum Coherence>
index_t ck_tile::buffer_view< address_space_enum::global, T, BufferSizeType, InvalidElementUseNumericalZeroValue, Coherence >::PackedSize = ck_tile::numeric_traits<remove_cvref_t<T>>::PackedSize
staticconstexpr
The documentation for this struct was generated from the following file: