DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB > Struct Template Reference

• ck
• tensor_operation
• device
• DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle

Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members

ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB > Struct Template Reference

#include <device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp>

Inheritance diagram for ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >:

Classes
struct	Argument
struct	Invoker

Public Types
using	DeviceOp = DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle
using	ADataType = OutDataType
using	BDataType = InDataType
using	EDataType = WeiDataType
using	AElementwiseOperation = OutElementwiseOperation
using	BElementwiseOperation = InElementwiseOperation
using	CDEElementwiseOperation = WeiElementwiseOperation
using	ABDataType = InDataType
using	ABCGridDescs = decltype(GetABCGridDesc< NDimSpatial >())
using	AGridDesc_K0_M_K1 = remove_cvref_t< decltype(ABCGridDescs{}[I0])>
using	BGridDesc_K0_N_K1 = remove_cvref_t< decltype(ABCGridDescs{}[I1])>
using	CGridDesc_M_N = remove_cvref_t< decltype(ABCGridDescs{}[I2])>
using	GridwiseGemm = GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_bwd_weight< BlockSize, ADataType, BDataType, AccDataType, AccDataType, InMemoryDataOperationEnum::AtomicAdd, AGridDesc_K0_M_K1, BGridDesc_K0_N_K1, CGridDesc_M_N, AElementwiseOperation, BElementwiseOperation, element_wise::PassThrough, MPerBlock, NPerBlock, K0PerBlock, MPerXdl, NPerXdl, K1, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, ABlockLdsM1PerBlock, ABlockLdsM0PerBlock, ABlockLdsM1Padding, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, BBlockLdsN1PerBlock, BBlockLdsN0PerBlock, BBlockLdsN1Padding, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, WorkspaceInOutScalarPerVector, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, true, true, 1, PipelineVersion::v1, ComputeTypeA, ComputeTypeB >
using	DsGridDesc_M_N = decltype(MakeDsGridDescriptor_M_N< NDimSpatial >({}, {}))
using	CDGridDesc_M_N = decltype(concat_tuple(Tuple< CGridDesc_M_N >{}, DsGridDesc_M_N{}))
using	DsGridPointerTuple = decltype(GetDsGridPointerTuple())
using	CDDataTypes = decltype(concat_tuple(Tuple< const AccDataType * >{}, DsGridPointerTuple{}))
using	EGridDesc_M_N = CGridDesc_M_N
using	Block2TileMapElementwise = BlockToCTileMap_M00_N0_M01Adapt< MPerBlock, NPerBlock >
using	GridwiseElementwise = GridwiseElementwise< CDGridDesc_M_N, Tuple< EGridDesc_M_N >, CDDataTypes, Tuple< EDataType * >, Block2TileMapElementwise, CDEElementwiseOperation, BlockSize, MPerBlock, NPerBlock, MPerBlock/ClusterLengthMPerBlock, NPerBlock/ClusterLengthNPerBlock, Sequence< 0, 1 >, decltype(MakeElementwiseInputSequence()), Sequence< CBlockTransferScalarPerVector_NWaveNPerXdl >, I1, I1 >
using	CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock = decltype(GridwiseGemm::MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(CGridDesc_M_N{}))
using	Block2CTileMap = decltype(GridwiseGemm::MakeCBlockClusterAdaptor(CGridDesc_M_N{}, 1, 1, 1))

Public Member Functions
bool	IsSupportedArgument (const BaseArgument *p_arg) override
std::unique_ptr< BaseArgument >	MakeArgumentPointer (const void *p_in_grid, void *p_wei_grid, const void *p_out_grid, const std::array< const void *, NumDTensor > &p_ds, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_k_c_xs_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &input_left_pads, const std::array< ck::index_t, NDimSpatial > &input_right_pads, InElementwiseOperation in_element_op, WeiElementwiseOperation wei_element_op, OutElementwiseOperation out_element_op, const ck::index_t split_k) override
std::unique_ptr< BaseInvoker >	MakeInvokerPointer () override
std::string	GetTypeString () const override
size_t	GetWorkSpaceSize (const BaseArgument *p_arg) const override
void	SetWorkSpacePointer (BaseArgument *p_arg, void *p_workspace, const StreamConfig &=StreamConfig{}) 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	GetTypeIdName () const
virtual std::optional< std::string >	GetObjectName () const
virtual std::optional< std::string >	GetTemplateInfo () const
virtual std::string	GetTypeIdHashCode () const
virtual	~BaseOperator ()

