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 | |
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
◆ IsSupportedArgument()
template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
template<typename InputStrides >
|
inlinestatic |
Validates if the given arguments are supported by the 2D 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()()
template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
template<typename InputShape , typename InputStrides , typename KeptDim , typename ReduceDims >
|
inline |
Member Data Documentation
◆ kBlockSize
template<typename Problem_ , typename Policy_ = Reduce2dDefaultPolicy>
|
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