GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ > Struct Template Reference

GroupedConvTraits&lt; NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ > Struct Template Reference
Classes | Public Types | Static Public Attributes | List of all members
ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ > Struct Template Reference

#include <grouped_convolution_utils.hpp>

Classes

struct  FixedGemmParams
 

Public Types

using InLayout = InLayout_
 
using WeiLayout = WeiLayout_
 
using DsLayout = DsLayout_
 
using OutLayout = OutLayout_
 
using AsLayoutFwd = ck_tile::tensor_layout::gemm::RowMajor
 
using BsLayoutFwd = ck_tile::tensor_layout::gemm::ColumnMajor
 
using CLayoutFwd = ck_tile::tensor_layout::gemm::RowMajor
 
using AsLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor
 
using BsLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor
 
using CLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor
 
using AsLayoutBwdWeight = ck_tile::tensor_layout::gemm::ColumnMajor
 
using BsLayoutBwdWeight = ck_tile::tensor_layout::gemm::RowMajor
 
using CLayoutBwdWeight = ck_tile::tensor_layout::gemm::RowMajor
 
template<ck_tile::index_t NumWaveGroups = 1>
using GroupedConvImplicitGemmTraitsFwd = TileGemmTraits< true, true, true, AsLayoutFwd, BsLayoutFwd, CLayoutFwd, NumWaveGroups >
 
template<ck_tile::index_t NumWaveGroups = 1>
using GroupedConvImplicitGemmTraitsBwdData = TileGemmTraits< true, true, true, AsLayoutBwdData, BsLayoutBwdData, CLayoutBwdData, NumWaveGroups >
 
template<ck_tile::index_t NumWaveGroups = 1>
using GroupedConvImplicitGemmTraitsBwdWeight = TileGemmTraits< true, true, true, AsLayoutBwdWeight, BsLayoutBwdWeight, CLayoutBwdWeight, NumWaveGroups >
 
using ImplicitGemmDsLayout = decltype(generate_implicit_gemm_layout())
 

Static Public Attributes

static constexpr bool EnableSplitImage = EnableSplitImage_
 
static constexpr index_t NumGroupsToMerge = NumGroupsToMerge_
 
static constexpr index_t NDimSpatial = NDimSpatial_
 
static constexpr ConvolutionSpecialization ConvSpecialization = ConvSpecialization_
 
static constexpr ck_tile::index_t VectorSizeA = VectorSizeA_
 
static constexpr ck_tile::index_t VectorSizeB = VectorSizeB_
 
static constexpr ck_tile::index_t VectorSizeC = VectorSizeC_
 
static constexpr ck_tile::index_t NumDTensor = DsLayout::size()
 

Member Typedef Documentation

◆ AsLayoutBwdData

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::AsLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor

◆ AsLayoutBwdWeight

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::AsLayoutBwdWeight = ck_tile::tensor_layout::gemm::ColumnMajor

◆ AsLayoutFwd

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::AsLayoutFwd = ck_tile::tensor_layout::gemm::RowMajor

◆ BsLayoutBwdData

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::BsLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor

◆ BsLayoutBwdWeight

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::BsLayoutBwdWeight = ck_tile::tensor_layout::gemm::RowMajor

◆ BsLayoutFwd

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::BsLayoutFwd = ck_tile::tensor_layout::gemm::ColumnMajor

◆ CLayoutBwdData

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::CLayoutBwdData = ck_tile::tensor_layout::gemm::RowMajor

◆ CLayoutBwdWeight

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::CLayoutBwdWeight = ck_tile::tensor_layout::gemm::RowMajor

◆ CLayoutFwd

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::CLayoutFwd = ck_tile::tensor_layout::gemm::RowMajor

◆ DsLayout

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::DsLayout = DsLayout_

◆ GroupedConvImplicitGemmTraitsBwdData

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
template<ck_tile::index_t NumWaveGroups = 1>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::GroupedConvImplicitGemmTraitsBwdData = TileGemmTraits<true, true, true, AsLayoutBwdData, BsLayoutBwdData, CLayoutBwdData, NumWaveGroups>

◆ GroupedConvImplicitGemmTraitsBwdWeight

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
template<ck_tile::index_t NumWaveGroups = 1>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::GroupedConvImplicitGemmTraitsBwdWeight = TileGemmTraits<true, true, true, AsLayoutBwdWeight, BsLayoutBwdWeight, CLayoutBwdWeight, NumWaveGroups>

◆ GroupedConvImplicitGemmTraitsFwd

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
template<ck_tile::index_t NumWaveGroups = 1>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::GroupedConvImplicitGemmTraitsFwd = TileGemmTraits<true, true, true, AsLayoutFwd, BsLayoutFwd, CLayoutFwd, NumWaveGroups>

◆ ImplicitGemmDsLayout

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::ImplicitGemmDsLayout = decltype(generate_implicit_gemm_layout())

◆ InLayout

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::InLayout = InLayout_

◆ OutLayout

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::OutLayout = OutLayout_

◆ WeiLayout

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
using ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::WeiLayout = WeiLayout_

Member Data Documentation

◆ ConvSpecialization

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr ConvolutionSpecialization ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::ConvSpecialization = ConvSpecialization_
staticconstexpr

◆ EnableSplitImage

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr bool ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::EnableSplitImage = EnableSplitImage_
staticconstexpr

◆ NDimSpatial

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::NDimSpatial = NDimSpatial_
staticconstexpr

◆ NumDTensor

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr ck_tile::index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::NumDTensor = DsLayout::size()
staticconstexpr

◆ NumGroupsToMerge

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::NumGroupsToMerge = NumGroupsToMerge_
staticconstexpr

◆ VectorSizeA

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr ck_tile::index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::VectorSizeA = VectorSizeA_
staticconstexpr

◆ VectorSizeB

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr ck_tile::index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::VectorSizeB = VectorSizeB_
staticconstexpr

◆ VectorSizeC

template<index_t NDimSpatial_, ConvolutionSpecialization ConvSpecialization_, typename InLayout_ , typename WeiLayout_ , typename DsLayout_ , typename OutLayout_ , index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1, index_t VectorSizeC_ = 1, index_t NumGroupsToMerge_ = 1, bool EnableSplitImage_ = false>
constexpr ck_tile::index_t ck_tile::GroupedConvTraits< NDimSpatial_, ConvSpecialization_, InLayout_, WeiLayout_, DsLayout_, OutLayout_, VectorSizeA_, VectorSizeB_, VectorSizeC_, NumGroupsToMerge_, EnableSplitImage_ >::VectorSizeC = VectorSizeC_
staticconstexpr
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp