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 
124 #endif
125 
126 typedef float2 hipFloatComplex;
127 
128 __device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
129 
130 __device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
131 
132 __device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
133  hipFloatComplex z;
134  z.x = a;
135  z.y = b;
136  return z;
137 }
138 
139 __device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
140  hipFloatComplex ret;
141  ret.x = z.x;
142  ret.y = -z.y;
143  return ret;
144 }
145 
146 __device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
147  return z.x * z.x + z.y * z.y;
148 }
149 
150 __device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
151  return make_hipFloatComplex(p.x + q.x, p.y + q.y);
152 }
153 
154 __device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
155  return make_hipFloatComplex(p.x - q.x, p.y - q.y);
156 }
157 
158 __device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
159  return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
160 }
161 
162 __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
163  float sqabs = hipCsqabsf(q);
164  hipFloatComplex ret;
165  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
166  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
167  return ret;
168 }
169 
170 __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
171 
172 
173 typedef double2 hipDoubleComplex;
174 
175 __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
176 
177 __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
178 
179 __device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
180  hipDoubleComplex z;
181  z.x = a;
182  z.y = b;
183  return z;
184 }
185 
186 __device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
187  hipDoubleComplex ret;
188  ret.x = z.x;
189  ret.y = z.y;
190  return ret;
191 }
192 
193 __device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
194  return z.x * z.x + z.y * z.y;
195 }
196 
197 __device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
198  return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
199 }
200 
201 __device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
202  return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
203 }
204 
205 __device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
206  return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
207 }
208 
209 __device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
210  double sqabs = hipCsqabs(q);
211  hipDoubleComplex ret;
212  ret.x = (p.x * q.x + p.y * q.y) / sqabs;
213  ret.y = (p.y * q.x - p.x * q.y) / sqabs;
214  return ret;
215 }
216 
217 __device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
218 
219 
220 #if __cplusplus
221 
222 COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
223 COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
224 COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
225 COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
226 COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
227 COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
228 COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
229 COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
230 COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
231 COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
232 COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
233 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
234 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
235 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
236 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
237 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
238 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
239 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
240 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
241 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
242 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
243 
244 COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
245 COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
246 COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
247 COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
248 COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
249 COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
250 COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
251 COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
252 COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
253 COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
254 COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
255 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
256 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
257 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
258 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
259 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
260 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
261 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
262 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
263 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
264 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
265 
266 #endif
267 
268 
269 typedef hipFloatComplex hipComplex;
270 
271 __device__ __host__ static inline hipComplex make_hipComplex(float x, float y) {
272  return make_hipFloatComplex(x, y);
273 }
274 
275 __device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
276  return make_hipFloatComplex((float)z.x, (float)z.y);
277 }
278 
279 __device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
280  return make_hipDoubleComplex((double)z.x, (double)z.y);
281 }
282 
283 __device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
284  float real = (p.x * q.x) + r.x;
285  float imag = (q.x * p.y) + r.y;
286 
287  real = -(p.y * q.y) + real;
288  imag = (p.x * q.y) + imag;
289 
290  return make_hipComplex(real, imag);
291 }
292 
293 __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
294  hipDoubleComplex r) {
295  double real = (p.x * q.x) + r.x;
296  double imag = (q.x * p.y) + r.y;
297 
298  real = -(p.y * q.y) + real;
299  imag = (p.x * q.y) + imag;
300 
301  return make_hipDoubleComplex(real, imag);
302 }
303 
304 // Complex functions returning real numbers.
305 #define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \
306 __device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \
307 __device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); }
308 
309 __DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
310 __DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal)
311 __DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag)
312 
313 // Complex functions returning complex numbers.
314 #define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \
315 __device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \
316 __device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); }
317 
318 __DEFINE_HIP_COMPLEX_FUN(conj, hipConj)
319 
320 #endif
#define __host__
Definition: host_defines.h:41
Defines the different newt vector types for HIP runtime.