From b71b24def238a32e05cee6d760148194d74398f0 Mon Sep 17 00:00:00 2001 From: carryyu <569782149@qq.com> Date: Wed, 20 Jul 2022 13:12:28 +0000 Subject: [PATCH 1/3] [PFCC] SeluKernel Optimization --- paddle/phi/kernels/funcs/activation_functor.h | 42 +++++++++++++++++++ paddle/phi/kernels/gpu/activation_kernel.cu | 2 + paddle/phi/kernels/gpu/selu_kernel.cu | 21 ---------- 3 files changed, 44 insertions(+), 21 deletions(-) delete mode 100644 paddle/phi/kernels/gpu/selu_kernel.cu diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h index bd5e3dec3d6b0..9ec2f6c264d96 100644 --- a/paddle/phi/kernels/funcs/activation_functor.h +++ b/paddle/phi/kernels/funcs/activation_functor.h @@ -2132,6 +2132,48 @@ struct CudaExpFunctor : public BaseActivationFunctor { } }; +template +struct CudaSeluFunctor : public BaseActivationFunctor { + float scale; + float alpha; + T zero = static_cast(0.0f); + + typename BaseActivationFunctor::AttrPair GetAttrs() { + return {{"scale", &scale}, {"alpha", &alpha}}; + } + + __device__ __forceinline__ T operator()(const T x) const { + T res = x; + if (res <= zero) { + res = alpha * exp(res) - alpha; + } + res *= scale; + return res; + } +}; + +template <> +struct CudaSeluFunctor : public BaseActivationFunctor { + float scale; + float alpha; + double zero = static_cast(0.0f); + + typename BaseActivationFunctor::AttrPair GetAttrs() { + return {{"scale", &scale}, {"alpha", &alpha}}; + } + + __device__ __forceinline__ double operator()(const double x) const { + double res = x; + double alpha_cast = static_cast(alpha); + double scale_cast = static_cast(scale); + if (res <= zero) { + res = alpha_cast * exp(res) - alpha_cast; + } + res *= scale_cast; + return res; + } +}; + template struct CudaSquareFunctor : public BaseActivationFunctor { // square(x) = x * x diff --git a/paddle/phi/kernels/gpu/activation_kernel.cu b/paddle/phi/kernels/gpu/activation_kernel.cu index b7ff76f744645..d229ae0cfaa62 100644 --- a/paddle/phi/kernels/gpu/activation_kernel.cu +++ b/paddle/phi/kernels/gpu/activation_kernel.cu @@ -131,6 +131,7 @@ DEFINE_GPU_ACT_KERNEL_WITH_TWO_ATTRS(HardSigmoid, CudaHardSigmoidFunctor, slope, offset) +DEFINE_GPU_ACT_KERNEL_WITH_TWO_ATTRS(Selu, CudaSeluFunctor, scale, alpha) template void HardSwishKernel(const Context& dev_ctx, @@ -263,3 +264,4 @@ PD_REGISTER_KERNEL(pow, int, int64_t, phi::dtype::float16) {} +PD_REGISTER_KERNEL(selu, GPU, ALL_LAYOUT, phi::SeluKernel, float, double) {} diff --git a/paddle/phi/kernels/gpu/selu_kernel.cu b/paddle/phi/kernels/gpu/selu_kernel.cu deleted file mode 100644 index 99303d8c18a97..0000000000000 --- a/paddle/phi/kernels/gpu/selu_kernel.cu +++ /dev/null @@ -1,21 +0,0 @@ -// Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#include "paddle/phi/kernels/selu_kernel.h" - -#include "paddle/phi/backends/gpu/gpu_context.h" -#include "paddle/phi/core/kernel_registry.h" -#include "paddle/phi/kernels/impl/selu_kernel_impl.h" - -PD_REGISTER_KERNEL(selu, GPU, ALL_LAYOUT, phi::SeluKernel, float, double) {} From 58fd1b686243d7568fce32188f4b1b13b12b8b61 Mon Sep 17 00:00:00 2001 From: carryyu <569782149@qq.com> Date: Wed, 20 Jul 2022 13:20:30 +0000 Subject: [PATCH 2/3] selu kernel optimization --- paddle/phi/kernels/funcs/activation_functor.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h index 9ec2f6c264d96..58f047c0a0dcb 100644 --- a/paddle/phi/kernels/funcs/activation_functor.h +++ b/paddle/phi/kernels/funcs/activation_functor.h @@ -2145,7 +2145,7 @@ struct CudaSeluFunctor : public BaseActivationFunctor { __device__ __forceinline__ T operator()(const T x) const { T res = x; if (res <= zero) { - res = alpha * exp(res) - alpha; + res = alpha * expf(res) - alpha; } res *= scale; return res; From 330e6e75df57a4cc0877d16f051178c28c9ecd6b Mon Sep 17 00:00:00 2001 From: carryyu <> Date: Mon, 1 Aug 2022 11:10:00 +0000 Subject: [PATCH 3/3] add private --- paddle/phi/kernels/funcs/activation_functor.h | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h index 58f047c0a0dcb..5bbfda0da5c6f 100644 --- a/paddle/phi/kernels/funcs/activation_functor.h +++ b/paddle/phi/kernels/funcs/activation_functor.h @@ -2134,10 +2134,6 @@ struct CudaExpFunctor : public BaseActivationFunctor { template struct CudaSeluFunctor : public BaseActivationFunctor { - float scale; - float alpha; - T zero = static_cast(0.0f); - typename BaseActivationFunctor::AttrPair GetAttrs() { return {{"scale", &scale}, {"alpha", &alpha}}; } @@ -2150,14 +2146,15 @@ struct CudaSeluFunctor : public BaseActivationFunctor { res *= scale; return res; } -}; -template <> -struct CudaSeluFunctor : public BaseActivationFunctor { + private: float scale; float alpha; - double zero = static_cast(0.0f); + T zero = static_cast(0.0f); +}; +template <> +struct CudaSeluFunctor : public BaseActivationFunctor { typename BaseActivationFunctor::AttrPair GetAttrs() { return {{"scale", &scale}, {"alpha", &alpha}}; } @@ -2172,6 +2169,11 @@ struct CudaSeluFunctor : public BaseActivationFunctor { res *= scale_cast; return res; } + + private: + float scale; + float alpha; + double zero = static_cast(0.0f); }; template