tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize > Struct Template Reference

tile_distribution_encoding_pattern_aq_transposed_c&lt; BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize &gt; Struct Template Reference#

Composable Kernel: ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize > Struct Template Reference
ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize > Struct Template Reference

#include <gemm_group_quant_utils.hpp>

Inheritance diagram for ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >:
ck_tile::tile_distribution_encoding_pattern

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto make_2d_static_tile_distribution ()

Static Public Attributes

static constexpr index_t warp_size = get_warp_size()
static constexpr index_t num_warps = BlockSize / get_warp_size()
static constexpr index_t MWarps = BlockGemmShape::BlockWarps::at(number<0>{})
static constexpr index_t NWarps = BlockGemmShape::BlockWarps::at(number<1>{})
static constexpr index_t KWarps = BlockGemmShape::BlockWarps::at(number<2>{})
static constexpr index_t MIterPerWarp = BlockGemmShape::kM / (MWarps * WarpGemm::kM)
static constexpr index_t X = XPerTile
static constexpr index_t XR = 2
static constexpr index_t Y0 = MIterPerWarp
static constexpr index_t Y1 = MWarps
static constexpr index_t Y2 = WarpGemm::kM

Member Function Documentation

◆ make_2d_static_tile_distribution()

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::make_2d_static_tile_distribution ( )
inlinestaticconstexpr

Member Data Documentation

◆ KWarps

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::KWarps = BlockGemmShape::BlockWarps::at(number<2>{})
staticconstexpr

◆ MIterPerWarp

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::MIterPerWarp = BlockGemmShape::kM / (MWarps * WarpGemm::kM)
staticconstexpr

◆ MWarps

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::MWarps = BlockGemmShape::BlockWarps::at(number<0>{})
staticconstexpr

◆ num_warps

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::num_warps = BlockSize / get_warp_size()
staticconstexpr

◆ NWarps

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::NWarps = BlockGemmShape::BlockWarps::at(number<1>{})
staticconstexpr

◆ warp_size

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::warp_size = get_warp_size()
staticconstexpr

◆ X

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::X = XPerTile
staticconstexpr

◆ XR

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::XR = 2
staticconstexpr

◆ Y0

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::Y0 = MIterPerWarp
staticconstexpr

◆ Y1

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::Y1 = MWarps
staticconstexpr

◆ Y2

template<typename BlockGemmShape, typename WarpGemm, index_t BlockSize, index_t YPerTile, index_t XPerTile, index_t VecSize>
index_t ck_tile::tile_distribution_encoding_pattern_aq_transposed_c< BlockGemmShape, WarpGemm, BlockSize, YPerTile, XPerTile, VecSize >::Y2 = WarpGemm::kM
staticconstexpr

The documentation for this struct was generated from the following file: