23 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H 24 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H 38 #define COMPLEX_NEG_OP_OVERLOAD(type) \ 39 __device__ __host__ static inline type operator-(const type& op) { \ 46 #define COMPLEX_EQ_OP_OVERLOAD(type) \ 47 __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \ 48 return lhs.x == rhs.x && lhs.y == rhs.y; \ 51 #define COMPLEX_NE_OP_OVERLOAD(type) \ 52 __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \ 53 return !(lhs == rhs); \ 56 #define COMPLEX_ADD_OP_OVERLOAD(type) \ 57 __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \ 59 ret.x = lhs.x + rhs.x; \ 60 ret.y = lhs.y + rhs.y; \ 64 #define COMPLEX_SUB_OP_OVERLOAD(type) \ 65 __device__ __host__ static inline type operator-(const type& lhs, const type& rhs) { \ 67 ret.x = lhs.x - rhs.x; \ 68 ret.y = lhs.y - rhs.y; \ 72 #define COMPLEX_MUL_OP_OVERLOAD(type) \ 73 __device__ __host__ static inline type operator*(const type& lhs, const type& rhs) { \ 75 ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \ 76 ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \ 80 #define COMPLEX_DIV_OP_OVERLOAD(type) \ 81 __device__ __host__ static inline type operator/(const type& lhs, const type& rhs) { \ 83 ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \ 84 ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \ 85 ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \ 86 ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \ 90 #define COMPLEX_ADD_PREOP_OVERLOAD(type) \ 91 __device__ __host__ static inline type& operator+=(type& lhs, const type& rhs) { \ 97 #define COMPLEX_SUB_PREOP_OVERLOAD(type) \ 98 __device__ __host__ static inline type& operator-=(type& lhs, const type& rhs) { \ 104 #define COMPLEX_MUL_PREOP_OVERLOAD(type) \ 105 __device__ __host__ static inline type& operator*=(type& lhs, const type& rhs) { \ 110 #define COMPLEX_DIV_PREOP_OVERLOAD(type) \ 111 __device__ __host__ static inline type& operator/=(type& lhs, const type& rhs) { \ 116 #define COMPLEX_SCALAR_PRODUCT(type, type1) \ 117 __device__ __host__ static inline type operator*(const type& lhs, type1 rhs) { \ 119 ret.x = lhs.x * rhs; \ 120 ret.y = lhs.y * rhs; \ 123 #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ 124 explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ 125 __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} 132 typedef float value_type;
134 explicit __device__
__host__ hipFloatComplex(
float x) : x(x), y(0.0f) {}
135 __device__
__host__ hipFloatComplex(
float x,
float y) : x(x), y(y) {}
136 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
unsigned short)
137 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
signed short)
138 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
unsigned int)
139 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
signed int)
140 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
double)
141 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
unsigned long)
142 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
signed long)
143 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
unsigned long long)
144 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex,
signed long long)
152 typedef double value_type;
154 explicit __device__
__host__ hipDoubleComplex(
double x) : x(x), y(0.0f) {}
155 __device__
__host__ hipDoubleComplex(
double x,
double y) : x(x), y(y) {}
156 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
unsigned short)
157 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
signed short)
158 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
unsigned int)
159 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
signed int)
160 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
float)
161 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
unsigned long)
162 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
signed long)
163 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
unsigned long long)
164 MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex,
signed long long)
236 return z.x * z.x + z.y * z.y;
240 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
244 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
248 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
252 float sqabs = hipCsqabsf(q);
254 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
255 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
280 return z.x * z.x + z.y * z.y;
284 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
288 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
292 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
296 double sqabs = hipCsqabs(q);
298 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
299 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
308 return make_hipFloatComplex(x, y);
312 return make_hipFloatComplex((
float)z.x, (
float)z.y);
316 return make_hipDoubleComplex((
double)z.x, (
double)z.y);
320 float real = (p.x * q.x) + r.x;
321 float imag = (q.x * p.y) + r.y;
323 real = -(p.y * q.y) + real;
324 imag = (p.x * q.y) + imag;
326 return make_hipComplex(real, imag);
331 double real = (p.x * q.x) + r.x;
332 double imag = (q.x * p.y) + r.y;
334 real = -(p.y * q.y) + real;
335 imag = (p.x * q.y) + imag;
337 return make_hipDoubleComplex(real, imag);
341 #define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \ 342 __device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \ 343 __device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); } 345 __DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
346 __DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal)
347 __DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag)
350 #define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \ 351 __device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \ 352 __device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); } 354 __DEFINE_HIP_COMPLEX_FUN(conj, hipConj)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Copies the memory address of symbol symbolName to devPtr.
Definition: hip_fp16_math_fwd.h:53
Definition: hip_complex.h:129
#define __host__
Definition: host_defines.h:41
Definition: hip_complex.h:149
Defines the different newt vector types for HIP runtime.