PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile > Struct Template Reference

PoolShape&lt; BlockWarps, BlockTile, WarpTile, ThreadTile &gt; Struct Template Reference#

Composable Kernel: ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile > Struct Template Reference
Static Public Attributes | List of all members
ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile > Struct Template Reference

#include <pool_shape.hpp>

Static Public Attributes

static constexpr index_t Block_M = BlockTile::at(number<0>{})
 
static constexpr index_t Block_N = BlockTile::at(number<1>{})
 
static constexpr index_t Warp_M = WarpTile::at(number<0>{})
 
static constexpr index_t Warp_N = WarpTile::at(number<1>{})
 
static constexpr index_t ThreadTile_M = ThreadTile::at(number<0>{})
 
static constexpr index_t ThreadTile_N = ThreadTile::at(number<1>{})
 
static constexpr index_t WarpPerBlock_M = BlockWarps::at(number<0>{})
 
static constexpr index_t WarpPerBlock_N = BlockWarps::at(number<1>{})
 
static constexpr index_t WarpSizeScaleFactor
 
static constexpr index_t WarpSizeScaleFactor_M
 
static constexpr index_t WarpSizeScaleFactor_N
 
static constexpr index_t ThreadPerWarp_M = Warp_M / ThreadTile_M / WarpSizeScaleFactor_M
 
static constexpr index_t ThreadPerWarp_N = Warp_N / ThreadTile_N / WarpSizeScaleFactor_N
 
static constexpr index_t Repeat_M = Block_M * WarpSizeScaleFactor_M / (WarpPerBlock_M * Warp_M)
 
static constexpr index_t Repeat_N = Block_N * WarpSizeScaleFactor_N / (WarpPerBlock_N * Warp_N)
 
static constexpr index_t BlockSize
 

Member Data Documentation

◆ Block_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Block_M = BlockTile::at(number<0>{})
staticconstexpr

◆ Block_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Block_N = BlockTile::at(number<1>{})
staticconstexpr

◆ BlockSize

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::BlockSize
staticconstexpr
Initial value:
=
ck_tile::get_warp_size() * reduce_on_sequence(BlockWarps{}, multiplies{}, number<1>{})
ck_tile::multiplies
__host__ __device__ multiplies() -> multiplies< void, void >
FIXME: create macro to replace 'host device' and nothing more.
ck_tile::reduce_on_sequence
constexpr CK_TILE_HOST_DEVICE index_t reduce_on_sequence(Seq, Reduce f, number< Init >)
Definition: sequence.hpp:982
ck::get_warp_size
constexpr __device__ index_t get_warp_size()
Definition: get_id.hpp:10

◆ Repeat_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Repeat_M = Block_M * WarpSizeScaleFactor_M / (WarpPerBlock_M * Warp_M)
staticconstexpr

◆ Repeat_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Repeat_N = Block_N * WarpSizeScaleFactor_N / (WarpPerBlock_N * Warp_N)
staticconstexpr

◆ ThreadPerWarp_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::ThreadPerWarp_M = Warp_M / ThreadTile_M / WarpSizeScaleFactor_M
staticconstexpr

◆ ThreadPerWarp_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::ThreadPerWarp_N = Warp_N / ThreadTile_N / WarpSizeScaleFactor_N
staticconstexpr

◆ ThreadTile_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::ThreadTile_M = ThreadTile::at(number<0>{})
staticconstexpr

◆ ThreadTile_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::ThreadTile_N = ThreadTile::at(number<1>{})
staticconstexpr

◆ Warp_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Warp_M = WarpTile::at(number<0>{})
staticconstexpr

◆ Warp_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Warp_N = WarpTile::at(number<1>{})
staticconstexpr

◆ WarpPerBlock_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpPerBlock_M = BlockWarps::at(number<0>{})
staticconstexpr

◆ WarpPerBlock_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpPerBlock_N = BlockWarps::at(number<1>{})
staticconstexpr

◆ WarpSizeScaleFactor

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpSizeScaleFactor
staticconstexpr
Initial value:
=
Warp_M * Warp_N / ThreadTile_M / ThreadTile_N / ck_tile::get_warp_size()
ck_tile::PoolShape::Warp_M
static constexpr index_t Warp_M
Definition: pool_shape.hpp:19
ck_tile::PoolShape::Warp_N
static constexpr index_t Warp_N
Definition: pool_shape.hpp:20
ck_tile::PoolShape::ThreadTile_M
static constexpr index_t ThreadTile_M
Definition: pool_shape.hpp:22
ck_tile::PoolShape::ThreadTile_N
static constexpr index_t ThreadTile_N
Definition: pool_shape.hpp:23

◆ WarpSizeScaleFactor_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpSizeScaleFactor_M
staticconstexpr
Initial value:
=
(Warp_M / ThreadTile_M > Warp_N / ThreadTile_N) ? WarpSizeScaleFactor : 1
ck_tile::PoolShape::WarpSizeScaleFactor
static constexpr index_t WarpSizeScaleFactor
Definition: pool_shape.hpp:35

◆ WarpSizeScaleFactor_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::PoolShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpSizeScaleFactor_N
staticconstexpr
Initial value:
=
(Warp_M / ThreadTile_M > Warp_N / ThreadTile_N) ? 1 : WarpSizeScaleFactor
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/pooling/pipeline/pool_shape.hpp