DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched > Struct Template Reference#
Classes |
Public Types |
Public Member Functions |
Static Public Member Functions |
Static Public Attributes |
List of all members
ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched > Struct Template Reference
#include <device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp>
Inheritance diagram for ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >:
Classes | |
| struct | Argument |
| struct | Invoker |
Public Types | |
| using | DeviceOp = DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle |
| using | ConvToGemmFwdTransformer = TransformConvFwdToGemm<NDimSpatial, ConvForwardSpecialization> |
| using | AGridDesc_M_K |
| using | BGridDesc_N_K |
| using | EGridDesc_M_N |
| using | RGridDesc_M = remove_cvref_t<decltype(MakeRGridDescriptor_M<RLayout>({}, {}))> |
| template<index_t NXdlPerWave_> | |
| using | GridwiseGemmBase |
| using | GridwiseGemm64 = GridwiseGemmBase<math::max(NXdlPerWave64, 1)> |
| using | GridwiseGemm32 = GridwiseGemmBase<NXdlPerWave32> |
| using | AGridDesc_AK0_M_AK1 |
| using | BGridDesc_BK0_N_BK1 |
| using | Block2ETileMap = typename GridwiseGemm64::DefaultBlock2ETileMap |
Public Member Functions | |
| bool | IsSupportedArgument (const BaseArgument *p_arg) override |
| std::unique_ptr< BaseArgument > | MakeArgumentPointer (const void *p_a, const void *p_b, const std::array< const void *, NumDTensor > &p_ds, void *p_e, std::array< void *, NumRTensor > p_rs, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_lengths, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_strides, const std::array< index_t, NDimSpatial > &conv_filter_strides, const std::array< index_t, NDimSpatial > &conv_filter_dilations, const std::array< index_t, NDimSpatial > &input_left_pads, const std::array< index_t, NDimSpatial > &input_right_pads, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CDEElementwiseOperation &cde_element_op, const QsElementwiseOperation &qs_element_op, const RsElementwiseOperation &rs_element_op) override |
| std::unique_ptr< BaseInvoker > | MakeInvokerPointer () override |
| std::string | GetTypeString () const override |
| Public Member Functions inherited from ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, DsDataType, EDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, RsElementwiseOperation, QsElementwiseOperation > | |
| virtual std::unique_ptr< BaseArgument > | MakeArgumentPointer (const void *p_a, const void *p_b, const std::array< const void *, NumDTensor > &p_ds, void *p_e, std::array< void *, NumRTensor > p_rs, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_lengths, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_strides, const std::array< index_t, NDimSpatial > &conv_filter_strides, const std::array< index_t, NDimSpatial > &conv_filter_dilations, const std::array< index_t, NDimSpatial > &input_left_pads, const std::array< index_t, NDimSpatial > &input_right_pads, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CDEElementwiseOperation &cde_element_op, const RsElementwiseOperation &qs_element_op, const QsElementwiseOperation &rs_element_op)=0 |
| Public Member Functions inherited from ck::tensor_operation::device::BaseOperator | |
| BaseOperator ()=default | |
| BaseOperator (const BaseOperator &)=default | |
| BaseOperator & | operator= (const BaseOperator &)=default |
| virtual std::string | GetInstanceString () const |
| virtual std::string | GetTypeIdName () const |
| virtual std::optional< std::string > | GetObjectName () const |
| virtual std::optional< std::string > | GetTemplateInfo () const |
| virtual std::string | GetTypeIdHashCode () const |
| virtual size_t | GetWorkSpaceSize (const BaseArgument *) const |
| virtual void | SetWorkSpacePointer (BaseArgument *p_arg, void *p_workspace, const StreamConfig &=StreamConfig{}) const |
| virtual | ~BaseOperator () |
Static Public Member Functions | |
| template<typename ALay> | |
| static auto | MakeAGridDescriptor_M_K (const ConvToGemmFwdTransformer &conv_to_gemm_transformer) |
| template<typename BLay> | |
| static auto | MakeBGridDescriptor_N_K (const ConvToGemmFwdTransformer &conv_to_gemm_transformer) |
| template<typename ELay> | |
| static auto | MakeEGridDescriptor_M_N (const ConvToGemmFwdTransformer &conv_to_gemm_transformer) |
| template<typename Descriptor> | |
| static auto | GetPaddedRGridDescriptor (Descriptor descriptor, index_t MRaw) |
| template<typename RLay, typename std::enable_if< is_same_v< RLay, tensor_layout::convolution::GNW >||is_same_v< RLay, tensor_layout::convolution::GNHW >||is_same_v< RLay, tensor_layout::convolution::GNDHW >, bool >::type = false> | |
| static auto | MakeRGridDescriptor_M (const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_lengths, const std::array< index_t, NDimSpatial+2 > &) |
| template<typename RLay, typename std::enable_if< is_same_v< RLay, tensor_layout::convolution::G_NW >||is_same_v< RLay, tensor_layout::convolution::G_NHW >||is_same_v< RLay, tensor_layout::convolution::G_NDHW >||is_same_v< RLay, tensor_layout::convolution::NWG >||is_same_v< RLay, tensor_layout::convolution::NHWG >||is_same_v< RLay, tensor_layout::convolution::NDHWG >, bool >::type = false> | |
| static auto | MakeRGridDescriptor_M (const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_lengths, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_strides) |
| static bool | IsSupportedArgument (const Argument &arg) |
| static auto | MakeArgument (const void *p_a, const void *p_b, const std::array< const void *, NumDTensor > &p_ds, void *p_e, std::array< void *, NumRTensor > p_rs, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_k_c_xs_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_n_k_wos_strides, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_lengths, const std::array< index_t, NDimSpatial+2 > &r_g_n_wos_strides, const std::array< index_t, NDimSpatial > &conv_filter_strides, const std::array< index_t, NDimSpatial > &conv_filter_dilations, const std::array< index_t, NDimSpatial > &input_left_pads, const std::array< index_t, NDimSpatial > &input_right_pads, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CDEElementwiseOperation &cde_element_op, const QsElementwiseOperation &qs_element_op, const RsElementwiseOperation &rs_element_op) |
| static auto | MakeInvoker () |
Static Public Attributes | |
| static GET_NXDL_PER_WAVE_IMPL constexpr auto | NXdlPerWave64 = GetNXdlPerWave<true>() |
| static constexpr auto | NXdlPerWave32 = GetNXdlPerWave<false>() |
| static constexpr index_t | NumDTensor = DsDataType::Size() |
| static constexpr index_t | NumRTensor = RsDataType::Size() |
| 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 | matrix_padder |
| static constexpr ConvToGemmFwdTransformer | dummy_conv_to_gemm_transformer |
| Static Public Attributes inherited from ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, DsDataType, EDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, RsElementwiseOperation, QsElementwiseOperation > | |
| static constexpr index_t | NumDTensor |
| static constexpr index_t | NumRTensor |
Member Typedef Documentation
◆ AGridDesc_AK0_M_AK1
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::AGridDesc_AK0_M_AK1 |
Initial value:
AGridDesc_M_K{}))>
decltype(MakeAGridDescriptor_M_K(1, 1, 1)) AGridDesc_M_K
Definition gridwise_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp:346
__host__ static __device__ constexpr auto MakeDefaultAGridDescriptor_AK0_M_AK1(const AGridDesc_M_K &a_grid_desc_m_k)
Definition gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp:174
◆ AGridDesc_M_K
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::AGridDesc_M_K |
Initial value:
static constexpr ConvToGemmFwdTransformer dummy_conv_to_gemm_transformer
Definition codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp:434
◆ BGridDesc_BK0_N_BK1
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::BGridDesc_BK0_N_BK1 |
Initial value:
BGridDesc_N_K{}))>
decltype(MakeBGridDescriptor_N_K(1, 1, 1)) BGridDesc_N_K
Definition gridwise_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp:347
__host__ static __device__ constexpr auto MakeDefaultBGridDescriptor_BK0_N_BK1(const BGridDesc_N_K &b_grid_desc_n_k)
Definition gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp:190
◆ BGridDesc_N_K
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::BGridDesc_N_K |
◆ Block2ETileMap
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::Block2ETileMap = typename GridwiseGemm64::DefaultBlock2ETileMap |
◆ ConvToGemmFwdTransformer
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::ConvToGemmFwdTransformer = TransformConvFwdToGemm<NDimSpatial, ConvForwardSpecialization> |
◆ DeviceOp
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::DeviceOp = DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle |
◆ EGridDesc_M_N
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::EGridDesc_M_N |
◆ GridwiseGemm32
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::GridwiseGemm32 = GridwiseGemmBase<NXdlPerWave32> |
◆ GridwiseGemm64
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::GridwiseGemm64 = GridwiseGemmBase<math::max(NXdlPerWave64, 1)> |
◆ GridwiseGemmBase
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<index_t NXdlPerWave_>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::GridwiseGemmBase |
◆ RGridDesc_M
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
| using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, DELayout, RLayout, ADataType, BDataType, AccDataType, CShuffleDataType, DsDataType, EDataType, ReduceAccDataType, RsDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, QsElementwiseOperation, RsElementwiseOperation, ThreadReduceOperations, RsGlobalMemoryDataOperation, ConvForwardSpecialization, GemmSpec, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, CDEBlockTransferScalarPerVector_NPerBlock, RThreadTransferDstScalarPerVector_MPerBlock, LoopSched >::RGridDesc_M = remove_cvref_t<decltype(MakeRGridDescriptor_M<RLayout>({}, {}))> |
Member Function Documentation
◆ GetPaddedRGridDescriptor()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename Descriptor>
|
inlinestatic |
◆ GetTypeString()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlineoverridevirtual |
Reimplemented from ck::tensor_operation::device::BaseOperator.
◆ IsSupportedArgument() [1/2]
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlinestatic |
◆ IsSupportedArgument() [2/2]
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlineoverridevirtual |
Reimplemented from ck::tensor_operation::device::BaseOperator.
◆ MakeAGridDescriptor_M_K()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename ALay>
|
inlinestatic |
◆ MakeArgument()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlinestatic |
◆ MakeArgumentPointer()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlineoverride |
◆ MakeBGridDescriptor_N_K()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename BLay>
|
inlinestatic |
◆ MakeEGridDescriptor_M_N()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename ELay>
|
inlinestatic |
◆ MakeInvoker()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlinestatic |
◆ MakeInvokerPointer()
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
inlineoverridevirtual |
◆ MakeRGridDescriptor_M() [1/2]
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename RLay, typename std::enable_if< is_same_v< RLay, tensor_layout::convolution::GNW >||is_same_v< RLay, tensor_layout::convolution::GNHW >||is_same_v< RLay, tensor_layout::convolution::GNDHW >, bool >::type = false>
|
inlinestatic |
◆ MakeRGridDescriptor_M() [2/2]
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
template<typename RLay, typename std::enable_if< is_same_v< RLay, tensor_layout::convolution::G_NW >||is_same_v< RLay, tensor_layout::convolution::G_NHW >||is_same_v< RLay, tensor_layout::convolution::G_NDHW >||is_same_v< RLay, tensor_layout::convolution::NWG >||is_same_v< RLay, tensor_layout::convolution::NHWG >||is_same_v< RLay, tensor_layout::convolution::NDHWG >, bool >::type = false>
|
inlinestatic |
Member Data Documentation
◆ dummy_conv_to_gemm_transformer
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ I0
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ I1
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ I2
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ I3
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ matrix_padder
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
Initial value:
=
MatrixPadder<GemmSpec, index_t, index_t, index_t>{MPerBlock, NPerBlock, KPerBlock}
Definition matrix_padder.hpp:180
◆ NumDTensor
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ NumRTensor
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ NXdlPerWave32
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
◆ NXdlPerWave64
template<index_t NDimSpatial, typename ALayout, typename BLayout, typename DELayout, typename RLayout, typename ADataType, typename BDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EDataType, typename ReduceAccDataType, typename RsDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, typename QsElementwiseOperation, typename RsElementwiseOperation, typename ThreadReduceOperations, typename RsGlobalMemoryDataOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1, index_t BK1, index_t MPerXDL, index_t NPerXDL, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CDRThreadTransferClusterLengths_MPerBlock_NPerBlock, index_t CDEBlockTransferScalarPerVector_NPerBlock, index_t RThreadTransferDstScalarPerVector_MPerBlock, LoopScheduler LoopSched = make_default_loop_scheduler()>
|
staticconstexpr |
The documentation for this struct was generated from the following file: