QUDA  0.9.0
color_spinor_util.cu
Go to the documentation of this file.
1 #include <color_spinor_field.h>
3 #include <index_helper.cuh>
4 
5 namespace quda {
6 
7  using namespace colorspinor;
8 
12  template <class T>
13  void random(T &t) {
14  for (int parity=0; parity<t.Nparity(); parity++) {
15  for (int x_cb=0; x_cb<t.VolumeCB(); x_cb++) {
16  for (int s=0; s<t.Nspin(); s++) {
17  for (int c=0; c<t.Ncolor(); c++) {
18  t(parity,x_cb,s,c).real(comm_drand());
19  t(parity,x_cb,s,c).imag(comm_drand());
20  }
21  }
22  }
23  }
24  }
25 
29  template <class T>
30  void point(T &t, int x, int s, int c) { t(x%2, x/2, s, c) = 1.0; }
31 
36  template <class T>
37  void constant(T &t, int k, int s, int c) {
38  for (int parity=0; parity<t.Nparity(); parity++) {
39  for (int x_cb=0; x_cb<t.VolumeCB(); x_cb++) {
40  // set all color-spin components to zero
41  for (int s2=0; s2<t.Nspin(); s2++) {
42  for (int c2=0; c2<t.Ncolor(); c2++) {
43  t(parity,x_cb,s2,c2) = 0.0;
44  }
45  }
46  t(parity,x_cb,s,c) = k; // now set the one we want
47  }
48  }
49  }
50 
54  template <class P>
55  void sin(P &p, int d, int n, int offset) {
56  int coord[4];
57  int X[4] = { p.X(0), p.X(1), p.X(2), p.X(3)};
58  X[0] *= (p.Nparity() == 1) ? 2 : 1; // need full lattice dims
59 
60  for (int parity=0; parity<p.Nparity(); parity++) {
61  for (int x_cb=0; x_cb<p.VolumeCB(); x_cb++) {
62  getCoords(coord, x_cb, X, parity);
63 
64  double mode = n * (double)coord[d] / X[d];
65  double k = (double)offset + sin (M_PI * mode);
66 
67  for (int s=0; s<p.Nspin(); s++)
68  for (int c=0; c<p.Ncolor(); c++)
69  p(parity, x_cb, s, c) = k;
70  }
71  }
72  }
73 
74  // print out the vector at volume point x
75  template <typename Float, int nSpin, int nColor, QudaFieldOrder order>
76  void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c) {
78  if (sourceType == QUDA_RANDOM_SOURCE) random(A);
79  else if (sourceType == QUDA_POINT_SOURCE) point(A, x, s, c);
80  else if (sourceType == QUDA_CONSTANT_SOURCE) constant(A, x, s, c);
81  else if (sourceType == QUDA_SINUSOIDAL_SOURCE) sin(A, x, s, c);
82  else errorQuda("Unsupported source type %d", sourceType);
83  }
84 
85  template <typename Float, int nSpin, QudaFieldOrder order>
86  void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c) {
87  if (a.Ncolor() == 2) {
88  genericSource<Float,nSpin,2,order>(a,sourceType, x, s, c);
89  } else if (a.Ncolor() == 3) {
90  genericSource<Float,nSpin,3,order>(a,sourceType, x, s, c);
91  } else if (a.Ncolor() == 4) {
92  genericSource<Float,nSpin,4,order>(a,sourceType, x, s, c);
93  } else if (a.Ncolor() == 8) {
94  genericSource<Float,nSpin,8,order>(a,sourceType, x, s, c);
95  } else if (a.Ncolor() == 12) {
96  genericSource<Float,nSpin,12,order>(a,sourceType, x, s, c);
97  } else if (a.Ncolor() == 16) {
98  genericSource<Float,nSpin,16,order>(a,sourceType, x, s, c);
99  } else if (a.Ncolor() == 20) {
100  genericSource<Float,nSpin,20,order>(a,sourceType, x, s, c);
101  } else if (a.Ncolor() == 24) {
102  genericSource<Float,nSpin,24,order>(a,sourceType, x, s, c);
103  } else if (a.Ncolor() == 32) {
104  genericSource<Float,nSpin,32,order>(a,sourceType, x, s, c);
105  } else {
106  errorQuda("Unsupported nColor=%d\n", a.Ncolor());
107  }
108  }
109 
110  template <typename Float, QudaFieldOrder order>
111  void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c) {
112  if (a.Nspin() == 1) {
113  genericSource<Float,1,order>(a,sourceType, x, s, c);
114  } else if (a.Nspin() == 2) {
115  genericSource<Float,2,order>(a,sourceType, x, s, c);
116  } else if (a.Nspin() == 4) {
117  genericSource<Float,4,order>(a,sourceType, x, s, c);
118  } else {
119  errorQuda("Unsupported nSpin=%d\n", a.Nspin());
120  }
121  }
122 
123  template <typename Float>
124  void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c) {
125  if (a.FieldOrder() == QUDA_SPACE_SPIN_COLOR_FIELD_ORDER) {
126  genericSource<Float,QUDA_SPACE_SPIN_COLOR_FIELD_ORDER>(a,sourceType, x, s, c);
127  } else {
128  errorQuda("Unsupported field order %d\n", a.FieldOrder());
129  }
130 
131  }
132 
133  void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c) {
134 
135  if (a.Precision() == QUDA_DOUBLE_PRECISION) {
136  genericSource<double>(a,sourceType, x, s, c);
137  } else if (a.Precision() == QUDA_SINGLE_PRECISION) {
138  genericSource<float>(a,sourceType, x, s, c);
139  } else {
140  errorQuda("Precision not supported");
141  }
142 
143  }
144 
145 
146  template <class U, class V>
147  int compareSpinor(const U &u, const V &v, const int tol) {
148  int fail_check = 16*tol;
149  int *fail = new int[fail_check];
150  for (int f=0; f<fail_check; f++) fail[f] = 0;
151 
152  int N = 2*u.Nspin()*u.Ncolor();
153  int *iter = new int[N];
154  for (int i=0; i<N; i++) iter[i] = 0;
155 
156  for (int parity=0; parity<v.Nparity(); parity++) {
157  for (int x_cb=0; x_cb<u.VolumeCB(); x_cb++) {
158 
159  for (int s=0; s<u.Nspin(); s++) {
160  for (int c=0; c<u.Ncolor(); c++) {
161  for (int z=0; z<2; z++) {
162  int j = (s*u.Ncolor() + c)*2+z;
163 
164  double diff = z==0 ? fabs(u(parity,x_cb,s,c,z).real() - v(parity,x_cb,s,c,z).real()) :
165  fabs(u(parity,x_cb,s,c).imag() - v(parity,x_cb,s,c).imag());
166 
167  for (int f=0; f<fail_check; f++) {
168  if (diff > pow(10.0,-(f+1)/(double)tol)) {
169  fail[f]++;
170  }
171  }
172 
173  if (diff > 1e-3) iter[j]++;
174  }
175  }
176  }
177  }
178  }
179 
180  // reduce over all processes
181  for (int i=0; i<N; i++) comm_allreduce_int(&iter[i]);
182  for (int f=0; f<fail_check; f++) comm_allreduce_int(&fail[f]);
183 
184  for (int i=0; i<N; i++) printfQuda("%d fails = %d\n", i, iter[i]);
185 
186  int accuracy_level =0;
187  for (int f=0; f<fail_check; f++) {
188  if (fail[f] == 0) accuracy_level = f+1;
189  }
190 
191  size_t total = u.Nparity()*u.VolumeCB()*N*comm_size();
192  for (int f=0; f<fail_check; f++) {
193  printfQuda("%e Failures: %d / %lu = %e\n", pow(10.0,-(f+1)/(double)tol),
194  fail[f], total, fail[f] / (double)total);
195  }
196 
197  delete []iter;
198  delete []fail;
199 
200  return accuracy_level;
201  }
202 
203  template <typename oFloat, typename iFloat, QudaFieldOrder order>
205  int ret = 0;
206  if (a.Ncolor() == 3) {
207  const int Nc = 3;
208  if (a.Nspin() == 4) {
209  const int Ns = 4;
212  ret = compareSpinor(A, B, tol);
213  } else if (a.Nspin() == 1) {
214  const int Ns = 1;
217  ret = compareSpinor(A, B, tol);
218  }
219  } else {
220  errorQuda("Number of colors %d not supported", a.Ncolor());
221  }
222  return ret;
223  }
224 
225 
226  template <typename oFloat, typename iFloat>
227  int genericCompare(const cpuColorSpinorField &a, const cpuColorSpinorField &b, int tol) {
228  int ret = 0;
229  if (a.FieldOrder() == QUDA_SPACE_SPIN_COLOR_FIELD_ORDER &&
230  a.FieldOrder() == QUDA_SPACE_SPIN_COLOR_FIELD_ORDER) {
231  ret = genericCompare<oFloat,iFloat,QUDA_SPACE_SPIN_COLOR_FIELD_ORDER>(a, b, tol);
232  } else {
233  errorQuda("Unsupported field order %d\n", a.FieldOrder());
234  }
235  return ret;
236  }
237 
238 
239  template <typename oFloat>
240  int genericCompare(const cpuColorSpinorField &a, const cpuColorSpinorField &b, int tol) {
241  int ret = 0;
242  if (b.Precision() == QUDA_DOUBLE_PRECISION) {
243  ret = genericCompare<oFloat,double>(a, b, tol);
244  } else if (b.Precision() == QUDA_SINGLE_PRECISION) {
245  ret = genericCompare<oFloat,float>(a, b, tol);
246  } else {
247  errorQuda("Precision not supported");
248  }
249  return ret;
250  }
251 
252 
253  int genericCompare(const cpuColorSpinorField &a, const cpuColorSpinorField &b, int tol) {
254  int ret = 0;
255  if (a.Precision() == QUDA_DOUBLE_PRECISION) {
256  ret = genericCompare<double>(a, b, tol);
257  } else if (a.Precision() == QUDA_SINGLE_PRECISION) {
258  ret = genericCompare<float>(a, b, tol);
259  } else {
260  errorQuda("Precision not supported");
261  }
262  return ret;
263  }
264 
265 
266  template <class Order>
267  void print_vector(const Order &o, unsigned int x) {
268 
269  int x_cb = x / o.Nparity();
270  int parity = x%o.Nparity();
271 
272  for (int s=0; s<o.Nspin(); s++) {
273  std::cout << "x = " << x << ", s = " << s << ", { ";
274  for (int c=0; c<o.Ncolor(); c++) {
275  std::cout << o(parity, x_cb, s, c) ;
276  std::cout << ((c<o.Ncolor()-1) ? " , " : " " ) ;
277  }
278  std::cout << "}" << std::endl;
279  }
280 
281  }
282 
283  // print out the vector at volume point x
284  template <typename Float, QudaFieldOrder order>
285  void genericPrintVector(cpuColorSpinorField &a, unsigned int x) {
286  if (a.Ncolor() == 3 && a.Nspin() == 4) {
288  print_vector(A, x);
289  }
290  else if (a.Ncolor() == 2 && a.Nspin() == 2) {
292  print_vector(A, x);
293  }
294  else if (a.Ncolor() == 24 && a.Nspin() == 2) {
296  print_vector(A, x);
297  }
298  else if (a.Ncolor() == 6 && a.Nspin() == 4) {
300  print_vector(A, x);
301  }
302  else if (a.Ncolor() == 72 && a.Nspin() == 4) {
304  print_vector(A, x);
305  }
306  else if (a.Ncolor() == 576 && a.Nspin() == 2) {
308  print_vector(A, x);
309  }
310  else {
311  errorQuda("Not supported Ncolor = %d, Nspin = %d", a.Ncolor(), a.Nspin());
312  }
313  }
314 
315  // print out the vector at volume point x
316  template <typename Float>
317  void genericPrintVector(cpuColorSpinorField &a, unsigned int x) {
318  if (a.FieldOrder() == QUDA_SPACE_SPIN_COLOR_FIELD_ORDER) {
319  genericPrintVector<Float,QUDA_SPACE_SPIN_COLOR_FIELD_ORDER>(a,x);
320  } else {
321  errorQuda("Unsupported field order %d\n", a.FieldOrder());
322  }
323  }
324 
325  // print out the vector at volume point x
326  void genericPrintVector(cpuColorSpinorField &a, unsigned int x) {
327  if (a.Precision() == QUDA_DOUBLE_PRECISION) {
328  genericPrintVector<double>(a,x);
329  } else if (a.Precision() == QUDA_SINGLE_PRECISION) {
330  genericPrintVector<float>(a,x);
331  } else {
332  errorQuda("Precision %d not implemented", a.Precision());
333  }
334  }
335 
336 } // namespace quda
int genericCompare(const cpuColorSpinorField &a, const cpuColorSpinorField &b, int tol)
void print_vector(const Order &o, unsigned int x)
#define errorQuda(...)
Definition: util_quda.h:90
static __inline__ enum cudaRoundMode mode
int compareSpinor(const U &u, const V &v, const int tol)
size_t * total
size_t size_t offset
enum QudaSourceType_s QudaSourceType
#define b
int comm_size(void)
Definition: comm_mpi.cpp:126
double tol
Definition: test_util.cpp:1647
__host__ __device__ ValueType sin(ValueType x)
Definition: complex_quda.h:40
static __inline__ size_t p
void random(T &t)
int V
Definition: test_util.cpp:28
__host__ __device__ ValueType pow(ValueType x, ExponentType e)
Definition: complex_quda.h:100
int int int enum cudaChannelFormatKind f
double comm_drand(void)
Definition: comm_common.cpp:82
void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c)
void genericPrintVector(cpuColorSpinorField &a, unsigned int x)
void point(T &t, int x, int s, int c)
#define printfQuda(...)
Definition: util_quda.h:84
double fabs(double)
void comm_allreduce_int(int *data)
Definition: comm_mpi.cpp:305
const void * c
static __inline__ size_t size_t d
QudaParity parity
Definition: covdev_test.cpp:53
#define a
void constant(T &t, int k, int s, int c)
static __device__ __host__ void getCoords(int x[], int cb_index, const I X[], int parity)