QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
complex_quda.h
Go to the documentation of this file.
1 /*
2  * Copyright 2008-2009 NVIDIA Corporation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
21 #pragma once
22 
23 //#include <math.h>
24 
25 #include <cmath>
26 #include <complex>
27 #include <sstream>
28 //#include <cuComplex.h>
29 
30 namespace quda {
31  namespace gauge {
32  template<typename Float, typename storeFloat> struct fieldorder_wrapper;
33  }
34 
35  namespace colorspinor {
36  template<typename Float, typename storeFloat> struct fieldorder_wrapper;
37  }
38 }
39 
40 // We need this to make sure code inside quda:: that calls sqrt() using real numbers
41 // doesn't try to call the complex sqrt, but the standard sqrt
42 namespace quda
43 {
44  template <typename ValueType>
45  __host__ __device__
46  inline ValueType cos(ValueType x){
47  return std::cos(x);
48  }
49  template <typename ValueType>
50  __host__ __device__
51  inline ValueType sin(ValueType x){
52  return std::sin(x);
53  }
54  template <typename ValueType>
55  __host__ __device__
56  inline ValueType tan(ValueType x){
57  return std::tan(x);
58  }
59  template <typename ValueType>
60  __host__ __device__
61  inline ValueType acos(ValueType x){
62  return std::acos(x);
63  }
64  template <typename ValueType>
65  __host__ __device__
66  inline ValueType asin(ValueType x){
67  return std::asin(x);
68  }
69  template <typename ValueType>
70  __host__ __device__
71  inline ValueType atan(ValueType x){
72  return std::atan(x);
73  }
74  template <typename ValueType>
75  __host__ __device__
76  inline ValueType atan2(ValueType x,ValueType y){
77  return std::atan2(x,y);
78  }
79  template <typename ValueType>
80  __host__ __device__
81  inline ValueType cosh(ValueType x){
82  return std::cosh(x);
83  }
84  template <typename ValueType>
85  __host__ __device__
86  inline ValueType sinh(ValueType x){
87  return std::sinh(x);
88  }
89  template <typename ValueType>
90  __host__ __device__
91  inline ValueType tanh(ValueType x){
92  return std::tanh(x);
93  }
94  template <typename ValueType>
95  __host__ __device__
96  inline ValueType exp(ValueType x){
97  return std::exp(x);
98  }
99  template <typename ValueType>
100  __host__ __device__
101  inline ValueType log(ValueType x){
102  return std::log(x);
103  }
104  template <typename ValueType>
105  __host__ __device__
106  inline ValueType log10(ValueType x){
107  return std::log10(x);
108  }
109  template <typename ValueType, typename ExponentType>
110  __host__ __device__
111  inline ValueType pow(ValueType x, ExponentType e){
112 #if (CUDA_VERSION < 7050)
113  return std::pow(x,static_cast<ValueType>(e));
114 #else
115  return std::pow(x,e);
116 #endif
117  }
118  template <typename ValueType>
119  __host__ __device__
120  inline ValueType sqrt(ValueType x){
121  return std::sqrt(x);
122  }
123  template <typename ValueType>
124  __host__ __device__
125  inline ValueType abs(ValueType x){
126  return std::abs(x);
127  }
128  template <typename ValueType>
129  __host__ __device__
130  inline ValueType conj(ValueType x){
131  return x;
132  }
133 
134  template <typename ValueType> struct complex;
135  //template <> struct complex<float>;
136  //template <> struct complex<double>;
137 
138 
140  template<typename ValueType>
141  __host__ __device__
142  ValueType abs(const complex<ValueType>& z);
144  template<typename ValueType>
145  __host__ __device__
146  ValueType arg(const complex<ValueType>& z);
148  template<typename ValueType>
149  __host__ __device__
150  ValueType norm(const complex<ValueType>& z);
151 
153  template<typename ValueType>
154  __host__ __device__
155  complex<ValueType> conj(const complex<ValueType>& z);
157 
158  template<typename ValueType>
159  __host__ __device__
160  complex<ValueType> polar(const ValueType& m, const ValueType& theta = 0);
161 
162  // Arithmetic operators:
163  // Multiplication
164  template <typename ValueType>
165  __host__ __device__
166  inline complex<ValueType> operator*(const complex<ValueType>& lhs, const complex<ValueType>& rhs);
167  template <typename ValueType>
168  __host__ __device__
169  inline complex<ValueType> operator*(const complex<ValueType>& lhs, const ValueType & rhs);
170  template <typename ValueType>
171  __host__ __device__
172  inline complex<ValueType> operator*(const ValueType& lhs, const complex<ValueType>& rhs);
173  // Division
174  template <typename ValueType>
175  __host__ __device__
176  inline complex<ValueType> operator/(const complex<ValueType>& lhs, const complex<ValueType>& rhs);
177  template <>
178  __host__ __device__
179  inline complex<float> operator/(const complex<float>& lhs, const complex<float>& rhs);
180  template <>
181  __host__ __device__
182  inline complex<double> operator/(const complex<double>& lhs, const complex<double>& rhs);
183 
184  // Addition
185  template <typename ValueType>
186  __host__ __device__
187  inline complex<ValueType> operator+(const complex<ValueType>& lhs, const complex<ValueType>& rhs);
188  template <typename ValueType>
189  __host__ __device__
190  inline complex<ValueType> operator+(const complex<ValueType>& lhs, const ValueType & rhs);
191  template <typename ValueType>
192  __host__ __device__
193  inline complex<ValueType> operator+(const ValueType& lhs, const complex<ValueType>& rhs);
194  // Subtraction
195  template <typename ValueType>
196  __host__ __device__
197  inline complex<ValueType> operator-(const complex<ValueType>& lhs, const complex<ValueType>& rhs);
198  template <typename ValueType>
199  __host__ __device__
200  inline complex<ValueType> operator-(const complex<ValueType>& lhs, const ValueType & rhs);
201  template <typename ValueType>
202  __host__ __device__
203  inline complex<ValueType> operator-(const ValueType& lhs, const complex<ValueType>& rhs);
204 
205  // Unary plus and minus
206  template <typename ValueType>
207  __host__ __device__
208  inline complex<ValueType> operator+(const complex<ValueType>& rhs);
209  template <typename ValueType>
210  __host__ __device__
211  inline complex<ValueType> operator-(const complex<ValueType>& rhs);
212 
213  // Transcendentals:
214  // Returns the complex cosine of z.
215  template <typename ValueType>
216  __host__ __device__
217  complex<ValueType> cos(const complex<ValueType>& z);
218  // Returns the complex hyperbolic cosine of z.
219  template <typename ValueType>
220  __host__ __device__
221  complex<ValueType> cosh(const complex<ValueType>& z);
222  // Returns the complex base e exponential of z.
223  template <typename ValueType>
224  __host__ __device__
225  complex<ValueType> exp(const complex<ValueType>& z);
226  // Returns the complex natural logarithm of z.
227  template <typename ValueType>
228  __host__ __device__
229  complex<ValueType> log(const complex<ValueType>& z);
230  // Returns the complex base 10 logarithm of z.
231  template <typename ValueType>
232  __host__ __device__
233  complex<ValueType> log10(const complex<ValueType>& z);
234  // Returns z to the n'th power.
235  template <typename ValueType>
236  __host__ __device__
237  complex<ValueType> pow(const complex<ValueType>& z, const int& n);
238  // Returns z to the x'th power.
239  template <typename ValueType>
240  __host__ __device__
241  complex<ValueType> pow(const complex<ValueType>&z, const ValueType&x);
242  // Returns z to the z2'th power.
243  template <typename ValueType>
244  __host__ __device__
245  complex<ValueType> pow(const complex<ValueType>&z, const complex<ValueType>&z2);
246  // Returns x to the z'th power.
247  template <typename ValueType>
248  __host__ __device__
249  complex<ValueType> pow(const ValueType& x, const complex<ValueType>& z);
250  // Returns the complex sine of z.
251  template <typename ValueType>
252  __host__ __device__
253  complex<ValueType> sin(const complex<ValueType>&z);
254  // Returns the complex hyperbolic sine of z.
255  template <typename ValueType>
256  __host__ __device__
257  complex<ValueType> sinh(const complex<ValueType>&z);
258  // Returns the complex square root of z.
259  template <typename ValueType>
260  __host__ __device__
261  complex<ValueType> sqrt(const complex<ValueType>&z);
262  // Returns the complex tangent of z.
263  template <typename ValueType>
264  __host__ __device__
265  complex<ValueType> tan(const complex<ValueType>&z);
266  // Returns the complex hyperbolic tangent of z.
267  template <typename ValueType>
268  __host__ __device__
269  complex<ValueType> tanh(const complex<ValueType>&z);
270 
271 
272  // Inverse Trigonometric:
273  // Returns the complex arc cosine of z.
274  template <typename ValueType>
275  __host__ __device__
276  complex<ValueType> acos(const complex<ValueType>& z);
277  // Returns the complex arc sine of z.
278  template <typename ValueType>
279  __host__ __device__
280  complex<ValueType> asin(const complex<ValueType>& z);
281  // Returns the complex arc tangent of z.
282  template <typename ValueType>
283  __host__ __device__
284  complex<ValueType> atan(const complex<ValueType>& z);
285  // Returns the complex hyperbolic arc cosine of z.
286  template <typename ValueType>
287  __host__ __device__
288  complex<ValueType> acosh(const complex<ValueType>& z);
289  // Returns the complex hyperbolic arc sine of z.
290  template <typename ValueType>
291  __host__ __device__
292  complex<ValueType> asinh(const complex<ValueType>& z);
293  // Returns the complex hyperbolic arc tangent of z.
294  template <typename ValueType>
295  __host__ __device__
296  complex<ValueType> atanh(const complex<ValueType>& z);
297 
298 
299 
300  // Stream operators:
301  template<typename ValueType,class charT, class traits>
302  std::basic_ostream<charT, traits>& operator<<(std::basic_ostream<charT, traits>& os, const complex<ValueType>& z);
303  template<typename ValueType, typename charT, class traits>
304  std::basic_istream<charT, traits>&
305  operator>>(std::basic_istream<charT, traits>& is, complex<ValueType>& z);
306 
307 
308  // Stream operators
309  template<typename ValueType,class charT, class traits>
310  std::basic_ostream<charT, traits>& operator<<(std::basic_ostream<charT, traits>& os, const complex<ValueType>& z)
311  {
312  os << '(' << z.real() << ',' << z.imag() << ')';
313  return os;
314  };
315 
316  template<typename ValueType, typename charT, class traits>
317  std::basic_istream<charT, traits>&
318  operator>>(std::basic_istream<charT, traits>& is, complex<ValueType>& z)
319  {
320  ValueType re, im;
321 
322  charT ch;
323  is >> ch;
324 
325  if(ch == '(')
326  {
327  is >> re >> ch;
328  if (ch == ',')
329  {
330  is >> im >> ch;
331  if (ch == ')')
332  {
333  z = complex<ValueType>(re, im);
334  }
335  else
336  {
337  is.setstate(std::ios_base::failbit);
338  }
339  }
340  else if (ch == ')')
341  {
342  z = re;
343  }
344  else
345  {
346  is.setstate(std::ios_base::failbit);
347  }
348  }
349  else
350  {
351  is.putback(ch);
352  is >> re;
353  z = re;
354  }
355  return is;
356  }
357 
358 template <typename T>
359  struct norm_type {
360  typedef T type;
361  };
362  template <typename T>
363  struct norm_type< complex<T> > {
364  typedef T type;
365  };
366 
367 template <typename ValueType>
368 struct complex
369 {
370 public:
371  typedef ValueType value_type;
372 
373  // Constructors
374  __host__ __device__
375  inline complex<ValueType>(const ValueType & re = ValueType(), const ValueType& im = ValueType())
376  {
377  real(re);
378  imag(im);
379  }
380 
381  template <class X>
382  __host__ __device__
383  inline complex<ValueType>(const complex<X> & z)
384  {
385  real(z.real());
386  imag(z.imag());
387  }
388 
389  template <class X>
390  __host__ __device__
391  inline complex<ValueType>(const std::complex<X> & z)
392  {
393  real(z.real());
394  imag(z.imag());
395  }
396 
397  template <typename T>
398  __host__ __device__
399  inline complex<ValueType>& operator=(const complex<T> z)
400  {
401  real(z.real());
402  imag(z.imag());
403  return *this;
404  }
405 
406  __host__ __device__
407  inline complex<ValueType>& operator+=(const complex<ValueType> z)
408  {
409  real(real()+z.real());
410  imag(imag()+z.imag());
411  return *this;
412  }
413 
414  __host__ __device__
415  inline complex<ValueType>& operator-=(const complex<ValueType> z)
416  {
417  real(real()-z.real());
418  imag(imag()-z.imag());
419  return *this;
420  }
421 
422  __host__ __device__
423  inline complex<ValueType>& operator*=(const complex<ValueType> z)
424  {
425  *this = *this * z;
426  return *this;
427  }
428 
429  __host__ __device__
430  inline complex<ValueType>& operator/=(const complex<ValueType> z)
431  {
432  *this = *this / z;
433  return *this;
434  }
435 
436  __host__ __device__
437  inline complex<ValueType>& operator*=(const ValueType z)
438  {
439  this->x *= z;
440  this->y *= z;
441  return *this;
442  }
443 
444  __host__ __device__ inline ValueType real() const volatile;
445  __host__ __device__ inline ValueType imag() const volatile;
446  __host__ __device__ inline ValueType real() const;
447  __host__ __device__ inline ValueType imag() const;
448  __host__ __device__ inline void real(ValueType) volatile;
449  __host__ __device__ inline void imag(ValueType) volatile;
450  __host__ __device__ inline void real(ValueType);
451  __host__ __device__ inline void imag(ValueType);
452 };
453 
454 // TODO make cuFloatComplex and cuDoubleComplex protected
455 // TODO see if returning references is a perf hazard
456 
457 template<>
458  struct complex <float> : public float2 //cuFloatComplex
459 {
460 public:
461  typedef float value_type;
462  __host__ __device__
463  inline complex<float>(){};
464  __host__ __device__
465  inline complex<float>(const float & re, const float& im = float())
466  {
467  real(re);
468  imag(im);
469  }
470 
471  // For some reason having the following constructor
472  // explicitly makes things faster with at least g++
473  __host__ __device__
475  : float2(z){}
476 
477  __host__ __device__
478  complex<float>(float2 z)
479  : float2(z){}
480 
481  template <class X>
482  inline complex<float>(const std::complex<X> & z)
483  {
484  real(z.real());
485  imag(z.imag());
486  }
487 
488  // Member operators
489  template <typename T>
490  __host__ __device__
491  inline volatile complex<float>& operator=(const complex<T> z) volatile
492  {
493  real(z.real());
494  imag(z.imag());
495  return *this;
496  }
497 
498  template <typename T>
499  __host__ __device__
500  inline complex<float>& operator=(const complex<T> z)
501  {
502  real(z.real());
503  imag(z.imag());
504  return *this;
505  }
506 
507  __host__ __device__
509  {
510  real(real()+z.real());
511  imag(imag()+z.imag());
512  return *this;
513  }
514 
515  __host__ __device__
517  {
518  real(real()-z.real());
519  imag(imag()-z.imag());
520  return *this;
521  }
522 
523  __host__ __device__
525  {
526  *this = *this * z;
527  return *this;
528  }
529 
530  __host__ __device__
532  {
533  *this = *this / z;
534  return *this;
535  }
536 
537  __host__ __device__
538  inline complex<float>& operator*=(const float z)
539  {
540  this->x *= z;
541  this->y *= z;
542  return *this;
543  }
544 
545  // Let the compiler synthesize the copy and assignment operators.
546  __host__ __device__ inline complex<float>(const volatile complex<float> & z)
547  {
548  real(z.real());
549  imag(z.imag());
550  }
551 
552  __host__ __device__ inline float real() const volatile{ return x; }
553  __host__ __device__ inline float imag() const volatile{ return y; }
554  __host__ __device__ inline float real() const{ return x; }
555  __host__ __device__ inline float imag() const{ return y; }
556  __host__ __device__ inline void real(float re)volatile{ x = re; }
557  __host__ __device__ inline void imag(float im)volatile{ y = im; }
558  __host__ __device__ inline void real(float re){ x = re; }
559  __host__ __device__ inline void imag(float im){ y = im; }
560 
561  // cast operators
562  inline operator std::complex<float>() const { return std::complex<float>(real(),imag()); }
563  template <typename T>
564  inline __host__ __device__ operator complex<T>() const { return complex<T>(static_cast<T>(real()),static_cast<T>(imag())); }
565 
566  template<typename otherFloat, typename storeFloat>
567  __host__ __device__ inline void operator=(const gauge::fieldorder_wrapper<otherFloat,storeFloat> &a);
568 
569  template<typename otherFloat, typename storeFloat>
570  __host__ __device__ inline complex<float>(const gauge::fieldorder_wrapper<otherFloat,storeFloat> &a);
571 
572  template<typename otherFloat, typename storeFloat>
573  __host__ __device__ inline void operator=(const colorspinor::fieldorder_wrapper<otherFloat,storeFloat> &a);
574 
575  template<typename otherFloat, typename storeFloat>
576  __host__ __device__ inline complex<float>(const colorspinor::fieldorder_wrapper<otherFloat,storeFloat> &a);
577 };
578 
579 template<>
580  struct complex <double> : public double2 //cuDoubleComplex
581 {
582 public:
583  typedef double value_type;
584  __host__ __device__
585  inline complex<double>(){};
586  __host__ __device__
587  inline complex<double>(const double & re, const double& im = double())
588  {
589  real(re);
590  imag(im);
591  }
592 
593  // For some reason having the following constructor
594  // explicitly makes things faster with at least g++
595  __host__ __device__
597  : double2(z) {}
598 
599  __host__ __device__
600  inline complex<double>(double2 z)
601  : double2(z) {}
602 
603  template <class X>
604  inline complex<double>(const std::complex<X> & z)
605  {
606  real(z.real());
607  imag(z.imag());
608  }
609 
610  // Member operators
611  template <typename T>
612  __host__ __device__
613  inline volatile complex<double>& operator=(const complex<T> z) volatile
614  {
615  real(z.real());
616  imag(z.imag());
617  return *this;
618  }
619 
620  template <typename T>
621  __host__ __device__
622  inline complex<double>& operator=(const complex<T> z)
623  {
624  real(z.real());
625  imag(z.imag());
626  return *this;
627  }
628 
629  __host__ __device__
631  {
632  real(real()+z.real());
633  imag(imag()+z.imag());
634  return *this;
635  }
636 
637  __host__ __device__
638  inline complex<double>& operator+=(const complex<float> z)
639  {
640  real(real()+z.real());
641  imag(imag()+z.imag());
642  return *this;
643  }
644 
645  __host__ __device__
647  {
648  real(real()-z.real());
649  imag(imag()-z.imag());
650  return *this;
651  }
652 
653  __host__ __device__
655  {
656  *this = *this * z;
657  return *this;
658  }
659 
660  __host__ __device__
662  {
663  *this = *this / z;
664  return *this;
665  }
666 
667  __host__ __device__
668  inline complex<double>& operator*=(const double z)
669  {
670  this->x *= z;
671  this->y *= z;
672  return *this;
673  }
674 
675  __host__ __device__ inline complex<double>(const volatile complex<double> & z)
676  {
677  real(z.real());
678  imag(z.imag());
679  }
680 
681  // Let the compiler synthesize the copy and assignment operators.
682  __host__ __device__ inline double real() const volatile { return x; }
683  __host__ __device__ inline double imag() const volatile { return y; }
684  __host__ __device__ inline double real() const { return x; }
685  __host__ __device__ inline double imag() const { return y; }
686  __host__ __device__ inline void real(double re)volatile{ x = re; }
687  __host__ __device__ inline void imag(double im)volatile{ y = im; }
688  __host__ __device__ inline void real(double re){ x = re; }
689  __host__ __device__ inline void imag(double im){ y = im; }
690 
691  // cast operators
692  inline operator std::complex<double>() const { return std::complex<double>(real(),imag()); }
693  template <typename T>
694  inline __host__ __device__ operator complex<T>() const { return complex<T>(static_cast<T>(real()),static_cast<T>(imag())); }
695 
696  template<typename otherFloat, typename storeFloat>
697  __host__ __device__ inline void operator=(const gauge::fieldorder_wrapper<otherFloat,storeFloat> &a);
698 
699  template<typename otherFloat, typename storeFloat>
700  __host__ __device__ inline complex<double>(const gauge::fieldorder_wrapper<otherFloat,storeFloat> &a);
701 
702  template<typename otherFloat, typename storeFloat>
703  __host__ __device__ inline void operator=(const colorspinor::fieldorder_wrapper<otherFloat,storeFloat> &a);
704 
705  template<typename otherFloat, typename storeFloat>
706  __host__ __device__ inline complex<double>(const colorspinor::fieldorder_wrapper<otherFloat,storeFloat> &a);
707 };
708 
709 template<>
710 struct complex <char> : public char2
711 {
712 public:
713  typedef char value_type;
714 
715  __host__ __device__ inline complex<char>(){};
716 
717  __host__ __device__ inline complex<char>(const char & re, const char& im = float())
718  {
719  real(re);
720  imag(im);
721  }
722 
723  __host__ __device__ inline complex<char>(const complex<char> & z) : char2(z){}
724 
725  __host__ __device__ inline complex<char>& operator+=(const complex<char> z)
726  {
727  real(real()+z.real());
728  imag(imag()+z.imag());
729  return *this;
730  }
731 
732  __host__ __device__ inline complex<char>& operator-=(const complex<char> z)
733  {
734  real(real()-z.real());
735  imag(imag()-z.imag());
736  return *this;
737  }
738 
739  __host__ __device__ inline char real() const volatile{ return x; }
740  __host__ __device__ inline char imag() const volatile{ return y; }
741  __host__ __device__ inline char real() const{ return x; }
742  __host__ __device__ inline char imag() const{ return y; }
743  __host__ __device__ inline void real(char re)volatile{ x = re; }
744  __host__ __device__ inline void imag(char im)volatile{ y = im; }
745  __host__ __device__ inline void real(char re){ x = re; }
746  __host__ __device__ inline void imag(char im){ y = im; }
747 
748  // cast operators
749  inline operator std::complex<char>() const { return std::complex<char>(real(),imag()); }
750  template <typename T>
751  inline __host__ __device__ operator complex<T>() const { return complex<T>(static_cast<T>(real()),static_cast<T>(imag())); }
752 
753 };
754 
755 
756 template<>
757 struct complex <short> : public short2
758 {
759 public:
760  typedef short value_type;
761 
762  __host__ __device__ inline complex<short>(){};
763 
764  __host__ __device__ inline complex<short>(const short & re, const short& im = float())
765  {
766  real(re);
767  imag(im);
768  }
769 
770  __host__ __device__ inline complex<short>(const complex<short> & z) : short2(z){}
771 
772  __host__ __device__ inline complex<short>& operator+=(const complex<short> z)
773  {
774  real(real()+z.real());
775  imag(imag()+z.imag());
776  return *this;
777  }
778 
779  __host__ __device__ inline complex<short>& operator-=(const complex<short> z)
780  {
781  real(real()-z.real());
782  imag(imag()-z.imag());
783  return *this;
784  }
785 
786  __host__ __device__ inline short real() const volatile{ return x; }
787  __host__ __device__ inline short imag() const volatile{ return y; }
788  __host__ __device__ inline short real() const{ return x; }
789  __host__ __device__ inline short imag() const{ return y; }
790  __host__ __device__ inline void real(short re)volatile{ x = re; }
791  __host__ __device__ inline void imag(short im)volatile{ y = im; }
792  __host__ __device__ inline void real(short re){ x = re; }
793  __host__ __device__ inline void imag(short im){ y = im; }
794 
795  // cast operators
796  inline operator std::complex<short>() const { return std::complex<short>(real(),imag()); }
797  template <typename T>
798  inline __host__ __device__ operator complex<T>() const { return complex<T>(static_cast<T>(real()),static_cast<T>(imag())); }
799 
800 };
801 
802 template<>
803 struct complex <int> : public int2
804 {
805 public:
806  typedef int value_type;
807 
808  __host__ __device__ inline complex<int>(){};
809 
810  __host__ __device__ inline complex<int>(const int& re, const int& im = float())
811  {
812  real(re);
813  imag(im);
814  }
815 
816  __host__ __device__ inline complex<int>(const complex<int> & z) : int2(z){}
817 
818  __host__ __device__ inline complex<int>& operator+=(const complex<int> z)
819  {
820  real(real()+z.real());
821  imag(imag()+z.imag());
822  return *this;
823  }
824 
825  __host__ __device__ inline complex<int>& operator-=(const complex<int> z)
826  {
827  real(real()-z.real());
828  imag(imag()-z.imag());
829  return *this;
830  }
831 
832  __host__ __device__ inline int real() const volatile{ return x; }
833  __host__ __device__ inline int imag() const volatile{ return y; }
834  __host__ __device__ inline int real() const{ return x; }
835  __host__ __device__ inline int imag() const{ return y; }
836  __host__ __device__ inline void real(int re)volatile{ x = re; }
837  __host__ __device__ inline void imag(int im)volatile{ y = im; }
838  __host__ __device__ inline void real(int re){ x = re; }
839  __host__ __device__ inline void imag(int im){ y = im; }
840 
841  // cast operators
842  inline operator std::complex<int>() const { return std::complex<int>(real(),imag()); }
843  template <typename T>
844  inline __host__ __device__ operator complex<T>() const { return complex<T>(static_cast<T>(real()),static_cast<T>(imag())); }
845 
846 };
847 
848  // Binary arithmetic operations
849  // At the moment I'm implementing the basic functions, and the
850  // corresponding cuComplex calls are commented.
851 
852  template<typename ValueType>
853  __host__ __device__
854  inline complex<ValueType> operator+(const complex<ValueType>& lhs,
855 const complex<ValueType>& rhs){
856  return complex<ValueType>(lhs.real()+rhs.real(),lhs.imag()+rhs.imag());
857  // return cuCaddf(lhs,rhs);
858  }
859 
860  template<typename ValueType>
861  __host__ __device__
862  inline complex<ValueType> operator+(const volatile complex<ValueType>& lhs,
863 const volatile complex<ValueType>& rhs){
864  return complex<ValueType>(lhs.real()+rhs.real(),lhs.imag()+rhs.imag());
865  // return cuCaddf(lhs,rhs);
866  }
867 
868  template <typename ValueType>
869  __host__ __device__
870  inline complex<ValueType> operator+(const complex<ValueType>& lhs, const ValueType & rhs){
871  return complex<ValueType>(lhs.real()+rhs,lhs.imag());
872  // return cuCaddf(lhs,complex<ValueType>(rhs));
873  }
874  template <typename ValueType>
875  __host__ __device__
876  inline complex<ValueType> operator+(const ValueType& lhs, const complex<ValueType>& rhs){
877  return complex<ValueType>(rhs.real()+lhs,rhs.imag());
878  // return cuCaddf(complex<float>(lhs),rhs);
879  }
880 
881  template <typename ValueType>
882  __host__ __device__
883  inline complex<ValueType> operator-(const complex<ValueType>& lhs, const complex<ValueType>& rhs){
884  return complex<ValueType>(lhs.real()-rhs.real(),lhs.imag()-rhs.imag());
885  // return cuCsubf(lhs,rhs);
886  }
887  template <typename ValueType>
888  __host__ __device__
889  inline complex<ValueType> operator-(const complex<ValueType>& lhs, const ValueType & rhs){
890  return complex<ValueType>(lhs.real()-rhs,lhs.imag());
891  // return cuCsubf(lhs,complex<float>(rhs));
892  }
893  template <typename ValueType>
894  __host__ __device__
895  inline complex<ValueType> operator-(const ValueType& lhs, const complex<ValueType>& rhs){
896  return complex<ValueType>(lhs-rhs.real(),-rhs.imag());
897  // return cuCsubf(complex<float>(lhs),rhs);
898  }
899 
900  template <typename ValueType>
901  __host__ __device__
902  inline complex<ValueType> operator*(const complex<ValueType>& lhs,
903 const complex<ValueType>& rhs){
904  return complex<ValueType>(lhs.real()*rhs.real()-lhs.imag()*rhs.imag(),
905 lhs.real()*rhs.imag()+lhs.imag()*rhs.real());
906  // return cuCmulf(lhs,rhs);
907  }
908 
909  template <typename ValueType>
910  __host__ __device__
911  inline complex<ValueType> operator*(const complex<ValueType>& lhs, const ValueType & rhs){
912  return complex<ValueType>(lhs.real()*rhs,lhs.imag()*rhs);
913  // return cuCmulf(lhs,complex<float>(rhs));
914  }
915 
916  template <typename ValueType>
917  __host__ __device__
918  inline complex<ValueType> operator*(const ValueType& lhs, const complex<ValueType>& rhs){
919  return complex<ValueType>(rhs.real()*lhs,rhs.imag()*lhs);
920  // return cuCmulf(complex<float>(lhs),rhs);
921  }
922 
923 
924  template <typename ValueType>
925  __host__ __device__
926  inline complex<ValueType> operator/(const complex<ValueType>& lhs, const complex<ValueType>& rhs){
927  const ValueType cross_norm = lhs.real() * rhs.real() + lhs.imag() * rhs.imag();
928  const ValueType rhs_norm = norm(rhs);
929  return complex<ValueType>(cross_norm/rhs_norm,
930 (lhs.imag() * rhs.real() - lhs.real() * rhs.imag()) / rhs_norm);
931  }
932 
933  template <>
934  __host__ __device__
935  inline complex<float> operator/(const complex<float>& lhs, const complex<float>& rhs){
936 
937  complex<float> quot;
938  float s = fabsf(rhs.real()) + fabsf(rhs.imag());
939  float oos = 1.0f / s;
940  float ars = lhs.real() * oos;
941  float ais = lhs.imag() * oos;
942  float brs = rhs.real() * oos;
943  float bis = rhs.imag() * oos;
944  s = (brs * brs) + (bis * bis);
945  oos = 1.0f / s;
946  return complex<float>(((ars * brs) + (ais * bis)) * oos,
947  ((ais * brs) - (ars * bis)) * oos);
948  }
949 
950  template <>
951  __host__ __device__
952  inline complex<double> operator/(const complex<double>& lhs, const complex<double>& rhs){
953 
954  complex<double> quot;
955  double s = fabs(rhs.real()) + fabs(rhs.imag());
956  double oos = 1.0 / s;
957  double ars = lhs.real() * oos;
958  double ais = lhs.imag() * oos;
959  double brs = rhs.real() * oos;
960  double bis = rhs.imag() * oos;
961  s = (brs * brs) + (bis * bis);
962  oos = 1.0 / s;
963  return complex<double>(((ars * brs) + (ais * bis)) * oos,
964  ((ais * brs) - (ars * bis)) * oos);
965  }
966 
967  template <typename ValueType>
968  __host__ __device__
969  inline complex<ValueType> operator/(const complex<ValueType>& lhs, const ValueType & rhs){
970  return complex<ValueType>(lhs.real()/rhs,lhs.imag()/rhs);
971  // return cuCdivf(lhs,complex<float>(rhs));
972  }
973 
974  template <typename ValueType>
975  __host__ __device__
976  inline complex<ValueType> operator/(const ValueType& lhs, const complex<ValueType>& rhs){
977  const ValueType cross_norm = lhs * rhs.real();
978  const ValueType rhs_norm = norm(rhs);
979  return complex<ValueType>(cross_norm/rhs_norm,(-lhs.real() * rhs.imag()) / rhs_norm);
980  }
981 
982  template <>
983  __host__ __device__
984  inline complex<float> operator/(const float& lhs, const complex<float>& rhs){
985  return complex<float>(lhs) / rhs;
986  }
987  template <>
988  __host__ __device__
989  inline complex<double> operator/(const double& lhs, const complex<double>& rhs){
990  return complex<double>(lhs) / rhs;
991  }
992 
993  // Unary arithmetic operations
994  template <typename ValueType>
995  __host__ __device__
996  inline complex<ValueType> operator+(const complex<ValueType>& rhs){
997  return rhs;
998  }
999  template <typename ValueType>
1000  __host__ __device__
1001  inline complex<ValueType> operator-(const complex<ValueType>& rhs){
1002  return rhs*-ValueType(1);
1003  }
1004 
1005  // Equality operators
1006  template <typename ValueType>
1007  __host__ __device__
1008  inline bool operator==(const complex<ValueType>& lhs, const complex<ValueType>& rhs){
1009  if(lhs.real() == rhs.real() && lhs.imag() == rhs.imag()){
1010  return true;
1011  }
1012  return false;
1013  }
1014 
1015  template <typename ValueType>
1016  __host__ __device__
1017  inline bool operator==(const ValueType & lhs, const complex<ValueType>& rhs){
1018  if(lhs == rhs.real() && rhs.imag() == 0){
1019  return true;
1020  }
1021  return false;
1022  }
1023  template <typename ValueType>
1024  __host__ __device__
1025  inline bool operator==(const complex<ValueType> & lhs, const ValueType& rhs){
1026  if(lhs.real() == rhs && lhs.imag() == 0){
1027  return true;
1028  }
1029  return false;
1030  }
1031 
1032 
1033  template <typename ValueType>
1034  __host__ __device__
1035  inline bool operator!=(const complex<ValueType>& lhs, const complex<ValueType>& rhs){
1036  return !(lhs == rhs);
1037  }
1038 
1039  template <typename ValueType>
1040  __host__ __device__
1041  inline bool operator!=(const ValueType & lhs, const complex<ValueType>& rhs){
1042  return !(lhs == rhs);
1043  }
1044 
1045  template <typename ValueType>
1046  __host__ __device__
1047  inline bool operator!=(const complex<ValueType> & lhs, const ValueType& rhs){
1048  return !(lhs == rhs);
1049  }
1050 
1051 
1052  template <typename ValueType>
1053  __host__ __device__
1054  inline complex<ValueType> conj(const complex<ValueType>& z){
1055  return complex<ValueType>(z.real(),-z.imag());
1056  }
1057 
1058  template <typename ValueType>
1059  __host__ __device__
1060  inline ValueType abs(const complex<ValueType>& z){
1061  return ::hypot(z.real(),z.imag());
1062  }
1063  template <>
1064  __host__ __device__
1065  inline float abs(const complex<float>& z){
1066  return ::hypotf(z.real(),z.imag());
1067  }
1068  template<>
1069  __host__ __device__
1070  inline double abs(const complex<double>& z){
1071  return ::hypot(z.real(),z.imag());
1072  }
1073 
1074  template <typename ValueType>
1075  __host__ __device__
1076  inline ValueType arg(const complex<ValueType>& z){
1077  return atan2(z.imag(),z.real());
1078  }
1079  template<>
1080  __host__ __device__
1081  inline float arg(const complex<float>& z){
1082  return atan2f(z.imag(),z.real());
1083  }
1084  template<>
1085  __host__ __device__
1086  inline double arg(const complex<double>& z){
1087  return atan2(z.imag(),z.real());
1088  }
1089 
1090  template <typename ValueType>
1091  __host__ __device__
1092  inline ValueType norm(const complex<ValueType>& z){
1093  return z.real()*z.real() + z.imag()*z.imag();
1094  }
1095 
1096  template <typename ValueType>
1097  __host__ __device__
1098  inline complex<ValueType> polar(const ValueType & m, const ValueType & theta){
1099  return complex<ValueType>(m * ::cos(theta),m * ::sin(theta));
1100  }
1101 
1102  template <>
1103  __host__ __device__
1104  inline complex<float> polar(const float & magnitude, const float & angle){
1105  return complex<float>(magnitude * ::cosf(angle),magnitude * ::sinf(angle));
1106  }
1107 
1108  template <>
1109  __host__ __device__
1110  inline complex<double> polar(const double & magnitude, const double & angle){
1111  return complex<double>(magnitude * ::cos(angle),magnitude * ::sin(angle));
1112  }
1113 
1114  // Transcendental functions implementation
1115  template <typename ValueType>
1116  __host__ __device__
1117  inline complex<ValueType> cos(const complex<ValueType>& z){
1118  const ValueType re = z.real();
1119  const ValueType im = z.imag();
1120  return complex<ValueType>(::cos(re) * ::cosh(im), -::sin(re) * ::sinh(im));
1121  }
1122 
1123  template <>
1124  __host__ __device__
1125  inline complex<float> cos(const complex<float>& z){
1126  const float re = z.real();
1127  const float im = z.imag();
1128  return complex<float>(cosf(re) * coshf(im), -sinf(re) * sinhf(im));
1129  }
1130 
1131  template <typename ValueType>
1132  __host__ __device__
1133  inline complex<ValueType> cosh(const complex<ValueType>& z){
1134  const ValueType re = z.real();
1135  const ValueType im = z.imag();
1136  return complex<ValueType>(::cosh(re) * ::cos(im), ::sinh(re) * ::sin(im));
1137  }
1138 
1139  template <>
1140  __host__ __device__
1141  inline complex<float> cosh(const complex<float>& z){
1142  const float re = z.real();
1143  const float im = z.imag();
1144  return complex<float>(::coshf(re) * ::cosf(im), ::sinhf(re) * ::sinf(im));
1145  }
1146 
1147 
1148  template <typename ValueType>
1149  __host__ __device__
1150  inline complex<ValueType> exp(const complex<ValueType>& z){
1151  return polar(::exp(z.real()),z.imag());
1152  }
1153 
1154  template <>
1155  __host__ __device__
1156  inline complex<float> exp(const complex<float>& z){
1157  return polar(::expf(z.real()),z.imag());
1158  }
1159 
1160  template <typename ValueType>
1161  __host__ __device__
1162  inline complex<ValueType> log(const complex<ValueType>& z){
1163  return complex<ValueType>(::log(abs(z)),arg(z));
1164  }
1165 
1166  template <>
1167  __host__ __device__
1168  inline complex<float> log(const complex<float>& z){
1169  return complex<float>(::logf(abs(z)),arg(z));
1170  }
1171 
1172 
1173  template <typename ValueType>
1174  __host__ __device__
1175  inline complex<ValueType> log10(const complex<ValueType>& z){
1176  // Using the explicit literal prevents compile time warnings in
1177  // devices that don't support doubles
1178  return log(z)/ValueType(2.30258509299404568402);
1179  // return log(z)/ValueType(::log(10.0));
1180  }
1181 
1182  template <typename ValueType>
1183  __host__ __device__
1184  inline complex<ValueType> pow(const complex<ValueType>& z, const ValueType & exponent){
1185  return exp(log(z)*exponent);
1186  }
1187 
1188  template <typename ValueType>
1189  __host__ __device__
1190  inline complex<ValueType> pow(const complex<ValueType>& z, const complex<ValueType> & exponent){
1191  return exp(log(z)*exponent);
1192  }
1193 
1194  template <typename ValueType>
1195  __host__ __device__
1196  inline complex<ValueType> pow(const ValueType & x, const complex<ValueType> & exponent){
1197  return exp(::log(x)*exponent);
1198  }
1199 
1200  template <>
1201  __host__ __device__
1202  inline complex<float> pow(const float & x, const complex<float> & exponent){
1203  return exp(::logf(x)*exponent);
1204  }
1205 
1206  template <typename ValueType>
1207  __host__ __device__
1208  inline complex<ValueType> pow(const complex<ValueType>& z,const int & exponent){
1209  return exp(log(z)*ValueType(exponent));
1210  }
1211 
1212  template <typename ValueType>
1213  __host__ __device__
1214  inline complex<ValueType> sin(const complex<ValueType>& z){
1215  const ValueType re = z.real();
1216  const ValueType im = z.imag();
1217  return complex<ValueType>(::sin(re) * ::cosh(im), ::cos(re) * ::sinh(im));
1218  }
1219 
1220  template <>
1221  __host__ __device__
1222  inline complex<float> sin(const complex<float>& z){
1223  const float re = z.real();
1224  const float im = z.imag();
1225  return complex<float>(::sinf(re) * ::coshf(im), ::cosf(re) * ::sinhf(im));
1226  }
1227 
1228  template <typename ValueType>
1229  __host__ __device__
1230  inline complex<ValueType> sinh(const complex<ValueType>& z){
1231  const ValueType re = z.real();
1232  const ValueType im = z.imag();
1233  return complex<ValueType>(::sinh(re) * ::cos(im), ::cosh(re) * ::sin(im));
1234  }
1235 
1236  template <>
1237  __host__ __device__
1238  inline complex<float> sinh(const complex<float>& z){
1239  const float re = z.real();
1240  const float im = z.imag();
1241  return complex<float>(::sinhf(re) * ::cosf(im), ::coshf(re) * ::sinf(im));
1242  }
1243 
1244  template <typename ValueType>
1245  __host__ __device__
1246  inline complex<ValueType> sqrt(const complex<ValueType>& z){
1247  return polar(::sqrt(abs(z)),arg(z)/ValueType(2));
1248  }
1249 
1250  template <>
1251  __host__ __device__
1252  inline complex<float> sqrt(const complex<float>& z){
1253  return polar(::sqrtf(abs(z)),arg(z)/float(2));
1254  }
1255 
1256  template <typename ValueType>
1257  __host__ __device__
1258  inline complex<ValueType> tan(const complex<ValueType>& z){
1259  return sin(z)/cos(z);
1260  }
1261 
1262  template <typename ValueType>
1263  __host__ __device__
1264  inline complex<ValueType> tanh(const complex<ValueType>& z){
1265  // This implementation seems better than the simple sin/cos
1266  return (exp(ValueType(2)*z)-ValueType(1))/(exp(ValueType(2)*z)+ValueType(1));
1267  // return sinh(z)/cosh(z);
1268  }
1269 
1270  // Inverse trigonometric functions implementation
1271 
1272  template <typename ValueType>
1273  __host__ __device__
1274  inline complex<ValueType> acos(const complex<ValueType>& z){
1275  const complex<ValueType> ret = asin(z);
1276  return complex<ValueType>(ValueType(M_PI/2.0) - ret.real(),-ret.imag());
1277  }
1278 
1279  template <typename ValueType>
1280  __host__ __device__
1281  inline complex<ValueType> asin(const complex<ValueType>& z){
1282  const complex<ValueType> i(0,1);
1283  return -i*asinh(i*z);
1284  }
1285 
1286  template <typename ValueType>
1287  __host__ __device__
1288  inline complex<ValueType> atan(const complex<ValueType>& z){
1289  const complex<ValueType> i(0,1);
1290  return -i*atanh(i*z);
1291  }
1292 
1293  template <typename ValueType>
1294  __host__ __device__
1295  inline complex<ValueType> acosh(const complex<ValueType>& z){
1296  quda::complex<ValueType> ret((z.real() - z.imag()) * (z.real() + z.imag()) - ValueType(1.0),
1297  ValueType(2.0) * z.real() * z.imag());
1298  ret = sqrt(ret);
1299  if (z.real() < ValueType(0.0)){
1300  ret = -ret;
1301  }
1302  ret += z;
1303  ret = log(ret);
1304  if (ret.real() < ValueType(0.0)){
1305  ret = -ret;
1306  }
1307  return ret;
1308 
1309  /*
1310  quda::complex<ValueType> ret = log(sqrt(z*z-ValueType(1))+z);
1311  if(ret.real() < 0){
1312  ret.real(-ret.real());
1313  }
1314  return ret;
1315  */
1316  }
1317 
1318  template <typename ValueType>
1319  __host__ __device__
1320  inline complex<ValueType> asinh(const complex<ValueType>& z){
1321  return log(sqrt(z*z+ValueType(1))+z);
1322  }
1323 
1324  template <typename ValueType>
1325  __host__ __device__
1326  inline complex<ValueType> atanh(const complex<ValueType>& z){
1327  ValueType imag2 = z.imag() * z.imag();
1328  ValueType n = ValueType(1.0) + z.real();
1329  n = imag2 + n * n;
1330 
1331  ValueType d = ValueType(1.0) - z.real();
1332  d = imag2 + d * d;
1333  complex<ValueType> ret(ValueType(0.25) * (::log(n) - ::log(d)),0);
1334 
1335  d = ValueType(1.0) - z.real() * z.real() - imag2;
1336 
1337  ret.imag(ValueType(0.5) * ::atan2(ValueType(2.0) * z.imag(), d));
1338  return ret;
1339  //return (log(ValueType(1)+z)-log(ValueType(1)-z))/ValueType(2);
1340  }
1341 
1342  template <typename ValueType>
1343  __host__ __device__
1344  inline complex<float> atanh(const complex<float>& z){
1345  float imag2 = z.imag() * z.imag();
1346  float n = float(1.0) + z.real();
1347  n = imag2 + n * n;
1348 
1349  float d = float(1.0) - z.real();
1350  d = imag2 + d * d;
1351  complex<float> ret(float(0.25) * (::logf(n) - ::logf(d)),0);
1352 
1353  d = float(1.0) - z.real() * z.real() - imag2;
1354 
1355  ret.imag(float(0.5) * ::atan2f(float(2.0) * z.imag(), d));
1356  return ret;
1357  //return (log(ValueType(1)+z)-log(ValueType(1)-z))/ValueType(2);
1358 
1359  }
1360 
1361  template <typename real> __host__ __device__ inline complex<real> cmul(const complex<real> &x, const complex<real> &y)
1362  {
1363  complex<real> w;
1364  w.x = x.real() * y.real();
1365  w.x -= x.imag() * y.imag();
1366  w.y = x.imag() * y.real();
1367  w.y += x.real() * y.imag();
1368  return w;
1369  }
1370 
1371  template <typename real>
1372  __host__ __device__ inline complex<real> cmac(const complex<real> &x, const complex<real> &y, const complex<real> &z)
1373  {
1374  complex<real> w = z;
1375  w.x += x.real() * y.real();
1376  w.x -= x.imag() * y.imag();
1377  w.y += x.imag() * y.real();
1378  w.y += x.real() * y.imag();
1379  return w;
1380  }
1381 
1382 } // end namespace quda
__host__ __device__ complex< float > sinh(const complex< float > &z)
__host__ __device__ ValueType tanh(ValueType x)
Definition: complex_quda.h:91
__host__ __device__ void real(short re) volatile
Definition: complex_quda.h:790
__host__ __device__ ValueType sinh(ValueType x)
Definition: complex_quda.h:86
__host__ __device__ void real(char re) volatile
Definition: complex_quda.h:743
__host__ __device__ void real(int re) volatile
Definition: complex_quda.h:836
__host__ __device__ complex< float > & operator-=(const complex< float > z)
Definition: complex_quda.h:516
__host__ __device__ complex< double > & operator=(const complex< T > z)
Definition: complex_quda.h:622
ValueType value_type
Definition: complex_quda.h:371
__host__ __device__ ValueType norm(const complex< ValueType > &z)
Returns the magnitude of z squared.
__host__ __device__ void imag(short im)
Definition: complex_quda.h:793
__host__ __device__ complex< double > & operator+=(const complex< float > z)
Definition: complex_quda.h:638
__host__ __device__ ValueType atan(ValueType x)
Definition: complex_quda.h:71
fieldorder_wrapper is an internal class that is used to wrap instances of FieldOrder accessors...
Definition: complex_quda.h:32
__host__ __device__ complex< double > & operator-=(const complex< double > z)
Definition: complex_quda.h:646
__host__ __device__ ValueType exp(ValueType x)
Definition: complex_quda.h:96
__host__ __device__ complex< float > & operator*=(const complex< float > z)
Definition: complex_quda.h:524
__host__ __device__ complex< char > & operator-=(const complex< char > z)
Definition: complex_quda.h:732
__host__ __device__ ValueType cosh(ValueType x)
Definition: complex_quda.h:81
__host__ __device__ void real(double re) volatile
Definition: complex_quda.h:686
__host__ __device__ ValueType sqrt(ValueType x)
Definition: complex_quda.h:120
__host__ __device__ complex< real > cmac(const complex< real > &x, const complex< real > &y, const complex< real > &z)
__host__ __device__ complex< float > cosh(const complex< float > &z)
__host__ __device__ void real(float re)
Definition: complex_quda.h:558
__host__ __device__ void imag(short im) volatile
Definition: complex_quda.h:791
__host__ __device__ void real(short re)
Definition: complex_quda.h:792
__host__ __device__ char real() const volatile
Definition: complex_quda.h:739
__host__ __device__ complex< ValueType > tanh(const complex< ValueType > &z)
__host__ __device__ void imag(float im)
Definition: complex_quda.h:559
__host__ __device__ int real() const
Definition: complex_quda.h:834
std::basic_istream< charT, traits > & operator>>(std::basic_istream< charT, traits > &is, complex< ValueType > &z)
Definition: complex_quda.h:318
__host__ __device__ double imag() const
Definition: complex_quda.h:685
__host__ __device__ complex< int > & operator+=(const complex< int > z)
Definition: complex_quda.h:818
__host__ __device__ ValueType tan(ValueType x)
Definition: complex_quda.h:56
__host__ __device__ void imag(double im) volatile
Definition: complex_quda.h:687
__host__ __device__ char real() const
Definition: complex_quda.h:741
fieldorder_wrapper is an internal class that is used to wrap instances of FieldOrder accessors...
__host__ __device__ complex< ValueType > & operator+=(const complex< ValueType > z)
Definition: complex_quda.h:407
__host__ __device__ void real(float re) volatile
Definition: complex_quda.h:556
__host__ __device__ complex< float > & operator=(const complex< T > z)
Definition: complex_quda.h:500
__host__ __device__ complex< ValueType > tan(const complex< ValueType > &z)
__host__ __device__ complex< ValueType > & operator-=(const complex< ValueType > z)
Definition: complex_quda.h:415
__host__ __device__ complex< ValueType > & operator*=(const complex< ValueType > z)
Definition: complex_quda.h:423
__host__ __device__ complex< ValueType > asin(const complex< ValueType > &z)
__host__ __device__ void imag(float im) volatile
Definition: complex_quda.h:557
__host__ __device__ ValueType sin(ValueType x)
Definition: complex_quda.h:51
__host__ __device__ char imag() const
Definition: complex_quda.h:742
__host__ __device__ complex< ValueType > acos(const complex< ValueType > &z)
__host__ __device__ complex< float > sin(const complex< float > &z)
__host__ __device__ int imag() const volatile
Definition: complex_quda.h:833
__host__ __device__ int real() const volatile
Definition: complex_quda.h:832
__host__ __device__ void real(double re)
Definition: complex_quda.h:688
__host__ __device__ ValueType atan2(ValueType x, ValueType y)
Definition: complex_quda.h:76
__host__ __device__ short imag() const volatile
Definition: complex_quda.h:787
__host__ __device__ void real(char re)
Definition: complex_quda.h:745
__host__ __device__ void real(int re)
Definition: complex_quda.h:838
__device__ __host__ ColorSpinor< Float, Nc, Ns > operator-(const ColorSpinor< Float, Nc, Ns > &x, const ColorSpinor< Float, Nc, Ns > &y)
ColorSpinor subtraction operator.
__host__ __device__ complex< float > cos(const complex< float > &z)
__host__ __device__ complex< int > & operator-=(const complex< int > z)
Definition: complex_quda.h:825
__host__ __device__ ValueType pow(ValueType x, ExponentType e)
Definition: complex_quda.h:111
__host__ __device__ void imag(double im)
Definition: complex_quda.h:689
__host__ __device__ float imag() const volatile
Definition: complex_quda.h:553
__host__ __device__ void imag(int im)
Definition: complex_quda.h:839
__host__ __device__ complex< float > & operator/=(const complex< float > z)
Definition: complex_quda.h:531
__host__ __device__ float real() const
Definition: complex_quda.h:554
__host__ __device__ short real() const volatile
Definition: complex_quda.h:786
__host__ __device__ complex< float > exp(const complex< float > &z)
__host__ __device__ complex< float > log(const complex< float > &z)
__host__ __device__ short real() const
Definition: complex_quda.h:788
__host__ __device__ complex< double > & operator/=(const complex< double > z)
Definition: complex_quda.h:661
__device__ __host__ complex< Float > operator*(const Float &a, const fieldorder_wrapper< Float, storeFloat > &b)
__host__ __device__ complex< ValueType > log10(const complex< ValueType > &z)
__host__ __device__ complex< ValueType > & operator/=(const complex< ValueType > z)
Definition: complex_quda.h:430
__host__ __device__ complex< ValueType > asinh(const complex< ValueType > &z)
__host__ __device__ complex< char > & operator+=(const complex< char > z)
Definition: complex_quda.h:725
__host__ __device__ complex< ValueType > & operator*=(const ValueType z)
Definition: complex_quda.h:437
__host__ __device__ complex< real > cmul(const complex< real > &x, const complex< real > &y)
__host__ __device__ complex< ValueType > & operator=(const complex< T > z)
Definition: complex_quda.h:399
__host__ __device__ ValueType log(ValueType x)
Definition: complex_quda.h:101
__host__ __device__ double imag() const volatile
Definition: complex_quda.h:683
__host__ __device__ complex< float > & operator*=(const float z)
Definition: complex_quda.h:538
__host__ __device__ complex< ValueType > operator/(const complex< ValueType > &lhs, const complex< ValueType > &rhs)
Definition: complex_quda.h:926
__host__ __device__ volatile complex< float > & operator=(const complex< T > z) volatile
Definition: complex_quda.h:491
__host__ __device__ ValueType log10(ValueType x)
Definition: complex_quda.h:106
__host__ __device__ char imag() const volatile
Definition: complex_quda.h:740
__host__ __device__ short imag() const
Definition: complex_quda.h:789
__host__ __device__ double abs(const complex< double > &z)
__host__ __device__ complex< float > sqrt(const complex< float > &z)
__shared__ float s[]
__host__ __device__ void imag(char im)
Definition: complex_quda.h:746
__host__ __device__ complex< double > & operator*=(const double z)
Definition: complex_quda.h:668
__host__ __device__ complex< ValueType > atan(const complex< ValueType > &z)
__host__ __device__ complex< short > & operator+=(const complex< short > z)
Definition: complex_quda.h:772
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
__host__ __device__ double real() const
Definition: complex_quda.h:684
__host__ __device__ void imag(int im) volatile
Definition: complex_quda.h:837
__host__ __device__ ValueType cos(ValueType x)
Definition: complex_quda.h:46
__host__ __device__ complex< ValueType > polar(const ValueType &m, const ValueType &theta=0)
Returns the complex with magnitude m and angle theta in radians.
__host__ __device__ ValueType abs(ValueType x)
Definition: complex_quda.h:125
__host__ __device__ int imag() const
Definition: complex_quda.h:835
__host__ __device__ ValueType acos(ValueType x)
Definition: complex_quda.h:61
__host__ __device__ complex< float > pow(const float &x, const complex< float > &exponent)
__host__ __device__ ValueType conj(ValueType x)
Definition: complex_quda.h:130
__host__ __device__ complex< double > & operator*=(const complex< double > z)
Definition: complex_quda.h:654
__device__ __host__ complex< Float > operator+(const fieldorder_wrapper< Float, storeFloat > &a, const complex< Float > &b)
__host__ __device__ bool operator!=(const complex< ValueType > &lhs, const complex< ValueType > &rhs)
__host__ __device__ complex< ValueType > atanh(const complex< ValueType > &z)
__host__ __device__ float imag() const
Definition: complex_quda.h:555
__host__ __device__ complex< float > & operator+=(const complex< float > z)
Definition: complex_quda.h:508
__host__ __device__ complex< short > & operator-=(const complex< short > z)
Definition: complex_quda.h:779
__host__ __device__ volatile complex< double > & operator=(const complex< T > z) volatile
Definition: complex_quda.h:613
__host__ __device__ double real() const volatile
Definition: complex_quda.h:682
__host__ __device__ void imag(char im) volatile
Definition: complex_quda.h:744
__host__ __device__ float real() const volatile
Definition: complex_quda.h:552
__host__ __device__ complex< ValueType > acosh(const complex< ValueType > &z)
__host__ __device__ bool operator==(const complex< ValueType > &lhs, const complex< ValueType > &rhs)
__host__ __device__ ValueType asin(ValueType x)
Definition: complex_quda.h:66
__host__ __device__ complex< double > & operator+=(const complex< double > z)
Definition: complex_quda.h:630