diff --git a/include/hcc_detail/hipComplex.h b/include/hcc_detail/hipComplex.h index 66d14d191e..910cee946d 100644 --- a/include/hcc_detail/hipComplex.h +++ b/include/hcc_detail/hipComplex.h @@ -26,45 +26,45 @@ typedef struct{ float y; }hipFloatComplex; -__device__ inline float hipCrealf(hipFloatComplex z){ +__device__ static inline float hipCrealf(hipFloatComplex z){ return z.x; } -__device__ inline float hipCimagf(hipFloatComplex z){ +__device__ static inline float hipCimagf(hipFloatComplex z){ return z.y; } -__device__ inline hipFloatComplex make_hipFloatComplex(float a, float b){ +__device__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){ hipFloatComplex z; z.x = a; z.y = b; return z; } -__device__ inline hipFloatComplex hipConjf(hipFloatComplex z){ +__device__ static inline hipFloatComplex hipConjf(hipFloatComplex z){ hipFloatComplex ret; ret.x = z.x; - ret.y = z.y; + ret.y = -z.y; return ret; } -__device__ inline float hipCsqabsf(hipFloatComplex z){ +__device__ static inline float hipCsqabsf(hipFloatComplex z){ return z.x * z.x + z.y * z.y; } -__device__ hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){ +__device__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x + q.x, p.y + q.y); } -__device__ hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){ +__device__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x - q.x, p.y - q.y); } -__device__ hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){ +__device__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); } -__device__ hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ +__device__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ float sqabs = hipCsqabsf(q); hipFloatComplex ret; ret.x = (p.x * q.x + p.y * q.y)/sqabs; @@ -72,7 +72,7 @@ __device__ hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ return ret; } -__device__ float hipCabsf(hipFloatComplex z){ +__device__ static inline float hipCabsf(hipFloatComplex z){ return sqrtf(hipCsqabsf(z)); } @@ -82,45 +82,45 @@ typedef struct{ double y; }hipDoubleComplex; -__device__ inline double hipCreal(hipDoubleComplex z){ +__device__ static inline double hipCreal(hipDoubleComplex z){ return z.x; } -__device__ inline double hipCimag(hipDoubleComplex z){ +__device__ static inline double hipCimag(hipDoubleComplex z){ return z.y; } -__device__ inline hipDoubleComplex make_hipDoubleComplex(double a, double b){ +__device__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){ hipDoubleComplex z; z.x = a; z.y = b; return z; } -__device__ inline hipDoubleComplex hipConj(hipDoubleComplex z){ +__device__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){ hipDoubleComplex ret; ret.x = z.x; ret.y = z.y; return ret; } -__device__ inline double hipCsqabs(hipDoubleComplex z){ +__device__ static inline double hipCsqabs(hipDoubleComplex z){ return z.x * z.x + z.y * z.y; } -__device__ hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){ +__device__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x + q.x, p.y + q.y); } -__device__ hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){ +__device__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x - q.x, p.y - q.y); } -__device__ hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){ +__device__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); } -__device__ hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ +__device__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ double sqabs = hipCsqabs(q); hipDoubleComplex ret; ret.x = (p.x * q.x + p.y * q.y)/sqabs; @@ -128,28 +128,28 @@ __device__ hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ return ret; } -__device__ inline double hipCabs(hipDoubleComplex z){ +__device__ static inline double hipCabs(hipDoubleComplex z){ return sqrtf(hipCsqabs(z)); } typedef hipFloatComplex hipComplex; -__device__ inline hipComplex make_hipComplex(float x, +__device__ static inline hipComplex make_hipComplex(float x, float y){ return make_hipFloatComplex(x, y); } -__device__ inline hipFloatComplex hipComplexDoubleToFloat +__device__ static inline hipFloatComplex hipComplexDoubleToFloat (hipDoubleComplex z){ return make_hipFloatComplex((float)z.x, (float)z.y); } -__device__ inline hipDoubleComplex hipComplexFloatToDouble +__device__ static inline hipDoubleComplex hipComplexFloatToDouble (hipFloatComplex z){ return make_hipDoubleComplex((double)z.x, (double)z.y); } -__device__ inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ +__device__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ float real = (p.x * q.x) + r.x; float imag = (q.x * p.y) + r.y; @@ -159,7 +159,7 @@ __device__ inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ return make_hipComplex(real, imag); } -__device__ inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){ +__device__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){ float real = (p.x * q.x) + r.x; float imag = (q.x * p.y) + r.y; @@ -169,6 +169,4 @@ __device__ inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex return make_hipDoubleComplex(real, imag); } - - #endif diff --git a/include/hipComplex.h b/include/hipComplex.h index 0fdafb8263..27281a2df4 100644 --- a/include/hipComplex.h +++ b/include/hipComplex.h @@ -24,7 +24,7 @@ THE SOFTWARE. #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include #elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) -#include "cuComplex.h" +#include #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif diff --git a/include/nvcc_detail/hipComplex.h b/include/nvcc_detail/hipComplex.h new file mode 100644 index 0000000000..b5c182bd4d --- /dev/null +++ b/include/nvcc_detail/hipComplex.h @@ -0,0 +1,108 @@ +#ifndef HIPCOMPLEX_H +#define HIPCOMPLEX_H + +#include"cuComplex.h" + +typedef cuFloatComplex hipFloatComplex; + +__device__ __host__ static inline float hipCrealf(hipFloatComplex z){ + return cuCrealf(z); +} + +__device__ __host__ static inline float hipCimagf(hipFloatComplex z){ + return cuCimagf(z); +} + +__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){ + return make_cuFloatComplex(a, b); +} + +__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z){ + return cuConjf(z); +} + +__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z){ + return cuCabsf(z) * cuCabsf(z); +} + +__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){ + return cuCaddf(p, q); +} + +__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){ + return cuCsubf(p, q); +} + +__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){ + return cuCmulf(p, q); +} + +__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ + return cuCdivf(p, q); +} + +__device__ __host__ static inline float hipCabsf(hipFloatComplex z){ + return cuCabsf(p, q); +} + +typedef cuDoubleComplex hipDoubleComplex; + +__device__ __host__ static inline double hipCreal(hipDoubleComplex z){ + return cuCreal(z); +} + +__device__ __host__ static inline double hipCimag(hipDoubleComplex z){ + return cuCimag(z); +} + +__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){ + return make_cuDoubleComplex(a, b); +} + +__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){ + return cuConj(z); +} + +__device__ __host__ static inline hipDoubleComplex hipCsqabs(hipDoubleComplex z){ + return cuCabs(z) * cuCabs(z); +} + +__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){ + return cuCadd(p, q); +} + +__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){ + return cuCsub(p, q); +} + +__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ + return cuCdiv(p, q); +} + +__device__ __host__ static inline double hipCabs(hipDoubleComplex z){ + return cuCabs(z); +} + +typedef cuFloatComplex hipComplex; + +__device__ __host__ static inline hipComplex make_Complex(float x, float y){ + return make_cuComplex(x, y); +} + +__device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z){ + return cuComplexDoubleToFloat(z); +} + +__device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z){ + return cuComplexFloatToDouble(z); +} + +__device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ + return cuCfmaf(p, q, r); +} + +__device__ __host__ static inline hipDoubleComplex hipCfma(hipComplex p, hipComplex q, hipComplex r){ + return cuCfma(p, q, r); +} + +#endif