Skip to content

P
Paddle

Crayon鑫 / Paddle
与 Fork 源项目一致

Fork自 PaddlePaddle / Paddle

通知 1
Star 1
Fork 0
P
Paddle
未验证 提交 7b9fcaca 编写于 作者: C Chen Weihang 提交者： GitHub
浏览文件
操作

add marco cond for multi function (#32239)

上级 63abd500
隐藏空白更改
内联 并排
Showing with 68 addition and 32 deletion
paddle/fluid/platform/complex128.h
浏览文件 @ 7b9fcaca
...@@ -47,6 +47,10 @@ ...@@ -47,6 +47,10 @@
#define HOST #define HOST
#endif #endif
#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP)
#define PADDLE_WITH_CUDA_OR_HIP_COMPLEX128
#endif
namespace paddle { namespace paddle {
namespace platform { namespace platform {
...@@ -217,7 +221,8 @@ struct PADDLE_ALIGN(16) complex128 { ...@@ -217,7 +221,8 @@ struct PADDLE_ALIGN(16) complex128 {
HOSTDEVICE inline complex128 operator+(const complex128& a, HOSTDEVICE inline complex128 operator+(const complex128& a,
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::complex<double>(a.real, a.imag) + return complex128(thrust::complex<double>(a.real, a.imag) +
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
#else #else
...@@ -227,7 +232,8 @@ HOSTDEVICE inline complex128 operator+(const complex128& a, ...@@ -227,7 +232,8 @@ HOSTDEVICE inline complex128 operator+(const complex128& a,
HOSTDEVICE inline complex128 operator-(const complex128& a, HOSTDEVICE inline complex128 operator-(const complex128& a,
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::complex<double>(a.real, a.imag) - return complex128(thrust::complex<double>(a.real, a.imag) -
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
#else #else
...@@ -237,7 +243,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a, ...@@ -237,7 +243,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a,
HOSTDEVICE inline complex128 operator*(const complex128& a, HOSTDEVICE inline complex128 operator*(const complex128& a,
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::complex<double>(a.real, a.imag) * return complex128(thrust::complex<double>(a.real, a.imag) *
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
#else #else
...@@ -248,7 +255,8 @@ HOSTDEVICE inline complex128 operator*(const complex128& a, ...@@ -248,7 +255,8 @@ HOSTDEVICE inline complex128 operator*(const complex128& a,
HOSTDEVICE inline complex128 operator/(const complex128& a, HOSTDEVICE inline complex128 operator/(const complex128& a,
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::complex<double>(a.real, a.imag) / return complex128(thrust::complex<double>(a.real, a.imag) /
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
#else #else
...@@ -259,7 +267,8 @@ HOSTDEVICE inline complex128 operator/(const complex128& a, ...@@ -259,7 +267,8 @@ HOSTDEVICE inline complex128 operator/(const complex128& a,
} }
HOSTDEVICE inline complex128 operator-(const complex128& a) { HOSTDEVICE inline complex128 operator-(const complex128& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(-thrust::complex<double>(a.real, a.imag)); return complex128(-thrust::complex<double>(a.real, a.imag));
#else #else
complex128 res; complex128 res;
...@@ -271,7 +280,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a) { ...@@ -271,7 +280,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a) {
HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex128(thrust::complex<double>(a.real, a.imag) += a = complex128(thrust::complex<double>(a.real, a.imag) +=
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
return a; return a;
...@@ -284,7 +294,8 @@ HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT ...@@ -284,7 +294,8 @@ HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT
HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex128(thrust::complex<double>(a.real, a.imag) -= a = complex128(thrust::complex<double>(a.real, a.imag) -=
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
return a; return a;
...@@ -297,7 +308,8 @@ HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT ...@@ -297,7 +308,8 @@ HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT
HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex128(thrust::complex<double>(a.real, a.imag) *= a = complex128(thrust::complex<double>(a.real, a.imag) *=
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
return a; return a;
...@@ -310,7 +322,8 @@ HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT ...@@ -310,7 +322,8 @@ HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT
HOSTDEVICE inline complex128& operator/=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator/=(complex128& a, // NOLINT
const complex128& b) { const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex128(thrust::complex<double>(a.real, a.imag) /= a = complex128(thrust::complex<double>(a.real, a.imag) /=
thrust::complex<double>(b.real, b.imag)); thrust::complex<double>(b.real, b.imag));
return a; return a;
...@@ -353,7 +366,7 @@ HOSTDEVICE inline bool operator>=(const complex128& a, const complex128& b) { ...@@ -353,7 +366,7 @@ HOSTDEVICE inline bool operator>=(const complex128& a, const complex128& b) {
} }
HOSTDEVICE inline bool(isnan)(const complex128& a) { HOSTDEVICE inline bool(isnan)(const complex128& a) {
#if defined(__CUDA_ARCH__) #if defined(PADDLE_WITH_CUDA) && defined(__CUDA_ARCH__)
// __isnanf not supported on HIP platform // __isnanf not supported on HIP platform
return __isnan(a.real) || __isnan(a.imag); return __isnan(a.real) || __isnan(a.imag);
#else #else
...@@ -362,7 +375,7 @@ HOSTDEVICE inline bool(isnan)(const complex128& a) { ...@@ -362,7 +375,7 @@ HOSTDEVICE inline bool(isnan)(const complex128& a) {
} }
HOSTDEVICE inline bool(isinf)(const complex128& a) { HOSTDEVICE inline bool(isinf)(const complex128& a) {
#if defined(__CUDA_ARCH__) #if defined(PADDLE_WITH_CUDA) && defined(__CUDA_ARCH__)
// __isinf not supported on HIP platform // __isinf not supported on HIP platform
return __isinf(a.real) || __isinf(a.imag); return __isinf(a.real) || __isinf(a.imag);
#else #else
...@@ -375,7 +388,8 @@ HOSTDEVICE inline bool(isfinite)(const complex128& a) { ...@@ -375,7 +388,8 @@ HOSTDEVICE inline bool(isfinite)(const complex128& a) {
} }
HOSTDEVICE inline double(abs)(const complex128& a) { HOSTDEVICE inline double(abs)(const complex128& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return thrust::abs(thrust::complex<double>(a.real, a.imag)); return thrust::abs(thrust::complex<double>(a.real, a.imag));
#else #else
return std::abs(std::complex<double>(a.real, a.imag)); return std::abs(std::complex<double>(a.real, a.imag));
...@@ -383,7 +397,8 @@ HOSTDEVICE inline double(abs)(const complex128& a) { ...@@ -383,7 +397,8 @@ HOSTDEVICE inline double(abs)(const complex128& a) {
} }
HOSTDEVICE inline complex128(pow)(const complex128& a, const complex128& b) { HOSTDEVICE inline complex128(pow)(const complex128& a, const complex128& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::pow(thrust::complex<double>(a.real, a.imag), return complex128(thrust::pow(thrust::complex<double>(a.real, a.imag),
thrust::complex<double>(b.real, b.imag))); thrust::complex<double>(b.real, b.imag)));
#else #else
...@@ -392,7 +407,8 @@ HOSTDEVICE inline complex128(pow)(const complex128& a, const complex128& b) { ...@@ -392,7 +407,8 @@ HOSTDEVICE inline complex128(pow)(const complex128& a, const complex128& b) {
} }
HOSTDEVICE inline complex128(sqrt)(const complex128& a) { HOSTDEVICE inline complex128(sqrt)(const complex128& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::sqrt(thrust::complex<double>(a.real, a.imag))); return complex128(thrust::sqrt(thrust::complex<double>(a.real, a.imag)));
#else #else
return std::sqrt(std::complex<double>(a)); return std::sqrt(std::complex<double>(a));
...@@ -400,7 +416,8 @@ HOSTDEVICE inline complex128(sqrt)(const complex128& a) { ...@@ -400,7 +416,8 @@ HOSTDEVICE inline complex128(sqrt)(const complex128& a) {
} }
HOSTDEVICE inline complex128(tanh)(const complex128& a) { HOSTDEVICE inline complex128(tanh)(const complex128& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::tanh(thrust::complex<double>(a.real, a.imag))); return complex128(thrust::tanh(thrust::complex<double>(a.real, a.imag)));
#else #else
return std::tanh(std::complex<double>(a)); return std::tanh(std::complex<double>(a));
...@@ -408,7 +425,8 @@ HOSTDEVICE inline complex128(tanh)(const complex128& a) { ...@@ -408,7 +425,8 @@ HOSTDEVICE inline complex128(tanh)(const complex128& a) {
} }
HOSTDEVICE inline complex128(log)(const complex128& a) { HOSTDEVICE inline complex128(log)(const complex128& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX128) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex128(thrust::log(thrust::complex<double>(a.real, a.imag))); return complex128(thrust::log(thrust::complex<double>(a.real, a.imag)));
#else #else
return complex128(std::log(std::complex<double>(a))); return complex128(std::log(std::complex<double>(a)));
......
paddle/fluid/platform/complex64.h
浏览文件 @ 7b9fcaca
...@@ -47,6 +47,10 @@ ...@@ -47,6 +47,10 @@
#define HOST #define HOST
#endif #endif
#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP)
#define PADDLE_WITH_CUDA_OR_HIP_COMPLEX64
#endif
#include "complex128.h" // NOLINT #include "complex128.h" // NOLINT
namespace paddle { namespace paddle {
...@@ -224,7 +228,8 @@ struct PADDLE_ALIGN(8) complex64 { ...@@ -224,7 +228,8 @@ struct PADDLE_ALIGN(8) complex64 {
}; };
HOSTDEVICE inline complex64 operator+(const complex64& a, const complex64& b) { HOSTDEVICE inline complex64 operator+(const complex64& a, const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::complex<float>(a.real, a.imag) + return complex64(thrust::complex<float>(a.real, a.imag) +
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
#else #else
...@@ -233,7 +238,8 @@ HOSTDEVICE inline complex64 operator+(const complex64& a, const complex64& b) { ...@@ -233,7 +238,8 @@ HOSTDEVICE inline complex64 operator+(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline complex64 operator-(const complex64& a, const complex64& b) { HOSTDEVICE inline complex64 operator-(const complex64& a, const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::complex<float>(a.real, a.imag) - return complex64(thrust::complex<float>(a.real, a.imag) -
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
#else #else
...@@ -242,7 +248,8 @@ HOSTDEVICE inline complex64 operator-(const complex64& a, const complex64& b) { ...@@ -242,7 +248,8 @@ HOSTDEVICE inline complex64 operator-(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline complex64 operator*(const complex64& a, const complex64& b) { HOSTDEVICE inline complex64 operator*(const complex64& a, const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::complex<float>(a.real, a.imag) * return complex64(thrust::complex<float>(a.real, a.imag) *
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
#else #else
...@@ -252,7 +259,8 @@ HOSTDEVICE inline complex64 operator*(const complex64& a, const complex64& b) { ...@@ -252,7 +259,8 @@ HOSTDEVICE inline complex64 operator*(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline complex64 operator/(const complex64& a, const complex64& b) { HOSTDEVICE inline complex64 operator/(const complex64& a, const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::complex<float>(a.real, a.imag) / return complex64(thrust::complex<float>(a.real, a.imag) /
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
#else #else
...@@ -263,7 +271,8 @@ HOSTDEVICE inline complex64 operator/(const complex64& a, const complex64& b) { ...@@ -263,7 +271,8 @@ HOSTDEVICE inline complex64 operator/(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline complex64 operator-(const complex64& a) { HOSTDEVICE inline complex64 operator-(const complex64& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(-thrust::complex<float>(a.real, a.imag)); return complex64(-thrust::complex<float>(a.real, a.imag));
#else #else
complex64 res; complex64 res;
...@@ -275,7 +284,8 @@ HOSTDEVICE inline complex64 operator-(const complex64& a) { ...@@ -275,7 +284,8 @@ HOSTDEVICE inline complex64 operator-(const complex64& a) {
HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT
const complex64& b) { const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex64(thrust::complex<float>(a.real, a.imag) += a = complex64(thrust::complex<float>(a.real, a.imag) +=
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
return a; return a;
...@@ -288,7 +298,8 @@ HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT ...@@ -288,7 +298,8 @@ HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT
HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT
const complex64& b) { const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex64(thrust::complex<float>(a.real, a.imag) -= a = complex64(thrust::complex<float>(a.real, a.imag) -=
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
return a; return a;
...@@ -301,7 +312,8 @@ HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT ...@@ -301,7 +312,8 @@ HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT
HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT
const complex64& b) { const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex64(thrust::complex<float>(a.real, a.imag) *= a = complex64(thrust::complex<float>(a.real, a.imag) *=
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
return a; return a;
...@@ -314,7 +326,8 @@ HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT ...@@ -314,7 +326,8 @@ HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT
HOSTDEVICE inline complex64& operator/=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator/=(complex64& a, // NOLINT
const complex64& b) { const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
a = complex64(thrust::complex<float>(a.real, a.imag) /= a = complex64(thrust::complex<float>(a.real, a.imag) /=
thrust::complex<float>(b.real, b.imag)); thrust::complex<float>(b.real, b.imag));
return a; return a;
...@@ -357,7 +370,7 @@ HOSTDEVICE inline bool operator>=(const complex64& a, const complex64& b) { ...@@ -357,7 +370,7 @@ HOSTDEVICE inline bool operator>=(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline bool(isnan)(const complex64& a) { HOSTDEVICE inline bool(isnan)(const complex64& a) {
#if defined(__CUDA_ARCH__) #if defined(PADDLE_WITH_CUDA) && defined(__CUDA_ARCH__)
// __isnanf not supported on HIP platform // __isnanf not supported on HIP platform
return __isnanf(a.real) || __isnanf(a.imag); return __isnanf(a.real) || __isnanf(a.imag);
#else #else
...@@ -366,7 +379,7 @@ HOSTDEVICE inline bool(isnan)(const complex64& a) { ...@@ -366,7 +379,7 @@ HOSTDEVICE inline bool(isnan)(const complex64& a) {
} }
HOSTDEVICE inline bool(isinf)(const complex64& a) { HOSTDEVICE inline bool(isinf)(const complex64& a) {
#if defined(__CUDA_ARCH__) #if defined(PADDLE_WITH_CUDA) && defined(__CUDA_ARCH__)
// __isinff not supported on HIP platform // __isinff not supported on HIP platform
return __isinff(a.real) || __isinff(a.imag); return __isinff(a.real) || __isinff(a.imag);
#else #else
...@@ -379,7 +392,8 @@ HOSTDEVICE inline bool(isfinite)(const complex64& a) { ...@@ -379,7 +392,8 @@ HOSTDEVICE inline bool(isfinite)(const complex64& a) {
} }
HOSTDEVICE inline float(abs)(const complex64& a) { HOSTDEVICE inline float(abs)(const complex64& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::abs(thrust::complex<float>(a.real, a.imag))); return complex64(thrust::abs(thrust::complex<float>(a.real, a.imag)));
#else #else
return std::abs(std::complex<float>(a.real, a.imag)); return std::abs(std::complex<float>(a.real, a.imag));
...@@ -387,7 +401,8 @@ HOSTDEVICE inline float(abs)(const complex64& a) { ...@@ -387,7 +401,8 @@ HOSTDEVICE inline float(abs)(const complex64& a) {
} }
HOSTDEVICE inline complex64(pow)(const complex64& a, const complex64& b) { HOSTDEVICE inline complex64(pow)(const complex64& a, const complex64& b) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::pow(thrust::complex<float>(a.real, a.imag), return complex64(thrust::pow(thrust::complex<float>(a.real, a.imag),
thrust::complex<float>(b.real, b.imag))); thrust::complex<float>(b.real, b.imag)));
#else #else
...@@ -396,7 +411,8 @@ HOSTDEVICE inline complex64(pow)(const complex64& a, const complex64& b) { ...@@ -396,7 +411,8 @@ HOSTDEVICE inline complex64(pow)(const complex64& a, const complex64& b) {
} }
HOSTDEVICE inline complex64(sqrt)(const complex64& a) { HOSTDEVICE inline complex64(sqrt)(const complex64& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::sqrt(thrust::complex<float>(a.real, a.imag))); return complex64(thrust::sqrt(thrust::complex<float>(a.real, a.imag)));
#else #else
return std::sqrt(std::complex<float>(a)); return std::sqrt(std::complex<float>(a));
...@@ -404,7 +420,8 @@ HOSTDEVICE inline complex64(sqrt)(const complex64& a) { ...@@ -404,7 +420,8 @@ HOSTDEVICE inline complex64(sqrt)(const complex64& a) {
} }
HOSTDEVICE inline complex64(tanh)(const complex64& a) { HOSTDEVICE inline complex64(tanh)(const complex64& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::tanh(thrust::complex<float>(a.real, a.imag))); return complex64(thrust::tanh(thrust::complex<float>(a.real, a.imag)));
#else #else
return std::tanh(std::complex<float>(a)); return std::tanh(std::complex<float>(a));
...@@ -412,7 +429,8 @@ HOSTDEVICE inline complex64(tanh)(const complex64& a) { ...@@ -412,7 +429,8 @@ HOSTDEVICE inline complex64(tanh)(const complex64& a) {
} }
HOSTDEVICE inline complex64(log)(const complex64& a) { HOSTDEVICE inline complex64(log)(const complex64& a) {
#if defined(__CUDA_ARCH__) || defined(__HIPCC__) #if defined(PADDLE_WITH_CUDA_OR_HIP_COMPLEX64) && \
(defined(__CUDA_ARCH__) || defined(__HIPCC__))
return complex64(thrust::log(thrust::complex<float>(a.real, a.imag))); return complex64(thrust::log(thrust::complex<float>(a.real, a.imag)));
#else #else
return std::log(std::complex<float>(a)); return std::log(std::complex<float>(a));
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或