Static Public Member Functions
template<ck::index_t NDim, typename ck::enable_if< NDim==1, bool >::type = false>
static auto	GetABCGridDesc ()
template<ck::index_t NDim, typename ck::enable_if< NDim==2, bool >::type = false>
static auto	GetABCGridDesc ()
template<ck::index_t NDim, typename ck::enable_if< NDim==3, bool >::type = false>
static auto	GetABCGridDesc ()
static constexpr auto	MakeElementwiseInputSequence ()
static constexpr auto	GetDsGridPointerTuple ()
template<index_t NDim, typename ck::enable_if< NDim==1, bool >::type = false>
static auto	MakeDsGridDescriptor_M_N (const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_strides)
template<index_t NDim, typename ck::enable_if< NDim==2, bool >::type = false>
static auto	MakeDsGridDescriptor_M_N (const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_strides)
template<index_t NDim, typename ck::enable_if< NDim==3, bool >::type = false>
static auto	MakeDsGridDescriptor_M_N (const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &ds_g_k_c_xs_strides)
template<typename ComputePtrOffsetOfBatch >
static void	InitElementwiseBatchStrides (const ComputePtrOffsetOfBatch &compute_ptr_offset_of_batch_, std::array< index_t, NumDTensor+I1 > &input_batch_strides, std::array< index_t, I1 > &output_batch_strides)
static constexpr bool	IsValidCompilationParameter ()
static bool	IsSupportedArgument (const Argument &arg)
static auto	MakeArgument (const InDataType *p_in_grid, WeiDataType *p_wei_grid, const OutDataType *p_out_grid, const std::array< const void *, NumDTensor > &p_ds, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &ds_g_k_c_xs_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &input_left_pads, const std::array< ck::index_t, NDimSpatial > &input_right_pads, InElementwiseOperation in_element_op, WeiElementwiseOperation wei_element_op, OutElementwiseOperation out_element_op, const ck::index_t split_k)
static auto	MakeInvoker ()

Static Public Attributes
static constexpr index_t	NumDTensor = DsLayout::Size()
static constexpr auto	I0 = Number<0>{}
static constexpr auto	I1 = Number<1>{}
static constexpr auto	I2 = Number<2>{}
static constexpr auto	I3 = Number<3>{}
static constexpr auto	I4 = Number<4>{}
static constexpr auto	I5 = Number<5>{}
static constexpr auto	K1Number = Number<K1>{}
static constexpr auto	conv_to_gemm_transformer
static constexpr index_t	MaxScalarPerVectorFP32 = 4
static constexpr index_t	WorkspaceInOutScalarPerVector
static constexpr auto	BankLength = 128
static constexpr auto	ElePerBank = BankLength / sizeof(ADataType)
static constexpr auto	ABlockLdsM1PerBlock = ElePerBank / K1
static constexpr auto	ABlockLdsM0PerBlock = MPerBlock / ABlockLdsM1PerBlock
static constexpr auto	ABlockLdsM1Padding = 4
static constexpr auto	BBlockLdsN1PerBlock = ElePerBank / K1
static constexpr auto	BBlockLdsN0PerBlock = NPerBlock / BBlockLdsN1PerBlock
static constexpr auto	BBlockLdsN1Padding = 4
static constexpr index_t	ClusterLengthMPerBlock
static constexpr index_t	ClusterLengthNPerBlock
Static Public Attributes inherited from ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, InDataType, InDataType >
static constexpr index_t	NumDTensor

Member Typedef Documentation

◆ ABCGridDescs

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ABCGridDescs = decltype(GetABCGridDesc<NDimSpatial>())

◆ ABDataType

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ABDataType = InDataType

◆ ADataType

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ADataType = OutDataType

◆ AElementwiseOperation

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::AElementwiseOperation = OutElementwiseOperation

◆ AGridDesc_K0_M_K1

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::AGridDesc_K0_M_K1 = remove_cvref_t<decltype(ABCGridDescs{}[I0])>

◆ BDataType

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BDataType = InDataType

◆ BElementwiseOperation

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BElementwiseOperation = InElementwiseOperation

◆ BGridDesc_K0_N_K1

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BGridDesc_K0_N_K1 = remove_cvref_t<decltype(ABCGridDescs{}[I1])>

◆ Block2CTileMap

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::Block2CTileMap = decltype(GridwiseGemm::MakeCBlockClusterAdaptor(CGridDesc_M_N{}, 1, 1, 1))

◆ Block2TileMapElementwise

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::Block2TileMapElementwise = BlockToCTileMap_M00_N0_M01Adapt<MPerBlock, NPerBlock>

◆ CDDataTypes

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::CDDataTypes = decltype(concat_tuple(Tuple<const AccDataType*>{}, DsGridPointerTuple{}))

◆ CDEElementwiseOperation

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::CDEElementwiseOperation = WeiElementwiseOperation

◆ CDGridDesc_M_N

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::CDGridDesc_M_N = decltype(concat_tuple(Tuple<CGridDesc_M_N>{}, DsGridDesc_M_N{}))

◆ CGridDesc_M_N

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::CGridDesc_M_N = remove_cvref_t<decltype(ABCGridDescs{}[I2])>

◆ CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock = decltype(GridwiseGemm::MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(CGridDesc_M_N{}))

◆ DeviceOp

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::DeviceOp = DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle

◆ DsGridDesc_M_N

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::DsGridDesc_M_N = decltype(MakeDsGridDescriptor_M_N<NDimSpatial>({}, {}))

◆ DsGridPointerTuple

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::DsGridPointerTuple = decltype(GetDsGridPointerTuple())

◆ EDataType

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::EDataType = WeiDataType

◆ EGridDesc_M_N

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::EGridDesc_M_N = CGridDesc_M_N

◆ GridwiseElementwise

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GridwiseElementwise = GridwiseElementwise<CDGridDesc_M_N, Tuple<EGridDesc_M_N>, CDDataTypes, Tuple<EDataType*>, Block2TileMapElementwise, CDEElementwiseOperation, BlockSize, MPerBlock, NPerBlock, MPerBlock / ClusterLengthMPerBlock, NPerBlock / ClusterLengthNPerBlock, Sequence<0, 1>, decltype(MakeElementwiseInputSequence()), Sequence<CBlockTransferScalarPerVector_NWaveNPerXdl>, I1, I1>

◆ GridwiseGemm

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

using ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GridwiseGemm = GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_bwd_weight< BlockSize, ADataType, BDataType, AccDataType, AccDataType, InMemoryDataOperationEnum::AtomicAdd, AGridDesc_K0_M_K1, BGridDesc_K0_N_K1, CGridDesc_M_N, AElementwiseOperation, BElementwiseOperation, element_wise::PassThrough, MPerBlock, NPerBlock, K0PerBlock, MPerXdl, NPerXdl, K1, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, ABlockLdsM1PerBlock, ABlockLdsM0PerBlock, ABlockLdsM1Padding, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, BBlockLdsN1PerBlock, BBlockLdsN0PerBlock, BBlockLdsN1Padding, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, WorkspaceInOutScalarPerVector, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, true, true, 1, PipelineVersion::v1, ComputeTypeA, ComputeTypeB>

Member Function Documentation

