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

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

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

#include <reduce2d_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 ThreadPerWarp_M = Warp_M / ThreadTile_M
 
static constexpr index_t ThreadPerWarp_N = Warp_N / ThreadTile_N
 
static constexpr index_t Repeat_M = Block_M / (WarpPerBlock_M * Warp_M)
 
static constexpr index_t Repeat_N = Block_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::Reduce2dShape< 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::Reduce2dShape< 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::Reduce2dShape< 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:979
ck::get_warp_size
__host__ constexpr __device__ index_t get_warp_size()
Definition: get_id.hpp:42

◆ Repeat_M

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

◆ Repeat_N

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

◆ ThreadPerWarp_M

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

◆ ThreadPerWarp_N

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

◆ ThreadTile_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::Reduce2dShape< 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::Reduce2dShape< 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::Reduce2dShape< 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::Reduce2dShape< 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::Reduce2dShape< 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::Reduce2dShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::WarpPerBlock_N = BlockWarps::at(number<1>{})
staticconstexpr
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/reduce/pipeline/reduce2d_shape.hpp