BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ > Struct Template Reference

BatchedTransposeLdsProblem&lt; DataType_, BlockTile, NumWarps, kPadM_, kPadN_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ > Struct Template Reference

#include <batched_transpose_lds_problem.hpp>

Public Types

using DataType = remove_cvref_t<DataType_>

Static Public Attributes

static constexpr index_t kRowWarps_ = NumWarps::at(number<0>{})
static constexpr index_t kColWarps_ = NumWarps::at(number<1>{})
static constexpr index_t kRowPerBlock_ = BlockTile::at(number<0>{})
static constexpr index_t kColPerBlock_ = BlockTile::at(number<1>{})
static constexpr index_t kBlockSize = get_warp_size() * kRowWarps_ * kColWarps_
static constexpr index_t kLeadNumWarps = kColWarps_
static constexpr index_t kSecondNumWarps = kRowWarps_
static constexpr index_t kLeadSizePerBlock = kColPerBlock_
static constexpr index_t kSecondSizePerBlock = kRowPerBlock_
static constexpr index_t kQuadrantLeadDim = LaneGroupTransposeTraits<DataType>::kleadDim
static constexpr index_t kQuadrantSecondDim = LaneGroupTransposeTraits<DataType>::ksecondDim
static constexpr index_t kLeadSizePerWarp = kLeadSizePerBlock / kLeadNumWarps
static constexpr index_t kSecondSizePerWarp = kSecondSizePerBlock / kSecondNumWarps
static constexpr index_t kQuadNumPerLeadDim = kLeadSizePerWarp / kQuadrantLeadDim
static constexpr index_t kQuadNumPerSecondDim = kSecondSizePerWarp / kQuadrantSecondDim
static constexpr index_t kIterationsInSecondDim
static constexpr bool kPadM = kPadM_
static constexpr bool kPadN = kPadN_
static constexpr auto kMPerBlock = kSecondSizePerBlock
static constexpr auto kNPerBlock = kLeadSizePerBlock
static constexpr index_t MaxLoadStoreSize = 16
static constexpr auto VectorSizeInput = kPadN ? 1 : MaxLoadStoreSize / sizeof(DataType)
static constexpr auto VectorSizeOutput = kPadM ? 1 : MaxLoadStoreSize / sizeof(DataType)
static constexpr auto LDSVectorSize = MaxLoadStoreSize / sizeof(DataType)

Member Typedef Documentation

◆ DataType

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
using ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::DataType = remove_cvref_t<DataType_>

Member Data Documentation

◆ kBlockSize

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kBlockSize = get_warp_size() * kRowWarps_ * kColWarps_
staticconstexpr

◆ kColPerBlock_

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kColPerBlock_ = BlockTile::at(number<1>{})
staticconstexpr

◆ kColWarps_

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kColWarps_ = NumWarps::at(number<1>{})
staticconstexpr

◆ kIterationsInSecondDim

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kIterationsInSecondDim
staticconstexpr
Initial value:
=
kQuadNumPerLeadDim * kQuadNumPerSecondDim * 16 / get_warp_size()
ck_tile::get_warp_size
CK_TILE_HOST_DEVICE constexpr index_t get_warp_size()
Definition arch.hpp:63
ck_tile::BatchedTransposeLdsProblem::kQuadNumPerLeadDim
static constexpr index_t kQuadNumPerLeadDim
Definition batched_transpose_lds_problem.hpp:50
ck_tile::BatchedTransposeLdsProblem::kQuadNumPerSecondDim
static constexpr index_t kQuadNumPerSecondDim
Definition batched_transpose_lds_problem.hpp:51

◆ kLeadNumWarps

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kLeadNumWarps = kColWarps_
staticconstexpr

◆ kLeadSizePerBlock

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kLeadSizePerBlock = kColPerBlock_
staticconstexpr

◆ kLeadSizePerWarp

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kLeadSizePerWarp = kLeadSizePerBlock / kLeadNumWarps
staticconstexpr

◆ kMPerBlock

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
auto ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kMPerBlock = kSecondSizePerBlock
staticconstexpr

◆ kNPerBlock

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
auto ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kNPerBlock = kLeadSizePerBlock
staticconstexpr

◆ kPadM

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
bool ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kPadM = kPadM_
staticconstexpr

◆ kPadN

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
bool ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kPadN = kPadN_
staticconstexpr

◆ kQuadNumPerLeadDim

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kQuadNumPerLeadDim = kLeadSizePerWarp / kQuadrantLeadDim
staticconstexpr

◆ kQuadNumPerSecondDim

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kQuadNumPerSecondDim = kSecondSizePerWarp / kQuadrantSecondDim
staticconstexpr

◆ kQuadrantLeadDim

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kQuadrantLeadDim = LaneGroupTransposeTraits<DataType>::kleadDim
staticconstexpr

◆ kQuadrantSecondDim

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kQuadrantSecondDim = LaneGroupTransposeTraits<DataType>::ksecondDim
staticconstexpr

◆ kRowPerBlock_

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kRowPerBlock_ = BlockTile::at(number<0>{})
staticconstexpr

◆ kRowWarps_

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kRowWarps_ = NumWarps::at(number<0>{})
staticconstexpr

◆ kSecondNumWarps

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kSecondNumWarps = kRowWarps_
staticconstexpr

◆ kSecondSizePerBlock

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kSecondSizePerBlock = kRowPerBlock_
staticconstexpr

◆ kSecondSizePerWarp

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::kSecondSizePerWarp = kSecondSizePerBlock / kSecondNumWarps
staticconstexpr

◆ LDSVectorSize

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
auto ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::LDSVectorSize = MaxLoadStoreSize / sizeof(DataType)
staticconstexpr

◆ MaxLoadStoreSize

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
index_t ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::MaxLoadStoreSize = 16
staticconstexpr

◆ VectorSizeInput

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
auto ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::VectorSizeInput = kPadN ? 1 : MaxLoadStoreSize / sizeof(DataType)
staticconstexpr

◆ VectorSizeOutput

template<typename DataType_, typename BlockTile, typename NumWarps, bool kPadM_, bool kPadN_>
auto ck_tile::BatchedTransposeLdsProblem< DataType_, BlockTile, NumWarps, kPadM_, kPadN_ >::VectorSizeOutput = kPadM ? 1 : MaxLoadStoreSize / sizeof(DataType)
staticconstexpr
The documentation for this struct was generated from the following file: