HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_hip_complex.h
1/*
2Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
3
4Permission is hereby granted, free of charge, to any person obtaining a copy
5of this software and associated documentation files (the "Software"), to deal
6in the Software without restriction, including without limitation the rights
7to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8copies of the Software, and to permit persons to whom the Software is
9furnished to do so, subject to the following conditions:
10
11The above copyright notice and this permission notice shall be included in
12all copies or substantial portions of the Software.
13
14THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20THE SOFTWARE.
21*/
22
23#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
24#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
25
26#include "hip/amd_detail/amd_hip_vector_types.h"
27
28#if defined(__HIPCC_RTC__)
29#define __HOST_DEVICE__ __device__
30#else
31#define __HOST_DEVICE__ __host__ __device__
32// TODO: Clang has a bug which allows device functions to call std functions
33// when std functions are introduced into default namespace by using statement.
34// math.h may be included after this bug is fixed.
35#if __cplusplus
36#include <cmath>
37#else
38#include "math.h"
39#endif
40#endif // !defined(__HIPCC_RTC__)
41
42#if __cplusplus
43#define COMPLEX_NEG_OP_OVERLOAD(type) \
44 __HOST_DEVICE__ static inline type operator-(const type& op) { \
45 type ret; \
46 ret.x = -op.x; \
47 ret.y = -op.y; \
48 return ret; \
49 }
50
51#define COMPLEX_EQ_OP_OVERLOAD(type) \
52 __HOST_DEVICE__ static inline bool operator==(const type& lhs, const type& rhs) { \
53 return lhs.x == rhs.x && lhs.y == rhs.y; \
54 }
55
56#define COMPLEX_NE_OP_OVERLOAD(type) \
57 __HOST_DEVICE__ static inline bool operator!=(const type& lhs, const type& rhs) { \
58 return !(lhs == rhs); \
59 }
60
61#define COMPLEX_ADD_OP_OVERLOAD(type) \
62 __HOST_DEVICE__ static inline type operator+(const type& lhs, const type& rhs) { \
63 type ret; \
64 ret.x = lhs.x + rhs.x; \
65 ret.y = lhs.y + rhs.y; \
66 return ret; \
67 }
68
69#define COMPLEX_SUB_OP_OVERLOAD(type) \
70 __HOST_DEVICE__ static inline type operator-(const type& lhs, const type& rhs) { \
71 type ret; \
72 ret.x = lhs.x - rhs.x; \
73 ret.y = lhs.y - rhs.y; \
74 return ret; \
75 }
76
77#define COMPLEX_MUL_OP_OVERLOAD(type) \
78 __HOST_DEVICE__ static inline type operator*(const type& lhs, const type& rhs) { \
79 type ret; \
80 ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \
81 ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \
82 return ret; \
83 }
84
85#define COMPLEX_DIV_OP_OVERLOAD(type) \
86 __HOST_DEVICE__ static inline type operator/(const type& lhs, const type& rhs) { \
87 type ret; \
88 ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \
89 ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \
90 ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \
91 ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \
92 return ret; \
93 }
94
95#define COMPLEX_ADD_PREOP_OVERLOAD(type) \
96 __HOST_DEVICE__ static inline type& operator+=(type& lhs, const type& rhs) { \
97 lhs.x += rhs.x; \
98 lhs.y += rhs.y; \
99 return lhs; \
100 }
101
102#define COMPLEX_SUB_PREOP_OVERLOAD(type) \
103 __HOST_DEVICE__ static inline type& operator-=(type& lhs, const type& rhs) { \
104 lhs.x -= rhs.x; \
105 lhs.y -= rhs.y; \
106 return lhs; \
107 }
108
109#define COMPLEX_MUL_PREOP_OVERLOAD(type) \
110 __HOST_DEVICE__ static inline type& operator*=(type& lhs, const type& rhs) { \
111 type temp{lhs}; \
112 lhs.x = rhs.x * temp.x - rhs.y * temp.y; \
113 lhs.y = rhs.y * temp.x + rhs.x * temp.y; \
114 return lhs; \
115 }
116
117#define COMPLEX_DIV_PREOP_OVERLOAD(type) \
118 __HOST_DEVICE__ static inline type& operator/=(type& lhs, const type& rhs) { \
119 type temp; \
120 temp.x = (lhs.x*rhs.x + lhs.y * rhs.y) / (rhs.x*rhs.x + rhs.y*rhs.y); \
121 temp.y = (lhs.y * rhs.x - lhs.x * rhs.y) / (rhs.x*rhs.x + rhs.y*rhs.y); \
122 lhs = temp; \
123 return lhs; \
124 }
125
126#define COMPLEX_SCALAR_PRODUCT(type, type1) \
127 __HOST_DEVICE__ static inline type operator*(const type& lhs, type1 rhs) { \
128 type ret; \
129 ret.x = lhs.x * rhs; \
130 ret.y = lhs.y * rhs; \
131 return ret; \
132 }
133
134#endif
135
136typedef float2 hipFloatComplex;
137
138__HOST_DEVICE__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
139
140__HOST_DEVICE__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
141
142__HOST_DEVICE__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
144 z.x = a;
145 z.y = b;
146 return z;
147}
148
149__HOST_DEVICE__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
150 hipFloatComplex ret;
151 ret.x = z.x;
152 ret.y = -z.y;
153 return ret;
154}
155
156__HOST_DEVICE__ static inline float hipCsqabsf(hipFloatComplex z) {
157 return z.x * z.x + z.y * z.y;
158}
159
160__HOST_DEVICE__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
161 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
162}
163
164__HOST_DEVICE__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
165 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
166}
167
168__HOST_DEVICE__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
169 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
170}
171
172__HOST_DEVICE__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
173 float sqabs = hipCsqabsf(q);
174 hipFloatComplex ret;
175 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
176 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
177 return ret;
178}
179
180__HOST_DEVICE__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
181
182
184
185__HOST_DEVICE__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
186
187__HOST_DEVICE__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
188
189__HOST_DEVICE__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
191 z.x = a;
192 z.y = b;
193 return z;
194}
195
196__HOST_DEVICE__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
198 ret.x = z.x;
199 ret.y = -z.y;
200 return ret;
201}
202
203__HOST_DEVICE__ static inline double hipCsqabs(hipDoubleComplex z) {
204 return z.x * z.x + z.y * z.y;
205}
206
207__HOST_DEVICE__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
208 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
209}
210
211__HOST_DEVICE__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
212 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
213}
214
215__HOST_DEVICE__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
216 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
217}
218
219__HOST_DEVICE__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
220 double sqabs = hipCsqabs(q);
222 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
223 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
224 return ret;
225}
226
227__HOST_DEVICE__ static inline double hipCabs(hipDoubleComplex z) { return sqrt(hipCsqabs(z)); }
228
229
230#if __cplusplus
231
232COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
233COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
234COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
235COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
236COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
237COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
238COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
239COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
240COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
241COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
242COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
243COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
244COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
245COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
246COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
247COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
248COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
249COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
250COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
251COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
252COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
253
254COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
255COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
256COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
257COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
258COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
259COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
260COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
261COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
262COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
263COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
264COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
265COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
266COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
267COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
268COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
269COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
270COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
271COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
272COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
273COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
274COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
275
276#endif
277
278
280
281__HOST_DEVICE__ static inline hipComplex make_hipComplex(float x, float y) {
282 return make_hipFloatComplex(x, y);
283}
284
285__HOST_DEVICE__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
286 return make_hipFloatComplex((float)z.x, (float)z.y);
287}
288
289__HOST_DEVICE__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
290 return make_hipDoubleComplex((double)z.x, (double)z.y);
291}
292
293__HOST_DEVICE__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
294 float real = (p.x * q.x) + r.x;
295 float imag = (q.x * p.y) + r.y;
296
297 real = -(p.y * q.y) + real;
298 imag = (p.x * q.y) + imag;
299
300 return make_hipComplex(real, imag);
301}
302
303__HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
305 double real = (p.x * q.x) + r.x;
306 double imag = (q.x * p.y) + r.y;
307
308 real = -(p.y * q.y) + real;
309 imag = (p.x * q.y) + imag;
310
311 return make_hipDoubleComplex(real, imag);
312}
313
314#endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:2023