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, 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 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> | |
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 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> | |
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 , 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_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> | |
constexpr CK_TILE_HOST_DEVICE 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> | |
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 >={}) |
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_ |
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>
|
constexprdefault |
◆ tensor_view() [2/2]
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
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>
|
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>
|
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>
|
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>
|
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>
|
inlineconstexpr |
◆ get_buffer_view() [1/2]
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
inlineconstexpr |
◆ get_buffer_view() [2/2]
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
inlineconstexpr |
◆ get_num_of_dimension()
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
inlinestaticconstexpr |
◆ get_tensor_descriptor()
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
inline |
◆ init_raw()
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
staticconstexpr |
◆ PackedSize
template<typename BufferView_ , typename TensorDesc_ , memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
|
staticconstexpr |
Initial value:
=
Definition: numeric.hpp:81
The documentation for this struct was generated from the following file:
- /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tensor_view.hpp