25#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
26#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
28#if !defined(__HIPCC_RTC__)
29#include "hip/amd_detail/amd_hip_vector_types.h"
32#if defined(__HIPCC_RTC__)
33#define __HOST_DEVICE__ __device__
35#define __HOST_DEVICE__ __host__ __device__
46typedef float2 hipFloatComplex;
48__HOST_DEVICE__
static inline float hipCrealf(hipFloatComplex z) {
return z.x; }
50__HOST_DEVICE__
static inline float hipCimagf(hipFloatComplex z) {
return z.y; }
52__HOST_DEVICE__
static inline hipFloatComplex make_hipFloatComplex(
float a,
float b) {
59__HOST_DEVICE__
static inline hipFloatComplex hipConjf(hipFloatComplex z) {
66__HOST_DEVICE__
static inline float hipCsqabsf(hipFloatComplex z) {
67 return z.x * z.x + z.y * z.y;
70__HOST_DEVICE__
static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
71 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
74__HOST_DEVICE__
static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
75 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
78__HOST_DEVICE__
static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
79 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
82__HOST_DEVICE__
static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
83 float sqabs = hipCsqabsf(q);
85 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
86 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
90__HOST_DEVICE__
static inline float hipCabsf(hipFloatComplex z) {
return sqrtf(hipCsqabsf(z)); }
93typedef double2 hipDoubleComplex;
95__HOST_DEVICE__
static inline double hipCreal(hipDoubleComplex z) {
return z.x; }
97__HOST_DEVICE__
static inline double hipCimag(hipDoubleComplex z) {
return z.y; }
99__HOST_DEVICE__
static inline hipDoubleComplex make_hipDoubleComplex(
double a,
double b) {
106__HOST_DEVICE__
static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
107 hipDoubleComplex ret;
113__HOST_DEVICE__
static inline double hipCsqabs(hipDoubleComplex z) {
114 return z.x * z.x + z.y * z.y;
117__HOST_DEVICE__
static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
118 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
121__HOST_DEVICE__
static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
122 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
125__HOST_DEVICE__
static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
126 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
129__HOST_DEVICE__
static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
130 double sqabs = hipCsqabs(q);
131 hipDoubleComplex ret;
132 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
133 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
137__HOST_DEVICE__
static inline double hipCabs(hipDoubleComplex z) {
return sqrt(hipCsqabs(z)); }
139typedef hipFloatComplex hipComplex;
141__HOST_DEVICE__
static inline hipComplex make_hipComplex(
float x,
float y) {
142 return make_hipFloatComplex(x, y);
145__HOST_DEVICE__
static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
146 return make_hipFloatComplex((
float)z.x, (
float)z.y);
149__HOST_DEVICE__
static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
150 return make_hipDoubleComplex((
double)z.x, (
double)z.y);
153__HOST_DEVICE__
static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
154 float real = (p.x * q.x) + r.x;
155 float imag = (q.x * p.y) + r.y;
157 real = -(p.y * q.y) + real;
158 imag = (p.x * q.y) + imag;
160 return make_hipComplex(real, imag);
163__HOST_DEVICE__
static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
164 hipDoubleComplex r) {
165 double real = (p.x * q.x) + r.x;
166 double imag = (q.x * p.y) + r.y;
168 real = -(p.y * q.y) + real;
169 imag = (p.x * q.y) + imag;
171 return make_hipDoubleComplex(real, imag);
Definition amd_hip_vector_types.h:2035
Definition amd_hip_vector_types.h:2072