GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ > Struct Template Reference

GroupedConvBwdDataKernelArgs&lt; GroupedConvTraitsType_, TilePartitioner_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ > Struct Template Reference
Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ > Struct Template Reference

The Grouped Convolution kernel device arguments. More...

#include <grouped_convolution_backward_data_kernel.hpp>

Public Types

using TilePartitioner = remove_cvref_t<TilePartitioner_>
using ConvToGemmTransformer
using ABCGridDescs
using AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>
using BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>
using CGridDescMN = remove_cvref_t<decltype(ABCGridDescs{}[number<2>{}])>

Public Member Functions

template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NWGK >, bool >::type = false>
CK_TILE_HOST GroupedConvBwdDataKernelArgs (const GroupedConvBwdDataHostArgs &args)
template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NHWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKYXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NHWGK >, bool >::type = false>
CK_TILE_HOST GroupedConvBwdDataKernelArgs (const GroupedConvBwdDataHostArgs &args)
template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NDHWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKZYXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NDHWGK >, bool >::type = false>
CK_TILE_HOST GroupedConvBwdDataKernelArgs (const GroupedConvBwdDataHostArgs &args)

Public Attributes

array< index_t, NonSpatialDims+GroupedConvTraitsType_::NDimSpatial > in_g_n_c_wis_lengths
array< index_t, NonSpatialDims+GroupedConvTraitsType_::NDimSpatial > wei_g_k_c_xs_lengths
array< index_t, NonSpatialDims+GroupedConvTraitsType_::NDimSpatial > out_g_n_k_wos_lengths
array< index_t, GroupedConvTraitsType_::NDimSpatial > conv_filter_strides
array< index_t, GroupedConvTraitsType_::NDimSpatial > conv_filter_dilations
array< index_t, GroupedConvTraitsType_::NDimSpatial > input_left_pads
array< index_t, GroupedConvTraitsType_::NDimSpatial > input_right_pads
array< index_t, GroupedConvTraitsType_::NDimSpatial > tildes
index_t k_batch
index_t GemmBatch
index_t grid_size_ = 0
index_t gemm_count = 0
const void * out_ptr
void * in_ptr
std::array< const void *, NumDTensor > ds_ptr
const void * wei_ptr
array< AGridDescMK, MaxGroupedGemmGroupsNum > a_grid_descs_m_k
array< BGridDescNK, MaxGroupedGemmGroupsNum > b_grid_descs_n_k
array< CGridDescMN, MaxGroupedGemmGroupsNum > c_grid_descs_m_n
array< index_t, MaxGroupedGemmGroupsNum > block_starts
array< index_t, MaxGroupedGemmGroupsNum > block_ends
long_index_t group_stride_a
long_index_t group_stride_b
long_index_t group_stride_c
index_t n_splits = 1
index_t n_per_split = 1
index_t original_n = 1
index_t input_batch_stride = 0
index_t output_batch_stride = 0

Static Public Attributes

static constexpr index_t NumDTensor = GroupedConvTraitsType_::NumDTensor
static constexpr auto I0 = number<0>()
static constexpr auto I1 = number<1>()
static constexpr index_t MaxGroupedGemmGroupsNum = 128
static constexpr index_t NonSpatialDims = 3

Detailed Description

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
struct ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >

The Grouped Convolution kernel device arguments.

Member Typedef Documentation

◆ ABCGridDescs

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::ABCGridDescs
Initial value:
remove_cvref_t<
decltype(ConvToGemmTransformer{}.MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(1))>
ck::remove_cvref_t
remove_cv_t< remove_reference_t< T > > remove_cvref_t
Definition type.hpp:297
ck_tile::GroupedConvBwdDataKernelArgs::ConvToGemmTransformer
TransformConvBwdDataToGemm< GroupedConvTraitsType_::NDimSpatial, GroupedConvTraitsType_::ConvSpecialization, GroupedConvTraitsType_::VectorSizeA, GroupedConvTraitsType_::VectorSizeB, GroupedConvTraitsType_::VectorSizeC, true > ConvToGemmTransformer
Definition grouped_convolution_backward_data_kernel.hpp:25

◆ AGridDescMK

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>

◆ BGridDescNK

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>

◆ CGridDescMN

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::CGridDescMN = remove_cvref_t<decltype(ABCGridDescs{}[number<2>{}])>

◆ ConvToGemmTransformer

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::ConvToGemmTransformer
Initial value:
TransformConvBwdDataToGemm<GroupedConvTraitsType_::NDimSpatial,
GroupedConvTraitsType_::ConvSpecialization,
GroupedConvTraitsType_::VectorSizeA,
GroupedConvTraitsType_::VectorSizeB,
GroupedConvTraitsType_::VectorSizeC,
true>
ck_tile::TransformConvBwdDataToGemm
Definition transform_conv_bwd_data_to_gemm.hpp:22

◆ TilePartitioner

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
using ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::TilePartitioner = remove_cvref_t<TilePartitioner_>

Constructor & Destructor Documentation

◆ GroupedConvBwdDataKernelArgs() [1/3]

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NWGK >, bool >::type = false>
CK_TILE_HOST ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::GroupedConvBwdDataKernelArgs ( const GroupedConvBwdDataHostArgs & args)
inline

◆ GroupedConvBwdDataKernelArgs() [2/3]

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NHWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKYXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NHWGK >, bool >::type = false>
CK_TILE_HOST ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::GroupedConvBwdDataKernelArgs ( const GroupedConvBwdDataHostArgs & args)
inline

◆ GroupedConvBwdDataKernelArgs() [3/3]

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
template<typename InLay = typename GroupedConvTraitsType_::InLayout, typename WeiLay = typename GroupedConvTraitsType_::WeiLayout, typename OutLay = typename GroupedConvTraitsType_::OutLayout, typename std::enable_if< std::is_same_v< InLay, tensor_layout::convolution::NDHWGC > &&std::is_same_v< WeiLay, tensor_layout::convolution::GKZYXC > &&std::is_same_v< OutLay, tensor_layout::convolution::NDHWGK >, bool >::type = false>
CK_TILE_HOST ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::GroupedConvBwdDataKernelArgs ( const GroupedConvBwdDataHostArgs & args)
inline

Member Data Documentation

◆ a_grid_descs_m_k

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<AGridDescMK, MaxGroupedGemmGroupsNum> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::a_grid_descs_m_k

◆ b_grid_descs_n_k

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<BGridDescNK, MaxGroupedGemmGroupsNum> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::b_grid_descs_n_k

◆ block_ends

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, MaxGroupedGemmGroupsNum> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::block_ends

◆ block_starts

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, MaxGroupedGemmGroupsNum> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::block_starts

◆ c_grid_descs_m_n

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<CGridDescMN, MaxGroupedGemmGroupsNum> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::c_grid_descs_m_n

◆ conv_filter_dilations

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::conv_filter_dilations

◆ conv_filter_strides

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::conv_filter_strides

◆ ds_ptr

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
std::array<const void*, NumDTensor> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::ds_ptr

◆ gemm_count

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::gemm_count = 0

◆ GemmBatch

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::GemmBatch

◆ grid_size_

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::grid_size_ = 0

◆ group_stride_a

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
long_index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::group_stride_a

◆ group_stride_b

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
long_index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::group_stride_b

◆ group_stride_c

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
long_index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::group_stride_c

◆ I0

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
auto ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::I0 = number<0>()
staticconstexpr

◆ I1

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
auto ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::I1 = number<1>()
staticconstexpr

◆ in_g_n_c_wis_lengths

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, NonSpatialDims + GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::in_g_n_c_wis_lengths

◆ in_ptr

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
void* ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::in_ptr

◆ input_batch_stride

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::input_batch_stride = 0

◆ input_left_pads

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::input_left_pads

◆ input_right_pads

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::input_right_pads

◆ k_batch

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::k_batch

◆ MaxGroupedGemmGroupsNum

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::MaxGroupedGemmGroupsNum = 128
staticconstexpr

◆ n_per_split

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::n_per_split = 1

◆ n_splits

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::n_splits = 1

◆ NonSpatialDims

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::NonSpatialDims = 3
staticconstexpr

◆ NumDTensor

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::NumDTensor = GroupedConvTraitsType_::NumDTensor
staticconstexpr

◆ original_n

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::original_n = 1

◆ out_g_n_k_wos_lengths

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, NonSpatialDims + GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::out_g_n_k_wos_lengths

◆ out_ptr

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
const void* ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::out_ptr

◆ output_batch_stride

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
index_t ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::output_batch_stride = 0

◆ tildes

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::tildes

◆ wei_g_k_c_xs_lengths

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
array<index_t, NonSpatialDims + GroupedConvTraitsType_::NDimSpatial> ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::wei_g_k_c_xs_lengths

◆ wei_ptr

template<typename GroupedConvTraitsType_, typename TilePartitioner_>
const void* ck_tile::GroupedConvBwdDataKernelArgs< GroupedConvTraitsType_, TilePartitioner_ >::wei_ptr
The documentation for this struct was generated from the following file: