QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
gauge_qcharge.cu
Go to the documentation of this file.
1 #include <quda_internal.h>
2 #include <tune_quda.h>
3 #include <gauge_field.h>
4 
5 #include <launch_kernel.cuh>
6 #include <jitify_helper.cuh>
8 
9 namespace quda
10 {
11 
12 #ifdef GPU_GAUGE_TOOLS
13 
14  template <typename Float, typename Arg> class QChargeCompute : TunableLocalParity
15  {
16  Arg &arg;
17  const GaugeField &meta;
18 
19 private:
20  bool tuneGridDim() const { return true; }
21  unsigned int minThreads() const { return arg.threads; }
22 
23 public:
24  QChargeCompute(Arg &arg, const GaugeField &meta) : arg(arg), meta(meta)
25  {
26 #ifdef JITIFY
27  create_jitify_program("kernels/gauge_qcharge.cuh");
28 #endif
29  }
30  virtual ~QChargeCompute() {}
31 
32  void apply(const cudaStream_t &stream)
33  {
34  if (meta.Location() == QUDA_CUDA_FIELD_LOCATION) {
35  arg.result_h[0] = 0.;
36  TuneParam tp = tuneLaunch(*this, getTuning(), getVerbosity());
37 #ifdef JITIFY
38  using namespace jitify::reflection;
39  jitify_error = program->kernel("quda::qChargeComputeKernel")
40  .instantiate((int)tp.block.x, Type<Float>(), Type<Arg>())
41  .configure(tp.grid, tp.block, tp.shared_bytes, stream)
42  .launch(arg);
43 #else
44  LAUNCH_KERNEL(qChargeComputeKernel, tp, stream, arg, Float);
45 #endif
47  } else { // run the CPU code
48  errorQuda("qChargeComputeKernel not supported on CPU");
49  }
50  }
51 
52  TuneKey tuneKey() const
53  {
54  std::stringstream aux;
55  aux << "threads=" << arg.threads << ",prec=" << sizeof(Float);
56  return TuneKey(meta.VolString(), typeid(*this).name(), aux.str().c_str());
57  }
58 
59  long long flops() const { return 2 * arg.threads * (3 * 198 + 9); }
60  long long bytes() const { return 2 * arg.threads * ((6 * 18) + Arg::density) * sizeof(Float); }
61  }; // QChargeCompute
62 
63  template <typename Float, typename Gauge, bool density>
64  void computeQCharge(const Gauge data, const GaugeField &Fmunu, Float *qDensity, Float &qChg)
65  {
66  QChargeArg<Float, Gauge, density> arg(data, Fmunu, qDensity);
67  QChargeCompute<Float, decltype(arg)> qChargeCompute(arg, Fmunu);
68  qChargeCompute.apply(0);
70  comm_allreduce((double *)arg.result_h);
71  qChg = arg.result_h[0];
72  }
73 
74  template <typename Float, bool density> Float computeQCharge(const GaugeField &Fmunu, Float *qDensity = nullptr)
75  {
76  Float qChg = 0.0;
77 
78  if (!Fmunu.isNative()) errorQuda("Topological charge computation only supported on native ordered fields");
79 
80  if (Fmunu.Reconstruct() == QUDA_RECONSTRUCT_NO) {
82  computeQCharge<Float, Gauge, density>(Gauge(Fmunu), Fmunu, qDensity, qChg);
83  } else if (Fmunu.Reconstruct() == QUDA_RECONSTRUCT_12) {
85  computeQCharge<Float, Gauge, density>(Gauge(Fmunu), Fmunu, qDensity, qChg);
86  } else if (Fmunu.Reconstruct() == QUDA_RECONSTRUCT_8) {
87  typedef typename gauge_mapper<Float, QUDA_RECONSTRUCT_8>::type Gauge;
88  computeQCharge<Float, Gauge, density>(Gauge(Fmunu), Fmunu, qDensity, qChg);
89  } else {
90  errorQuda("Reconstruction type %d of gauge field not supported", Fmunu.Reconstruct());
91  }
92 
93  return qChg;
94  }
95 #endif // GPU_GAUGE_TOOLS
96 
97  double computeQCharge(const GaugeField &Fmunu)
98  {
99  double qChg = 0.0;
100 #ifdef GPU_GAUGE_TOOLS
101  if (!Fmunu.isNative()) errorQuda("Order %d with %d reconstruct not supported", Fmunu.Order(), Fmunu.Reconstruct());
102 
103  if (Fmunu.Precision() == QUDA_SINGLE_PRECISION) {
104  qChg = computeQCharge<float, false>(Fmunu);
105  } else if (Fmunu.Precision() == QUDA_DOUBLE_PRECISION) {
106  qChg = computeQCharge<double, false>(Fmunu);
107  } else {
108  errorQuda("Precision %d not supported", Fmunu.Precision());
109  }
110 #else
111  errorQuda("Gauge tools are not built");
112 #endif // GPU_GAUGE_TOOLS
113  return qChg;
114  }
115 
116  double computeQChargeDensity(const GaugeField &Fmunu, void *qDensity)
117  {
118  double qChg = 0.0;
119 #ifdef GPU_GAUGE_TOOLS
120  if (!Fmunu.isNative()) errorQuda("Order %d with %d reconstruct not supported", Fmunu.Order(), Fmunu.Reconstruct());
121 
122  if (Fmunu.Precision() == QUDA_SINGLE_PRECISION) {
123  qChg = computeQCharge<float, true>(Fmunu, (float *)qDensity);
124  } else if (Fmunu.Precision() == QUDA_DOUBLE_PRECISION) {
125  qChg = computeQCharge<double, true>(Fmunu, (double *)qDensity);
126  } else {
127  errorQuda("Precision %d not supported", Fmunu.Precision());
128  }
129 #else
130  errorQuda("Gauge tools are not built");
131 #endif // GPU_GAUGE_TOOLS
132  return qChg;
133  }
134 } // namespace quda
double computeQCharge(const GaugeField &Fmunu)
Compute the topological charge.
QudaVerbosity getVerbosity()
Definition: util_quda.cpp:21
#define errorQuda(...)
Definition: util_quda.h:121
Helper file when using jitify run-time compilation. This file should be included in source code...
cudaStream_t * stream
const char * VolString() const
double computeQChargeDensity(const GaugeField &Fmunu, void *result)
Compute the topological charge density per lattice site.
#define qudaDeviceSynchronize()
__global__ void qChargeComputeKernel(Arg arg)
TuneParam & tuneLaunch(Tunable &tunable, QudaTune enabled, QudaVerbosity verbosity)
Definition: tune.cpp:643
#define LAUNCH_KERNEL(kernel, tp, stream, arg,...)
QudaFieldLocation Location() const
unsigned long long flops
Definition: blas_quda.cu:22
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
QudaReconstructType Reconstruct() const
Definition: gauge_field.h:250
QudaGaugeFieldOrder Order() const
Definition: gauge_field.h:251
#define checkCudaError()
Definition: util_quda.h:161
void comm_allreduce(double *data)
Definition: comm_mpi.cpp:242
QudaTune getTuning()
Query whether autotuning is enabled or not. Default is enabled but can be overridden by setting QUDA_...
Definition: util_quda.cpp:52
QudaPrecision Precision() const
bool isNative() const
unsigned long long bytes
Definition: blas_quda.cu:23