ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType > Struct Template Reference

ElementWiseShape&lt; BlockWarps, BlockTile, WarpTile, ComputeDataType &gt; Struct Template Reference#

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

#include <elementwise_shape.hpp>

Static Public Attributes

static constexpr index_t kBlockM = BlockTile::at(number<0>{})
 
static constexpr index_t kWarpM = WarpTile::at(number<0>{})
 
static constexpr index_t kVectorM
 
static constexpr index_t kWarpPerBlockM = BlockWarps::at(number<0>{})
 
static constexpr index_t kThreadPerWarpM = get_warp_size()
 
static constexpr index_t kRepeatM = kBlockM / (kWarpPerBlockM * kVectorM * kThreadPerWarpM)
 
static constexpr index_t kBlockSize
 

Member Data Documentation

◆ kBlockM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kBlockM = BlockTile::at(number<0>{})
staticconstexpr

◆ kBlockSize

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kBlockSize
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

◆ kRepeatM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kRepeatM = kBlockM / (kWarpPerBlockM * kVectorM * kThreadPerWarpM)
staticconstexpr

◆ kThreadPerWarpM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kThreadPerWarpM = get_warp_size()
staticconstexpr

◆ kVectorM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kVectorM
staticconstexpr
Initial value:
=
min(static_cast<index_t>(16 / sizeof(ComputeDataType)), kWarpM / get_warp_size())
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
ck_tile::min
constexpr CK_TILE_HOST_DEVICE T min(T x)
Definition: math.hpp:210
ck_tile::ElementWiseShape::kWarpM
static constexpr index_t kWarpM
Definition: elementwise_shape.hpp:15

◆ kWarpM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kWarpM = WarpTile::at(number<0>{})
staticconstexpr

◆ kWarpPerBlockM

template<typename BlockWarps , typename BlockTile , typename WarpTile , typename ComputeDataType >
constexpr index_t ck_tile::ElementWiseShape< BlockWarps, BlockTile, WarpTile, ComputeDataType >::kWarpPerBlockM = BlockWarps::at(number<0>{})
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/elementwise/pipeline/elementwise_shape.hpp