MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ > Struct Template Reference

• ck_tile
• MultiReduce2d

Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members

ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ > Struct Template Reference

#include <multi_reduce2d_kernel.hpp>

Public Types
using	Problem = ck_tile::remove_cvref_t< Problem_ >
using	Policy = ck_tile::remove_cvref_t< Policy_ >
using	XDataType = ck_tile::remove_cvref_t< typename Problem::XDataType >
using	ComputeDataType = ck_tile::remove_cvref_t< typename Problem::ComputeDataType >
using	YDataType = ck_tile::remove_cvref_t< typename Problem::YDataType >
using	TilePartitioner = Reduce2dTilePartitioner< typename Problem::BlockShape, ForceMultiBlock_ >

Public Member Functions
template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims , typename ElementwiseOps , typename AccumulatorOps >
CK_TILE_DEVICE void	operator() (const XDataType *p_x, YDataType *p_y_tuple, InputShape input_shape, InputStrides input_strides, KeptDim kept_dim, ReduceDims reduce_dims, index_t output_tensor_offset, ElementwiseOps elementwise_ops, AccumulatorOps accumulator_ops) const
template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims , typename ElementwiseOps , typename AccumulatorOps , typename InterblockReduceOps >
CK_TILE_DEVICE void	operator() (const XDataType *p_x, YDataType *p_y_tuple, InputShape input_shape, InputStrides input_strides, KeptDim kept_dim, ReduceDims reduce_dims, index_t output_tensor_offset, ElementwiseOps elementwise_ops, AccumulatorOps accumulator_ops, InterblockReduceOps interblock_reduce_ops) const

Static Public Member Functions
static constexpr CK_TILE_HOST auto	BlockSize ()
template<typename InputStrides >
static CK_TILE_HOST bool	IsSupportedArgument (index_t y_continous_dim, InputStrides input_strides)
	Validates if the given arguments are supported by the 2D multi reduction kernel. More...

Static Public Attributes
static constexpr bool	ForceMultiBlock = ForceMultiBlock_
static constexpr index_t	kBlockSize = Problem::BlockShape::BlockSize

Member Typedef Documentation

◆ ComputeDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ Policy

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::Policy = ck_tile::remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::Problem = ck_tile::remove_cvref_t<Problem_>

◆ TilePartitioner

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::TilePartitioner = Reduce2dTilePartitioner<typename Problem::BlockShape, ForceMultiBlock_>

◆ XDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>

◆ YDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

using ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::YDataType = ck_tile::remove_cvref_t<typename Problem::YDataType>

Member Function Documentation

◆ BlockSize()

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

static constexpr CK_TILE_HOST auto ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::BlockSize ( )

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

template<typename InputStrides >

static CK_TILE_HOST bool ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::IsSupportedArgument ( index_t  y_continous_dim, InputStrides  input_strides  )

inlinestatic

Validates if the given arguments are supported by the 2D multi reduction kernel.

Parameters

y_continous_dim	Size of the continuous dimension of the output tensor. Must be a multiple of ThreadTile_N for proper thread mapping.
input_strides	The stride configuration of the input tensor. The last stride must be 1 to ensure contiguous memory access and enable efficient vectorized loads.

Returns

true if the arguments are supported, false otherwise. Error messages are logged when CK_TILE_LOGGING is enabled.

Note

Requirements:

• y_continous_dim % ThreadTile_N == 0 (for proper thread distribution)
• input_strides[-1] == 1 (for contiguous memory access)

◆ operator()() [1/2]

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims , typename ElementwiseOps , typename AccumulatorOps >

CK_TILE_DEVICE void ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::operator() ( const XDataType *  p_x, YDataType *  p_y_tuple, InputShape  input_shape, InputStrides  input_strides, KeptDim  kept_dim, ReduceDims  reduce_dims, index_t  output_tensor_offset, ElementwiseOps  elementwise_ops, AccumulatorOps  accumulator_ops  ) const

inline

◆ operator()() [2/2]

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims , typename ElementwiseOps , typename AccumulatorOps , typename InterblockReduceOps >

CK_TILE_DEVICE void ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::operator() ( const XDataType *  p_x, YDataType *  p_y_tuple, InputShape  input_shape, InputStrides  input_strides, KeptDim  kept_dim, ReduceDims  reduce_dims, index_t  output_tensor_offset, ElementwiseOps  elementwise_ops, AccumulatorOps  accumulator_ops, InterblockReduceOps  interblock_reduce_ops  ) const

inline

Member Data Documentation

◆ ForceMultiBlock

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

constexpr bool ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::ForceMultiBlock = ForceMultiBlock_

staticconstexpr

◆ kBlockSize

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy, bool ForceMultiBlock_ = false>

constexpr index_t ck_tile::MultiReduce2d< Problem_, Policy_, ForceMultiBlock_ >::kBlockSize = Problem::BlockShape::BlockSize

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/kernel/multi_reduce2d_kernel.hpp