QUDA  v0.5.0
A library for QCD on GPUs
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
Classes | Namespaces | Macros | Enumerations | Functions | Variables
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 <blas_quda.h>
#include <face_quda.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 <tm_ndeg_dslash_def.h>
#include "misc_helpers.cu"

Go to the source code of this file.

Classes

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

Namespaces

namespace  quda
 

Macros

#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 NDEGTM_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__)
 
#define MORE_GENERIC_ASYM_DSLASH(FUNC, DAG, X, kernel_type, gridDim, blockDim, shared, stream, param,...)
 
#define GENERIC_ASYM_DSLASH(FUNC, DAG, X, gridDim, blockDim, shared, stream, param,...)
 
#define ASYM_DSLASH(FUNC, gridDim, blockDim, shared, stream, param,...)
 

Enumerations

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

Functions

void quda::setDslashTuning (QudaTune tune, QudaVerbosity verbose)
 
void quda::setKernelPackT (bool pack)
 
bool quda::getKernelPackT ()
 
void quda::setFace (const FaceBuffer &face)
 
void quda::createDslashEvents ()
 
void quda::destroyDslashEvents ()
 
void quda::setTwistParam (double &a, double &b, const double &kappa, const double &mu, const int dagger, const QudaTwistGamma5Type twist)
 
void quda::initDslashCommsPattern ()
 
void quda::dslashCuda (DslashCuda &dslash, const size_t regSize, const int parity, const int dagger, const int volume, const int *faceVolumeCB)
 
void quda::wilsonDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int oddBit, const int daggerBit, const cudaColorSpinorField *x, const double &k, const int *commDim)
 
void quda::cloverDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover cloverInv, const cudaColorSpinorField *in, const int oddBit, const int daggerBit, const cudaColorSpinorField *x, const double &k, const int *commDim)
 
void quda::asymCloverDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover cloverInv, const cudaColorSpinorField *in, const int oddBit, const int daggerBit, const cudaColorSpinorField *x, const double &k, const int *commDim)
 
void quda::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 &epsilon, const int *commDim)
 ndeg tm:
 
void quda::domainWallDslashCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const double &m_f, const double &k, const int *commDim)
 
void quda::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 *commDim)
 NEW:extra argument.
 
void quda::cloverCuda (cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover clover, const cudaColorSpinorField *in, const int oddBit)
 
void quda::twistGamma5Cuda (cudaColorSpinorField *out, const cudaColorSpinorField *in, const int dagger, const double &kappa, const double &mu, const double &epsilon, const QudaTwistGamma5Type)
 ndeg tm:
 

Variables

DslashParam quda::dslashParam
 
int quda::Vspatial
 
int quda::gatherCompleted [Nstream]
 
int quda::previousDir [Nstream]
 
int quda::commsCompleted [Nstream]
 
int quda::dslashCompleted [Nstream]
 
int quda::commDimTotal
 

Macro Definition Documentation

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

Definition at line 402 of file dslash_quda.cu.

#define CLOVER_SHARED_FLOATS_PER_THREAD   0

Definition at line 207 of file dslash_quda.cu.

#define CUDA_EVENT_RECORD (   a,
  b 
)

Definition at line 124 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 346 of file dslash_quda.cu.

#define DSLASH_SHARED_FLOATS_PER_THREAD   0

Definition at line 203 of file dslash_quda.cu.

#define DSLASH_TIME_PROFILE ( )

Definition at line 125 of file dslash_quda.cu.

#define GENERIC_ASYM_DSLASH (   FUNC,
  DAG,
  X,
  gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)
Value:
switch(param.kernel_type) { \
case INTERIOR_KERNEL: \
MORE_GENERIC_ASYM_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 369 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 313 of file dslash_quda.cu.

#define MORE_GENERIC_ASYM_DSLASH (   FUNC,
  DAG,
  X,
  kernel_type,
  gridDim,
  blockDim,
  shared,
  stream,
  param,
  ... 
)
Value:
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 358 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 292 of file dslash_quda.cu.

#define NDEGTM_SHARED_FLOATS_PER_THREAD   0

Definition at line 211 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 354 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 53 of file dslash_quda.cu.

Generated on Wed Mar 20 2013 12:52:19 for QUDA by   doxygen 1.8.2