MAGMA  magma-1.4.0
Matrix Algebra on GPU and Multicore Architectures
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
operators.h
Go to the documentation of this file.
1 /*
2  -- MAGMA (version 1.4.0) --
3  Univ. of Tennessee, Knoxville
4  Univ. of California, Berkeley
5  Univ. of Colorado, Denver
6  August 2013
7 
8  @author Mathieu Faverge
9  @author Mark Gates
10 */
11 
12 #ifndef MAGMA_OPERATORS_H
13 #define MAGMA_OPERATORS_H
14 
15 // __host__ and __device__ are defined in CUDA headers.
16 #include "magma.h"
17 
18 /* names to match C++ std complex functions */
19 __host__ __device__ static inline double real(const magmaDoubleComplex &x) { return MAGMA_Z_REAL(x); }
20 __host__ __device__ static inline float real(const magmaFloatComplex &x) { return MAGMA_C_REAL(x); }
21 __host__ __device__ static inline double real(const double &x) { return x; }
22 __host__ __device__ static inline float real(const float &x) { return x; }
23 
24 __host__ __device__ static inline double imag(const magmaDoubleComplex &x) { return MAGMA_Z_IMAG(x); }
25 __host__ __device__ static inline float imag(const magmaFloatComplex &x) { return MAGMA_C_IMAG(x); }
26 __host__ __device__ static inline double imag(const double &x) { return 0.; }
27 __host__ __device__ static inline float imag(const float &x) { return 0.; }
28 
29 __host__ __device__ static inline magmaDoubleComplex conj(const magmaDoubleComplex &x) { return MAGMA_Z_CNJG(x); }
30 __host__ __device__ static inline magmaFloatComplex conj(const magmaFloatComplex &x) { return MAGMA_C_CNJG(x); }
31 __host__ __device__ static inline double conj(const double &x) { return x; }
32 __host__ __device__ static inline float conj(const float &x) { return x; }
33 
34 
35 /*************************************************************
36  * magmaDoubleComplex
37  */
38 
39 // ---------- negate
40 __host__ __device__ static inline magmaDoubleComplex
41 operator - (const magmaDoubleComplex &a)
42 {
43  return MAGMA_Z_MAKE( -real(a),
44  -imag(a) );
45 }
46 
47 
48 // ---------- add
49 __host__ __device__ static inline magmaDoubleComplex
50 operator + (const magmaDoubleComplex a, const magmaDoubleComplex b)
51 {
52  return MAGMA_Z_MAKE( real(a) + real(b),
53  imag(a) + imag(b) );
54 }
55 
56 __host__ __device__ static inline magmaDoubleComplex
57 operator + (const magmaDoubleComplex a, const double s)
58 {
59  return MAGMA_Z_MAKE( real(a) + s,
60  imag(a) );
61 }
62 
63 __host__ __device__ static inline magmaDoubleComplex
64 operator + (const double s, const magmaDoubleComplex b)
65 {
66  return MAGMA_Z_MAKE( s + real(b),
67  imag(b) );
68 }
69 
70 __host__ __device__ static inline magmaDoubleComplex&
71 operator += (magmaDoubleComplex &a, const magmaDoubleComplex b)
72 {
73  a = MAGMA_Z_MAKE( real(a) + real(b),
74  imag(a) + imag(b) );
75  return a;
76 }
77 
78 __host__ __device__ static inline magmaDoubleComplex&
79 operator += (magmaDoubleComplex &a, const double s)
80 {
81  a = MAGMA_Z_MAKE( real(a) + s,
82  imag(a) );
83  return a;
84 }
85 
86 
87 // ---------- subtract
88 __host__ __device__ static inline magmaDoubleComplex
89 operator - (const magmaDoubleComplex a, const magmaDoubleComplex b)
90 {
91  return MAGMA_Z_MAKE( real(a) - real(b),
92  imag(a) - imag(b) );
93 }
94 
95 __host__ __device__ static inline magmaDoubleComplex
96 operator - (const magmaDoubleComplex a, const double s)
97 {
98  return MAGMA_Z_MAKE( real(a) - s,
99  imag(a) );
100 }
101 
102 __host__ __device__ static inline magmaDoubleComplex
103 operator - (const double s, const magmaDoubleComplex b)
104 {
105  return MAGMA_Z_MAKE( s - real(b),
106  - imag(b) );
107 }
108 
109 __host__ __device__ static inline magmaDoubleComplex&
110 operator -= (magmaDoubleComplex &a, const magmaDoubleComplex b)
111 {
112  a = MAGMA_Z_MAKE( real(a) - real(b),
113  imag(a) - imag(b) );
114  return a;
115 }
116 
117 __host__ __device__ static inline magmaDoubleComplex&
118 operator -= (magmaDoubleComplex &a, const double s)
119 {
120  a = MAGMA_Z_MAKE( real(a) - s,
121  imag(a) );
122  return a;
123 }
124 
125 
126 // ---------- multiply
127 __host__ __device__ static inline magmaDoubleComplex
128 operator * (const magmaDoubleComplex a, const magmaDoubleComplex b)
129 {
130  return MAGMA_Z_MAKE( real(a)*real(b) - imag(a)*imag(b),
131  imag(a)*real(b) + real(a)*imag(b) );
132 }
133 
134 __host__ __device__ static inline magmaDoubleComplex
135 operator * (const magmaDoubleComplex a, const double s)
136 {
137  return MAGMA_Z_MAKE( real(a)*s,
138  imag(a)*s );
139 }
140 
141 __host__ __device__ static inline magmaDoubleComplex
142 operator * (const double s, const magmaDoubleComplex a)
143 {
144  return MAGMA_Z_MAKE( real(a)*s,
145  imag(a)*s );
146 }
147 
148 __host__ __device__ static inline magmaDoubleComplex&
149 operator *= (magmaDoubleComplex &a, const magmaDoubleComplex b)
150 {
151  a = MAGMA_Z_MAKE( real(a)*real(b) - imag(a)*imag(b),
152  imag(a)*real(b) + real(a)*imag(b) );
153  return a;
154 }
155 
156 __host__ __device__ static inline magmaDoubleComplex&
157 operator *= (magmaDoubleComplex &a, const double s)
158 {
159  a = MAGMA_Z_MAKE( real(a)*s,
160  imag(a)*s );
161  return a;
162 }
163 
164 
165 // ---------- divide
166 /* From LAPACK DLADIV
167  * Performs complex division in real arithmetic, avoiding unnecessary overflow.
168  *
169  * a + i*b
170  * p + i*q = ---------
171  * c + i*d
172  */
173 __host__ __device__ static inline magmaDoubleComplex
174 operator / (const magmaDoubleComplex x, const magmaDoubleComplex y)
175 {
176  double a = real(x);
177  double b = imag(x);
178  double c = real(y);
179  double d = imag(y);
180  double e, f, p, q;
181  if ( fabs( d ) < fabs( c ) ) {
182  e = d / c;
183  f = c + d*e;
184  p = ( a + b*e ) / f;
185  q = ( b - a*e ) / f;
186  }
187  else {
188  e = c / d;
189  f = d + c*e;
190  p = ( b + a*e ) / f;
191  q = ( -a + b*e ) / f;
192  }
193  return MAGMA_Z_MAKE( p, q );
194 }
195 
196 __host__ __device__ static inline magmaDoubleComplex
197 operator / (const magmaDoubleComplex a, const double s)
198 {
199  return MAGMA_Z_MAKE( real(a)/s,
200  imag(a)/s );
201 }
202 
203 __host__ __device__ static inline magmaDoubleComplex
204 operator / (const double a, const magmaDoubleComplex y)
205 {
206  double c = real(y);
207  double d = imag(y);
208  double e, f, p, q;
209  if ( fabs( d ) < fabs( c ) ) {
210  e = d / c;
211  f = c + d*e;
212  p = a / f;
213  q = -a*e / f;
214  }
215  else {
216  e = c / d;
217  f = d + c*e;
218  p = a*e / f;
219  q = -a / f;
220  }
221  return MAGMA_Z_MAKE( p, q );
222 }
223 
224 __host__ __device__ static inline magmaDoubleComplex&
225 operator /= (magmaDoubleComplex &a, const magmaDoubleComplex b)
226 {
227  a = a/b;
228  return a;
229 }
230 
231 __host__ __device__ static inline magmaDoubleComplex&
232 operator /= (magmaDoubleComplex &a, const double s)
233 {
234  a = MAGMA_Z_MAKE( real(a)/s,
235  imag(a)/s );
236  return a;
237 }
238 
239 
240 // ---------- equality
241 __host__ __device__ static inline bool
242 operator == (const magmaDoubleComplex a, const magmaDoubleComplex b)
243 {
244  return ( real(a) == real(b) &&
245  imag(a) == imag(b) );
246 }
247 
248 __host__ __device__ static inline bool
249 operator == (const magmaDoubleComplex a, const double s)
250 {
251  return ( real(a) == s &&
252  imag(a) == 0. );
253 }
254 
255 __host__ __device__ static inline bool
256 operator == (const double s, const magmaDoubleComplex a)
257 {
258  return ( real(a) == s &&
259  imag(a) == 0. );
260 }
261 
262 
263 // ---------- not equality
264 __host__ __device__ static inline bool
265 operator != (const magmaDoubleComplex a, const magmaDoubleComplex b)
266 {
267  return ! (a == b);
268 }
269 
270 __host__ __device__ static inline bool
271 operator != (const magmaDoubleComplex a, const double s)
272 {
273  return ! (a == s);
274 }
275 
276 __host__ __device__ static inline bool
277 operator != (const double s, const magmaDoubleComplex a)
278 {
279  return ! (a == s);
280 }
281 
282 
283 /*************************************************************
284  * magmaFloatComplex
285  */
286 
287 // ---------- negate
288 __host__ __device__ static inline magmaFloatComplex
289 operator - (const magmaFloatComplex &a)
290 {
291  return MAGMA_C_MAKE( -real(a),
292  -imag(a) );
293 }
294 
295 
296 // ---------- add
297 __host__ __device__ static inline magmaFloatComplex
298 operator + (const magmaFloatComplex a, const magmaFloatComplex b)
299 {
300  return MAGMA_C_MAKE( real(a) + real(b),
301  imag(a) + imag(b) );
302 }
303 
304 __host__ __device__ static inline magmaFloatComplex
305 operator + (const magmaFloatComplex a, const float s)
306 {
307  return MAGMA_C_MAKE( real(a) + s,
308  imag(a) );
309 }
310 
311 __host__ __device__ static inline magmaFloatComplex
312 operator + (const float s, const magmaFloatComplex b)
313 {
314  return MAGMA_C_MAKE( s + real(b),
315  imag(b) );
316 }
317 
318 __host__ __device__ static inline magmaFloatComplex&
319 operator += (magmaFloatComplex &a, const magmaFloatComplex b)
320 {
321  a = MAGMA_C_MAKE( real(a) + real(b),
322  imag(a) + imag(b) );
323  return a;
324 }
325 
326 __host__ __device__ static inline magmaFloatComplex&
327 operator += (magmaFloatComplex &a, const float s)
328 {
329  a = MAGMA_C_MAKE( real(a) + s,
330  imag(a) );
331  return a;
332 }
333 
334 
335 // ---------- subtract
336 __host__ __device__ static inline magmaFloatComplex
337 operator - (const magmaFloatComplex a, const magmaFloatComplex b)
338 {
339  return MAGMA_C_MAKE( real(a) - real(b),
340  imag(a) - imag(b) );
341 }
342 
343 __host__ __device__ static inline magmaFloatComplex
344 operator - (const magmaFloatComplex a, const float s)
345 {
346  return MAGMA_C_MAKE( real(a) - s,
347  imag(a) );
348 }
349 
350 __host__ __device__ static inline magmaFloatComplex
351 operator - (const float s, const magmaFloatComplex b)
352 {
353  return MAGMA_C_MAKE( s - real(b),
354  - imag(b) );
355 }
356 
357 __host__ __device__ static inline magmaFloatComplex&
358 operator -= (magmaFloatComplex &a, const magmaFloatComplex b)
359 {
360  a = MAGMA_C_MAKE( real(a) - real(b),
361  imag(a) - imag(b) );
362  return a;
363 }
364 
365 __host__ __device__ static inline magmaFloatComplex&
366 operator -= (magmaFloatComplex &a, const float s)
367 {
368  a = MAGMA_C_MAKE( real(a) - s,
369  imag(a) );
370  return a;
371 }
372 
373 
374 // ---------- multiply
375 __host__ __device__ static inline magmaFloatComplex
376 operator * (const magmaFloatComplex a, const magmaFloatComplex b)
377 {
378  return MAGMA_C_MAKE( real(a)*real(b) - imag(a)*imag(b),
379  imag(a)*real(b) + real(a)*imag(b) );
380 }
381 
382 __host__ __device__ static inline magmaFloatComplex
383 operator * (const magmaFloatComplex a, const float s)
384 {
385  return MAGMA_C_MAKE( real(a)*s,
386  imag(a)*s );
387 }
388 
389 __host__ __device__ static inline magmaFloatComplex
390 operator * (const float s, const magmaFloatComplex a)
391 {
392  return MAGMA_C_MAKE( real(a)*s,
393  imag(a)*s );
394 }
395 
396 __host__ __device__ static inline magmaFloatComplex&
397 operator *= (magmaFloatComplex &a, const magmaFloatComplex b)
398 {
399  a = MAGMA_C_MAKE( real(a)*real(b) - imag(a)*imag(b),
400  imag(a)*real(b) + real(a)*imag(b) );
401  return a;
402 }
403 
404 __host__ __device__ static inline magmaFloatComplex&
405 operator *= (magmaFloatComplex &a, const float s)
406 {
407  a = MAGMA_C_MAKE( real(a)*s,
408  imag(a)*s );
409  return a;
410 }
411 
412 
413 // ---------- divide
414 /* From LAPACK DLADIV
415  * Performs complex division in real arithmetic, avoiding unnecessary overflow.
416  *
417  * a + i*b
418  * p + i*q = ---------
419  * c + i*d
420  */
421 __host__ __device__ static inline magmaFloatComplex
422 operator / (const magmaFloatComplex x, const magmaFloatComplex y)
423 {
424  float a = real(x);
425  float b = imag(x);
426  float c = real(y);
427  float d = imag(y);
428  float e, f, p, q;
429  if ( fabs( d ) < fabs( c ) ) {
430  e = d / c;
431  f = c + d*e;
432  p = ( a + b*e ) / f;
433  q = ( b - a*e ) / f;
434  }
435  else {
436  e = c / d;
437  f = d + c*e;
438  p = ( b + a*e ) / f;
439  q = ( -a + b*e ) / f;
440  }
441  return MAGMA_C_MAKE( p, q );
442 }
443 
444 __host__ __device__ static inline magmaFloatComplex
445 operator / (const magmaFloatComplex a, const float s)
446 {
447  return MAGMA_C_MAKE( real(a)/s,
448  imag(a)/s );
449 }
450 
451 __host__ __device__ static inline magmaFloatComplex
452 operator / (const float a, const magmaFloatComplex y)
453 {
454  float c = real(y);
455  float d = imag(y);
456  float e, f, p, q;
457  if ( fabs( d ) < fabs( c ) ) {
458  e = d / c;
459  f = c + d*e;
460  p = a / f;
461  q = -a*e / f;
462  }
463  else {
464  e = c / d;
465  f = d + c*e;
466  p = a*e / f;
467  q = -a / f;
468  }
469  return MAGMA_C_MAKE( p, q );
470 }
471 
472 __host__ __device__ static inline magmaFloatComplex&
473 operator /= (magmaFloatComplex &a, const magmaFloatComplex b)
474 {
475  a = a/b;
476  return a;
477 }
478 
479 __host__ __device__ static inline magmaFloatComplex&
480 operator /= (magmaFloatComplex &a, const float s)
481 {
482  a = MAGMA_C_MAKE( real(a)/s,
483  imag(a)/s );
484  return a;
485 }
486 
487 
488 // ---------- equality
489 __host__ __device__ static inline bool
490 operator == (const magmaFloatComplex a, const magmaFloatComplex b)
491 {
492  return ( real(a) == real(b) &&
493  imag(a) == imag(b) );
494 }
495 
496 __host__ __device__ static inline bool
497 operator == (const magmaFloatComplex a, const float s)
498 {
499  return ( real(a) == s &&
500  imag(a) == 0. );
501 }
502 
503 __host__ __device__ static inline bool
504 operator == (const float s, const magmaFloatComplex a)
505 {
506  return ( real(a) == s &&
507  imag(a) == 0. );
508 }
509 
510 
511 // ---------- not equality
512 __host__ __device__ static inline bool
513 operator != (const magmaFloatComplex a, const magmaFloatComplex b)
514 {
515  return ! (a == b);
516 }
517 
518 __host__ __device__ static inline bool
519 operator != (const magmaFloatComplex a, const float s)
520 {
521  return ! (a == s);
522 }
523 
524 __host__ __device__ static inline bool
525 operator != (const float s, const magmaFloatComplex a)
526 {
527  return ! (a == s);
528 }
529 
530 #endif /* MAGMA_OPERATORS_H */
#define MAGMA_Z_MAKE(r, i)
Definition: magma.h:123
__host__ static __device__ magmaDoubleComplex & operator+=(magmaDoubleComplex &a, const magmaDoubleComplex b)
Definition: operators.h:71
__host__ static __device__ magmaDoubleComplex & operator*=(magmaDoubleComplex &a, const magmaDoubleComplex b)
Definition: operators.h:149
#define MAGMA_C_IMAG(a)
Definition: magma.h:147
#define MAGMA_C_CNJG(v, t)
Definition: magma.h:142
__host__ static __device__ magmaDoubleComplex conj(const magmaDoubleComplex &x)
Definition: operators.h:29
#define MAGMA_Z_IMAG(a)
Definition: magma.h:125
__host__ static __device__ magmaDoubleComplex operator/(const magmaDoubleComplex x, const magmaDoubleComplex y)
Definition: operators.h:174
__host__ static __device__ magmaDoubleComplex operator+(const magmaDoubleComplex a, const magmaDoubleComplex b)
Definition: operators.h:50
__host__ static __device__ magmaDoubleComplex & operator-=(magmaDoubleComplex &a, const magmaDoubleComplex b)
Definition: operators.h:110
__host__ static __device__ magmaDoubleComplex operator*(const magmaDoubleComplex a, const magmaDoubleComplex b)
Definition: operators.h:128
__host__ static __device__ double imag(const magmaDoubleComplex &x)
Definition: operators.h:24
#define MAGMA_C_REAL(a)
Definition: magma.h:146
__host__ static __device__ magmaDoubleComplex & operator/=(magmaDoubleComplex &a, const magmaDoubleComplex b)
Definition: operators.h:225
__host__ static __device__ double real(const magmaDoubleComplex &x)
Definition: operators.h:19
__host__ static __device__ magmaDoubleComplex operator-(const magmaDoubleComplex &a)
Definition: operators.h:41
__host__ static __device__ bool operator==(const magmaDoubleComplex a, const magmaDoubleComplex b)
Definition: operators.h:242
#define MAGMA_Z_CNJG(v, t)
Definition: magma.h:120
#define MAGMA_C_MAKE(r, i)
Definition: magma.h:145
__host__ static __device__ bool operator!=(const magmaDoubleComplex a, const magmaDoubleComplex b)
Definition: operators.h:265
#define MAGMA_Z_REAL(a)
Definition: magma.h:124