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 RepeatInWarp
 
static constexpr index_t RepeatInWarp_M
 
static constexpr index_t RepeatInWarp_N
 
static constexpr index_t ThreadPerWarp_M = Warp_M / ThreadTile_M / RepeatInWarp_M
 
static constexpr index_t ThreadPerWarp_N = Warp_N / ThreadTile_N / RepeatInWarp_N
 
static constexpr index_t Repeat_M = Block_M * RepeatInWarp_M / (WarpPerBlock_M * Warp_M)
 
static constexpr index_t Repeat_N = Block_N * RepeatInWarp_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: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::Reduce2dShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::Repeat_M = Block_M * RepeatInWarp_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 * RepeatInWarp_N / (WarpPerBlock_N * Warp_N)
staticconstexpr

◆ RepeatInWarp

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

◆ RepeatInWarp_M

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::Reduce2dShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::RepeatInWarp_M
staticconstexpr
Initial value:
=
(Warp_M / ThreadTile_M > Warp_N / ThreadTile_N) ? RepeatInWarp : 1
ck_tile::Reduce2dShape::RepeatInWarp
static constexpr index_t RepeatInWarp
Definition: reduce2d_shape.hpp:28

◆ RepeatInWarp_N

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ThreadTile >
constexpr index_t ck_tile::Reduce2dShape< BlockWarps, BlockTile, WarpTile, ThreadTile >::RepeatInWarp_N
staticconstexpr
Initial value:
=
(Warp_M / ThreadTile_M > Warp_N / ThreadTile_N) ? 1 : RepeatInWarp

◆ 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 / RepeatInWarp_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 / RepeatInWarp_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