TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec > Struct Template Reference

TransformBatchedContractionContractionToBatchedGemmGemm_Wmma&lt; NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec &gt; Struct Template Reference#

Composable Kernel: ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec > Struct Template Reference
Static Public Member Functions | Static Public Attributes | List of all members
ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec > Struct Template Reference

#include <transform_contraction_to_gemm_arraybase.hpp>

Static Public Member Functions

__host__ static __device__ auto MakeAGridDescriptorPair (const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_strides_vec)
__host__ static __device__ auto MakeAGridDescriptor_G_M_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_strides_vec)
__host__ static __device__ auto MakeAGridDescriptor_M_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &a_gs_ms_ks_strides_vec)
template<typename AGridDesc_M_K, typename Number>
__host__ static __device__ constexpr auto MakeAGridDescriptor_AK0_M_AK1 (const AGridDesc_M_K &a_grid_desc_m_k, const Number &AK1)
template<typename AGridDesc_M_K, typename WmmaK, typename MRepeat, typename MWaves, typename MPerWmma, typename AK1>
__host__ static __device__ constexpr auto MakeAGridDescriptor_AKWmma_MBlockRepeat_MWaves_AK0PerWmma_AKRow_MPerWmma_AK1 (const AGridDesc_M_K &a_grid_desc_m_k, const WmmaK &, const MRepeat &, const MWaves &, const MPerWmma &, const AK1 &)
__host__ static __device__ auto MakeB0GridDescriptorPair (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_strides_vec)
__host__ static __device__ auto MakeB0GridDescriptor_G_N_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_strides_vec)
__host__ static __device__ auto MakeB0GridDescriptor_N_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b0_gs_ns_ks_strides_vec)
template<typename BGridDesc_N_K, typename Number>
__host__ static __device__ constexpr auto MakeB0GridDescriptor_BK0_N_BK1 (const BGridDesc_N_K &b_grid_desc_n_k, const Number &BK1)
template<typename BGridDesc_L_K, typename WmmaK, typename LRepeat, typename LWaves, typename LPerWmma, typename BK1>
__host__ static __device__ constexpr auto MakeB0GridDescriptor_BKWmma_LBlockRepeat_LWaves_BK0PerWmma_BKRow_LPerWmma_BK1 (const BGridDesc_L_K &b_grid_desc_l_k, const WmmaK &, const LRepeat &, const LWaves &, const LPerWmma &, const BK1 &)
__host__ static __device__ auto MakeB1GridDescriptorPair (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_strides_vec)
__host__ static __device__ auto MakeB1GridDescriptor_G_N_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_strides_vec)
__host__ static __device__ auto MakeB1GridDescriptor_N_K (const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &b1_gs_os_ns_strides_vec)
template<typename B1GridDesc_N_K, typename Number>
__host__ static __device__ constexpr auto MakeB1GridDescriptor_BK0_N_BK1 (const B1GridDesc_N_K &b1_grid_desc_n_k, const Number &B1K1)
template<typename BGridDesc_N_L, typename WmmaL, typename NRepeat, typename NWaves, typename NPerWmma, typename BL1>
__host__ static __device__ constexpr auto MakeB1GridDescriptor_BLWmma_NBlockRepeat_NWaves__BL0PerWmma_BLRow_NPerWmma_BL1 (const BGridDesc_N_L &b_grid_desc_n_l, const WmmaL &, const NRepeat &, const NWaves &, const NPerWmma &, const BL1 &)
__host__ static __device__ auto MakeCGridDescriptorPair (const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_strides_vec)
__host__ static __device__ auto MakeCGridDescriptor_G_M_N (const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_strides_vec)
__host__ static __device__ auto MakeCGridDescriptor_M_N (const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_lengths_vec, const std::array< index_t, NumDimG+NumDimM+NumDimN > &c_gs_ms_os_strides_vec)

Static Public Attributes

