QUDA  v1.1.0
A library for QCD on GPUs
checksum.cu
Go to the documentation of this file.
1 #include <gauge_field_order.h>
2 
3 namespace quda {
4 
5  template <typename T, QudaGaugeFieldOrder order, int Nc>
6  struct ChecksumArg {
7  static constexpr int nColor = Nc;
8  typedef typename mapper<T>::type real;
9  typedef typename gauge_order_mapper<T,order,Nc>::type G;
10  const G U;
11  const int volumeCB;
12  ChecksumArg(const GaugeField &U, bool mini) : U(U), volumeCB(mini ? 1 : U.VolumeCB()) { }
13  };
14 
15  template <typename Arg>
16  __device__ __host__ inline uint64_t siteChecksum(const Arg &arg, int d, int parity, int x_cb) {
17  const Matrix<complex<typename Arg::real>,Arg::nColor> u = arg.U(d, x_cb, parity);
18  return u.checksum();
19  }
20 
21  template <typename Arg>
22  uint64_t ChecksumCPU(const Arg &arg)
23  {
24  uint64_t checksum_ = 0;
25  for (int parity=0; parity<2; parity++)
26  for (int x_cb=0; x_cb<arg.volumeCB; x_cb++)
27  for (int d=0; d<arg.U.geometry; d++)
28  checksum_ ^= siteChecksum(arg, d, parity, x_cb);
29  return checksum_;
30  }
31 
32  template <typename T, int Nc>
33  uint64_t Checksum(const GaugeField &u, bool mini)
34  {
35  uint64_t checksum = 0;
36  if (u.Order() == QUDA_QDP_GAUGE_ORDER) {
37  ChecksumArg<T,QUDA_QDP_GAUGE_ORDER,Nc> arg(u,mini);
38  checksum = ChecksumCPU(arg);
39  } else if (u.Order() == QUDA_QDPJIT_GAUGE_ORDER) {
40  ChecksumArg<T,QUDA_QDPJIT_GAUGE_ORDER,Nc> arg(u,mini);
41  checksum = ChecksumCPU(arg);
42  } else if (u.Order() == QUDA_MILC_GAUGE_ORDER) {
43  ChecksumArg<T,QUDA_MILC_GAUGE_ORDER,Nc> arg(u,mini);
44  checksum = ChecksumCPU(arg);
45  } else if (u.Order() == QUDA_BQCD_GAUGE_ORDER) {
46  ChecksumArg<T,QUDA_BQCD_GAUGE_ORDER,Nc> arg(u,mini);
47  checksum = ChecksumCPU(arg);
48  } else if (u.Order() == QUDA_TIFR_GAUGE_ORDER) {
49  ChecksumArg<T,QUDA_TIFR_GAUGE_ORDER,Nc> arg(u,mini);
50  checksum = ChecksumCPU(arg);
51  } else if (u.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
52  ChecksumArg<T,QUDA_TIFR_PADDED_GAUGE_ORDER,Nc> arg(u,mini);
53  checksum = ChecksumCPU(arg);
54  } else {
55  errorQuda("Checksum not implemented");
56  }
57 
58  return checksum;
59  }
60 
61  template <typename T>
62  uint64_t Checksum(const GaugeField &u, bool mini)
63  {
64  uint64_t checksum = 0;
65  switch (u.Ncolor()) {
66  case 3: checksum = Checksum<T,3>(u,mini); break;
67  default: errorQuda("Unsupported nColor = %d", u.Ncolor());
68  }
69  return checksum;
70  }
71 
72  uint64_t Checksum(const GaugeField &u, bool mini)
73  {
74  uint64_t checksum = 0;
75  switch (u.Precision()) {
76  case QUDA_DOUBLE_PRECISION: checksum = Checksum<double>(u,mini); break;
77  case QUDA_SINGLE_PRECISION: checksum = Checksum<float>(u,mini); break;
78  default: errorQuda("Unsupported precision = %d", u.Precision());
79  }
80 
81  comm_allreduce_xor(&checksum);
82 
83  return checksum;
84  }
85 
86 }