GroupedConvFwdKernelArgs< GroupedConvTraitsType_, CDElementwise_ > Struct Template Reference

GroupedConvFwdKernelArgs< GroupedConvTraitsType_, CDElementwise_ > Struct Template Reference#

Composable Kernel: ck_tile::GroupedConvFwdKernelArgs< GroupedConvTraitsType_, CDElementwise_ > Struct Template Reference

The Grouped Convolution kernel device arguments. More...

#include <grouped_convolution_forward_kernel.hpp>

Classes
struct	SplitImageInfo

Public Types
using	ConvToGemmFwdTransformer
using	CDElementwise = CDElementwise_
using	AGridDescMK
using	BGridDescNK
using	CGridDescMN
using	ConvToGemmFwdTransformer_t = ConvToGemmFwdTransformer
using	AGridDescMK_t = AGridDescMK
using	CGridDescMN_t = CGridDescMN

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	GroupedConvFwdKernelArgs (const GroupedConvFwdHostArgs< CDElementwise > &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	GroupedConvFwdKernelArgs (const GroupedConvFwdHostArgs< CDElementwise > &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	GroupedConvFwdKernelArgs (const GroupedConvFwdHostArgs< CDElementwise > &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
index_t	k_batch
index_t	GemmM
index_t	GemmN
index_t	GemmK
index_t	GemmBatch
const void *	in_ptr
const void *	wei_ptr
std::array< const void *, NumDTensor >	ds_ptr
const CDElementwise	elfunc
void *	out_ptr
AGridDescMK	a_grid_desc_m_k
BGridDescNK	b_grid_desc_n_k
CGridDescMN	c_grid_desc_m_n
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
long_index_t	spatial_offset_in = 0
long_index_t	spatial_offset_out = 0
ConvToGemmFwdTransformer	transformer_
index_t	num_spatial_pieces = 1
SplitImageInfo	split_image

Static Public Attributes
static constexpr index_t	NumDTensor = GroupedConvTraitsType_::NumDTensor
static constexpr index_t	NonSpatialDims = 3

Detailed Description

template<typename GroupedConvTraitsType_, typename CDElementwise_>
struct ck_tile::GroupedConvFwdKernelArgs< GroupedConvTraitsType_, CDElementwise_ >

The Grouped Convolution kernel device arguments.