HIP: Heterogenous-computing Interface for Portability
hip_complex.h
1 /*
2 Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
24 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
25 
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  lhs = lhs * rhs; \
112  return lhs; \
113  }
114 
115 #define COMPLEX_DIV_PREOP_OVERLOAD(type) \
116  __HOST_DEVICE__ static inline type& operator/=(type& lhs, const type& rhs) { \
117  lhs = lhs / rhs; \
118  return lhs; \
119  }
120 
121 #define COMPLEX_SCALAR_PRODUCT(type, type1) \
122  __HOST_DEVICE__ static inline type operator*(const type& lhs, type1 rhs) { \
123  type ret; \
124  ret.x = lhs.x * rhs; \
125  ret.y = lhs.y * rhs; \
126  return ret; \
127  }
128 
129 #endif
130 
131 typedef float2 hipFloatComplex;
132 
133 __HOST_DEVICE__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
134 
135 __HOST_DEVICE__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
136 
137 __HOST_DEVICE__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
138  hipFloatComplex z;
139  z.x = a;
140  z.y = b;
141  return z;
142 }
143 
144 __HOST_DEVICE__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
145  hipFloatComplex ret;
146  ret.x = z.x;
147  ret.y = -z.y;
148  return ret;
149 }
150 
151 __HOST_DEVICE__ static inline float hipCsqabsf(hipFloatComplex z) {
152  return z.x * z.x + z.y * z.y;
153 }
154 
155 __HOST_DEVICE__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
156  return make_hipFloatComplex(p.x + q.x, p.y + q.y);
157 }
158 
159 __HOST_DEVICE__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
160  return make_hipFloatComplex(p.x - q.x, p.y - q.y);
161 }
162 
163 __HOST_DEVICE__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
164  return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
165 }
166 
167 __HOST_DEVICE__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
168  float sqabs = hipCsqabsf(q);
169  hipFloatComplex ret;
170  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
171  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
172  return ret;
173 }
174 
175 __HOST_DEVICE__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
176 
177 
178 typedef double2 hipDoubleComplex;
179 
180 __HOST_DEVICE__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
181 
182 __HOST_DEVICE__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
183 
184 __HOST_DEVICE__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
186  z.x = a;
187  z.y = b;
188  return z;
189 }
190 
191 __HOST_DEVICE__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
192  hipDoubleComplex ret;
193  ret.x = z.x;
194  ret.y = -z.y;
195  return ret;
196 }
197 
198 __HOST_DEVICE__ static inline double hipCsqabs(hipDoubleComplex z) {
199  return z.x * z.x + z.y * z.y;
200 }
201 
202 __HOST_DEVICE__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
203  return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
204 }
205 
206 __HOST_DEVICE__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
207  return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
208 }
209 
210 __HOST_DEVICE__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
211  return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
212 }
213 
214 __HOST_DEVICE__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
215  double sqabs = hipCsqabs(q);
216  hipDoubleComplex ret;
217  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
218  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
219  return ret;
220 }
221 
222 __HOST_DEVICE__ static inline double hipCabs(hipDoubleComplex z) { return sqrt(hipCsqabs(z)); }
223 
224 
225 #if __cplusplus
226 
227 COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
228 COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
229 COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
230 COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
231 COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
232 COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
233 COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
234 COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
235 COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
236 COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
237 COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
238 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
239 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
240 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
241 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
242 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
243 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
244 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
245 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
246 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
247 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
248 
249 COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
250 COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
251 COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
252 COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
253 COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
254 COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
255 COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
256 COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
257 COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
258 COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
259 COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
260 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
261 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
262 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
263 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
264 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
265 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
266 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
267 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
268 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
269 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
270 
271 #endif
272 
273 
275 
276 __HOST_DEVICE__ static inline hipComplex make_hipComplex(float x, float y) {
277  return make_hipFloatComplex(x, y);
278 }
279 
280 __HOST_DEVICE__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
281  return make_hipFloatComplex((float)z.x, (float)z.y);
282 }
283 
284 __HOST_DEVICE__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
285  return make_hipDoubleComplex((double)z.x, (double)z.y);
286 }
287 
288 __HOST_DEVICE__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
289  float real = (p.x * q.x) + r.x;
290  float imag = (q.x * p.y) + r.y;
291 
292  real = -(p.y * q.y) + real;
293  imag = (p.x * q.y) + imag;
294 
295  return make_hipComplex(real, imag);
296 }
297 
298 __HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
299  hipDoubleComplex r) {
300  double real = (p.x * q.x) + r.x;
301  double imag = (q.x * p.y) + r.y;
302 
303  real = -(p.y * q.y) + real;
304  imag = (p.x * q.y) + imag;
305 
306  return make_hipDoubleComplex(real, imag);
307 }
308 
309 #endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
double2
Definition: hip_vector_types.h:1589
hip_vector_types.h
Defines the different newt vector types for HIP runtime.
float2
Definition: hip_vector_types.h:1582