[cherry-pick2.3]fix compile bug of windows cuda11.5 (#41464)

cherry-pick fix compile bug of windows cuda11.5 #41433
PaddlePaddle · Apr 7, 2022 · 7f502b7 · 7f502b7
1 parent 5b85f3d
commit 7f502b7
Showing 1 changed file with 9 additions and 4 deletions.
diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h
@@ -1878,12 +1878,17 @@ struct CudaCosGradFunctor : public BaseActivationFunctor<T> {
 
 template <typename T>
 struct CudaExpFunctor : public BaseActivationFunctor<T> {
-  using MPType = typename phi::dtype::MPTypeTrait<T>::Type;
+  // exp(x) = expf(x)
+  __device__ __forceinline__ T operator()(const T x) const {
+    return static_cast<T>(expf(static_cast<float>(x)));
+  }
+};
 
+template <>
+struct CudaExpFunctor<double> : public BaseActivationFunctor<double> {
   // exp(x) = exp(x)
-  __device__ __forceinline__ T operator()(const T arg_x) const {
-    MPType x = static_cast<MPType>(arg_x);
-    return static_cast<T>(exp(x));
+  __device__ __forceinline__ double operator()(const double x) const {
+    return exp(x);
   }
 };