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_HCC_DETAIL_HIP_COMPLEX_H
24 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
25 
27 
28 // TODO: Clang has a bug which allows device functions to call std functions
29 // when std functions are introduced into default namespace by using statement.
30 // math.h may be included after this bug is fixed.
31 #if __cplusplus
32 #include <cmath>
33 #else
34 #include "math.h"
35 #endif
36 
37 #if __cplusplus
38 #define COMPLEX_NEG_OP_OVERLOAD(type) \
39  __device__ __host__ static inline type operator-(const type& op) { \
40  type ret; \
41  ret.x = -op.x; \
42  ret.y = -op.y; \
43  return ret; \
44  }
45 
46 #define COMPLEX_EQ_OP_OVERLOAD(type) \
47  __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \
48  return lhs.x == rhs.x && lhs.y == rhs.y; \
49  }
50 
51 #define COMPLEX_NE_OP_OVERLOAD(type) \
52  __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \
53  return !(lhs == rhs); \
54  }
55 
56 #define COMPLEX_ADD_OP_OVERLOAD(type) \
57  __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \
58  type ret; \
59  ret.x = lhs.x + rhs.x; \
60  ret.y = lhs.y + rhs.y; \
61  return ret; \
62  }
63 
64 #define COMPLEX_SUB_OP_OVERLOAD(type) \
65  __device__ __host__ static inline type operator-(const type& lhs, const type& rhs) { \
66  type ret; \
67  ret.x = lhs.x - rhs.x; \
68  ret.y = lhs.y - rhs.y; \
69  return ret; \
70  }
71 
72 #define COMPLEX_MUL_OP_OVERLOAD(type) \
73  __device__ __host__ static inline type operator*(const type& lhs, const type& rhs) { \
74  type ret; \
75  ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \
76  ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \
77  return ret; \
78  }
79 
80 #define COMPLEX_DIV_OP_OVERLOAD(type) \
81  __device__ __host__ static inline type operator/(const type& lhs, const type& rhs) { \
82  type ret; \
83  ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \
84  ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \
85  ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \
86  ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \
87  return ret; \
88  }
89 
90 #define COMPLEX_ADD_PREOP_OVERLOAD(type) \
91  __device__ __host__ static inline type& operator+=(type& lhs, const type& rhs) { \
92  lhs.x += rhs.x; \
93  lhs.y += rhs.y; \
94  return lhs; \
95  }
96 
97 #define COMPLEX_SUB_PREOP_OVERLOAD(type) \
98  __device__ __host__ 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_MUL_PREOP_OVERLOAD(type) \
105  __device__ __host__ static inline type& operator*=(type& lhs, const type& rhs) { \
106  lhs = lhs * rhs; \
107  return lhs; \
108  }
109 
110 #define COMPLEX_DIV_PREOP_OVERLOAD(type) \
111  __device__ __host__ static inline type& operator/=(type& lhs, const type& rhs) { \
112  lhs = lhs / rhs; \
113  return lhs; \
114  }
115 
116 #define COMPLEX_SCALAR_PRODUCT(type, type1) \
117  __device__ __host__ static inline type operator*(const type& lhs, type1 rhs) { \
118  type ret; \
119  ret.x = lhs.x * rhs; \
120  ret.y = lhs.y * rhs; \
121  return ret; \
122  }
123 #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \
124  explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \
125  __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
126 
127 #endif
128 
130 #ifdef __cplusplus
131  public:
132  typedef float value_type;
133  __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
134  explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
135  __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
136  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short)
137  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short)
138  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned int)
139  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed int)
140  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, double)
141  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long)
142  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long)
143  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long long)
144  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long long)
145 #endif
146  float x, y;
147 } __attribute__((aligned(8)));
148 
150 #ifdef __cplusplus
151  public:
152  typedef double value_type;
153  __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
154  explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
155  __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
156  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short)
157  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short)
158  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned int)
159  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed int)
160  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, float)
161  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long)
162  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long)
163  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long long)
164  MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long long)
165 #endif
166  double x, y;
167 } __attribute__((aligned(16)));
168 
169 #if __cplusplus
170 
171 COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
172 COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
173 COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
174 COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
175 COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
176 COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
177 COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
178 COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
179 COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
180 COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
181 COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
182 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
183 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
184 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
185 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
186 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
187 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
188 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
189 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
190 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
191 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
192 
193 COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
194 COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
195 COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
196 COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
197 COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
198 COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
199 COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
200 COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
201 COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
202 COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
203 COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
204 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
205 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
206 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
207 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
208 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
209 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
210 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
211 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
212 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
213 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
214 
215 #endif
216 
217 __device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
218 
219 __device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
220 
221 __device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
222  hipFloatComplex z;
223  z.x = a;
224  z.y = b;
225  return z;
226 }
227 
228 __device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
229  hipFloatComplex ret;
230  ret.x = z.x;
231  ret.y = -z.y;
232  return ret;
233 }
234 
235 __device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
236  return z.x * z.x + z.y * z.y;
237 }
238 
239 __device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
240  return make_hipFloatComplex(p.x + q.x, p.y + q.y);
241 }
242 
243 __device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
244  return make_hipFloatComplex(p.x - q.x, p.y - q.y);
245 }
246 
247 __device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
248  return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
249 }
250 
251 __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
252  float sqabs = hipCsqabsf(q);
253  hipFloatComplex ret;
254  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
255  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
256  return ret;
257 }
258 
259 __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
260 
261 __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
262 
263 __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
264 
265 __device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
267  z.x = a;
268  z.y = b;
269  return z;
270 }
271 
272 __device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
273  hipDoubleComplex ret;
274  ret.x = z.x;
275  ret.y = z.y;
276  return ret;
277 }
278 
279 __device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
280  return z.x * z.x + z.y * z.y;
281 }
282 
283 __device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
284  return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
285 }
286 
287 __device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
288  return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
289 }
290 
291 __device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
292  return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
293 }
294 
295 __device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
296  double sqabs = hipCsqabs(q);
297  hipDoubleComplex ret;
298  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
299  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
300  return ret;
301 }
302 
303 __device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
304 
306 
307 __device__ __host__ static inline hipComplex make_hipComplex(float x, float y) {
308  return make_hipFloatComplex(x, y);
309 }
310 
311 __device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
312  return make_hipFloatComplex((float)z.x, (float)z.y);
313 }
314 
315 __device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
316  return make_hipDoubleComplex((double)z.x, (double)z.y);
317 }
318 
319 __device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
320  float real = (p.x * q.x) + r.x;
321  float imag = (q.x * p.y) + r.y;
322 
323  real = -(p.y * q.y) + real;
324  imag = (p.x * q.y) + imag;
325 
326  return make_hipComplex(real, imag);
327 }
328 
329 __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
330  hipDoubleComplex r) {
331  float real = (p.x * q.x) + r.x;
332  float imag = (q.x * p.y) + r.y;
333 
334  real = -(p.y * q.y) + real;
335  imag = (p.x * q.y) + imag;
336 
337  return make_hipDoubleComplex(real, imag);
338 }
339 
340 // Complex functions returning real numbers.
341 #define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \
342 __device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \
343 __device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); }
344 
345 __DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
346 __DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal)
347 __DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag)
348 
349 // Complex functions returning complex numbers.
350 #define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \
351 __device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \
352 __device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); }
353 
354 __DEFINE_HIP_COMPLEX_FUN(conj, hipConj)
355 
356 #endif
Definition: hip_complex.h:129
#define __host__
Definition: host_defines.h:41
Definition: hip_complex.h:149
Defines the different newt vector types for HIP runtime.