diff --git a/paddle/fluid/platform/complex128.h b/paddle/fluid/platform/complex128.h index d6fddd672a0f890b5cae6a2c4c453f9da761da68..da2f83c3497cce7b162336360690e1e76bce8b19 100644 --- a/paddle/fluid/platform/complex128.h +++ b/paddle/fluid/platform/complex128.h @@ -47,6 +47,10 @@ #define HOST #endif +#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP) +#define PADDLE_WITH_CUDA_OR_HIP_COMPLEX128 +#endif + namespace paddle { namespace platform { @@ -217,7 +221,8 @@ struct PADDLE_ALIGN(16) complex128 { HOSTDEVICE inline complex128 operator+(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::complex(a.real, a.imag) + thrust::complex(b.real, b.imag)); #else @@ -227,7 +232,8 @@ HOSTDEVICE inline complex128 operator+(const complex128& a, HOSTDEVICE inline complex128 operator-(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::complex(a.real, a.imag) - thrust::complex(b.real, b.imag)); #else @@ -237,7 +243,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a, HOSTDEVICE inline complex128 operator*(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::complex(a.real, a.imag) * thrust::complex(b.real, b.imag)); #else @@ -248,7 +255,8 @@ HOSTDEVICE inline complex128 operator*(const complex128& a, HOSTDEVICE inline complex128 operator/(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::complex(a.real, a.imag) / thrust::complex(b.real, b.imag)); #else @@ -259,7 +267,8 @@ 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(a.real, a.imag)); #else complex128 res; @@ -271,7 +280,8 @@ HOSTDEVICE inline complex128 operator-(const complex128& a) { HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT 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(a.real, a.imag) += thrust::complex(b.real, b.imag)); return a; @@ -284,7 +294,8 @@ HOSTDEVICE inline complex128& operator+=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT 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(a.real, a.imag) -= thrust::complex(b.real, b.imag)); return a; @@ -297,7 +308,8 @@ HOSTDEVICE inline complex128& operator-=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT 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(a.real, a.imag) *= thrust::complex(b.real, b.imag)); return a; @@ -310,7 +322,8 @@ HOSTDEVICE inline complex128& operator*=(complex128& a, // NOLINT HOSTDEVICE inline complex128& operator/=(complex128& a, // NOLINT 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(a.real, a.imag) /= thrust::complex(b.real, b.imag)); return a; @@ -353,7 +366,7 @@ HOSTDEVICE inline bool operator>=(const complex128& a, const complex128& b) { } 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 return __isnan(a.real) || __isnan(a.imag); #else @@ -362,7 +375,7 @@ HOSTDEVICE inline bool(isnan)(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 return __isinf(a.real) || __isinf(a.imag); #else @@ -375,7 +388,8 @@ HOSTDEVICE inline bool(isfinite)(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(a.real, a.imag)); #else return std::abs(std::complex(a.real, a.imag)); @@ -383,7 +397,8 @@ HOSTDEVICE inline double(abs)(const complex128& a) { } 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(a.real, a.imag), thrust::complex(b.real, b.imag))); #else @@ -392,7 +407,8 @@ HOSTDEVICE inline complex128(pow)(const complex128& a, const complex128& b) { } 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(a.real, a.imag))); #else return std::sqrt(std::complex(a)); @@ -400,7 +416,8 @@ HOSTDEVICE inline complex128(sqrt)(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(a.real, a.imag))); #else return std::tanh(std::complex(a)); @@ -408,7 +425,8 @@ HOSTDEVICE inline complex128(tanh)(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(a.real, a.imag))); #else return complex128(std::log(std::complex(a))); diff --git a/paddle/fluid/platform/complex64.h b/paddle/fluid/platform/complex64.h index 9d55ba19105a617266f4aa77594c9da1b54145fe..0aad7bd9dd2a8f1d59833720b442e34afa176ca6 100644 --- a/paddle/fluid/platform/complex64.h +++ b/paddle/fluid/platform/complex64.h @@ -47,6 +47,10 @@ #define HOST #endif +#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP) +#define PADDLE_WITH_CUDA_OR_HIP_COMPLEX64 +#endif + #include "complex128.h" // NOLINT namespace paddle { @@ -224,7 +228,8 @@ struct PADDLE_ALIGN(8) complex64 { }; 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(a.real, a.imag) + thrust::complex(b.real, b.imag)); #else @@ -233,7 +238,8 @@ 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(a.real, a.imag) - thrust::complex(b.real, b.imag)); #else @@ -242,7 +248,8 @@ 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(a.real, a.imag) * thrust::complex(b.real, b.imag)); #else @@ -252,7 +259,8 @@ 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(a.real, a.imag) / thrust::complex(b.real, b.imag)); #else @@ -263,7 +271,8 @@ HOSTDEVICE inline complex64 operator/(const complex64& a, const complex64& b) { } 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(a.real, a.imag)); #else complex64 res; @@ -275,7 +284,8 @@ HOSTDEVICE inline complex64 operator-(const complex64& a) { HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT 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(a.real, a.imag) += thrust::complex(b.real, b.imag)); return a; @@ -288,7 +298,8 @@ HOSTDEVICE inline complex64& operator+=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT 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(a.real, a.imag) -= thrust::complex(b.real, b.imag)); return a; @@ -301,7 +312,8 @@ HOSTDEVICE inline complex64& operator-=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT 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(a.real, a.imag) *= thrust::complex(b.real, b.imag)); return a; @@ -314,7 +326,8 @@ HOSTDEVICE inline complex64& operator*=(complex64& a, // NOLINT HOSTDEVICE inline complex64& operator/=(complex64& a, // NOLINT 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(a.real, a.imag) /= thrust::complex(b.real, b.imag)); return a; @@ -357,7 +370,7 @@ HOSTDEVICE inline bool operator>=(const complex64& a, const complex64& b) { } 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 return __isnanf(a.real) || __isnanf(a.imag); #else @@ -366,7 +379,7 @@ HOSTDEVICE inline bool(isnan)(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 return __isinff(a.real) || __isinff(a.imag); #else @@ -379,7 +392,8 @@ HOSTDEVICE inline bool(isfinite)(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(a.real, a.imag))); #else return std::abs(std::complex(a.real, a.imag)); @@ -387,7 +401,8 @@ HOSTDEVICE inline float(abs)(const complex64& a) { } 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(a.real, a.imag), thrust::complex(b.real, b.imag))); #else @@ -396,7 +411,8 @@ HOSTDEVICE inline complex64(pow)(const complex64& a, const complex64& b) { } 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(a.real, a.imag))); #else return std::sqrt(std::complex(a)); @@ -404,7 +420,8 @@ HOSTDEVICE inline complex64(sqrt)(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(a.real, a.imag))); #else return std::tanh(std::complex(a)); @@ -412,7 +429,8 @@ HOSTDEVICE inline complex64(tanh)(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(a.real, a.imag))); #else return std::log(std::complex(a));