static constexpr auto I0 = Number<0>{}
static constexpr auto I1 = Number<1>{}
static constexpr auto I2 = Number<2>{}
static constexpr auto I3 = Number<3>{}
static constexpr auto I4 = Number<4>{}
static constexpr index_t NumDimG = NumDims_G_M_N_K_O::At(I0)
static constexpr index_t NumDimM = NumDims_G_M_N_K_O::At(I1)
static constexpr index_t NumDimN = NumDims_G_M_N_K_O::At(I2)
static constexpr index_t NumDimK = NumDims_G_M_N_K_O::At(I3)
static constexpr index_t NumDimO = NumDims_G_M_N_K_O::At(I4)
static constexpr index_t MPerBlock = PerBlock_M_N_K_O::At(I0)
static constexpr index_t NPerBlock = PerBlock_M_N_K_O::At(I1)
static constexpr index_t KPerBlock = PerBlock_M_N_K_O::At(I2)
static constexpr index_t OPerBlock = PerBlock_M_N_K_O::At(I3)
static constexpr auto matrix_padder

Member Function Documentation

◆ MakeAGridDescriptor_AK0_M_AK1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename AGridDesc_M_K, typename Number>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeAGridDescriptor_AK0_M_AK1 ( const AGridDesc_M_K & a_grid_desc_m_k,
const Number & AK1 )
inlinestaticconstexpr

◆ MakeAGridDescriptor_AKWmma_MBlockRepeat_MWaves_AK0PerWmma_AKRow_MPerWmma_AK1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename AGridDesc_M_K, typename WmmaK, typename MRepeat, typename MWaves, typename MPerWmma, typename AK1>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeAGridDescriptor_AKWmma_MBlockRepeat_MWaves_AK0PerWmma_AKRow_MPerWmma_AK1 ( const AGridDesc_M_K & a_grid_desc_m_k,
const WmmaK & ,
const MRepeat & ,
const MWaves & ,
const MPerWmma & ,
const AK1 &  )
inlinestaticconstexpr

◆ MakeAGridDescriptor_G_M_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeAGridDescriptor_G_M_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_strides_vec )
inlinestatic

◆ MakeAGridDescriptor_M_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeAGridDescriptor_M_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_strides_vec )
inlinestatic

◆ MakeAGridDescriptorPair()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeAGridDescriptorPair ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & a_gs_ms_ks_strides_vec )
inlinestatic

◆ MakeB0GridDescriptor_BK0_N_BK1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename BGridDesc_N_K, typename Number>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB0GridDescriptor_BK0_N_BK1 ( const BGridDesc_N_K & b_grid_desc_n_k,
const Number & BK1 )
inlinestaticconstexpr

◆ MakeB0GridDescriptor_BKWmma_LBlockRepeat_LWaves_BK0PerWmma_BKRow_LPerWmma_BK1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename BGridDesc_L_K, typename WmmaK, typename LRepeat, typename LWaves, typename LPerWmma, typename BK1>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB0GridDescriptor_BKWmma_LBlockRepeat_LWaves_BK0PerWmma_BKRow_LPerWmma_BK1 ( const BGridDesc_L_K & b_grid_desc_l_k,
const WmmaK & ,
const LRepeat & ,
const LWaves & ,
const LPerWmma & ,
const BK1 &  )
inlinestaticconstexpr

◆ MakeB0GridDescriptor_G_N_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB0GridDescriptor_G_N_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_strides_vec )
inlinestatic

◆ MakeB0GridDescriptor_N_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB0GridDescriptor_N_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_strides_vec )
inlinestatic

◆ MakeB0GridDescriptorPair()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB0GridDescriptorPair ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b0_gs_ns_ks_strides_vec )
inlinestatic

◆ MakeB1GridDescriptor_BK0_N_BK1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename B1GridDesc_N_K, typename Number>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB1GridDescriptor_BK0_N_BK1 ( const B1GridDesc_N_K & b1_grid_desc_n_k,
const Number & B1K1 )
inlinestaticconstexpr

