include/ck_tile/core/utility/reduce_operator.hpp Source File

include/ck_tile/core/utility/reduce_operator.hpp Source File#

Composable Kernel: include/ck_tile/core/utility/reduce_operator.hpp Source File
reduce_operator.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck_tile/core/config.hpp"
7 
8 namespace ck_tile {
9 
10 namespace ReduceOp {
11 // y = ReduceOp(y, x);
12 struct Add
13 {
14  template <typename T>
15  CK_TILE_HOST_DEVICE static constexpr T GetIdentityValue()
16  {
17  return type_convert<T>(0.0f);
18  };
19 
20  template <typename T,
21  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
22  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
23  CK_TILE_HOST_DEVICE constexpr T operator()(const T& y, const T x) const
24  {
25  return y + x;
26  }
27 
28  template <typename T,
29  typename = std::enable_if_t<std::is_same_v<T, half_t> || std::is_same_v<T, bf16_t>>>
30  CK_TILE_HOST_DEVICE constexpr T operator()(T& y, T x) const
31  {
32  float y_ = type_convert<float>(y);
33  float x_ = type_convert<float>(x);
34 
35  return type_convert<T>(y_ + x_);
36  }
37 };
38 
39 struct SquareAdd
40 {
41  template <typename T>
42  CK_TILE_HOST_DEVICE static constexpr T GetIdentityValue()
43  {
44  return type_convert<T>(0.0f);
45  };
46 
47  template <typename T,
48  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
49  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
50  CK_TILE_HOST_DEVICE constexpr T operator()(const T& y, const T x) const
51  {
52  return y + (x * x);
53  }
54 };
55 
56 struct Max
57 {
58  template <typename T,
59  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
60  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
61  CK_TILE_HOST_DEVICE static constexpr T GetIdentityValue()
62  {
63  return numeric<T>::min();
64  };
65 
66  template <typename T,
67  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
68  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
69  CK_TILE_HOST_DEVICE constexpr T operator()(const T& y, const T x) const
70  {
71  return max(y, x);
72  }
73 };
74 
75 struct AbsMax
76 {
77  template <typename T,
78  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
79  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
80  CK_TILE_HOST_DEVICE static constexpr T GetIdentityValue()
81  {
82  return numeric<T>::min();
83  };
84 
85  template <typename T,
86  typename = std::enable_if_t<std::is_same_v<T, float> || std::is_same_v<T, double> ||
87  std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t>>>
88  CK_TILE_HOST_DEVICE constexpr T operator()(const T& y, const T x) const
89  {
90  return max(y, abs(x));
91  }
92 };
93 
94 } // namespace ReduceOp
95 } // namespace ck_tile
CK_TILE_HOST_DEVICE
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:41
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::abs
CK_TILE_HOST_DEVICE bfloat16_t abs(const bfloat16_t &x)
Definition: bfloat16.hpp:395
ck_tile::max
constexpr CK_TILE_HOST_DEVICE T max(T x)
Definition: math.hpp:161
ck_tile::ReduceOp::AbsMax
Definition: reduce_operator.hpp:76
ck_tile::ReduceOp::AbsMax::operator()
constexpr CK_TILE_HOST_DEVICE T operator()(const T &y, const T x) const
Definition: reduce_operator.hpp:88
ck_tile::ReduceOp::AbsMax::GetIdentityValue
static constexpr CK_TILE_HOST_DEVICE T GetIdentityValue()
Definition: reduce_operator.hpp:80
ck_tile::ReduceOp::Add
Definition: reduce_operator.hpp:13
ck_tile::ReduceOp::Add::operator()
constexpr CK_TILE_HOST_DEVICE T operator()(T &y, T x) const
Definition: reduce_operator.hpp:30
ck_tile::ReduceOp::Add::GetIdentityValue
static constexpr CK_TILE_HOST_DEVICE T GetIdentityValue()
Definition: reduce_operator.hpp:15
ck_tile::ReduceOp::Add::operator()
constexpr CK_TILE_HOST_DEVICE T operator()(const T &y, const T x) const
Definition: reduce_operator.hpp:23
ck_tile::ReduceOp::Max
Definition: reduce_operator.hpp:57
ck_tile::ReduceOp::Max::GetIdentityValue
static constexpr CK_TILE_HOST_DEVICE T GetIdentityValue()
Definition: reduce_operator.hpp:61
ck_tile::ReduceOp::Max::operator()
constexpr CK_TILE_HOST_DEVICE T operator()(const T &y, const T x) const
Definition: reduce_operator.hpp:69
ck_tile::ReduceOp::SquareAdd
Definition: reduce_operator.hpp:40
ck_tile::ReduceOp::SquareAdd::operator()
constexpr CK_TILE_HOST_DEVICE T operator()(const T &y, const T x) const
Definition: reduce_operator.hpp:50
ck_tile::ReduceOp::SquareAdd::GetIdentityValue
static constexpr CK_TILE_HOST_DEVICE T GetIdentityValue()
Definition: reduce_operator.hpp:42
ck_tile::numeric::min
static constexpr CK_TILE_HOST_DEVICE T min()
Definition: numeric.hpp:20