/******************************************************************************* * * MIT License * * Copyright (c) 2020 Advanced Micro Devices, Inc. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal * in the Software without restriction, including without limitation the rights * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell * copies of the Software, and to permit persons to whom the Software is * furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice shall be included in all * copies or substantial portions of the Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. * *******************************************************************************/ #ifndef CK_REDUCTION_OPERATOR_HPP #define CK_REDUCTION_OPERATOR_HPP #include "reduction_common.hpp" namespace ck { namespace reduce { // Every binary operator used in reduction is represented by a templated functor class. Each functor // class must provide at least // three members: // 1) GetZeroVal() -- the interface to return the "identity element" for the binary operator, // "identity element" is the unique // element in the algebraic space that doesn't affect the value of other elements // when operated with any of them. // 2) indexable -- boolean value indicating whether indices of the operated elements could be // recorded. Usually, Min/Max operator could // need to record the indices of elements. For operator like Add/Mul, no need to // record the indices. // 3) operator() -- the first argument of the operator must be both an input & output, and the // corresponding variable usually stores // the accumulated result of many operator() calls; the second argument is only an // input. For indexable binary // operator, the second version of operator() has third argument (which is an // output) to indicate whether the // accumulated value (the first argument) has changed, in which case the recorded // accumulated index also need be // changed. template struct Add { using dataType = T; __device__ static T GetZeroVal() { return type_convert{}(0.0f); }; __device__ inline constexpr void operator()(T& a, T b) const { a = a + b; } static constexpr bool indexable = false; }; template struct Mul { using dataType = T; __device__ static T GetZeroVal() { return type_convert{}(1.0f); }; __device__ inline constexpr void operator()(T& a, T b) const { a = a * b; } static constexpr bool indexable = false; }; template struct Max { using dataType = T; __device__ static T GetZeroVal() { return std::numeric_limits::min(); }; __device__ inline constexpr void operator()(T& a, T b) const { if(a < b) a = b; } __device__ inline constexpr void operator()(T& a, T b, bool& changed) const { if(a < b) { a = b; changed = true; } } static constexpr bool indexable = true; }; template struct Min { using dataType = T; __device__ static T GetZeroVal() { return std::numeric_limits::max(); }; __device__ inline constexpr void operator()(T& a, T b) const { if(a > b) a = b; } __device__ inline constexpr void operator()(T& a, T b, bool& changed) const { if(a > b) { a = b; changed = true; } } static constexpr bool indexable = true; }; template <> __device__ half_t Max::GetZeroVal() { return type_convert{}(std::numeric_limits::min()); }; template <> __device__ half_t Min::GetZeroVal() { return type_convert{}(std::numeric_limits::max()); }; // Unary operators are usually called element-wisely before the reduction is executed on the // elements. // They are needed for easy implementation of reduction types of AVG, NRM1, NRM2 template struct unary_identic { __device__ unary_identic(const int divider = 1) { scaler = 1.0f / static_cast(divider); }; __device__ inline constexpr T operator()(T a) const { return a * type_convert{}(scaler); }; float scaler = 1.0f; }; template struct unary_identic { __device__ unary_identic(const int divider = 1) { (void)divider; }; __device__ inline constexpr T operator()(T a) const { return a; }; }; template struct unary_square { __device__ unary_square(const int divider = 1) { scaler = 1.0f / static_cast(divider); }; __device__ inline constexpr T operator()(T a) const { a = a * a; return a * type_convert{}(scaler); }; float scaler = 1.0f; }; template struct unary_square { __device__ unary_square(const int divider = 1) { (void)divider; }; __device__ inline constexpr T operator()(T a) const { return a * a; }; }; template struct unary_abs { __device__ unary_abs(const int divider = 1) { scaler = 1.0f / static_cast(divider); }; __device__ inline constexpr T operator()(T a) const { a = abs(a); return a * type_convert{}(scaler); }; float scaler = 1.0f; }; template struct unary_abs { __device__ unary_abs(const int divider = 1) { (void)divider; }; __device__ inline constexpr T operator()(T a) const { return abs(a); }; }; // We know for sure that 4.0 has __habs(), but 3.0 does not have it. // Let's assume that __habs() exists since 3.5. #if HIP_PACKAGE_VERSION_FLAT < 3005000000 inline __device__ __half __habs(__half x) { union { __half half; unsigned short u16; } val; val.half = x; val.u16 = val.u16 & 0x7fff; return val.half; } #endif template struct unary_abs { __device__ unary_abs(const int divider = 1) { scaler = 1.0f / static_cast(divider); }; __device__ inline half_t operator()(half_t a) const { a = static_cast(__habs(a)); return a * type_convert{}(scaler); }; float scaler = 1.0f; }; template <> struct unary_abs { __device__ unary_abs(const int divider = 1) { (void)divider; }; __device__ inline half_t operator()(half_t a) const { return static_cast(__habs(a)); }; }; template struct unary_sqrt { __device__ unary_sqrt(const int divider = 1) { (void)divider; }; __device__ inline T operator()(T a) const { return sqrtf(a); }; }; template <> struct unary_sqrt { __device__ unary_sqrt(const int divider = 1) { (void)divider; }; __device__ inline half_t operator()(half_t a) const { return static_cast(hsqrt(a)); }; }; }; // end of namespace reduce // The templated struct reduce_binary_operator maps the enum Ids of binary operators to their // respective functor classes. // The "GetZeroVal()" interface and boolean member "indexable" are also provided in // reduce_binary_operactor for // easier checking by the upper-layer codes in the kernels. template struct reduce_binary_operator; template struct reduce_binary_operator { using opType = reduce::Add; using dataType = T; __device__ static T GetZeroVal() { return reduce::Add::GetZeroVal(); }; static constexpr bool indexable = reduce::Add::indexable; }; template struct reduce_binary_operator { using opType = reduce::Mul; using dataType = T; __device__ static T GetZeroVal() { return reduce::Mul::GetZeroVal(); }; static constexpr bool indexable = reduce::Mul::indexable; }; template struct reduce_binary_operator { using opType = reduce::Min; using dataType = T; __device__ static T GetZeroVal() { return reduce::Min::GetZeroVal(); }; static constexpr bool indexable = reduce::Min::indexable; }; template struct reduce_binary_operator { using opType = reduce::Max; using dataType = T; __device__ static T GetZeroVal() { return reduce::Max::GetZeroVal(); }; static constexpr bool indexable = reduce::Max::indexable; }; template struct reduce_binary_operator { using opType = reduce::Max; using dataType = T; __device__ static T GetZeroVal() { return reduce::Max::GetZeroVal(); }; static constexpr bool indexable = reduce::Max::indexable; }; template struct reduce_binary_operator { using opType = reduce::Add; using dataType = T; __device__ static T GetZeroVal() { return reduce::Add::GetZeroVal(); }; static constexpr bool indexable = reduce::Add::indexable; }; template struct reduce_binary_operator { using opType = reduce::Add; using dataType = T; __device__ static T GetZeroVal() { return reduce::Add::GetZeroVal(); }; static constexpr bool indexable = reduce::Add::indexable; }; template struct reduce_binary_operator { using opType = reduce::Add; using dataType = T; __device__ static T GetZeroVal() { return reduce::Add::GetZeroVal(); }; static constexpr bool indexable = reduce::Add::indexable; }; // The templated struct reduce_unary_operator maps the enum Ids of Reduce operators to two unary // functor classes. // The two unary functors are called before and afer the Reduction is executed respectively template struct reduce_unary_operator { using preUnaryOp = reduce::unary_identic; using posUnaryOp = reduce::unary_identic; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_identic; using posUnaryOp = reduce::unary_identic; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_abs; using posUnaryOp = reduce::unary_identic; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_abs; using posUnaryOp = reduce::unary_identic; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_square; using posUnaryOp = reduce::unary_identic; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_square; using posUnaryOp = reduce::unary_sqrt; }; template struct reduce_unary_operator { using preUnaryOp = reduce::unary_identic; using posUnaryOp = reduce::unary_sqrt; }; } // end of namespace ck #endif