Composable Kernel: ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, > Struct Template Reference

DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, > Struct Template Reference

DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, > Struct Template Reference#

Composable Kernel: ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, > Struct Template Reference

ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, > Struct Template Reference

#include <device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp>

Inheritance diagram for ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >:

Classes
struct	Argument
struct	Invoker

Public Types
using	DeviceOp = DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK
using	ConvToGemmFwdTransformer = TransformConvFwdToGemm<NDimSpatial, ConvForwardSpecialization>
using	AGridDesc_AK0_M_AK1
using	BGridDesc_BK0_N_BK1
using	CGridDesc_M_N
using	GridwiseGemm
using	AGridDesc_K0_M0_M1_K1
using	BGridDesc_K0_N0_N1_K1
using	CGridDesc_M0_M10_M11_N0_N10_N11
using	DefaultBlock2CTileMap

Public Member Functions
bool	IsSupportedArgument (const BaseArgument *p_arg) override
std::unique_ptr< BaseArgument >	MakeArgumentPointer (const void p_a, const void p_b, void *p_c, 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< index_t, NDimSpatial+3 > &c_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &c_g_n_k_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 CElementwiseOperation &c_element_op) override
std::unique_ptr< BaseInvoker >	MakeInvokerPointer () override
std::string	GetTypeString () const override
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_AK0_M_AK1 (const ConvToGemmFwdTransformer &conv_to_gemm_transformer)
template<typename BLay>
static auto	MakeBGridDescriptor_BK0_N_BK1 (const ConvToGemmFwdTransformer &conv_to_gemm_transformer)
template<typename CLay>
static auto	MakeCGridDescriptor_M_N (const ConvToGemmFwdTransformer &conv_to_gemm_transformer)
static bool	IsSupportedArgument (const Argument &arg)
static auto	MakeArgument (const void p_a, const void p_b, void *p_c, 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< index_t, NDimSpatial+3 > &c_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &c_g_n_k_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 CElementwiseOperation &c_element_op)
static auto	MakeInvoker ()

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	matrix_padder
static constexpr ConvToGemmFwdTransformer	dummy_conv_to_gemm_transformer

Member Typedef Documentation

◆ AGridDesc_AK0_M_AK1

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::AGridDesc_AK0_M_AK1

Initial value:

remove_cvref_t<decltype(MakeAGridDescriptor_AK0_M_AK1<ALayout>(

dummy_conv_to_gemm_transformer))>

ck::remove_cvref_t

remove_cv_t< remove_reference_t< T > > remove_cvref_t

Definition type.hpp:297

ck::tensor_operation::device::DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK::dummy_conv_to_gemm_transformer

static constexpr ConvToGemmFwdTransformer dummy_conv_to_gemm_transformer

Definition device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp:316

◆ AGridDesc_K0_M0_M1_K1

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::AGridDesc_K0_M0_M1_K1

Initial value:

decltype(GridwiseGemm::MakeAGridDescriptor_K0_M0_M1_K1(AGridDesc_AK0_M_AK1{}))

ck::GridwiseGemmDl_km_kn_mn_v1r3< BlockSize, ADataType, AccDataType, CDataType, InMemoryDataOperationEnum::Set, AGridDesc_AK0_M_AK1, BGridDesc_BK0_N_BK1, CGridDesc_M_N, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector >::MakeAGridDescriptor_K0_M0_M1_K1

__host__ static __device__ constexpr auto MakeAGridDescriptor_K0_M0_M1_K1(const AGridDesc_AK0_M_AK1 &a_grid_desc_k0_m_k1)

Definition gridwise_gemm_dl_v1r3.hpp:168

ck::tensor_operation::device::DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK::AGridDesc_AK0_M_AK1

remove_cvref_t< decltype(MakeAGridDescriptor_AK0_M_AK1< ALayout >( dummy_conv_to_gemm_transformer))> AGridDesc_AK0_M_AK1

Definition device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp:317

◆ BGridDesc_BK0_N_BK1

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::BGridDesc_BK0_N_BK1

Initial value:

remove_cvref_t<decltype(MakeBGridDescriptor_BK0_N_BK1<BLayout>(

dummy_conv_to_gemm_transformer))>

◆ BGridDesc_K0_N0_N1_K1

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::BGridDesc_K0_N0_N1_K1

Initial value:

decltype(GridwiseGemm::MakeBGridDescriptor_K0_N0_N1_K1(BGridDesc_BK0_N_BK1{}))

ck::GridwiseGemmDl_km_kn_mn_v1r3< BlockSize, ADataType, AccDataType, CDataType, InMemoryDataOperationEnum::Set, AGridDesc_AK0_M_AK1, BGridDesc_BK0_N_BK1, CGridDesc_M_N, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector >::MakeBGridDescriptor_K0_N0_N1_K1

__host__ static __device__ constexpr auto MakeBGridDescriptor_K0_N0_N1_K1(const BGridDesc_BK0_N_BK1 &b_grid_desc_k0_n_k1)

