【PFCC算子性能优化】 SeluKernel Optimization

paddle/phi/kernels/funcs/activation_functor.h

@@ -2158,6 +2158,48 @@ struct CudaExpFunctor<double> : public BaseActivationFunctor<double> {
   }
 };
 
+template <typename T>
+struct CudaSeluFunctor : public BaseActivationFunctor<T> {
+  float scale;
+  float alpha;
+  T zero = static_cast<T>(0.0f);
+
+  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;
+  }
+};
+
+template <>
+struct CudaSeluFunctor<double> : public BaseActivationFunctor<double> {
+  float scale;
+  float alpha;
+  double zero = static_cast<double>(0.0f);
+
+  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;
+  }
+};
+
 template <typename T>
 struct CudaSquareFunctor : public BaseActivationFunctor<T> {
   // square(x) = x * x

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) {}