tensor_view< BufferView_, TensorDesc_, DstInMemOp_ > Struct Template Reference

tensor_view&lt; BufferView_, TensorDesc_, DstInMemOp_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ > Struct Template Reference

#include <tensor_view.hpp>

Public Types

using buffer_view = remove_reference_t<BufferView_>
using DataType = typename buffer_view::type
using TensorDesc = remove_cvref_t<TensorDesc_>
using TensorIndex = array<index_t, TensorDesc::get_num_of_top_dimension()>
using TensorCoord = decltype(make_tensor_coordinate(TensorDesc{}, TensorIndex{}))

Public Member Functions

CK_TILE_HOST_DEVICE constexpr tensor_view ()=default
CK_TILE_HOST_DEVICE constexpr tensor_view (const buffer_view &buffer_view, const TensorDesc &desc)
CK_TILE_HOST_DEVICE void init_raw ()
CK_TILE_HOST_DEVICE constexpr auto & get_tensor_descriptor () const
CK_TILE_HOST_DEVICE constexpr const auto & get_buffer_view () const
CK_TILE_HOST_DEVICE constexpr auto & get_buffer_view ()
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > get_vectorized_elements (const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}) const
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > get_vectorized_elements (const TensorCoord &coord, 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_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE void get_vectorized_elements_raw (remove_cvref_t< X > &dst, const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE void get_vectorized_elements_raw (remove_cvref_t< X > &dst, const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void async_get_vectorized_elements (CK_TILE_LDS_ADDR remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}) const
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void async_get_vectorized_elements (CK_TILE_LDS_ADDR remove_cvref_t< DataType > *smem, const TensorCoord &coord, 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_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void async_get_vectorized_elements_raw (remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool_constant< pre_nop >={}) const
template<typename X, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void async_get_vectorized_elements_raw (remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t coord_extra_offset, index_t linear_offset, bool_constant< pre_nop >={}) const
template<typename X, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void async_get_vectorized_elements_raw (remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< pre_nop >={}) const
template<typename X, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > get_transpose_vectorized_elements (const TensorCoord &coord, index_t linear_offset) const
template<typename X, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > get_transpose_vectorized_elements (const TensorCoord &coord, index_t linear_offset, bool is_valid_element) const
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void set_vectorized_elements (const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void set_vectorized_elements (const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void set_vectorized_elements_raw (const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void set_vectorized_elements_raw (const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void update_vectorized_elements (const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void update_vectorized_elements (const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void update_vectorized_elements_raw (const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void update_vectorized_elements_raw (const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr index_t get_num_of_dimension ()

Public Attributes

buffer_view buf_
TensorDesc desc_

Static Public Attributes

static constexpr auto DstInMemOp = DstInMemOp_
static constexpr index_t PackedSize

Member Typedef Documentation

◆ buffer_view

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
using ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::buffer_view = remove_reference_t<BufferView_>

◆ DataType

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
using ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::DataType = typename buffer_view::type

◆ TensorCoord

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
using ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::TensorCoord = decltype(make_tensor_coordinate(TensorDesc{}, TensorIndex{}))

◆ TensorDesc

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
using ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::TensorDesc = remove_cvref_t<TensorDesc_>

◆ TensorIndex

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
using ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::TensorIndex = array<index_t, TensorDesc::get_num_of_top_dimension()>

Constructor & Destructor Documentation

◆ tensor_view() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::tensor_view ( )
constexprdefault

◆ tensor_view() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::tensor_view ( const buffer_view & buffer_view,
const TensorDesc & desc )
inlineconstexpr

Member Function Documentation

◆ async_get_vectorized_elements() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::async_get_vectorized_elements ( CK_TILE_LDS_ADDR remove_cvref_t< DataType > * smem,
const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ async_get_vectorized_elements() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::async_get_vectorized_elements ( CK_TILE_LDS_ADDR remove_cvref_t< DataType > * smem,
const TensorCoord & coord,
index_t linear_offset,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ async_get_vectorized_elements_raw() [1/3]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::async_get_vectorized_elements_raw ( remove_cvref_t< DataType > * smem,
const TensorCoord & coord,
index_t coord_extra_offset,
index_t linear_offset,
bool_constant< pre_nop > = {} ) const
inlineconstexpr

◆ async_get_vectorized_elements_raw() [2/3]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::async_get_vectorized_elements_raw ( remove_cvref_t< DataType > * smem,
const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
bool_constant< pre_nop > = {} ) const
inlineconstexpr

◆ async_get_vectorized_elements_raw() [3/3]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::async_get_vectorized_elements_raw ( remove_cvref_t< DataType > * smem,
const TensorCoord & coord,
index_t linear_offset,
bool_constant< pre_nop > = {} ) const
inlineconstexpr

◆ get_buffer_view() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr auto & ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_buffer_view ( )
inlineconstexpr

◆ get_buffer_view() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr const auto & ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_buffer_view ( ) const
inlineconstexpr

◆ get_num_of_dimension()

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_num_of_dimension ( )
inlinestaticconstexpr

◆ get_tensor_descriptor()

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE constexpr auto & ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_tensor_descriptor ( ) const
inlineconstexpr

◆ get_transpose_vectorized_elements() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_transpose_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset ) const
inlineconstexpr

◆ get_transpose_vectorized_elements() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_transpose_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element ) const
inlineconstexpr

◆ get_vectorized_elements() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ get_vectorized_elements() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr remove_cvref_t< X > ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
bool_constant< oob_conditional_check > = {} ) const
inlineconstexpr

◆ get_vectorized_elements_raw() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_vectorized_elements_raw ( remove_cvref_t< X > & dst,
const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} ) const
inline

◆ get_vectorized_elements_raw() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_vectorized_elements_raw ( remove_cvref_t< X > & dst,
const TensorCoord & coord,
index_t linear_offset,
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} ) const
inline

◆ init_raw()

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
CK_TILE_HOST_DEVICE void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::init_raw ( )
inline

◆ set_vectorized_elements() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::set_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ set_vectorized_elements() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::set_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ set_vectorized_elements_raw() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::set_vectorized_elements_raw ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ set_vectorized_elements_raw() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::set_vectorized_elements_raw ( const TensorCoord & coord,
index_t linear_offset,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ update_vectorized_elements() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::update_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ update_vectorized_elements() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::update_vectorized_elements ( const TensorCoord & coord,
index_t linear_offset,
const X & x,
bool_constant< oob_conditional_check > = {} )
inlineconstexpr

◆ update_vectorized_elements_raw() [1/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::update_vectorized_elements_raw ( const TensorCoord & coord,
index_t linear_offset,
bool is_valid_element,
const X & x,
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} )
inlineconstexpr

◆ update_vectorized_elements_raw() [2/2]

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
template<typename X, bool oob_conditional_check = true, bool pre_nop = false, typename std::enable_if< std::is_same_v< typename vector_traits< remove_cvref_t< X > >::scalar_type, typename vector_traits< remove_cvref_t< DataType > >::scalar_type >, bool >::type = false>
CK_TILE_HOST_DEVICE constexpr void ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::update_vectorized_elements_raw ( const TensorCoord & coord,
index_t linear_offset,
const X & x,
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} )
inlineconstexpr

Member Data Documentation

◆ buf_

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
buffer_view ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::buf_

◆ desc_

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
TensorDesc ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::desc_

◆ DstInMemOp

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
auto ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::DstInMemOp = DstInMemOp_
staticconstexpr

◆ PackedSize

template<typename BufferView_, typename TensorDesc_, memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
index_t ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::PackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<DataType>>::PackedSize
ck_tile::numeric_traits
Definition tile/core/numeric/numeric.hpp:81
ck_tile::static_distributed_tensor::PackedSize
static constexpr index_t PackedSize
Definition static_distributed_tensor.hpp:30
The documentation for this struct was generated from the following file: