Reduce< Problem_, Policy_ > Struct Template Reference

• ck_tile
• Reduce

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 CK_TILE_HOST constexpr 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.

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>

CK_TILE_HOST constexpr auto ck_tile::Reduce< Problem_, Policy_ >::BlockSize ( )

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename Problem_, typename Policy_ = Reduce2dDefaultPolicy>

template<typename InputStrides>

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_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>

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>

index_t ck_tile::Reduce< Problem_, Policy_ >::kBlockSize = Problem::BlockShape::BlockSize

staticconstexpr

---

The documentation for this struct was generated from the following file:

• reduce2d_kernel.hpp