/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp Source File#
warp_gemm_attribute_mfma.hpp
Go to the documentation of this file.
Definition: cluster_descriptor.hpp:13
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
typename impl::ext_vector< T, N >::type ext_vector_t
Definition: vector_type.hpp:83
Definition: warp_gemm_attribute_mfma.hpp:23
static constexpr auto get_warp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:48
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:38
typename Impl::BDataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:29
static constexpr index_t kCMLane
Definition: warp_gemm_attribute_mfma.hpp:40
typename Impl::AVecType AVecType
Definition: warp_gemm_attribute_mfma.hpp:32
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:34
decltype(get_warp_dstr_encoding< Impl::kAMLane >()) AWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:74
static constexpr auto AttrNumAccess
Definition: warp_gemm_attribute_mfma.hpp:25
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:30
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:97
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:42
static constexpr auto AttrNumAccessV
Definition: warp_gemm_attribute_mfma.hpp:26
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:36
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:24
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:88
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:39
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:37
typename Impl::ADataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:28
decltype(get_warp_dstr_encoding< Impl::kBNLane >()) BWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:75
typename Impl::BVecType BVecType
Definition: warp_gemm_attribute_mfma.hpp:33
Definition: warp_gemm_attribute_mfma.hpp:846
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:847
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, number< iKIter >, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:920
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:860
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:859
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:861
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:942
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:862
ext_vector_t< BDataType, vector_traits< typename Impl::BVecType >::vector_size *kKIter > BVecType
Definition: warp_gemm_attribute_mfma.hpp:856
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:901
typename Impl::BDataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:850
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:857
static constexpr index_t SFactor
Definition: warp_gemm_attribute_mfma.hpp:863
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:865
typename Impl::ADataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:849
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:851
ext_vector_t< ADataType, vector_traits< typename Impl::AVecType >::vector_size *kKIter > AVecType
Definition: warp_gemm_attribute_mfma.hpp:854
Definition: warp_gemm_attribute_mfma.hpp:702
typename Impl::BDataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:706
typename Impl::ADataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:707
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:719
ext_vector_t< ADataType, vector_traits< typename Impl::AVecType >::vector_size *kKIter > AVecType
Definition: warp_gemm_attribute_mfma.hpp:711
static constexpr index_t SFactor
Definition: warp_gemm_attribute_mfma.hpp:720
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:718
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:780
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:714
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:703
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:708
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:717
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:821
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, number< iKIter >, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:799
ext_vector_t< BDataType, vector_traits< typename Impl::BVecType >::vector_size *kKIter > BVecType
Definition: warp_gemm_attribute_mfma.hpp:713
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:716
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:722
Definition: warp_gemm_attribute_mfma.hpp:552
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:565
static constexpr CK_TILE_DEVICE auto get_cwarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:589
ext_vector_t< BDataType, vector_traits< typename Impl::BVecType >::vector_size *kKIter > BVecType
Definition: warp_gemm_attribute_mfma.hpp:564
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:570
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:559
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:572
typename Impl::BDataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:557
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:567
static constexpr CK_TILE_DEVICE auto get_awarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:577
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:569
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, number< iKIter >, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:655
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:568
decltype(get_awarp_dstr_encoding()) AWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:627
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:677
decltype(get_bwarp_dstr_encoding()) BWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:629
static constexpr CK_TILE_DEVICE auto get_bwarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:583
decltype(get_cwarp_dstr_encoding()) CWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:631
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:553
static constexpr auto AttrNumAccess
Definition: warp_gemm_attribute_mfma.hpp:554
ext_vector_t< ADataType, vector_traits< typename Impl::AVecType >::vector_size *kKIter > AVecType
Definition: warp_gemm_attribute_mfma.hpp:562
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:635
typename Impl::ADataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:558
Definition: warp_gemm_attribute_mfma.hpp:107
static constexpr auto AttrNumAccess
Definition: warp_gemm_attribute_mfma.hpp:111
static constexpr index_t kCMLane
Definition: warp_gemm_attribute_mfma.hpp:128
decltype(get_bwarp_dstr_encoding()) BWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:298
static constexpr CK_TILE_DEVICE auto get_cwarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:258
static constexpr CK_TILE_DEVICE auto get_bwarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:196
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:122
typename Impl::BDataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:115
ext_vector_t< ADataType, vector_traits< typename Impl::AVecType >::vector_size *kKIter > AVecType
Definition: warp_gemm_attribute_mfma.hpp:119
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, number< iKIter >, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:323
decltype(get_cwarp_dstr_encoding()) CWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:300
ext_vector_t< BDataType, vector_traits< typename Impl::BVecType >::vector_size *kKIter > BVecType
Definition: warp_gemm_attribute_mfma.hpp:121
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:345
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:130
static constexpr CK_TILE_DEVICE auto get_awarp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:135
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:116
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:304
typename Impl::ADataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:114
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:124
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:126
static constexpr auto AttrNumAccessV
Definition: warp_gemm_attribute_mfma.hpp:112
decltype(get_awarp_dstr_encoding()) AWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:296
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:125
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:110
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:127
Definition: warp_gemm_attribute_mfma.hpp:456
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:468
static constexpr index_t SFactor
Definition: warp_gemm_attribute_mfma.hpp:471
typename Impl::BVecType AVecType
Definition: warp_gemm_attribute_mfma.hpp:463
typename Impl::AVecType BVecType
Definition: warp_gemm_attribute_mfma.hpp:464
typename Impl::ADataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:460
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:541
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:457
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:531
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:470
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:461
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:467
typename Impl::BDataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:459
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:473
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:469
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:465
Definition: warp_gemm_attribute_mfma.hpp:372
CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const
Definition: warp_gemm_attribute_mfma.hpp:447
static constexpr index_t kCMLane
Definition: warp_gemm_attribute_mfma.hpp:389
typename Impl::BDataType ADataType
Definition: warp_gemm_attribute_mfma.hpp:377
typename Impl::AVecType BVecType
Definition: warp_gemm_attribute_mfma.hpp:382
static constexpr index_t kKPerThread
Definition: warp_gemm_attribute_mfma.hpp:388
static constexpr CK_TILE_HOST_DEVICE auto get_num_of_access()
Definition: warp_gemm_attribute_mfma.hpp:391
decltype(get_warp_dstr_encoding< Impl::kAMLane >()) BWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:424
typename Impl::BVecType AVecType
Definition: warp_gemm_attribute_mfma.hpp:381
static constexpr index_t kK
Definition: warp_gemm_attribute_mfma.hpp:387
typename Impl::ADataType BDataType
Definition: warp_gemm_attribute_mfma.hpp:378
static constexpr auto AttrNumAccessV
Definition: warp_gemm_attribute_mfma.hpp:375
static constexpr auto get_warp_dstr_encoding()
Definition: warp_gemm_attribute_mfma.hpp:397
typename Impl::CDataType CDataType
Definition: warp_gemm_attribute_mfma.hpp:379
static constexpr auto AttrNumAccess
Definition: warp_gemm_attribute_mfma.hpp:374
typename Impl::CVecType CVecType
Definition: warp_gemm_attribute_mfma.hpp:383
static constexpr index_t kN
Definition: warp_gemm_attribute_mfma.hpp:386
remove_cvref_t< WarpGemmAttributeMfmaImpl_ > Impl
Definition: warp_gemm_attribute_mfma.hpp:373
decltype(get_warp_dstr_encoding< Impl::kBNLane >()) AWarpDstrEncoding
Definition: warp_gemm_attribute_mfma.hpp:423
static constexpr index_t kM
Definition: warp_gemm_attribute_mfma.hpp:385
CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
Definition: warp_gemm_attribute_mfma.hpp:437
Definition: integral_constant.hpp:13
Definition: sequence.hpp:49
Definition: functional.hpp:43
Definition: debug.hpp:67
Definition: tile_distribution_encoding.hpp:26
Definition: tuple.hpp:192