From 859c4077173dbc28bcc5b874cb5e86745351ebdd Mon Sep 17 00:00:00 2001
From: carryyu <569782149@qq.com>
Date: Tue, 2 Aug 2022 21:40:08 +0800
Subject: [PATCH] =?UTF-8?q?=E3=80=90PFCC=E7=AE=97=E5=AD=90=E6=80=A7?=
 =?UTF-8?q?=E8=83=BD=E4=BC=98=E5=8C=96=E3=80=91=20SeluKernel=20Optimizatio?=
 =?UTF-8?q?n=20(#44490)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

* [PFCC] SeluKernel Optimization

* selu kernel optimization

* add private

Co-authored-by: carryyu <>
---
 paddle/phi/kernels/funcs/activation_functor.h | 44 +++++++++++++++++++
 paddle/phi/kernels/gpu/activation_kernel.cu   |  2 +
 paddle/phi/kernels/gpu/selu_kernel.cu         | 21 ---------
 3 files changed, 46 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 ee79cafd155..542c59bec1b 100644
--- a/paddle/phi/kernels/funcs/activation_functor.h
+++ b/paddle/phi/kernels/funcs/activation_functor.h
@@ -2158,6 +2158,50 @@ struct CudaExpFunctor<double> : public BaseActivationFunctor<double> {
   }
 };
 
+template <typename T>
+struct CudaSeluFunctor : public BaseActivationFunctor<T> {
+  typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
+    return {{"scale", &scale}, {"alpha", &alpha}};
+  }
+
+  __device__ __forceinline__ T operator()(const T x) const {
+    T res = x;
+    if (res <= zero) {
+      res = alpha * expf(res) - alpha;
+    }
+    res *= scale;
+    return res;
+  }
+
+ private:
+  float scale;
+  float alpha;
+  T zero = static_cast<T>(0.0f);
+};
+
+template <>
+struct CudaSeluFunctor<double> : public BaseActivationFunctor<double> {
+  typename BaseActivationFunctor<double>::AttrPair GetAttrs() {
+    return {{"scale", &scale}, {"alpha", &alpha}};
+  }
+
+  __device__ __forceinline__ double operator()(const double x) const {
+    double res = x;
+    double alpha_cast = static_cast<double>(alpha);
+    double scale_cast = static_cast<double>(scale);
+    if (res <= zero) {
+      res = alpha_cast * exp(res) - alpha_cast;
+    }
+    res *= scale_cast;
+    return res;
+  }
+
+ private:
+  float scale;
+  float alpha;
+  double zero = static_cast<double>(0.0f);
+};
+
 template <typename T>
 struct CudaSquareFunctor : public BaseActivationFunctor<T> {
   // square(x) = x * x
diff --git a/paddle/phi/kernels/gpu/activation_kernel.cu b/paddle/phi/kernels/gpu/activation_kernel.cu
index 6e116a3e157..0e9e754a997 100644
--- a/paddle/phi/kernels/gpu/activation_kernel.cu
+++ b/paddle/phi/kernels/gpu/activation_kernel.cu
@@ -132,6 +132,7 @@ DEFINE_GPU_ACT_KERNEL_WITH_TWO_ATTRS(HardSigmoid,
                                      CudaHardSigmoidFunctor,
                                      slope,
                                      offset)
+DEFINE_GPU_ACT_KERNEL_WITH_TWO_ATTRS(Selu, CudaSeluFunctor, scale, alpha)
 
 template <typename T, typename Context>
 void HardSwishKernel(const Context& dev_ctx,
@@ -265,3 +266,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 99303d8c18a..00000000000
--- 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) {}
-- 
GitLab