MfmaSelector< base_type, MPerXdlops, NPerXdlops, additional_type, is_single_rate_mfma, is_scale_mfma > Class Template Reference#
Public Member Functions |
Static Public Member Functions |
Static Public Attributes |
List of all members
ck::MfmaSelector< base_type, MPerXdlops, NPerXdlops, additional_type, is_single_rate_mfma, is_scale_mfma > Class Template Reference
Selects the appropriate MFMA instruction type and configuration for given data types and tile sizes on AMD GPUs. More...
#include <xdlops_gemm.hpp>
Public Member Functions | |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
template<> | |
constexpr auto | GetMfma () |
__host__ constexpr __device__ | MfmaSelector () |
Static Public Member Functions | |
template<typename base_type_ , index_t MPerXdlops_, index_t NPerXdlops_, typename additional_type_ = base_type_, bool is_single_rate_mfma_ = false, bool is_scale_mfma_ = false> | |
static constexpr auto | GetMfma () |
static constexpr bool | IsABroadcast () |
static constexpr index_t | GetKPerXdlops () |
static constexpr index_t | GetK1PerXdlops () |
Static Public Attributes | |
static constexpr auto | selected_mfma |
Detailed Description
template<typename base_type, index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
class ck::MfmaSelector< base_type, MPerXdlops, NPerXdlops, additional_type, is_single_rate_mfma, is_scale_mfma >
Selects the appropriate MFMA instruction type and configuration for given data types and tile sizes on AMD GPUs.
- Template Parameters
-
base_type The base data type for the matrix operation (e.g., float, half_t). MPerXdlops The number of rows per XDLops tile. NPerXdlops The number of columns per XDLops tile. additional_type (Optional) Additional data type for mixed-precision or special cases. Defaults to base_type. is_single_rate_mfma (Optional) Whether to use single-rate MFMA instructions. Defaults to false. is_scale_mfma (Optional) Whether to use scale MFMA instructions. Defaults to false.
Constructor & Destructor Documentation
◆ MfmaSelector()
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
|
inlineconstexpr |
Member Function Documentation
◆ GetK1PerXdlops()
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
|
inlinestaticconstexpr |
◆ GetKPerXdlops()
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
|
inlinestaticconstexpr |
◆ GetMfma() [1/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<typename base_type_ , index_t MPerXdlops_, index_t NPerXdlops_, typename additional_type_ = base_type_, bool is_single_rate_mfma_ = false, bool is_scale_mfma_ = false>
|
staticconstexpr |
◆ GetMfma() [2/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [3/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [4/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [5/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [6/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [7/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [8/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [9/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [10/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [11/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [12/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [13/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [14/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [15/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [16/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [17/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [18/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [19/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [20/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [21/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [22/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [23/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [24/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [25/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [26/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [27/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [28/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [29/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [30/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [31/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [32/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [33/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [34/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [35/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [36/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [37/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [38/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [39/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [40/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [41/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [42/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [43/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [44/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [45/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [46/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [47/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [48/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [49/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [50/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [51/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [52/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [53/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [54/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [55/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ GetMfma() [56/56]
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
template<>
|
inlineconstexpr |
◆ IsABroadcast()
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
|
inlinestaticconstexpr |
Member Data Documentation
◆ selected_mfma
template<typename base_type , index_t MPerXdlops, index_t NPerXdlops, typename additional_type = base_type, bool is_single_rate_mfma = false, bool is_scale_mfma = false>
|
staticconstexpr |
Initial value:
= mfma_type<GetMfma<element_type_t<base_type>,
MPerXdlops,
NPerXdlops,
element_type_t<additional_type>,
is_single_rate_mfma,
is_scale_mfma>()>{}
The documentation for this class was generated from the following file:
- /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/warp/xdlops_gemm.hpp