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

constexpr CK_TILE_HOST_DEVICE tensor_view ()=default
 
constexpr CK_TILE_HOST_DEVICE tensor_view (const buffer_view &buffer_view, const TensorDesc &desc)
 
CK_TILE_HOST_DEVICE void init_raw ()
 
constexpr CK_TILE_HOST_DEVICE auto & get_tensor_descriptor () const
 
constexpr CK_TILE_HOST_DEVICE const auto & get_buffer_view () const
 
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements (CK_TILE_LDS_ADDR remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset) 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>
constexpr CK_TILE_HOST_DEVICE 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) 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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 , 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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 >={})
 
CK_TILE_HOST_DEVICE void print () const
 

Static Public Member Functions

static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_dimension ()
 

Public Attributes

buffer_view buf_
 
TensorDesc desc_
 

Static Public Attributes

static constexpr auto DstInMemOp = DstInMemOp_
 

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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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 
) 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>
constexpr CK_TILE_HOST_DEVICE 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 
) const
inlineconstexpr

◆ async_get_vectorized_elements_raw() [1/2]

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>
constexpr CK_TILE_HOST_DEVICE 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() [2/2]

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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
static constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE auto& ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::get_tensor_descriptor ( ) 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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

◆ print()

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_ >::print ( ) const
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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr CK_TILE_HOST_DEVICE 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>
constexpr auto ck_tile::tensor_view< BufferView_, TensorDesc_, DstInMemOp_ >::DstInMemOp = DstInMemOp_
staticconstexpr
The documentation for this struct was generated from the following file: