QUDA  v1.1.0
A library for QCD on GPUs
dslash_quda.h
Go to the documentation of this file.
1 #ifndef _DSLASH_QUDA_H
2 #define _DSLASH_QUDA_H
3 
4 #include <quda_internal.h>
5 #include <tune_quda.h>
6 #include <color_spinor_field.h>
7 #include <gauge_field.h>
8 #include <clover_field.h>
9 #include <worker.h>
10 #ifdef NVSHMEM_COMMS
11 #include <cuda/atomic>
12 #endif
13 
14 namespace quda {
15 
19  void setKernelPackT(bool pack);
20 
25 
26  void pushKernelPackT(bool pack);
28 
35  void setPackComms(const int *dim_pack);
36 
38 
41 
42  namespace dslash
43  {
47  using shmem_sync_t = uint64_t;
48 
55 
63 
70 #ifdef NVSHMEM_COMMS
71  using shmem_retcount_intra_t = cuda::atomic<int, cuda::thread_scope_system>;
72  using shmem_retcount_inter_t = cuda::atomic<int, cuda::thread_scope_device>;
73  using shmem_interior_done_t = cuda::atomic<shmem_sync_t, cuda::thread_scope_device>;
74  using shmem_interior_count_t = cuda::atomic<int, cuda::thread_scope_block>;
75 
81  shmem_sync_t *get_shmem_sync_arr();
82 
88  shmem_retcount_intra_t *get_shmem_retcount_intra();
89 
95  shmem_retcount_inter_t *get_shmem_retcount_inter();
96 
102  shmem_interior_done_t *get_shmem_interior_done();
103 
109  shmem_interior_count_t *get_shmem_interior_count();
110 #endif
111  } // namespace dslash
112 
134  void ApplyWilson(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double kappa,
135  const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
136 
160  double kappa, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
161 
186  const CloverField &A, double kappa, double mu, const ColorSpinorField &x,
187  int parity, bool dagger, const int *comm_override, TimeProfile &profile);
188 
227  const CloverField &A, double kappa, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override,
228  TimeProfile &profile);
229 
276  const CloverField &A, double kappa, double mu,
277  const ColorSpinorField &x, int parity, bool dagger,
278  const int *comm_override, TimeProfile &profile);
279 
304  const GaugeField &U, const CloverField &A, double kappa, double mu,
305  const ColorSpinorField &x, int parity, bool dagger,
306  const int *comm_override, TimeProfile &profile);
307 
308  // old
309  void ApplyTwistedMass(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b,
310  const ColorSpinorField &x, int parity, bool dagger, const int *comm_override,
311  TimeProfile &profile);
312 
350  double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric,
351  const int *comm_override, TimeProfile &profile);
352 
380  void ApplyNdegTwistedMass(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b,
381  double c, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
382 
429  double a, double b, double c, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric,
430  const int *comm_override, TimeProfile &profile);
431 
456  double a, double b, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override,
457  TimeProfile &profile);
458 
497  const CloverField &C, double a, double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger,
498  const int *comm_override, TimeProfile &profile);
499 
524  void ApplyDomainWall5D(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double m_f,
525  const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
526 
553  void ApplyDomainWall4D(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double m_5,
554  const Complex *b_5, const Complex *c_5, const ColorSpinorField &x, int parity, bool dagger,
555  const int *comm_override, TimeProfile &profile);
556 
557  enum Dslash5Type {
565  M5INV_EOFA
566  };
567 
574  enum class MdwfFusedDslashType {
579  D5PRE,
580  };
581 
597  void ApplyDslash5(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, double m_f,
598  double m_5, const Complex *b_5, const Complex *c_5, double a, bool dagger, Dslash5Type type);
599 
600  // Tensor core functions for Mobius DWF
601  namespace mobius_tensor_core
602  {
604  const ColorSpinorField &x, double m_f, double m_5, const Complex *b_5, const Complex *c_5,
605  bool dagger, int parity, int shift[4], int halo_shift[4], MdwfFusedDslashType type);
606  }
607 
608  // The EOFA stuff
609  namespace mobius_eofa
610  {
611  void apply_dslash5(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, double m_f,
612  double m_5, const Complex *b_5, const Complex *c_5, double a, int eofa_pm, double inv,
613  double kappa, const double *eofa_u, const double *eofa_x, const double *eofa_y,
614  double sherman_morrison, bool dagger, Dslash5Type type);
615  }
616 
636  void ApplyLaplace(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, int dir, double a, double b,
637  const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
638 
658  void ApplyCovDev(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, int mu, int parity,
659  bool dagger, const int *comm_override, TimeProfile &profile);
660 
670  ColorSpinorField &out, const ColorSpinorField &in, const CloverField &clover, bool inverse, int parity);
671 
683  void ApplyStaggered(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a,
684  const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile);
685 
699  const GaugeField &L, double a, const ColorSpinorField &x, int parity, bool dagger,
700  const int *comm_override, TimeProfile &profile);
701 
710  bool dagger);
711 
723  void ApplyTwistGamma(ColorSpinorField &out, const ColorSpinorField &in, int d, double kappa, double mu,
724  double epsilon, int dagger, QudaTwistGamma5Type type);
725 
740  void ApplyTwistClover(ColorSpinorField &out, const ColorSpinorField &in, const CloverField &clover,
741  double kappa, double mu, double epsilon, int parity, int dagger, QudaTwistGamma5Type twist);
742 
757  void PackGhost(void *ghost[2 * QUDA_MAX_DIM], const ColorSpinorField &field, MemoryLocation location, int nFace,
758  bool dagger, int parity, bool spin_project, double a, double b, double c, int shmem,
759  const qudaStream_t &stream);
760 
766  void gamma5(ColorSpinorField &out, const ColorSpinorField &in);
767 
768 }
769 
770 #endif // _DSLASH_QUDA_H
double kappa
double epsilon
double mu
int eofa_pm
bool dagger
QudaParity parity
Definition: covdev_test.cpp:40
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
Definition: blas_quda.h:45
shmem_sync_t set_shmem_sync_counter(shmem_sync_t count)
Set the shmem sync counter to count.
uint64_t shmem_sync_t
type used for shmem signaling
Definition: dslash_quda.h:47
shmem_sync_t inc_shmem_sync_counter()
increase the shmem sync counter for the next dslash application
shmem_sync_t get_shmem_sync_counter()
Get the shmem sync counter.
void apply_dslash5(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, double m_f, double m_5, const Complex *b_5, const Complex *c_5, double a, int eofa_pm, double inv, double kappa, const double *eofa_u, const double *eofa_x, const double *eofa_y, double sherman_morrison, bool dagger, Dslash5Type type)
void apply_fused_dslash(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, ColorSpinorField &y, const ColorSpinorField &x, double m_f, double m_5, const Complex *b_5, const Complex *c_5, bool dagger, int parity, int shift[4], int halo_shift[4], MdwfFusedDslashType type)
void createDslashEvents()
void setKernelPackT(bool pack)
void ApplyStaggered(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Apply the staggered dslash operator to a color-spinor field.
void ApplyClover(ColorSpinorField &out, const ColorSpinorField &in, const CloverField &clover, bool inverse, int parity)
Apply clover-matrix field to a color-spinor field.
bool getDslashLaunch()
void ApplyTwistedCloverPreconditioned(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &C, double a, double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the preconditioned twisted-clover stencil.
void setPackComms(const int *dim_pack)
Helper function that sets which dimensions the packing kernel should be packing for.
void ApplyWilsonCloverHasenbuschTwistPCClovInv(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &A, double kappa, double mu, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the twisted-mass stencil.
void PackGhost(void *ghost[2 *QUDA_MAX_DIM], const ColorSpinorField &field, MemoryLocation location, int nFace, bool dagger, int parity, bool spin_project, double a, double b, double c, int shmem, const qudaStream_t &stream)
Dslash face packing routine.
void destroyDslashEvents()
MdwfFusedDslashType
Definition: dslash_quda.h:574
__device__ __host__ Matrix< T, 3 > inverse(const Matrix< T, 3 > &u)
Definition: quda_matrix.h:605
std::complex< double > Complex
Definition: quda_internal.h:86
void ApplyTwistClover(ColorSpinorField &out, const ColorSpinorField &in, const CloverField &clover, double kappa, double mu, double epsilon, int parity, int dagger, QudaTwistGamma5Type twist)
Apply twisted clover-matrix field to a color-spinor field.
void ApplyCovDev(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, int mu, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the covariant derivative.
void ApplyNdegTwistedMass(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, double c, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the non-degenerate twisted-mass stencil.
void ApplyTwistGamma(ColorSpinorField &out, const ColorSpinorField &in, int d, double kappa, double mu, double epsilon, int dagger, QudaTwistGamma5Type type)
Apply the twisted-mass gamma operator to a color-spinor field.
void ApplyLaplace(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, int dir, double a, double b, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Laplace stencil.
qudaStream_t * stream
void ApplyWilsonClover(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &A, double kappa, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Wilson-clover stencil.
void ApplyTwistedMass(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
void ApplyWilsonCloverPreconditioned(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &A, double kappa, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the preconditioned Wilson-clover stencil.
void ApplyStaggeredKahlerDiracInverse(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &Xinv, bool dagger)
Apply the (improved) staggered Kahler-Dirac inverse block to a color-spinor field.
void ApplyTwistedMassPreconditioned(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric, const int *comm_override, TimeProfile &profile)
Driver for applying the preconditioned twisted-mass stencil.
void ApplyTwistedClover(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &C, double a, double b, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the twisted-clover stencil.
bool getKernelPackT()
void pushKernelPackT(bool pack)
void ApplyDomainWall4D(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double m_5, const Complex *b_5, const Complex *c_5, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the batched Wilson 4-d stencil to a 5-d vector with 4-d preconditioned data order...
void popKernelPackT()
void gamma5(ColorSpinorField &out, const ColorSpinorField &in)
Applies a gamma5 matrix to a spinor (wrapper to ApplyGamma)
void ApplyWilsonCloverHasenbuschTwistPCNoClovInv(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &A, double kappa, double mu, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Wilson-clover stencil with thist for Hasenbusch.
void ApplyImprovedStaggered(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const GaugeField &L, double a, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Apply the improved staggered dslash operator to a color-spinor field.
void ApplyWilson(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double kappa, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Wilson stencil.
void ApplyNdegTwistedMassPreconditioned(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, double c, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric, const int *comm_override, TimeProfile &profile)
Driver for applying the preconditioned non-degenerate twisted-mass stencil.
Dslash5Type
Definition: dslash_quda.h:557
@ DSLASH5_DWF
Definition: dslash_quda.h:558
@ DSLASH5_MOBIUS_PRE
Definition: dslash_quda.h:559
@ DSLASH5_MOBIUS
Definition: dslash_quda.h:560
@ M5_INV_ZMOBIUS
Definition: dslash_quda.h:563
@ M5INV_EOFA
Definition: dslash_quda.h:565
@ M5_EOFA
Definition: dslash_quda.h:564
@ M5_INV_MOBIUS
Definition: dslash_quda.h:562
@ M5_INV_DWF
Definition: dslash_quda.h:561
void ApplyWilsonCloverHasenbuschTwist(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, const CloverField &A, double kappa, double mu, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Wilson-clover stencil.
void ApplyDslash5(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, double m_f, double m_5, const Complex *b_5, const Complex *c_5, double a, bool dagger, Dslash5Type type)
Apply either the domain-wall / mobius Dslash5 operator or the M5 inverse operator....
void ApplyDomainWall5D(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double m_f, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the Domain-wall 5-d stencil to a 5-d vector with 5-d preconditioned data order.
cudaStream_t qudaStream_t
Definition: quda_api.h:9
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5.