◆ MakeB1GridDescriptor_BLWmma_NBlockRepeat_NWaves__BL0PerWmma_BLRow_NPerWmma_BL1()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
template<typename BGridDesc_N_L, typename WmmaL, typename NRepeat, typename NWaves, typename NPerWmma, typename BL1>
__host__ static __device__ constexpr auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB1GridDescriptor_BLWmma_NBlockRepeat_NWaves__BL0PerWmma_BLRow_NPerWmma_BL1 ( const BGridDesc_N_L & b_grid_desc_n_l,
const WmmaL & ,
const NRepeat & ,
const NWaves & ,
const NPerWmma & ,
const BL1 &  )
inlinestaticconstexpr

◆ MakeB1GridDescriptor_G_N_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB1GridDescriptor_G_N_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_strides_vec )
inlinestatic

◆ MakeB1GridDescriptor_N_K()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB1GridDescriptor_N_K ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_strides_vec )
inlinestatic

◆ MakeB1GridDescriptorPair()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeB1GridDescriptorPair ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & b1_gs_os_ns_strides_vec )
inlinestatic

◆ MakeCGridDescriptor_G_M_N()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeCGridDescriptor_G_M_N ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_strides_vec )
inlinestatic

◆ MakeCGridDescriptor_M_N()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeCGridDescriptor_M_N ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_strides_vec )
inlinestatic

◆ MakeCGridDescriptorPair()

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
__host__ static __device__ auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MakeCGridDescriptorPair ( const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_lengths_vec,
const std::array< index_t, NumDimG+NumDimM+NumDimN > & c_gs_ms_os_strides_vec )
inlinestatic

Member Data Documentation

◆ I0

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::I2 = Number<2>{}
staticconstexpr

◆ I3

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::I3 = Number<3>{}
staticconstexpr

◆ I4

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::I4 = Number<4>{}
staticconstexpr

◆ KPerBlock

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::KPerBlock = PerBlock_M_N_K_O::At(I2)
staticconstexpr

◆ matrix_padder

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
auto ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::matrix_padder
staticconstexpr
Initial value:
=
device::GemmGemmPadder<GemmSpec, index_t, index_t, index_t, index_t>{
MPerBlock, NPerBlock, KPerBlock, OPerBlock}
ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma::KPerBlock
static constexpr index_t KPerBlock
Definition transform_contraction_to_gemm_arraybase.hpp:137
ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma::NPerBlock
static constexpr index_t NPerBlock
Definition transform_contraction_to_gemm_arraybase.hpp:136
ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma::MPerBlock
static constexpr index_t MPerBlock
Definition transform_contraction_to_gemm_arraybase.hpp:135
ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm::OPerBlock
static constexpr index_t OPerBlock
Definition transform_contraction_to_gemm.hpp:137
ck::tensor_operation::device::GemmGemmPadder
Definition matrix_padder.hpp:63

◆ MPerBlock

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::MPerBlock = PerBlock_M_N_K_O::At(I0)
staticconstexpr

◆ NPerBlock

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NPerBlock = PerBlock_M_N_K_O::At(I1)
staticconstexpr

◆ NumDimG

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NumDimG = NumDims_G_M_N_K_O::At(I0)
staticconstexpr

◆ NumDimK

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NumDimK = NumDims_G_M_N_K_O::At(I3)
staticconstexpr

◆ NumDimM

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NumDimM = NumDims_G_M_N_K_O::At(I1)
staticconstexpr

◆ NumDimN

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NumDimN = NumDims_G_M_N_K_O::At(I2)
staticconstexpr

◆ NumDimO

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::NumDimO = NumDims_G_M_N_K_O::At(I4)
staticconstexpr

◆ OPerBlock

template<typename NumDims_G_M_N_K_O, typename PerBlock_M_N_K_O, device::GemmSpecialization GemmSpec, device::TensorSpecialization ASpec, device::TensorSpecialization B0Spec, device::TensorSpecialization B1Spec, device::TensorSpecialization CSpec>
index_t ck::tensor_operation::TransformBatchedContractionContractionToBatchedGemmGemm_Wmma< NumDims_G_M_N_K_O, PerBlock_M_N_K_O, GemmSpec, ASpec, B0Spec, B1Spec, CSpec >::OPerBlock = PerBlock_M_N_K_O::At(I3)
staticconstexpr
The documentation for this struct was generated from the following file: