GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce > Struct Template Reference

GridwiseNormalizationBwdData_mk_to_mk&lt; DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce &gt; Struct Template Reference#

Composable Kernel: ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce > Struct Template Reference

#include <gridwise_normalization_bwd_data.hpp>

Public Types

using ThreadClusterLengths_M_K = Sequence<MThreadClusterSize, KThreadClusterSize>
using DYThreadBufferDimAccessOrder
using XThreadBufferDimAccessOrder
using GammaThreadBufferDimAccessOrder
using MeanInvStdThreadBufferDimAccessOrder
using DXThreadBufferDimAccessOrder
using ThreadClusterArrangeOrder = DYThreadBufferDimAccessOrder
using ThreadBufferLengths_M_K = Sequence<MThreadSliceSize, KThreadSliceSize>
using PassThroughOp = tensor_operation::element_wise::PassThrough
using BlockwiseSumReduce

Static Public Member Functions

static __device__ void Run (const GridDesc_M_K &dy_grid_desc_m_k, const GridDesc_M_K &x_grid_desc_m_k, const GridDesc_M_K &gamma_grid_desc_m_k, const GridDesc_M_K &mean_grid_desc_m_k, const GridDesc_M_K &inv_std_grid_desc_m_k, const GridDesc_M_K &dx_grid_desc_m_k, index_t num_k_block_tile_iteration, const DYDataType *const __restrict__ p_dy_global, const XDataType *const __restrict__ p_x_global, const GammaDataType *const __restrict__ p_gamma_global, const MeanInvStdDataType *const __restrict__ p_mean_global, const MeanInvStdDataType *const __restrict__ p_inv_std_global, DXDataType *const __restrict__ p_dx_global)

Static Public Attributes

static constexpr auto thread_cluster_desc
static constexpr auto thread_buffer_desc_m_k
static constexpr auto thread_buffer_desc_m
static constexpr auto I0 = Number<0>{}
static constexpr auto I1 = Number<1>{}
static constexpr auto I2 = Number<2>{}
static constexpr index_t M_BlockTileSize = MThreadClusterSize * MThreadSliceSize
static constexpr index_t K_BlockTileSize = KThreadClusterSize * KThreadSliceSize

Member Typedef Documentation

◆ BlockwiseSumReduce

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::BlockwiseSumReduce
Initial value:
PartitionedBlockwiseReduction<ComputeDataType,
BlockSize,
ThreadClusterLengths_M_K,
ThreadClusterArrangeOrder,
reduce::Add,
true>
ck::BlockwiseSoftmax::ThreadClusterLengths_M_K
decltype(ThreadClusterDesc_M_K{}.GetLengths()) ThreadClusterLengths_M_K
Definition blockwise_softmax.hpp:69
ck::GridwiseNormalizationBwdData_mk_to_mk::ThreadClusterArrangeOrder
DYThreadBufferDimAccessOrder ThreadClusterArrangeOrder
Definition gridwise_normalization_bwd_data.hpp:86
ck::PartitionedBlockwiseReduction
Definition reduction_functions_blockwise.hpp:28
ck::reduce::Add
Definition reduction_operator.hpp:37

◆ DXThreadBufferDimAccessOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::DXThreadBufferDimAccessOrder
Initial value:
typename conditional<DXDstVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type
ck::Sequence
Definition utility/sequence.hpp:43
ck::conditional
Definition utility/functional.hpp:100

◆ DYThreadBufferDimAccessOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::DYThreadBufferDimAccessOrder
Initial value:
typename conditional<DYSrcVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type

◆ GammaThreadBufferDimAccessOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::GammaThreadBufferDimAccessOrder
Initial value:
typename conditional<GammaSrcVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type

◆ MeanInvStdThreadBufferDimAccessOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::MeanInvStdThreadBufferDimAccessOrder
Initial value:
typename conditional<MeanInvStdSrcVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type

◆ PassThroughOp

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::PassThroughOp = tensor_operation::element_wise::PassThrough

◆ ThreadBufferLengths_M_K

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::ThreadBufferLengths_M_K = Sequence<MThreadSliceSize, KThreadSliceSize>

◆ ThreadClusterArrangeOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::ThreadClusterArrangeOrder = DYThreadBufferDimAccessOrder

◆ ThreadClusterLengths_M_K

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::ThreadClusterLengths_M_K = Sequence<MThreadClusterSize, KThreadClusterSize>

◆ XThreadBufferDimAccessOrder

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
using ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::XThreadBufferDimAccessOrder
Initial value:
typename conditional<XSrcVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type

Member Function Documentation

◆ Run()

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
__device__ void ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::Run ( const GridDesc_M_K & dy_grid_desc_m_k,
const GridDesc_M_K & x_grid_desc_m_k,
const GridDesc_M_K & gamma_grid_desc_m_k,
const GridDesc_M_K & mean_grid_desc_m_k,
const GridDesc_M_K & inv_std_grid_desc_m_k,
const GridDesc_M_K & dx_grid_desc_m_k,
index_t num_k_block_tile_iteration,
const DYDataType *const __restrict__ p_dy_global,
const XDataType *const __restrict__ p_x_global,
const GammaDataType *const __restrict__ p_gamma_global,
const MeanInvStdDataType *const __restrict__ p_mean_global,
const MeanInvStdDataType *const __restrict__ p_inv_std_global,
DXDataType *const __restrict__ p_dx_global )
inlinestatic

Member Data Documentation

◆ I0

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::I2 = Number<2>{}
staticconstexpr

◆ K_BlockTileSize

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
index_t ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::K_BlockTileSize = KThreadClusterSize * KThreadSliceSize
staticconstexpr

◆ M_BlockTileSize

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
index_t ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::M_BlockTileSize = MThreadClusterSize * MThreadSliceSize
staticconstexpr

◆ thread_buffer_desc_m

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::thread_buffer_desc_m
staticconstexpr
Initial value:
=
make_naive_tensor_descriptor_packed(make_tuple(Number<MThreadSliceSize>{}))
ck::Number
integral_constant< index_t, N > Number
Definition number.hpp:12
ck::make_naive_tensor_descriptor_packed
__host__ __device__ constexpr auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition tensor_descriptor_helper.hpp:101
ck::make_tuple
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211

◆ thread_buffer_desc_m_k

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::thread_buffer_desc_m_k
staticconstexpr
Initial value:
= make_naive_tensor_descriptor_packed(
make_tuple(Number<MThreadSliceSize>{}, Number<KThreadSliceSize>{}))

◆ thread_cluster_desc

template<typename DYDataType, typename XDataType, typename GammaDataType, typename MeanInvStdDataType, typename ComputeDataType, typename DXDataType, typename GridDesc_M_K, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t DYSrcVectorDim, index_t DYSrcVectorSize, index_t XSrcVectorDim, index_t XSrcVectorSize, index_t GammaSrcVectorDim, index_t GammaSrcVectorSize, index_t MeanInvStdSrcVectorDim, index_t MeanInvStdSrcVectorSize, index_t DXDstVectorDim, index_t DXDstVectorSize, bool SweepOnce>
auto ck::GridwiseNormalizationBwdData_mk_to_mk< DYDataType, XDataType, GammaDataType, MeanInvStdDataType, ComputeDataType, DXDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, DYSrcVectorDim, DYSrcVectorSize, XSrcVectorDim, XSrcVectorSize, GammaSrcVectorDim, GammaSrcVectorSize, MeanInvStdSrcVectorDim, MeanInvStdSrcVectorSize, DXDstVectorDim, DXDstVectorSize, SweepOnce >::thread_cluster_desc
staticconstexpr
Initial value:
=
make_cluster_descriptor(ThreadClusterLengths_M_K{}, ThreadClusterArrangeOrder{})
ck::make_cluster_descriptor
__host__ __device__ constexpr auto make_cluster_descriptor(const Lengths &lengths, ArrangeOrder order=typename arithmetic_sequence_gen< 0, Lengths::Size(), 1 >::type{})
Definition tensor_description/cluster_descriptor.hpp:13
ck::GridwiseNormalizationBwdData_mk_to_mk::ThreadClusterLengths_M_K
Sequence< MThreadClusterSize, KThreadClusterSize > ThreadClusterLengths_M_K
Definition gridwise_normalization_bwd_data.hpp:73
The documentation for this struct was generated from the following file: