QUDA  v0.5.0
A library for QCD on GPUs
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
float_vector.h
Go to the documentation of this file.
1 
9 #pragma once
10 
11 __host__ __device__ inline double2 operator+(const double2& x, const double2 &y) {
12  return make_double2(x.x + y.x, x.y + y.y);
13 }
14 
15 __host__ __device__ inline double2 operator-(const double2& x, const double2 &y) {
16  return make_double2(x.x - y.x, x.y - y.y);
17 }
18 
19 __host__ __device__ inline float2 operator-(const float2& x, const float2 &y) {
20  return make_float2(x.x - y.x, x.y - y.y);
21 }
22 
23 __host__ __device__ inline float4 operator-(const float4& x, const float4 &y) {
24  return make_float4(x.x - y.x, x.y - y.y, x.z - y.z, x.w - y.w);
25 }
26 
27 __host__ inline double3 operator+(const double3& x, const double3 &y) {
28  double3 z;
29  z.x = x.x + y.x; z.y = x.y + y.y; z.z = x.z + y.z;
30  return z;
31 }
32 
33 __device__ inline float4 operator*(const float a, const float4 x) {
34  float4 y;
35  y.x = a*x.x;
36  y.y = a*x.y;
37  y.z = a*x.z;
38  y.w = a*x.w;
39  return y;
40 }
41 
42 __device__ inline float2 operator*(const float a, const float2 x) {
43  float2 y;
44  y.x = a*x.x;
45  y.y = a*x.y;
46  return y;
47 }
48 
49 __device__ inline double2 operator*(const double a, const double2 x) {
50  double2 y;
51  y.x = a*x.x;
52  y.y = a*x.y;
53  return y;
54 }
55 
56 __device__ inline double4 operator*(const double a, const double4 x) {
57  double4 y;
58  y.x = a*x.x;
59  y.y = a*x.y;
60  y.z = a*x.z;
61  y.w = a*x.w;
62  return y;
63 }
64 
65 __device__ inline float2 operator+(const float2 x, const float2 y) {
66  float2 z;
67  z.x = x.x + y.x;
68  z.y = x.y + y.y;
69  return z;
70 }
71 
72 __device__ inline float4 operator+(const float4 x, const float4 y) {
73  float4 z;
74  z.x = x.x + y.x;
75  z.y = x.y + y.y;
76  z.z = x.z + y.z;
77  z.w = x.w + y.w;
78  return z;
79 }
80 
81 __device__ inline float4 operator+=(float4 &x, const float4 y) {
82  x.x += y.x;
83  x.y += y.y;
84  x.z += y.z;
85  x.w += y.w;
86  return x;
87 }
88 
89 __device__ inline float2 operator+=(float2 &x, const float2 y) {
90  x.x += y.x;
91  x.y += y.y;
92  return x;
93 }
94 
95 __host__ __device__ inline double2 operator+=(double2 &x, const double2 y) {
96  x.x += y.x;
97  x.y += y.y;
98  return x;
99 }
100 
101 __host__ __device__ inline double3 operator+=(double3 &x, const double3 y) {
102  x.x += y.x;
103  x.y += y.y;
104  x.z += y.z;
105  return x;
106 }
107 
108 __device__ inline float4 operator-=(float4 &x, const float4 y) {
109  x.x -= y.x;
110  x.y -= y.y;
111  x.z -= y.z;
112  x.w -= y.w;
113  return x;
114 }
115 
116 __device__ inline float2 operator-=(float2 &x, const float2 y) {
117  x.x -= y.x;
118  x.y -= y.y;
119  return x;
120 }
121 
122 __device__ inline double2 operator-=(double2 &x, const double2 y) {
123  x.x -= y.x;
124  x.y -= y.y;
125  return x;
126 }
127 
128 __device__ inline float2 operator*=(float2 &x, const float a) {
129  x.x *= a;
130  x.y *= a;
131  return x;
132 }
133 
134 __device__ inline float4 operator*=(float4 &a, const float &b) {
135  a.x *= b;
136  a.y *= b;
137  a.z *= b;
138  a.w *= b;
139  return a;
140 }
141 
142 __device__ inline double2 operator*=(double2 &a, const float &b) {
143  a.x *= b;
144  a.y *= b;
145  return a;
146 }
147 
148 __device__ inline double4 operator*=(double4 &a, const float &b) {
149  a.x *= b;
150  a.y *= b;
151  a.z *= b;
152  a.w *= b;
153  return a;
154 }
155 
156 __device__ inline float2 operator-(const float2 &x) {
157  return make_float2(-x.x, -x.y);
158 }
159 
160 __device__ inline double2 operator-(const double2 &x) {
161  return make_double2(-x.x, -x.y);
162 }
163 
164 
165 /*
166  Operations to return the maximium absolute value of a FloatN vector
167  */
168 
169 __forceinline__ __device__ float max_fabs(const float4 &c) {
170  float a = fmaxf(fabsf(c.x), fabsf(c.y));
171  float b = fmaxf(fabsf(c.z), fabsf(c.w));
172  return fmaxf(a, b);
173 };
174 
175 __forceinline__ __device__ float max_fabs(const float2 &b) {
176  return fmaxf(fabsf(b.x), fabsf(b.y));
177 };
178 
179 __forceinline__ __device__ double max_fabs(const double4 &c) {
180  double a = fmaxf(fabsf(c.x), fabsf(c.y));
181  double b = fmaxf(fabsf(c.z), fabsf(c.w));
182  return fmaxf(a, b);
183 };
184 
185 __forceinline__ __device__ double max_fabs(const double2 &b) {
186  return fmaxf(fabsf(b.x), fabsf(b.y));
187 };
188 
189 /*
190  Precision conversion routines for vector types
191  */
192 
193 __forceinline__ __device__ float2 make_FloatN(const double2 &a) {
194  return make_float2(a.x, a.y);
195 }
196 
197 __forceinline__ __device__ float4 make_FloatN(const double4 &a) {
198  return make_float4(a.x, a.y, a.z, a.w);
199 }
200 
201 __forceinline__ __device__ double2 make_FloatN(const float2 &a) {
202  return make_double2(a.x, a.y);
203 }
204 
205 __forceinline__ __device__ double4 make_FloatN(const float4 &a) {
206  return make_double4(a.x, a.y, a.z, a.w);
207 }
208 
209 __forceinline__ __device__ short4 make_shortN(const float4 &a) {
210  return make_short4(a.x, a.y, a.z, a.w);
211 }
212 
213 __forceinline__ __device__ short2 make_shortN(const float2 &a) {
214  return make_short2(a.x, a.y);
215 }
216 
217 __forceinline__ __device__ short4 make_shortN(const double4 &a) {
218  return make_short4(a.x, a.y, a.z, a.w);
219 }
220 
221 __forceinline__ __device__ short2 make_shortN(const double2 &a) {
222  return make_short2(a.x, a.y);
223 }
224