From 0b10ab2b7e4b7671d923bbdc98abc6e01deb2902 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 5 Jun 2018 08:46:20 -0400 Subject: [PATCH 1/6] Add missing macro MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT to hip_complex.h [ROCm/hip commit: 5eeb57b0a6b2113bf5ec5fdcb333f6834522698c] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 973d5f564b..22808ed576 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -94,6 +94,9 @@ THE SOFTWARE. ret.y = lhs.y * rhs; \ return ret; \ } +#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ + ComplexT(T val) : x(val), y(val) {} \ + ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif From cf6cdab0294a95f91448accf1407589a5b61cf0f Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 31 May 2018 13:58:44 -0400 Subject: [PATCH 2/6] Include cmath instead of math.h in hip_complex.h [ROCm/hip commit: 325cf3ccf0ec626bed2dccfacef6176e3e40cc35] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 22808ed576..3dc7a2ed56 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,7 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" -#include +#include #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From 0e29f327e2e2507986a4dca9231b87ce1b0654c7 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 6 Jun 2018 11:15:54 -0400 Subject: [PATCH 3/6] Includes or by __cplusplus in hip_complex.h [ROCm/hip commit: 2523c39a3706fc8d95d196a1eb4179497dda4a7e] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 3dc7a2ed56..72f139ba84 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,15 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" + +// TODO: Clang has a bug which allows device functions to call std functions +// when std functions are introduced into default namespace by using statement. +// math.h may be included after this bug is fixed. +#if __cplusplus #include +#else +#include "math.h" +#endif #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From fa9e73ccdc3be16d5c2564a7c238967647515317 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sun, 17 Jun 2018 12:18:05 -0400 Subject: [PATCH 4/6] Add missing __device__ __host__ to complex constructor Also add missing typedef value_type [ROCm/hip commit: 7a5605d006862c423710171f33641be7f54432ac] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 72f139ba84..783833dc7e 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -103,14 +103,15 @@ THE SOFTWARE. return ret; \ } #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ - ComplexT(T val) : x(val), y(val) {} \ - ComplexT(T val1, T val2) : x(val1), y(val2) {} + __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ + __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif struct hipFloatComplex { #ifdef __cplusplus public: + typedef float value_type; __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} @@ -130,6 +131,7 @@ struct hipFloatComplex { struct hipDoubleComplex { #ifdef __cplusplus public: + typedef double value_type; __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {} From a46f62a5c05aae9d2509d1aaa376711e5312ca7b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 11:57:57 -0400 Subject: [PATCH 5/6] Add abs/real/imag functions for hipFloatComplex/hipDoubleComplex [ROCm/hip commit: 9181fbb0b7e5f7db4df486b8600ea78092c4ffb1] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 783833dc7e..7f4dd840b7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -234,7 +234,6 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); } - __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; } __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; } @@ -314,4 +313,12 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } +#define __DEFINE_HIP_COMPLEX_FUN(func) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } + +__DEFINE_HIP_COMPLEX_FUN(abs) +__DEFINE_HIP_COMPLEX_FUN(real) +__DEFINE_HIP_COMPLEX_FUN(imag) + #endif From b8ae2784e955b798f241a5b01c52a24d3044a253 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 14:35:48 -0400 Subject: [PATCH 6/6] Add conj, operator-,==,!= for hipFloatComplex/hipDoubleComplex [ROCm/hip commit: 84da72dae8b204cd4c89e78b2d99219e33011290] --- .../hip/include/hip/hcc_detail/hip_complex.h | 44 ++++++++++++++++--- 1 file changed, 38 insertions(+), 6 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 7f4dd840b7..48a2852a5a 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -35,6 +35,24 @@ THE SOFTWARE. #endif #if __cplusplus +#define COMPLEX_NEG_OP_OVERLOAD(type) \ + __device__ __host__ static inline type operator-(const type& op) { \ + type ret; \ + ret.x = -op.x; \ + ret.y = -op.y; \ + return ret; \ + } + +#define COMPLEX_EQ_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \ + return lhs.x == rhs.x && lhs.y == rhs.y; \ + } + +#define COMPLEX_NE_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \ + return !(lhs == rhs); \ + } + #define COMPLEX_ADD_OP_OVERLOAD(type) \ __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \ type ret; \ @@ -150,6 +168,9 @@ struct hipDoubleComplex { #if __cplusplus +COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex) +COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex) +COMPLEX_NE_OP_OVERLOAD(hipFloatComplex) COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex) COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex) COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex) @@ -169,6 +190,9 @@ COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long) +COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex) COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex) COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex) COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex) @@ -313,12 +337,20 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } -#define __DEFINE_HIP_COMPLEX_FUN(func) \ -__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ -__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } +// Complex functions returning real numbers. +#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); } -__DEFINE_HIP_COMPLEX_FUN(abs) -__DEFINE_HIP_COMPLEX_FUN(real) -__DEFINE_HIP_COMPLEX_FUN(imag) +__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs) +__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal) +__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag) + +// Complex functions returning complex numbers. +#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \ +__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); } + +__DEFINE_HIP_COMPLEX_FUN(conj, hipConj) #endif