/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp Source File
reduction_operator_mapping.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
9 // FIXME: can it be replaced with ck::Tuple?
10 #include <tuple>
11 
12 namespace ck {
13 
14 // The templated struct reduce_binary_operator maps the enum Ids of binary operators to their
15 // respective functor classes.
16 // The boolean member "indexable" are also provided in reduce_binary_operactor for
17 // easier checking by the upper-layer codes in the kernels.
18 
19 template <ReduceTensorOp Op>
21 
22 template <>
24 {
26 
27  static constexpr bool indexable = false;
28 };
29 
30 template <>
32 {
34 
35  static constexpr bool indexable = false;
36 };
37 
38 template <>
40 {
42 
43  static constexpr bool indexable = true;
44 };
45 
46 template <>
48 {
50 
51  static constexpr bool indexable = true;
52 };
53 
54 template <>
56 {
58 
59  static constexpr bool indexable = true;
60 };
61 
62 template <>
64 {
66 
67  static constexpr bool indexable = false;
68 };
69 
70 template <>
72 {
74 
75  static constexpr bool indexable = false;
76 };
77 
78 template <>
80 {
82 
83  static constexpr bool indexable = false;
84 };
85 
86 // The templated struct reduce_unary_operator maps the enum Ids of Reduce operators to two unary
87 // functor classes.
88 // The two unary functors are called before and afer the Reduction is executed respectively
89 template <ReduceTensorOp Op, bool IsFirstReduce, bool IsLastReduce>
91 {
94 
95  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
96  GetElementwiseOperator(int32_t reduceLength)
97  {
98  (void)reduceLength;
100  };
101 };
102 
103 template <bool IsFirstReduce>
104 struct reduce_unary_operator<ReduceTensorOp::AVG, IsFirstReduce, true>
105 {
108 
109  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
110  GetElementwiseOperator(int32_t reduceLength)
111  {
113  };
114 };
115 
116 template <bool IsLastReduce>
117 struct reduce_unary_operator<ReduceTensorOp::NORM1, true, IsLastReduce>
118 {
121 
122  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
123  GetElementwiseOperator(int32_t reduceLength)
124  {
125  (void)reduceLength;
127  };
128 };
129 
130 template <bool IsLastReduce>
131 struct reduce_unary_operator<ReduceTensorOp::AMAX, true, IsLastReduce>
132 {
135 
136  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
137  GetElementwiseOperator(int32_t reduceLength)
138  {
139  (void)reduceLength;
141  };
142 };
143 
144 template <>
146 {
149 
150  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
151  GetElementwiseOperator(int32_t reduceLength)
152  {
153  (void)reduceLength;
155  };
156 };
157 
158 template <>
160 {
163 
164  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
165  GetElementwiseOperator(int32_t reduceLength)
166  {
167  (void)reduceLength;
169  };
170 };
171 
172 template <>
174 {
177 
178  static std::tuple<InElementwiseOperation, AccElementwiseOperation>
179  GetElementwiseOperator(int32_t reduceLength)
180  {
181  (void)reduceLength;
183  };
184 };
185 
186 } // namespace ck
Definition: ck.hpp:264
ReduceTensorOp
Definition: reduction_enums.hpp:9
__host__ constexpr __device__ auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:211
Definition: reduction_operator.hpp:409
Definition: reduction_operator.hpp:37
Definition: reduction_operator.hpp:163
Definition: reduction_operator.hpp:285
Definition: reduction_operator.hpp:114
Definition: reduction_operator_mapping.hpp:20
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:137
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:110
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:123
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:179
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:151
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:165
Definition: reduction_operator_mapping.hpp:91
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:96
Definition: unary_element_wise_operation.hpp:241
Definition: unary_element_wise_operation.hpp:629
Definition: unary_element_wise_operation.hpp:569
Definition: unary_element_wise_operation.hpp:650
Definition: unary_element_wise_operation.hpp:613