QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
blas_helper.cuh
Go to the documentation of this file.
1 #pragma once
2 
3 #include <color_spinor_field.h>
4 
5 //#define QUAD_SUM
6 #ifdef QUAD_SUM
7 #include <dbldbl.h>
8 #endif
9 
10 // these definitions are used to avoid calling
11 // std::complex<type>::real/imag which have C++11 ABI incompatibility
12 // issues with certain versions of GCC
13 
14 #define REAL(a) (*((double *)&a))
15 #define IMAG(a) (*((double *)&a + 1))
16 
17 namespace quda
18 {
19 
20  inline void checkSpinor(const ColorSpinorField &a, const ColorSpinorField &b)
21  {
22  if (a.Length() != b.Length()) errorQuda("lengths do not match: %lu %lu", a.Length(), b.Length());
23  if (a.Stride() != b.Stride()) errorQuda("strides do not match: %d %d", a.Stride(), b.Stride());
24  }
25 
26  inline void checkLength(const ColorSpinorField &a, const ColorSpinorField &b)
27  {
28  if (a.Length() != b.Length()) errorQuda("lengths do not match: %lu %lu", a.Length(), b.Length());
29  if (a.Stride() != b.Stride()) errorQuda("strides do not match: %d %d", a.Stride(), b.Stride());
30  }
31 
32 #ifdef QUAD_SUM
33 #define QudaSumFloat doubledouble
34 #define QudaSumFloat2 doubledouble2
35 #define QudaSumFloat3 doubledouble3
36  template <> struct scalar<doubledouble> {
37  typedef doubledouble type;
38  };
39  template <> struct scalar<doubledouble2> {
40  typedef doubledouble type;
41  };
42  template <> struct scalar<doubledouble3> {
43  typedef doubledouble type;
44  };
45  template <> struct scalar<doubledouble4> {
46  typedef doubledouble type;
47  };
48  template <> struct vector<doubledouble, 2> {
49  typedef doubledouble2 type;
50  };
51 #else
52 #define QudaSumFloat double
53 #define QudaSumFloat2 double2
54 #define QudaSumFloat3 double3
55 #define QudaSumFloat4 double4
56 #endif
57 
58  __host__ __device__ inline double set(double &x) { return x; }
59  __host__ __device__ inline double2 set(double2 &x) { return x; }
60  __host__ __device__ inline double3 set(double3 &x) { return x; }
61  __host__ __device__ inline double4 set(double4 &x) { return x; }
62  __host__ __device__ inline void sum(double &a, double &b) { a += b; }
63  __host__ __device__ inline void sum(double2 &a, double2 &b)
64  {
65  a.x += b.x;
66  a.y += b.y;
67  }
68  __host__ __device__ inline void sum(double3 &a, double3 &b)
69  {
70  a.x += b.x;
71  a.y += b.y;
72  a.z += b.z;
73  }
74  __host__ __device__ inline void sum(double4 &a, double4 &b)
75  {
76  a.x += b.x;
77  a.y += b.y;
78  a.z += b.z;
79  a.w += b.w;
80  }
81 
82 #ifdef QUAD_SUM
83  __host__ __device__ inline double set(doubledouble &a) { return a.head(); }
84  __host__ __device__ inline double2 set(doubledouble2 &a) { return make_double2(a.x.head(), a.y.head()); }
85  __host__ __device__ inline double3 set(doubledouble3 &a) { return make_double3(a.x.head(), a.y.head(), a.z.head()); }
86  __host__ __device__ inline void sum(double &a, doubledouble &b) { a += b.head(); }
87  __host__ __device__ inline void sum(double2 &a, doubledouble2 &b)
88  {
89  a.x += b.x.head();
90  a.y += b.y.head();
91  }
92  __host__ __device__ inline void sum(double3 &a, doubledouble3 &b)
93  {
94  a.x += b.x.head();
95  a.y += b.y.head();
96  a.z += b.z.head();
97  }
98 #endif
99 
100 } // namespace quda
doubledouble x
Definition: dbldbl.h:339
__device__ __host__ double head() const
Definition: dbldbl.h:303
#define errorQuda(...)
Definition: util_quda.h:121
__host__ __device__ void sum(double &a, double &b)
Definition: blas_helper.cuh:62
doubledouble x
Definition: dbldbl.h:357
doubledouble z
Definition: dbldbl.h:359
void checkLength(const ColorSpinorField &a, const ColorSpinorField &b)
Definition: blas_helper.cuh:26
doubledouble y
Definition: dbldbl.h:358
void checkSpinor(const ColorSpinorField &a, const ColorSpinorField &b)
Definition: blas_helper.cuh:20
doubledouble y
Definition: dbldbl.h:340