HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_hip_complex.h
1/*
2Copyright (c) 2015 - 2023 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#if !defined(__HIPCC_RTC__)
27#include "hip/amd_detail/amd_hip_vector_types.h"
28#endif
29
30#if defined(__HIPCC_RTC__)
31#define __HOST_DEVICE__ __device__
32#else
33#define __HOST_DEVICE__ __host__ __device__
34// TODO: Clang has a bug which allows device functions to call std functions
35// when std functions are introduced into default namespace by using statement.
36// math.h may be included after this bug is fixed.
37#if __cplusplus
38#include <cmath>
39#else
40#include "math.h"
41#endif
42#endif // !defined(__HIPCC_RTC__)
43
44#if __cplusplus
45#define COMPLEX_NEG_OP_OVERLOAD(type) \
46 __HOST_DEVICE__ static inline type operator-(const type& op) { \
47 type ret; \
48 ret.x = -op.x; \
49 ret.y = -op.y; \
50 return ret; \
51 }
52
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; \
56 }
57
58#define COMPLEX_NE_OP_OVERLOAD(type) \
59 __HOST_DEVICE__ static inline bool operator!=(const type& lhs, const type& rhs) { \
60 return !(lhs == rhs); \
61 }
62
63#define COMPLEX_ADD_OP_OVERLOAD(type) \
64 __HOST_DEVICE__ static inline type operator+(const type& lhs, const type& rhs) { \
65 type ret; \
66 ret.x = lhs.x + rhs.x; \
67 ret.y = lhs.y + rhs.y; \
68 return ret; \
69 }
70
71#define COMPLEX_SUB_OP_OVERLOAD(type) \
72 __HOST_DEVICE__ static inline type operator-(const type& lhs, const type& rhs) { \
73 type ret; \
74 ret.x = lhs.x - rhs.x; \
75 ret.y = lhs.y - rhs.y; \
76 return ret; \
77 }
78
79#define COMPLEX_MUL_OP_OVERLOAD(type) \
80 __HOST_DEVICE__ static inline type operator*(const type& lhs, const type& rhs) { \
81 type ret; \
82 ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \
83 ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \
84 return ret; \
85 }
86
87#define COMPLEX_DIV_OP_OVERLOAD(type) \
88 __HOST_DEVICE__ static inline type operator/(const type& lhs, const type& rhs) { \
89 type ret; \
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); \
94 return ret; \
95 }
96
97#define COMPLEX_ADD_PREOP_OVERLOAD(type) \
98 __HOST_DEVICE__ static inline type& operator+=(type& lhs, const type& rhs) { \
99 lhs.x += rhs.x; \
100 lhs.y += rhs.y; \
101 return lhs; \
102 }
103
104#define COMPLEX_SUB_PREOP_OVERLOAD(type) \
105 __HOST_DEVICE__ static inline type& operator-=(type& lhs, const type& rhs) { \
106 lhs.x -= rhs.x; \
107 lhs.y -= rhs.y; \
108 return lhs; \
109 }
110
111#define COMPLEX_MUL_PREOP_OVERLOAD(type) \
112 __HOST_DEVICE__ static inline type& operator*=(type& lhs, const type& rhs) { \
113 type temp{lhs}; \
114 lhs.x = rhs.x * temp.x - rhs.y * temp.y; \
115 lhs.y = rhs.y * temp.x + rhs.x * temp.y; \
116 return lhs; \
117 }
118
119#define COMPLEX_DIV_PREOP_OVERLOAD(type) \
120 __HOST_DEVICE__ static inline type& operator/=(type& lhs, const type& rhs) { \
121 type temp; \
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); \
124 lhs = temp; \
125 return lhs; \
126 }
127
128#define COMPLEX_SCALAR_PRODUCT(type, type1) \
129 __HOST_DEVICE__ static inline type operator*(const type& lhs, type1 rhs) { \
130 type ret; \
131 ret.x = lhs.x * rhs; \
132 ret.y = lhs.y * rhs; \
133 return ret; \
134 }
135
136#endif
137
138typedef float2 hipFloatComplex;
139
140__HOST_DEVICE__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
141
142__HOST_DEVICE__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
143
144__HOST_DEVICE__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
146 z.x = a;
147 z.y = b;
148 return z;
149}
150
151__HOST_DEVICE__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
152 hipFloatComplex ret;
153 ret.x = z.x;
154 ret.y = -z.y;
155 return ret;
156}
157
158__HOST_DEVICE__ static inline float hipCsqabsf(hipFloatComplex z) {
159 return z.x * z.x + z.y * z.y;
160}
161
162__HOST_DEVICE__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
163 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
164}
165
166__HOST_DEVICE__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
167 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
168}
169
170__HOST_DEVICE__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
171 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
172}
173
174__HOST_DEVICE__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
175 float sqabs = hipCsqabsf(q);
176 hipFloatComplex ret;
177 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
178 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
179 return ret;
180}
181
182__HOST_DEVICE__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
183
184
186
187__HOST_DEVICE__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
188
189__HOST_DEVICE__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
190
191__HOST_DEVICE__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
193 z.x = a;
194 z.y = b;
195 return z;
196}
197
198__HOST_DEVICE__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
200 ret.x = z.x;
201 ret.y = -z.y;
202 return ret;
203}
204
205__HOST_DEVICE__ static inline double hipCsqabs(hipDoubleComplex z) {
206 return z.x * z.x + z.y * z.y;
207}
208
209__HOST_DEVICE__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
210 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
211}
212
213__HOST_DEVICE__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
214 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
215}
216
217__HOST_DEVICE__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
218 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
219}
220
221__HOST_DEVICE__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
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;
226 return ret;
227}
228
229__HOST_DEVICE__ static inline double hipCabs(hipDoubleComplex z) { return sqrt(hipCsqabs(z)); }
230
231
232#if __cplusplus
233
234COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
235COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
236COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
237COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
238COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
239COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
240COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
241COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
242COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
243COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
244COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
245COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
246COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
247COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
248COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
249COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
250COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
251COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
252COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
253COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
254COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
255
256COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
257COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
258COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
259COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
260COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
261COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
262COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
263COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
264COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
265COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
266COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
267COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
268COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
269COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
270COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
271COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
272COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
273COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
274COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
275COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
276COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
277
278#endif
279
280
282
283__HOST_DEVICE__ static inline hipComplex make_hipComplex(float x, float y) {
284 return make_hipFloatComplex(x, y);
285}
286
287__HOST_DEVICE__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
288 return make_hipFloatComplex((float)z.x, (float)z.y);
289}
290
291__HOST_DEVICE__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
292 return make_hipDoubleComplex((double)z.x, (double)z.y);
293}
294
295__HOST_DEVICE__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
296 float real = (p.x * q.x) + r.x;
297 float imag = (q.x * p.y) + r.y;
298
299 real = -(p.y * q.y) + real;
300 imag = (p.x * q.y) + imag;
301
302 return make_hipComplex(real, imag);
303}
304
305__HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
307 double real = (p.x * q.x) + r.x;
308 double imag = (q.x * p.y) + r.y;
309
310 real = -(p.y * q.y) + real;
311 imag = (p.x * q.y) + imag;
312
313 return make_hipDoubleComplex(real, imag);
314}
315
316#endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:2023