Definition gridwise_gemm_dl_v1r3.hpp:188

ck::tensor_operation::device::DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK::BGridDesc_BK0_N_BK1

remove_cvref_t< decltype(MakeBGridDescriptor_BK0_N_BK1< BLayout >( dummy_conv_to_gemm_transformer))> BGridDesc_BK0_N_BK1

Definition device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp:319

◆ CGridDesc_M0_M10_M11_N0_N10_N11

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::CGridDesc_M0_M10_M11_N0_N10_N11

Initial value:

decltype(GridwiseGemm::MakeCGridDescriptor_M0_M10_M11_N0_N10_N11(CGridDesc_M_N{}))

ck::GridwiseGemmDl_km_kn_mn_v1r3< BlockSize, ADataType, AccDataType, CDataType, InMemoryDataOperationEnum::Set, AGridDesc_AK0_M_AK1, BGridDesc_BK0_N_BK1, CGridDesc_M_N, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector >::MakeCGridDescriptor_M0_M10_M11_N0_N10_N11

__host__ static __device__ constexpr auto MakeCGridDescriptor_M0_M10_M11_N0_N10_N11(const CGridDesc_M_N &c_grid_desc_m_n)

Definition gridwise_gemm_dl_v1r3.hpp:208

ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK::CGridDesc_M_N

remove_cvref_t< decltype(MakeCGridDescriptor_M_N< CLayout >(dummy_conv_to_gemm_transformer))> CGridDesc_M_N

Definition device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp:303

◆ CGridDesc_M_N

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::CGridDesc_M_N

Initial value:

remove_cvref_t<decltype(MakeCGridDescriptor_M_N<CLayout>(dummy_conv_to_gemm_transformer))>

ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK::dummy_conv_to_gemm_transformer

static constexpr ConvToGemmFwdTransformer dummy_conv_to_gemm_transformer

Definition device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp:298

◆ ConvToGemmFwdTransformer

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::ConvToGemmFwdTransformer = TransformConvFwdToGemm<NDimSpatial, ConvForwardSpecialization>

◆ DefaultBlock2CTileMap

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::DefaultBlock2CTileMap

Initial value:

decltype(GridwiseGemm::MakeDefaultBlock2CTileMap(CGridDesc_M_N{}))

ck::GridwiseGemmDl_km_kn_mn_v1r3< BlockSize, ADataType, AccDataType, CDataType, InMemoryDataOperationEnum::Set, AGridDesc_AK0_M_AK1, BGridDesc_BK0_N_BK1, CGridDesc_M_N, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector >::MakeDefaultBlock2CTileMap

__host__ static __device__ constexpr auto MakeDefaultBlock2CTileMap(const CGridDesc_M_N &c_grid_desc_m_n)

Definition gridwise_gemm_dl_v1r3.hpp:241

◆ DeviceOp

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::DeviceOp = DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK

◆ GridwiseGemm

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

using ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::GridwiseGemm

Member Function Documentation

◆ GetTypeString()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

std::string ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::GetTypeString ( ) const

inlineoverridevirtual

Reimplemented from ck::tensor_operation::device::BaseOperator.

◆ IsSupportedArgument() [1/2]

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

bool ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::IsSupportedArgument ( const Argument & arg )

inlinestatic

◆ IsSupportedArgument() [2/2]

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

bool ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::IsSupportedArgument ( const BaseArgument * p_arg )

inlineoverridevirtual

Reimplemented from ck::tensor_operation::device::BaseOperator.

◆ MakeAGridDescriptor_AK0_M_AK1()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

template<typename ALay>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeAGridDescriptor_AK0_M_AK1 ( const ConvToGemmFwdTransformer & conv_to_gemm_transformer )

inlinestatic

◆ MakeArgument()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeArgument	(	const void *	p_a,
		const void *	p_b,
		void *	p_c,
		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< index_t, NDimSpatial+3 > &	c_g_n_k_wos_lengths,
		const std::array< index_t, NDimSpatial+3 > &	c_g_n_k_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 CElementwiseOperation &	c_element_op )

inlinestatic

◆ MakeArgumentPointer()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

std::unique_ptr< BaseArgument > ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeArgumentPointer	(	const void *	p_a,
		const void *	p_b,
		void *	p_c,
		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< index_t, NDimSpatial+3 > &	c_g_n_k_wos_lengths,
		const std::array< index_t, NDimSpatial+3 > &	c_g_n_k_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 CElementwiseOperation &	c_element_op )

inlineoverridevirtual

Implements ck::tensor_operation::device::DeviceGroupedConvFwd< NDimSpatial, ALayout, BLayout, CLayout, ADataType, BDataType, CDataType, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation >.

◆ MakeBGridDescriptor_BK0_N_BK1()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

template<typename BLay>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeBGridDescriptor_BK0_N_BK1 ( const ConvToGemmFwdTransformer & conv_to_gemm_transformer )

inlinestatic

◆ MakeCGridDescriptor_M_N()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

