AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2.

Browse Source
This commit is contained in:
Zuhaib Khan
2019-05-28 16:57:51 -04:00
parent e17f94e080
commit cc0f3445bb
1 changed files with 92 additions and 128 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_complex.h
+92 -128
View File
@@ -120,51 +120,102 @@ THE SOFTWARE.
ret.y = lhs.y * rhs; \ ret.y = lhs.y * rhs; \
return ret; \ return ret; \
} }
#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \
explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \
__device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
#endif #endif
struct hipFloatComplex { typedef float2 hipFloatComplex;
#ifdef __cplusplus
public: __device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
typedef float value_type;
__device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} __device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} __device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short) hipFloatComplex z;
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short) z.x = a;
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned int) z.y = b;
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed int) return z;
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, double) }
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long) __device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long long) hipFloatComplex ret;
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long long) ret.x = z.x;
#endif ret.y = -z.y;
float x, y; return ret;
} __attribute__((aligned(8))); }
__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
return z.x * z.x + z.y * z.y;
}
__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
return make_hipFloatComplex(p.x + q.x, p.y + q.y);
}
__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
return make_hipFloatComplex(p.x - q.x, p.y - q.y);
}
__device__ __host__ 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__ __host__ 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;
ret.y = (p.y * q.x - p.x * q.y) / sqabs;
return ret;
}
__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
typedef double2 hipDoubleComplex;
__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
hipDoubleComplex z;
z.x = a;
z.y = b;
return z;
}
__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
hipDoubleComplex ret;
ret.x = z.x;
ret.y = z.y;
return ret;
}
__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
return z.x * z.x + z.y * z.y;
}
__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
}
__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
}
__device__ __host__ 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__ __host__ 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;
ret.y = (p.y * q.x - p.x * q.y) / sqabs;
return ret;
}
__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
struct hipDoubleComplex {
#ifdef __cplusplus
public:
typedef double value_type;
__device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, float)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long long)
#endif
double x, y;
} __attribute__((aligned(16)));
#if __cplusplus #if __cplusplus
@@ -214,93 +265,6 @@ COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
#endif #endif
__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
hipFloatComplex z;
z.x = a;
z.y = b;
return z;
}
__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
hipFloatComplex ret;
ret.x = z.x;
ret.y = -z.y;
return ret;
}
__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
return z.x * z.x + z.y * z.y;
}
__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
return make_hipFloatComplex(p.x + q.x, p.y + q.y);
}
__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
return make_hipFloatComplex(p.x - q.x, p.y - q.y);
}
__device__ __host__ 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__ __host__ 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;
ret.y = (p.y * q.x - p.x * q.y) / sqabs;
return ret;
}
__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; }
__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
hipDoubleComplex z;
z.x = a;
z.y = b;
return z;
}
__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
hipDoubleComplex ret;
ret.x = z.x;
ret.y = z.y;
return ret;
}
__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
return z.x * z.x + z.y * z.y;
}
__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
}
__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
}
__device__ __host__ 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__ __host__ 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;
ret.y = (p.y * q.x - p.x * q.y) / sqabs;
return ret;
}
__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
typedef hipFloatComplex hipComplex; typedef hipFloatComplex hipComplex;