◆ GetABCGridDesc() [1/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<ck::index_t NDim, typename ck::enable_if< NDim==1, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetABCGridDesc ( )

inlinestatic

◆ GetABCGridDesc() [2/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<ck::index_t NDim, typename ck::enable_if< NDim==2, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetABCGridDesc ( )

inlinestatic

◆ GetABCGridDesc() [3/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<ck::index_t NDim, typename ck::enable_if< NDim==3, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetABCGridDesc ( )

inlinestatic

◆ GetDsGridPointerTuple()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetDsGridPointerTuple ( )

inlinestaticconstexpr

◆ GetTypeString()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

std::string ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetTypeString ( ) const

inlineoverridevirtual

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

◆ GetWorkSpaceSize()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

size_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::GetWorkSpaceSize ( const BaseArgument *  p_arg ) const

inlineoverridevirtual

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

◆ InitElementwiseBatchStrides()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<typename ComputePtrOffsetOfBatch >

static void ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::InitElementwiseBatchStrides ( const ComputePtrOffsetOfBatch &  compute_ptr_offset_of_batch_, std::array< index_t, NumDTensor+I1 > &  input_batch_strides, std::array< index_t, I1 > &  output_batch_strides  )

inlinestatic

◆ IsSupportedArgument() [1/2]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static bool ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::IsSupportedArgument ( const Argument &  arg )

inlinestatic

◆ IsSupportedArgument() [2/2]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

bool ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::IsSupportedArgument ( const BaseArgument *  p_arg )

inlineoverridevirtual

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

◆ IsValidCompilationParameter()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static constexpr bool ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::IsValidCompilationParameter ( )

inlinestaticconstexpr

◆ MakeArgument()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeArgument ( const InDataType *  p_in_grid, WeiDataType *  p_wei_grid, const OutDataType *  p_out_grid, const std::array< const void *, NumDTensor > &  p_ds, const std::array< index_t, NDimSpatial+3 > &  b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &  b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &  e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &  e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &  a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &  a_g_n_k_wos_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &  ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &  ds_g_k_c_xs_strides, const std::array< ck::index_t, NDimSpatial > &  conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &  conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &  input_left_pads, const std::array< ck::index_t, NDimSpatial > &  input_right_pads, InElementwiseOperation  in_element_op, WeiElementwiseOperation  wei_element_op, OutElementwiseOperation  out_element_op, const ck::index_t  split_k  )

inlinestatic

◆ MakeArgumentPointer()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

std::unique_ptr<BaseArgument> ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeArgumentPointer ( const void *  p_in_grid, void *  p_wei_grid, const void *  p_out_grid, const std::array< const void *, NumDTensor > &  p_ds, const std::array< index_t, NDimSpatial+3 > &  b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &  b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &  e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &  e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &  a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &  a_g_n_k_wos_strides, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &  ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDimSpatial+3 >, NumDTensor > &  ds_g_k_c_xs_strides, const std::array< ck::index_t, NDimSpatial > &  conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &  conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &  input_left_pads, const std::array< ck::index_t, NDimSpatial > &  input_right_pads, InElementwiseOperation  in_element_op, WeiElementwiseOperation  wei_element_op, OutElementwiseOperation  out_element_op, const ck::index_t  split_k  )

inlineoverridevirtual

Implements ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, InDataType, InDataType >.

◆ MakeDsGridDescriptor_M_N() [1/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<index_t NDim, typename ck::enable_if< NDim==1, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeDsGridDescriptor_M_N ( const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_strides  )

inlinestatic

◆ MakeDsGridDescriptor_M_N() [2/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<index_t NDim, typename ck::enable_if< NDim==2, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeDsGridDescriptor_M_N ( const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_strides  )

inlinestatic

◆ MakeDsGridDescriptor_M_N() [3/3]

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

template<index_t NDim, typename ck::enable_if< NDim==3, bool >::type = false>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeDsGridDescriptor_M_N ( const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_lengths, const std::array< std::array< index_t, NDim+3 >, NumDTensor > &  ds_g_k_c_xs_strides  )

inlinestatic

◆ MakeElementwiseInputSequence()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeElementwiseInputSequence ( )

inlinestaticconstexpr

◆ MakeInvoker()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

static auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeInvoker ( )

inlinestatic

◆ MakeInvokerPointer()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

std::unique_ptr<BaseInvoker> ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MakeInvokerPointer ( )

inlineoverridevirtual

Implements ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, InDataType, InDataType >.

◆ SetWorkSpacePointer()

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

void ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::SetWorkSpacePointer ( BaseArgument *  p_arg, void *  p_workspace, const StreamConfig &  = StreamConfig{}  ) const

inlineoverridevirtual

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

Member Data Documentation

◆ ABlockLdsM0PerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ABlockLdsM0PerBlock = MPerBlock / ABlockLdsM1PerBlock

staticconstexpr

◆ ABlockLdsM1Padding

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ABlockLdsM1Padding = 4

staticconstexpr

◆ ABlockLdsM1PerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ABlockLdsM1PerBlock = ElePerBank / K1

staticconstexpr

◆ BankLength

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BankLength = 128

staticconstexpr

◆ BBlockLdsN0PerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BBlockLdsN0PerBlock = NPerBlock / BBlockLdsN1PerBlock

staticconstexpr

◆ BBlockLdsN1Padding

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BBlockLdsN1Padding = 4

staticconstexpr

◆ BBlockLdsN1PerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::BBlockLdsN1PerBlock = ElePerBank / K1

staticconstexpr

◆ ClusterLengthMPerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr index_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ClusterLengthMPerBlock

staticconstexpr

Initial value:

=

CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock::At(1)

◆ ClusterLengthNPerBlock

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr index_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ClusterLengthNPerBlock

staticconstexpr

Initial value:

=

CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock::At(3)

◆ conv_to_gemm_transformer

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::conv_to_gemm_transformer

staticconstexpr

Initial value:

=

TransformConvBwdWeightToGemm<NDimSpatial,

MPerBlock,

NPerBlock,

K1Number,

K0PerBlock,

ConvBackwardWeightSpecialization>{}

ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle::K1Number

static constexpr auto K1Number

Definition: device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp:189

◆ ElePerBank

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::ElePerBank = BankLength / sizeof(ADataType)

staticconstexpr

◆ I0

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I0 = Number<0>{}

staticconstexpr

◆ I1

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I1 = Number<1>{}

staticconstexpr

◆ I2

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I2 = Number<2>{}

staticconstexpr

◆ I3

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I3 = Number<3>{}

staticconstexpr

◆ I4

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I4 = Number<4>{}

staticconstexpr

◆ I5

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::I5 = Number<5>{}

staticconstexpr

◆ K1Number

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr auto ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::K1Number = Number<K1>{}

staticconstexpr

◆ MaxScalarPerVectorFP32

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr index_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::MaxScalarPerVectorFP32 = 4

staticconstexpr

◆ NumDTensor

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr index_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::NumDTensor = DsLayout::Size()

staticconstexpr

◆ WorkspaceInOutScalarPerVector

template<index_t NDimSpatial, typename InLayout , typename WeiLayout , typename OutLayout , typename DsLayout , typename InDataType , typename WeiDataType , typename OutDataType , typename AccDataType , typename DsDataType , typename InElementwiseOperation , typename WeiElementwiseOperation , typename OutElementwiseOperation , ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization, ck::index_t BlockSize, ck::index_t MPerBlock, ck::index_t NPerBlock, ck::index_t K0PerBlock, ck::index_t K1, ck::index_t MPerXdl, ck::index_t NPerXdl, ck::index_t MXdlPerWave, ck::index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_K0_M_K1 , typename ABlockTransferThreadClusterArrangeOrder , typename ABlockTransferSrcAccessOrder , ck::index_t ABlockTransferSrcVectorDim, ck::index_t ABlockTransferSrcScalarPerVector, ck::index_t ABlockTransferDstScalarPerVector_K1, bool ABlockLdsAddExtraM, typename BBlockTransferThreadClusterLengths_K0_N_K1 , typename BBlockTransferThreadClusterArrangeOrder , typename BBlockTransferSrcAccessOrder , ck::index_t BBlockTransferSrcVectorDim, ck::index_t BBlockTransferSrcScalarPerVector, ck::index_t BBlockTransferDstScalarPerVector_K1, bool BBlockLdsAddExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , index_t CBlockTransferScalarPerVector_NWaveNPerXdl, typename ComputeTypeA = InDataType, typename ComputeTypeB = ComputeTypeA>

constexpr index_t ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle< NDimSpatial, InLayout, WeiLayout, OutLayout, DsLayout, InDataType, WeiDataType, OutDataType, AccDataType, DsDataType, InElementwiseOperation, WeiElementwiseOperation, OutElementwiseOperation, ConvBackwardWeightSpecialization, BlockSize, MPerBlock, NPerBlock, K0PerBlock, K1, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, ComputeTypeA, ComputeTypeB >::WorkspaceInOutScalarPerVector

staticconstexpr

Initial value:

=

is_same_v<AccDataType, float>

? math::min(CBlockTransferScalarPerVector_NWaveNPerXdl, MaxScalarPerVectorFP32)

: CBlockTransferScalarPerVector_NWaveNPerXdl

ck::math::min

__host__ constexpr __device__ T min(T x)

Definition: math.hpp:116

ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle::MaxScalarPerVectorFP32

static constexpr index_t MaxScalarPerVectorFP32

Definition: device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp:199

---

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

• /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.1/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp