23#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
24#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
26#if !defined(__HIPCC_RTC__)
27#include "hip/amd_detail/amd_hip_vector_types.h"
30#if defined(__HIPCC_RTC__)
31#define __HOST_DEVICE__ __device__
33#define __HOST_DEVICE__ __host__ __device__
45#define COMPLEX_NEG_OP_OVERLOAD(type) \
46 __HOST_DEVICE__ static inline type operator-(const type& op) { \
53#define COMPLEX_EQ_OP_OVERLOAD(type) \
54 __HOST_DEVICE__ static inline bool operator==(const type& lhs, const type& rhs) { \
55 return lhs.x == rhs.x && lhs.y == rhs.y; \
58#define COMPLEX_NE_OP_OVERLOAD(type) \
59 __HOST_DEVICE__ static inline bool operator!=(const type& lhs, const type& rhs) { \
60 return !(lhs == rhs); \
63#define COMPLEX_ADD_OP_OVERLOAD(type) \
64 __HOST_DEVICE__ static inline type operator+(const type& lhs, const type& rhs) { \
66 ret.x = lhs.x + rhs.x; \
67 ret.y = lhs.y + rhs.y; \
71#define COMPLEX_SUB_OP_OVERLOAD(type) \
72 __HOST_DEVICE__ static inline type operator-(const type& lhs, const type& rhs) { \
74 ret.x = lhs.x - rhs.x; \
75 ret.y = lhs.y - rhs.y; \
79#define COMPLEX_MUL_OP_OVERLOAD(type) \
80 __HOST_DEVICE__ static inline type operator*(const type& lhs, const type& rhs) { \
82 ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \
83 ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \
87#define COMPLEX_DIV_OP_OVERLOAD(type) \
88 __HOST_DEVICE__ static inline type operator/(const type& lhs, const type& rhs) { \
90 ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \
91 ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \
92 ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \
93 ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \
97#define COMPLEX_ADD_PREOP_OVERLOAD(type) \
98 __HOST_DEVICE__ static inline type& operator+=(type& lhs, const type& rhs) { \
104#define COMPLEX_SUB_PREOP_OVERLOAD(type) \
105 __HOST_DEVICE__ static inline type& operator-=(type& lhs, const type& rhs) { \
111#define COMPLEX_MUL_PREOP_OVERLOAD(type) \
112 __HOST_DEVICE__ static inline type& operator*=(type& lhs, const type& rhs) { \
114 lhs.x = rhs.x * temp.x - rhs.y * temp.y; \
115 lhs.y = rhs.y * temp.x + rhs.x * temp.y; \
119#define COMPLEX_DIV_PREOP_OVERLOAD(type) \
120 __HOST_DEVICE__ static inline type& operator/=(type& lhs, const type& rhs) { \
122 temp.x = (lhs.x*rhs.x + lhs.y * rhs.y) / (rhs.x*rhs.x + rhs.y*rhs.y); \
123 temp.y = (lhs.y * rhs.x - lhs.x * rhs.y) / (rhs.x*rhs.x + rhs.y*rhs.y); \
128#define COMPLEX_SCALAR_PRODUCT(type, type1) \
129 __HOST_DEVICE__ static inline type operator*(const type& lhs, type1 rhs) { \
131 ret.x = lhs.x * rhs; \
132 ret.y = lhs.y * rhs; \
140__HOST_DEVICE__
static inline float hipCrealf(
hipFloatComplex z) {
return z.x; }
142__HOST_DEVICE__
static inline float hipCimagf(
hipFloatComplex z) {
return z.y; }
144__HOST_DEVICE__
static inline hipFloatComplex make_hipFloatComplex(
float a,
float b) {
159 return z.x * z.x + z.y * z.y;
163 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
167 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
171 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
175 float sqabs = hipCsqabsf(q);
177 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
178 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
182__HOST_DEVICE__
static inline float hipCabsf(
hipFloatComplex z) {
return sqrtf(hipCsqabsf(z)); }
187__HOST_DEVICE__
static inline double hipCreal(
hipDoubleComplex z) {
return z.x; }
189__HOST_DEVICE__
static inline double hipCimag(
hipDoubleComplex z) {
return z.y; }
191__HOST_DEVICE__
static inline hipDoubleComplex make_hipDoubleComplex(
double a,
double b) {
206 return z.x * z.x + z.y * z.y;
210 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
214 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
218 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
222 double sqabs = hipCsqabs(q);
224 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
225 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
229__HOST_DEVICE__
static inline double hipCabs(
hipDoubleComplex z) {
return sqrt(hipCsqabs(z)); }
283__HOST_DEVICE__
static inline hipComplex make_hipComplex(
float x,
float y) {
284 return make_hipFloatComplex(x, y);
288 return make_hipFloatComplex((
float)z.x, (
float)z.y);
292 return make_hipDoubleComplex((
double)z.x, (
double)z.y);
296 float real = (p.x * q.x) + r.x;
297 float imag = (q.x * p.y) + r.y;
299 real = -(p.y * q.y) + real;
300 imag = (p.x * q.y) + imag;
302 return make_hipComplex(real, imag);
307 double real = (p.x * q.x) + r.x;
308 double imag = (q.x * p.y) + r.y;
310 real = -(p.y * q.y) + real;
311 imag = (p.x * q.y) + imag;
313 return make_hipDoubleComplex(real, imag);
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:2023