template<typename CLay>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeCGridDescriptor_M_N ( const ConvToGemmFwdTransformer & conv_to_gemm_transformer )

inlinestatic

◆ MakeInvoker()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeInvoker ( )

inlinestatic

◆ MakeInvokerPointer()

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

std::unique_ptr< BaseInvoker > ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::MakeInvokerPointer ( )

inlineoverridevirtual

Implements ck::tensor_operation::device::DeviceGroupedConvFwd< NDimSpatial, ALayout, BLayout, CLayout, ADataType, BDataType, CDataType, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation >.

Member Data Documentation

◆ dummy_conv_to_gemm_transformer

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

ConvToGemmFwdTransformer ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::dummy_conv_to_gemm_transformer

staticconstexpr

◆ I0

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::I0 = Number<0>{}

staticconstexpr

◆ I1

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::I1 = Number<1>{}

staticconstexpr

◆ I2

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::I2 = Number<2>{}

staticconstexpr

◆ I3

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::I3 = Number<3>{}

staticconstexpr

◆ matrix_padder

template<index_t NDimSpatial, typename ADataType, typename BDataType, typename CDataType, typename AccDataType, typename ALayout, typename BLayout, typename CLayout, typename AElementwiseOperation, typename BElementwiseOperation, typename CElementwiseOperation, ConvolutionForwardSpecialization ConvForwardSpecialization, GemmSpecialization GemmSpec, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t K0PerBlock, index_t K1, index_t M1PerThread, index_t N1PerThread, index_t KPerThread, typename M1N1ThreadClusterM1Xs, typename M1N1ThreadClusterN1Xs, typename ABlockTransferThreadSliceLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterLengths_K0_M0_M1_K1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, typename ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, typename ABlockTransferSrcVectorTensorContiguousDimOrder, typename ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, typename BBlockTransferThreadSliceLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterLengths_K0_N0_N1_K1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, typename BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, typename BBlockTransferSrcVectorTensorContiguousDimOrder, typename BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, typename CThreadTransferSrcDstAccessOrder, index_t CThreadTransferSrcDstVectorDim, index_t CThreadTransferDstScalarPerVector, enable_if_t< is_same_v< AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< BElementwiseOperation, ck::tensor_operation::element_wise::PassThrough > &&is_same_v< CElementwiseOperation, ck::tensor_operation::element_wise::PassThrough >, bool > = false>

auto ck::tensor_operation::device::DeviceGroupedConvFwdDl_NHWC_KYXC_NHWK< NDimSpatial, ADataType, BDataType, CDataType, AccDataType, ALayout, BLayout, CLayout, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, ConvForwardSpecialization, GemmSpec, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, M1PerThread, N1PerThread, KPerThread, M1N1ThreadClusterM1Xs, M1N1ThreadClusterN1Xs, ABlockTransferThreadSliceLengths_K0_M0_M1_K1, ABlockTransferThreadClusterLengths_K0_M0_M1_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorTensorLengths_K0_M0_M1_K1, ABlockTransferSrcVectorTensorContiguousDimOrder, ABlockTransferDstVectorTensorLengths_K0_M0_M1_K1, BBlockTransferThreadSliceLengths_K0_N0_N1_K1, BBlockTransferThreadClusterLengths_K0_N0_N1_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorTensorLengths_K0_N0_N1_K1, BBlockTransferSrcVectorTensorContiguousDimOrder, BBlockTransferDstVectorTensorLengths_K0_N0_N1_K1, CThreadTransferSrcDstAccessOrder, CThreadTransferSrcDstVectorDim, CThreadTransferDstScalarPerVector, >::matrix_padder

staticconstexpr

Initial value:

=

MatrixPadder<GemmSpec, index_t, index_t, index_t>{MPerBlock, NPerBlock, K0PerBlock}

ck::tensor_operation::device::MatrixPadder

Definition matrix_padder.hpp:180

The documentation for this struct was generated from the following file:

device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp

Classes

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ AGridDesc_AK0_M_AK1

◆ AGridDesc_K0_M0_M1_K1

◆ BGridDesc_BK0_N_BK1

◆ BGridDesc_K0_N0_N1_K1

◆ CGridDesc_M0_M10_M11_N0_N10_N11

◆ CGridDesc_M_N

◆ ConvToGemmFwdTransformer

◆ DefaultBlock2CTileMap

◆ DeviceOp

◆ GridwiseGemm

Member Function Documentation

◆ GetTypeString()

◆ IsSupportedArgument() [1/2]

◆ IsSupportedArgument() [2/2]

◆ MakeAGridDescriptor_AK0_M_AK1()

◆ MakeArgument()

◆ MakeArgumentPointer()

◆ MakeBGridDescriptor_BK0_N_BK1()

◆ MakeCGridDescriptor_M_N()

◆ MakeInvoker()

◆ MakeInvokerPointer()

Member Data Documentation

◆ dummy_conv_to_gemm_transformer

◆ I0

◆ I1

◆ I2

◆ I3

◆ matrix_padder