1 /***************************************************************************************************
2 * Copyright (c) 2017 - 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
3 * SPDX-License-Identifier: BSD-3-Clause
4 *
5 * Redistribution and use in source and binary forms, with or without
6 * modification, are permitted provided that the following conditions are met:
7 *
8 * 1. Redistributions of source code must retain the above copyright notice, this
9 * list of conditions and the following disclaimer.
10 *
11 * 2. Redistributions in binary form must reproduce the above copyright notice,
12 * this list of conditions and the following disclaimer in the documentation
13 * and/or other materials provided with the distribution.
14 *
15 * 3. Neither the name of the copyright holder nor the names of its
16 * contributors may be used to endorse or promote products derived from
17 * this software without specific prior written permission.
18 *
19 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
20 * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
21 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
22 * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
23 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
24 * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
25 * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
26 * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
27 * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
28 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
29 *
30 **************************************************************************************************/
31 /*! \file
32 \brief Functor performing linear combination with a maximum operation used by epilogues.
33 */
34
35 #pragma once
36
37 #include <cutlass/array.h>
38 #include <cutlass/cutlass.h>
39 #include <cutlass/epilogue/thread/activation.h>
40 #include <cutlass/epilogue/thread/scale_type.h>
41 #include <cutlass/functional.h>
42 #include <cutlass/half.h>
43 #include <cutlass/numeric_conversion.h>
44 #include <cutlass/numeric_types.h>
45
46 /////////////////////////////////////////////////////////////////////////////////////////////////
47
48 namespace cutlass {
49 namespace epilogue {
50 namespace thread {
51
52 /////////////////////////////////////////////////////////////////////////////////////////////////
53
copysignf_pos(float a,float b)54 __forceinline__ __device__ float copysignf_pos(float a, float b)
55 {
56 float r;
57 r = __int_as_float(__float_as_int(a) | (__float_as_int(b) & 0x80000000));
58 return r;
59 }
60
tanh_opt(float x)61 __forceinline__ __device__ float tanh_opt(float x)
62 {
63 #if (__CUDACC_VER_MAJOR__ < 11) || (__CUDA_ARCH__ < 750)
64 const float exp_val = -1.f * fabs(2 * x);
65 return copysignf_pos((1.0f - __expf(exp_val)) / (__expf(exp_val) + 1.0f), x);
66 #else
67 return fast_tanh(x);
68 #endif
69 }
70
71 /////////////////////////////////////////////////////////////////////////////////////////////////
72 template<>
73 struct GELU_taylor<float> {
74 static const bool kIsHeavy = true;
75 CUTLASS_DEVICE
76 float operator()(float const& z) const
77 {
78
79 float k0 = float(0.7978845608028654);
80 float k1 = float(0.044715);
81
82 return float(
83 cutlass::constants::half<float>() * z
84 * (cutlass::constants::one<float>() + tanh_opt(k0 * z * (cutlass::constants::one<float>() + k1 * z * z))));
85 }
86
87 using Params = LinearCombinationGenericParams<float>;
88
89 CUTLASS_DEVICE
90 float operator()(float const& scalar, Params const& params_) const
91 {
92 return this->operator()(scalar);
93 }
94 };
95
96 } // namespace thread
97 } // namespace epilogue
98 } // namespace cutlass
99
100 /////////////////////////////////////////////////////////////////////////////////////////////////
101