Reduce< Problem_, Policy_ > Struct Template Reference

Reduce&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::Reduce< Problem_, Policy_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::Reduce< Problem_, Policy_ > Struct Template Reference

#include <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 >
 

Public Member Functions

template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims >
CK_TILE_DEVICE void operator() (const XDataType *p_x, YDataType *p_y, InputShape input_shape, InputStrides input_strides, KeptDim kept_dim, ReduceDims reduce_dims) 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 reduction kernel. More...
 

Static Public Attributes

static constexpr index_t kBlockSize = Problem::BlockShape::BlockSize
 

Member Typedef Documentation

◆ ComputeDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
using ck_tile::Reduce< Problem_, Policy_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ Policy

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
using ck_tile::Reduce< Problem_, Policy_ >::Policy = ck_tile::remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
using ck_tile::Reduce< Problem_, Policy_ >::Problem = ck_tile::remove_cvref_t<Problem_>

◆ XDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
using ck_tile::Reduce< Problem_, Policy_ >::XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>

◆ YDataType

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
using ck_tile::Reduce< Problem_, Policy_ >::YDataType = ck_tile::remove_cvref_t<typename Problem::YDataType>

Member Function Documentation

◆ BlockSize()

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
static constexpr CK_TILE_HOST auto ck_tile::Reduce< Problem_, Policy_ >::BlockSize ( )
inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
template<typename InputStrides >
static CK_TILE_HOST bool ck_tile::Reduce< Problem_, Policy_ >::IsSupportedArgument ( index_t  y_continous_dim,
InputStrides  input_strides 
)
inlinestatic

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

Parameters
y_continous_dimSize of the continuous dimension of the output tensor. Must be a multiple of ThreadTile_N for proper thread mapping.
input_stridesThe 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()()

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims >
CK_TILE_DEVICE void ck_tile::Reduce< Problem_, Policy_ >::operator() ( const XDataType *  p_x,
YDataType *  p_y,
InputShape  input_shape,
InputStrides  input_strides,
KeptDim  kept_dim,
ReduceDims  reduce_dims 
) const
inline

Member Data Documentation

◆ kBlockSize

template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
constexpr index_t ck_tile::Reduce< Problem_, Policy_ >::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/reduce2d_kernel.hpp