diff --git a/paddle/fluid/operators/kernel_primitives/functor_primitives.h b/paddle/fluid/operators/kernel_primitives/functor_primitives.h index 2bd8721b82fa2dcc52aa6f3acda7b262d532b829..5e3c1fc202d595cf7406841cae716d3ddcb59d02 100644 --- a/paddle/fluid/operators/kernel_primitives/functor_primitives.h +++ b/paddle/fluid/operators/kernel_primitives/functor_primitives.h @@ -53,7 +53,7 @@ struct ExpFunctor { HOSTDEVICE explicit inline ExpFunctor(int n) {} - HOSTDEVICE inline Ty operator()(const Tx& x) const { + HOSTDEVICE inline Ty operator()(const Tx x) const { return static_cast(details::Exp(x)); } }; @@ -67,7 +67,7 @@ struct IdentityFunctor { HOSTDEVICE explicit inline IdentityFunctor(int n) {} - HOSTDEVICE inline Ty operator()(const Tx& x) const { + HOSTDEVICE inline Ty operator()(const Tx x) const { return static_cast(x); } }; @@ -85,7 +85,7 @@ struct DivideFunctor { HOSTDEVICE explicit inline DivideFunctor(int n) : n_inv((MPType)(1.0 / n)) {} - HOSTDEVICE inline Ty operator()(const Tx& x) const { + HOSTDEVICE inline Ty operator()(const Tx x) const { return static_cast(static_cast(x) * n_inv); } @@ -102,7 +102,7 @@ struct InverseFunctor { HOSTDEVICE explicit inline InverseFunctor(int n) {} - HOSTDEVICE inline Ty operator()(const Tx& x) const { + HOSTDEVICE inline Ty operator()(const Tx x) const { return static_cast(-x); } }; @@ -116,7 +116,7 @@ struct SquareFunctor { HOSTDEVICE explicit inline SquareFunctor(int n) {} - HOSTDEVICE inline Ty operator()(const Tx& x) const { + HOSTDEVICE inline Ty operator()(const Tx x) const { return static_cast(x) * static_cast(x); } }; @@ -130,7 +130,7 @@ template struct MinFunctor { inline T initial() { return static_cast(std::numeric_limits::max()); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return (b < a) ? b : a; } }; @@ -144,7 +144,7 @@ struct MaxFunctor { return static_cast(std::numeric_limits::lowest()); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return (b > a) ? b : a; } }; @@ -156,7 +156,7 @@ template struct AddFunctor { inline T initial() { return static_cast(0.0f); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return b + a; } }; @@ -168,7 +168,7 @@ template struct MulFunctor { inline T initial() { return static_cast(1.0f); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return b * a; } }; @@ -180,7 +180,7 @@ template struct LogicalOrFunctor { inline T initial() { return static_cast(false); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return b || a; } }; @@ -192,7 +192,7 @@ template struct LogicalAndFunctor { inline T initial() { return static_cast(true); } - __device__ __forceinline__ T operator()(const T& a, const T& b) const { + __device__ __forceinline__ T operator()(const T a, const T b) const { return b && a; } }; @@ -204,7 +204,7 @@ template struct SubFunctor { inline T initial() { return static_cast(0.0f); } - inline HOSTDEVICE T operator()(const T& a, const T& b) const { return a - b; } + inline HOSTDEVICE T operator()(const T a, const T b) const { return a - b; } }; /** @@ -214,7 +214,7 @@ template struct DivFunctor { inline T initial() { return static_cast(1.0f); } - inline HOSTDEVICE T operator()(const T& a, const T& b) const { return a / b; } + inline HOSTDEVICE T operator()(const T a, const T b) const { return a / b; } }; template @@ -222,7 +222,7 @@ struct DivFunctor::value>::type> { inline T initial() { return static_cast(1.0f); } - inline HOSTDEVICE T operator()(const T& a, const T& b) const { + inline HOSTDEVICE T operator()(const T a, const T b) const { // For int32/int64, need to check whether the divison is zero. PADDLE_ENFORCE_NE(b, 0, platform::errors::InvalidArgument( @@ -239,7 +239,7 @@ template struct FloorDivFunctor { inline T initial() { return static_cast(1.0f); } - inline HOSTDEVICE T operator()(const T& a, const T& b) const { + inline HOSTDEVICE T operator()(const T a, const T b) const { PADDLE_ENFORCE_NE(b, 0, platform::errors::InvalidArgument( "Integer division by zero encountered " diff --git a/paddle/pten/kernels/gpu/math_kernel.cu b/paddle/pten/kernels/gpu/math_kernel.cu index f41934313d6740e44d198c411ccea05dabe1bda6..557080638038d4cbc7705e28b2ee3ff49340df9a 100644 --- a/paddle/pten/kernels/gpu/math_kernel.cu +++ b/paddle/pten/kernels/gpu/math_kernel.cu @@ -61,7 +61,7 @@ struct DivideFunctor { HOSTDEVICE explicit inline DivideFunctor(int n) : n_inv(static_cast(1.0 / n)) {} - HOSTDEVICE inline T operator()(const T& x) const { return x * n_inv; } + HOSTDEVICE inline T operator()(const T x) const { return x * n_inv; } private: T n_inv;