QUDA v0.4.0
A library for QCD on GPUs
Classes | Defines | Enumerations | Functions | Variables
quda/lib/dslash_quda.cu File Reference
#include <cstdlib>
#include <cstdio>
#include <string>
#include <iostream>
#include <color_spinor_field.h>
#include <clover_field.h>
#include <quda_internal.h>
#include <dslash_quda.h>
#include <sys/time.h>
#include <inline_ptx.h>
#include <dslash_textures.h>
#include <dslash_constants.h>
#include <pack_face_def.h>
#include <staggered_dslash_def.h>
#include <wilson_dslash_def.h>
#include <dw_dslash_def.h>
#include <tm_dslash_def.h>
#include <tm_core.h>
#include <clover_def.h>
#include <blas_quda.h>
#include <face_quda.h>
#include "misc_helpers.cu"

Go to the source code of this file.

Classes

struct  DslashParam
class  DslashCuda
class  SharedDslashCuda
class  WilsonDslashCuda< sFloat, gFloat >
class  CloverDslashCuda< sFloat, gFloat, cFloat >
class  TwistedDslashCuda< sFloat, gFloat >
class  DomainWallDslashCuda< sFloat, gFloat >
class  StaggeredDslashCuda< sFloat, fatGFloat, longGFloat >
class  CloverCuda< sFloat, cFloat >
class  TwistGamma5Cuda< sFloat >

Defines

#define BLOCK_DIM   64
#define DIRECT_ACCESS_FAT_LINK
#define CUDA_EVENT_RECORD(a, b)
#define DSLASH_TIME_PROFILE()
#define DSLASH_SHARED_FLOATS_PER_THREAD   0
#define CLOVER_SHARED_FLOATS_PER_THREAD   0
#define MORE_GENERIC_DSLASH(FUNC, DAG, X, kernel_type, gridDim, blockDim, shared, stream, param,...)
#define GENERIC_DSLASH(FUNC, DAG, X, gridDim, blockDim, shared, stream, param,...)
#define DSLASH(FUNC, gridDim, blockDim, shared, stream, param,...)
#define STAGGERED_DSLASH(gridDim, blockDim, shared, stream, param,...)   GENERIC_DSLASH(staggeredDslash, , Axpy, gridDim, blockDim, shared, stream, param, __VA_ARGS__)

Enumerations

enum  KernelType {
  INTERIOR_KERNEL = 5, EXTERIOR_KERNEL_X = 0, EXTERIOR_KERNEL_Y = 1, EXTERIOR_KERNEL_Z = 2,
  EXTERIOR_KERNEL_T = 3
}

Functions

void setDslashTuning (QudaTune tune, QudaVerbosity verbose)
__global__ void dummyKernel ()
void initCache ()
void setFace (const FaceBuffer &Face)
void setTwistParam (double &a, double &b, const double &kappa, const double &mu, const int dagger, const QudaTwistGamma5Type twist)
void initDslashCommsPattern ()
void dslashCuda (DslashCuda &dslash, const size_t regSize, const int parity, const int dagger, const int volume, const int *faceVolumeCB)
void wilsonDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &k, const int *commOverride)
void cloverDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover cloverInv, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &a, const int *commOverride)
void twistedMassDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &kappa, const double &mu, const double &a, const int *commOverride)
void domainWallDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &m_f, const double &k2)
void staggeredDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &fatGauge, const cudaGaugeField &longGauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &k, const int *commOverride)
void cloverCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover clover, const cudaColorSpinorField *in, const int parity)
void twistGamma5Cuda (cudaColorSpinorField *out, const cudaColorSpinorField *in, const int dagger, const double &kappa, const double &mu, const QudaTwistGamma5Type twist)

Variables

bool kernelPackT = false
DslashParam dslashParam
int Vspatial
FaceBuffer * face
cudaColorSpinorField * inSpinor
int gatherCompleted [Nstream]
int previousDir [Nstream]
int commsCompleted [Nstream]
int commDimTotal

Define Documentation

#define BLOCK_DIM   64

Definition at line 9 of file dslash_quda.cu.

#define CLOVER_SHARED_FLOATS_PER_THREAD   0

Definition at line 174 of file dslash_quda.cu.

#define CUDA_EVENT_RECORD (   a,
  b 
)

Definition at line 99 of file dslash_quda.cu.

#define DIRECT_ACCESS_FAT_LINK

Definition at line 28 of file dslash_quda.cu.

#define DSLASH (   FUNC,
  gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)
Value:
if (!dagger) {                                                  \
    GENERIC_DSLASH(FUNC, , Xpay, gridDim, blockDim, shared, stream, param, __VA_ARGS__) \
  } else {                                                              \
    GENERIC_DSLASH(FUNC, Dagger, Xpay, gridDim, blockDim, shared, stream, param, __VA_ARGS__) \
 }

Definition at line 258 of file dslash_quda.cu.

#define DSLASH_SHARED_FLOATS_PER_THREAD   0

Definition at line 170 of file dslash_quda.cu.

#define DSLASH_TIME_PROFILE ( )

Definition at line 100 of file dslash_quda.cu.

#define GENERIC_DSLASH (   FUNC,
  DAG,
  X,
  gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)
Value:
switch(param.kernel_type) {                                                                                   \
  case INTERIOR_KERNEL:                                                                                       \
    MORE_GENERIC_DSLASH(FUNC, DAG, X, INTERIOR_KERNEL, gridDim, blockDim, shared, stream, param, __VA_ARGS__) \
    break;                                                                                                    \
  default:                                                                                                    \
    errorQuda("KernelType %d not defined for single GPU", param.kernel_type);                                 \
  }

Definition at line 225 of file dslash_quda.cu.

#define MORE_GENERIC_DSLASH (   FUNC,
  DAG,
  X,
  kernel_type,
  gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)
Value:
if (x==0) {                                                                                                     \
    if (reconstruct == QUDA_RECONSTRUCT_NO) {                                                                     \
      FUNC ## 18 ## DAG ## Kernel<kernel_type><<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__ , param);     \
    } else if (reconstruct == QUDA_RECONSTRUCT_12) {                                                              \
      FUNC ## 12 ## DAG ## Kernel<kernel_type><<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__ , param);     \
    } else {                                                                                                      \
      FUNC ## 8 ## DAG ## Kernel<kernel_type><<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__, param);       \
    }                                                                                                             \
  } else {                                                                                                        \
    if (reconstruct == QUDA_RECONSTRUCT_NO) {                                                                     \
      FUNC ## 18 ## DAG ## X ## Kernel<kernel_type><<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__, param); \
    } else if (reconstruct == QUDA_RECONSTRUCT_12) {                                                              \
      FUNC ## 12 ## DAG ## X ## Kernel<kernel_type><<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__, param); \
    } else if (reconstruct == QUDA_RECONSTRUCT_8) {                                                               \
      FUNC ## 8 ## DAG ## X ## Kernel<kernel_type> <<<gridDim, blockDim, shared, stream>>> ( __VA_ARGS__, param); \
    }                                                                                                             \
  }

Definition at line 204 of file dslash_quda.cu.

#define STAGGERED_DSLASH (   gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)    GENERIC_DSLASH(staggeredDslash, , Axpy, gridDim, blockDim, shared, stream, param, __VA_ARGS__)

Definition at line 266 of file dslash_quda.cu.

Enumeration Type Documentation

enum KernelType
Enumerator:
INTERIOR_KERNEL 
EXTERIOR_KERNEL_X 
EXTERIOR_KERNEL_Y 
EXTERIOR_KERNEL_Z 
EXTERIOR_KERNEL_T 

Definition at line 41 of file dslash_quda.cu.

Function Documentation

void cloverCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  gauge,
const FullClover  clover,
const cudaColorSpinorField *  in,
const int  parity 
)

Definition at line 1551 of file dslash_quda.cu.

void cloverDslashCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  gauge,
const FullClover  cloverInv,
const cudaColorSpinorField *  in,
const int  parity,
const int  dagger,
const cudaColorSpinorField *  x,
const double &  a,
const int *  commOverride 
)

Definition at line 1207 of file dslash_quda.cu.

void domainWallDslashCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  gauge,
const cudaColorSpinorField *  in,
const int  parity,
const int  dagger,
const cudaColorSpinorField *  x,
const double &  m_f,
const double &  k2 
)

Definition at line 1336 of file dslash_quda.cu.

void dslashCuda ( DslashCuda &  dslash,
const size_t  regSize,
const int  parity,
const int  dagger,
const int  volume,
const int *  faceVolumeCB 
)

Definition at line 1035 of file dslash_quda.cu.

__global__ void dummyKernel ( )

Definition at line 181 of file dslash_quda.cu.

void initCache ( )

Definition at line 185 of file dslash_quda.cu.

void initDslashCommsPattern ( )

Initialize the arrays used for the dynamic scheduling.

Definition at line 1007 of file dslash_quda.cu.

void setDslashTuning ( QudaTune  tune,
QudaVerbosity  verbose 
)

Definition at line 119 of file dslash_quda.cu.

void setFace ( const FaceBuffer &  Face)

Definition at line 200 of file dslash_quda.cu.

void setTwistParam ( double &  a,
double &  b,
const double &  kappa,
const double &  mu,
const int  dagger,
const QudaTwistGamma5Type  twist 
)

Definition at line 647 of file dslash_quda.cu.

void staggeredDslashCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  fatGauge,
const cudaGaugeField &  longGauge,
const cudaColorSpinorField *  in,
const int  parity,
const int  dagger,
const cudaColorSpinorField *  x,
const double &  k,
const int *  commOverride 
)

Definition at line 1393 of file dslash_quda.cu.

void twistedMassDslashCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  gauge,
const cudaColorSpinorField *  in,
const int  parity,
const int  dagger,
const cudaColorSpinorField *  x,
const double &  kappa,
const double &  mu,
const double &  a,
const int *  commOverride 
)

Definition at line 1275 of file dslash_quda.cu.

void twistGamma5Cuda ( cudaColorSpinorField *  out,
const cudaColorSpinorField *  in,
const int  dagger,
const double &  kappa,
const double &  mu,
const QudaTwistGamma5Type  twist 
)

Definition at line 1667 of file dslash_quda.cu.

void wilsonDslashCuda ( cudaColorSpinorField *  out,
const cudaGaugeField &  gauge,
const cudaColorSpinorField *  in,
const int  parity,
const int  dagger,
const cudaColorSpinorField *  x,
const double &  k,
const int *  commOverride 
)

Definition at line 1149 of file dslash_quda.cu.

Variable Documentation

int commDimTotal

Definition at line 1002 of file dslash_quda.cu.

int commsCompleted[Nstream]

Definition at line 1001 of file dslash_quda.cu.

DslashParam dslashParam

Definition at line 63 of file dslash_quda.cu.

FaceBuffer* face

Definition at line 103 of file dslash_quda.cu.

int gatherCompleted[Nstream]

Definition at line 999 of file dslash_quda.cu.

cudaColorSpinorField* inSpinor

Definition at line 104 of file dslash_quda.cu.

bool kernelPackT = false

Definition at line 61 of file dslash_quda.cu.

int Ls

Definition at line 110 of file dslash_quda.cu.

int previousDir[Nstream]

Definition at line 1000 of file dslash_quda.cu.

int Vspatial

Definition at line 66 of file dslash_quda.cu.

int x[4]

Definition at line 109 of file dslash_quda.cu.

 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Defines
Generated on Wed Apr 4 2012 21:08:07 for QUDA by  doxygen 1.7.4