TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType > Struct Template Reference

TransformConvFwdToGemm&lt; NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType &gt; Struct Template Reference#

Composable Kernel: ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType > Struct Template Reference
Classes | Public Member Functions | Static Public Member Functions | List of all members
ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType > Struct Template Reference

#include <transform_conv_fwd_to_gemm.hpp>

Classes

struct  SplitImageInfo

Public Member Functions

CK_TILE_HOST constexpr IndexType GetN () const
CK_TILE_HOST constexpr IndexType GetOriginalN () const
CK_TILE_HOST constexpr TransformConvFwdToGemm ()
template<typename TransformConvFwdToGemmBase>
CK_TILE_HOST TransformConvFwdToGemm (const TransformConvFwdToGemmBase &transform_conv_fwd_to_gemm_base)
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==1, bool >::type = false>
CK_TILE_HOST TransformConvFwdToGemm (const ConvDimsType &a_g_n_c_wis_lengths, const ConvDimsType &b_g_k_c_xs_lengths, const ConvDimsType &c_g_n_k_wos_lengths, const ConvSpatialDimsType &conv_filter_strides, const ConvSpatialDimsType &conv_filter_dilations, const ConvSpatialDimsType &input_left_pads, const ConvSpatialDimsType &input_right_pads)
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==2, bool >::type = false>
CK_TILE_HOST TransformConvFwdToGemm (const ConvDimsType &a_g_n_c_wis_lengths, const ConvDimsType &b_g_k_c_xs_lengths, const ConvDimsType &c_g_n_k_wos_lengths, const ConvSpatialDimsType &conv_filter_strides, const ConvSpatialDimsType &conv_filter_dilations, const ConvSpatialDimsType &input_left_pads, const ConvSpatialDimsType &input_right_pads)
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==3, bool >::type = false>
CK_TILE_HOST TransformConvFwdToGemm (const ConvDimsType &a_g_n_c_wis_lengths, const ConvDimsType &b_g_k_c_xs_lengths, const ConvDimsType &c_g_n_k_wos_lengths, const ConvSpatialDimsType &conv_filter_strides, const ConvSpatialDimsType &conv_filter_dilations, const ConvSpatialDimsType &input_left_pads, const ConvSpatialDimsType &input_right_pads)
CK_TILE_HOST bool AreDescriptorsSmallerThan2GB () const
template<typename ALayout, typename std::enable_if< NDimSpatial==1 &&std::is_same_v< ALayout, tensor_layout::convolution::NWGC >, bool >::type = false>
CK_TILE_HOST auto MakeADescriptor_M_K () const
template<typename ALayout, typename std::enable_if< NDimSpatial==2 &&std::is_same_v< ALayout, tensor_layout::convolution::NHWGC >, bool >::type = false>
CK_TILE_HOST auto MakeADescriptor_M_K () const
template<typename ALayout, typename std::enable_if< NDimSpatial==3 &&std::is_same_v< ALayout, tensor_layout::convolution::NDHWGC >, bool >::type = false>
CK_TILE_HOST auto MakeADescriptor_M_K () const
template<typename BLayout, typename std::enable_if< std::is_same_v< BLayout, tensor_layout::convolution::GKXC >||std::is_same_v< BLayout, tensor_layout::convolution::GKYXC >||std::is_same_v< BLayout, tensor_layout::convolution::GKZYXC >, bool >::type = false>
CK_TILE_HOST auto MakeBDescriptor_N_K () const
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==1 &&std::is_same_v< CLayout, tensor_layout::convolution::NWGK >, bool >::type = false>
CK_TILE_HOST auto MakeCDescriptor_M_N () const
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==2 &&std::is_same_v< CLayout, tensor_layout::convolution::NHWGK >, bool >::type = false>
CK_TILE_HOST auto MakeCDescriptor_M_N () const
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==3 &&std::is_same_v< CLayout, tensor_layout::convolution::NDHWGK >, bool >::type = false>
CK_TILE_HOST auto MakeCDescriptor_M_N () const

Static Public Member Functions

static CK_TILE_HOST SplitImageInfo GetSplitImageInfo (index_t G, index_t N, index_t C, index_t K, index_t D_out, index_t H_out, index_t W_out)

Constructor & Destructor Documentation

◆ TransformConvFwdToGemm() [1/5]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
CK_TILE_HOST constexpr ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::TransformConvFwdToGemm ( )
inlineconstexpr

◆ TransformConvFwdToGemm() [2/5]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename TransformConvFwdToGemmBase>
CK_TILE_HOST ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::TransformConvFwdToGemm ( const TransformConvFwdToGemmBase & transform_conv_fwd_to_gemm_base)
inline

◆ TransformConvFwdToGemm() [3/5]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==1, bool >::type = false>
CK_TILE_HOST ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::TransformConvFwdToGemm ( const ConvDimsType & a_g_n_c_wis_lengths,
const ConvDimsType & b_g_k_c_xs_lengths,
const ConvDimsType & c_g_n_k_wos_lengths,
const ConvSpatialDimsType & conv_filter_strides,
const ConvSpatialDimsType & conv_filter_dilations,
const ConvSpatialDimsType & input_left_pads,
const ConvSpatialDimsType & input_right_pads )
inline

◆ TransformConvFwdToGemm() [4/5]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==2, bool >::type = false>
CK_TILE_HOST ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::TransformConvFwdToGemm ( const ConvDimsType & a_g_n_c_wis_lengths,
const ConvDimsType & b_g_k_c_xs_lengths,
const ConvDimsType & c_g_n_k_wos_lengths,
const ConvSpatialDimsType & conv_filter_strides,
const ConvSpatialDimsType & conv_filter_dilations,
const ConvSpatialDimsType & input_left_pads,
const ConvSpatialDimsType & input_right_pads )
inline

◆ TransformConvFwdToGemm() [5/5]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ConvDimsType, typename ConvSpatialDimsType, index_t NDim = NDimSpatial, typename std::enable_if< NDim==3, bool >::type = false>
CK_TILE_HOST ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::TransformConvFwdToGemm ( const ConvDimsType & a_g_n_c_wis_lengths,
const ConvDimsType & b_g_k_c_xs_lengths,
const ConvDimsType & c_g_n_k_wos_lengths,
const ConvSpatialDimsType & conv_filter_strides,
const ConvSpatialDimsType & conv_filter_dilations,
const ConvSpatialDimsType & input_left_pads,
const ConvSpatialDimsType & input_right_pads )
inline

Member Function Documentation

◆ AreDescriptorsSmallerThan2GB()

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
CK_TILE_HOST bool ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::AreDescriptorsSmallerThan2GB ( ) const
inline

◆ GetN()

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
CK_TILE_HOST constexpr IndexType ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::GetN ( ) const
inlineconstexpr

◆ GetOriginalN()

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
CK_TILE_HOST constexpr IndexType ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::GetOriginalN ( ) const
inlineconstexpr

◆ GetSplitImageInfo()

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
CK_TILE_HOST SplitImageInfo ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::GetSplitImageInfo ( index_t G,
index_t N,
index_t C,
index_t K,
index_t D_out,
index_t H_out,
index_t W_out )
inlinestatic

◆ MakeADescriptor_M_K() [1/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ALayout, typename std::enable_if< NDimSpatial==3 &&std::is_same_v< ALayout, tensor_layout::convolution::NDHWGC >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeADescriptor_M_K ( ) const
inline

◆ MakeADescriptor_M_K() [2/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ALayout, typename std::enable_if< NDimSpatial==2 &&std::is_same_v< ALayout, tensor_layout::convolution::NHWGC >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeADescriptor_M_K ( ) const
inline

◆ MakeADescriptor_M_K() [3/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename ALayout, typename std::enable_if< NDimSpatial==1 &&std::is_same_v< ALayout, tensor_layout::convolution::NWGC >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeADescriptor_M_K ( ) const
inline

◆ MakeBDescriptor_N_K()

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename BLayout, typename std::enable_if< std::is_same_v< BLayout, tensor_layout::convolution::GKXC >||std::is_same_v< BLayout, tensor_layout::convolution::GKYXC >||std::is_same_v< BLayout, tensor_layout::convolution::GKZYXC >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeBDescriptor_N_K ( ) const
inline

◆ MakeCDescriptor_M_N() [1/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==3 &&std::is_same_v< CLayout, tensor_layout::convolution::NDHWGK >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeCDescriptor_M_N ( ) const
inline

◆ MakeCDescriptor_M_N() [2/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==2 &&std::is_same_v< CLayout, tensor_layout::convolution::NHWGK >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeCDescriptor_M_N ( ) const
inline

◆ MakeCDescriptor_M_N() [3/3]

template<index_t NDimSpatial, ConvolutionSpecialization ConvSpecialization, index_t VectorSizeA, index_t VectorSizeB, index_t VectorSizeC, index_t NumGroupsToMerge = 1, bool SplitN = false, typename ADataType = float, typename CDataType = float, typename IndexType = index_t>
template<typename CLayout, index_t NDimSp = NDimSpatial, typename std::enable_if< NDimSp==1 &&std::is_same_v< CLayout, tensor_layout::convolution::NWGK >, bool >::type = false>
CK_TILE_HOST auto ck_tile::TransformConvFwdToGemm< NDimSpatial, ConvSpecialization, VectorSizeA, VectorSizeB, VectorSizeC, NumGroupsToMerge, SplitN, ADataType, CDataType, IndexType >::MakeCDescriptor_M_N ( ) const
inline
The documentation for this struct was generated from the following file: