18 static const double projector[8][4][4][2] = {
20 {{1,0}, {0,0}, {0,0}, {0,-1}},
21 {{0,0}, {1,0}, {0,-1}, {0,0}},
22 {{0,0}, {0,1}, {1,0}, {0,0}},
23 {{0,1}, {0,0}, {0,0}, {1,0}}
26 {{1,0}, {0,0}, {0,0}, {0,1}},
27 {{0,0}, {1,0}, {0,1}, {0,0}},
28 {{0,0}, {0,-1}, {1,0}, {0,0}},
29 {{0,-1}, {0,0}, {0,0}, {1,0}}
32 {{1,0}, {0,0}, {0,0}, {1,0}},
33 {{0,0}, {1,0}, {-1,0}, {0,0}},
34 {{0,0}, {-1,0}, {1,0}, {0,0}},
35 {{1,0}, {0,0}, {0,0}, {1,0}}
38 {{1,0}, {0,0}, {0,0}, {-1,0}},
39 {{0,0}, {1,0}, {1,0}, {0,0}},
40 {{0,0}, {1,0}, {1,0}, {0,0}},
41 {{-1,0}, {0,0}, {0,0}, {1,0}}
44 {{1,0}, {0,0}, {0,-1}, {0,0}},
45 {{0,0}, {1,0}, {0,0}, {0,1}},
46 {{0,1}, {0,0}, {1,0}, {0,0}},
47 {{0,0}, {0,-1}, {0,0}, {1,0}}
50 {{1,0}, {0,0}, {0,1}, {0,0}},
51 {{0,0}, {1,0}, {0,0}, {0,-1}},
52 {{0,-1}, {0,0}, {1,0}, {0,0}},
53 {{0,0}, {0,1}, {0,0}, {1,0}}
56 {{1,0}, {0,0}, {-1,0}, {0,0}},
57 {{0,0}, {1,0}, {0,0}, {-1,0}},
58 {{-1,0}, {0,0}, {1,0}, {0,0}},
59 {{0,0}, {-1,0}, {0,0}, {1,0}}
62 {{1,0}, {0,0}, {1,0}, {0,0}},
63 {{0,0}, {1,0}, {0,0}, {1,0}},
64 {{1,0}, {0,0}, {1,0}, {0,0}},
65 {{0,0}, {1,0}, {0,0}, {1,0}}
71 template <
typename Float>
73 for (
int i=0; i<4*3*2; i++) res[i] = 0.0;
75 for (
int s = 0; s < 4; s++) {
76 for (
int t = 0; t < 4; t++) {
80 for (
int m = 0; m < 3; m++) {
81 Float spinorRe = spinorIn[t*(3*2) + m*(2) + 0];
82 Float spinorIm = spinorIn[t*(3*2) + m*(2) + 1];
83 res[s*(3*2) + m*(2) + 0] += projRe*spinorRe - projIm*spinorIm;
84 res[s*(3*2) + m*(2) + 1] += projRe*spinorIm + projIm*spinorRe;
103 template <
typename sFloat,
typename gFloat>
104 void dslashReference(sFloat *res, gFloat **gaugeFull, sFloat *spinorField,
int oddBit,
int daggerBit)
108 gFloat *gaugeEven[4], *gaugeOdd[4];
109 for (
int dir = 0; dir < 4; dir++) {
110 gaugeEven[dir] = gaugeFull[dir];
114 for (
int i = 0; i <
Vh; i++) {
115 for (
int dir = 0; dir < 8; dir++) {
116 gFloat *gauge = gaugeLink(i, dir, oddBit, gaugeEven, gaugeOdd, 1);
117 sFloat *
spinor = spinorNeighbor(i, dir, oddBit, spinorField, 1);
120 int projIdx = 2*(dir/2)+(dir+daggerBit)%2;
123 for (
int s = 0; s < 4; s++) {
124 if (dir % 2 == 0) su3Mul(&gaugedSpinor[s*(3*2)], gauge, &projectedSpinor[s*(3*2)]);
125 else su3Tmul(&gaugedSpinor[s*(3*2)], gauge, &projectedSpinor[s*(3*2)]);
135 template <
typename sFloat,
typename gFloat>
136 void dslashReference(sFloat *res, gFloat **gaugeFull, gFloat **ghostGauge, sFloat *spinorField, sFloat **fwdSpinor,
137 sFloat **backSpinor,
int oddBit,
int daggerBit)
141 gFloat *gaugeEven[4], *gaugeOdd[4];
142 gFloat *ghostGaugeEven[4], *ghostGaugeOdd[4];
143 for (
int dir = 0; dir < 4; dir++) {
144 gaugeEven[dir] = gaugeFull[dir];
147 ghostGaugeEven[dir] = ghostGauge[dir];
151 for (
int i = 0; i <
Vh; i++) {
153 for (
int dir = 0; dir < 8; dir++) {
154 gFloat *gauge = gaugeLink_mg4dir(i, dir, oddBit, gaugeEven, gaugeOdd, ghostGaugeEven, ghostGaugeOdd, 1, 1);
155 sFloat *
spinor = spinorNeighbor_mg4dir(i, dir, oddBit, spinorField, fwdSpinor, backSpinor, 1, 1);
158 int projIdx = 2*(dir/2)+(dir+daggerBit)%2;
161 for (
int s = 0; s < 4; s++) {
162 if (dir % 2 == 0) su3Mul(&gaugedSpinor[s*(3*2)], gauge, &projectedSpinor[s*(3*2)]);
163 else su3Tmul(&gaugedSpinor[s*(3*2)], gauge, &projectedSpinor[s*(3*2)]);
175 void wil_dslash(
void *out,
void **gauge,
void *in,
int oddBit,
int daggerBit,
180 dslashReference((
double*)out, (
double**)gauge, (
double*)in, oddBit, daggerBit);
182 dslashReference((
float*)out, (
float**)gauge, (
float*)in, oddBit, daggerBit);
188 void **ghostGauge = (
void**)cpu.
Ghost();
197 for (
int d=0; d<4; d++)
csParam.
x[d] =
Z[d];
213 else errorQuda(
"ERROR: full parity not supported in function %s", __FUNCTION__);
222 dslashReference((
double*)out, (
double**)gauge, (
double**)ghostGauge, (
double*)in,
223 (
double**)fwd_nbr_spinor, (
double**)back_nbr_spinor, oddBit, daggerBit);
225 dslashReference((
float*)out, (
float**)gauge, (
float**)ghostGauge, (
float*)in,
226 (
float**)fwd_nbr_spinor, (
float**)back_nbr_spinor, oddBit, daggerBit);
234 template <
typename sFloat>
243 a = -2.0 *
kappa *
mu * flavor;
244 b = 1.0 / (1.0 + a*a);
246 printf(
"Twist type %d not defined\n", twist);
252 for(
int i = 0; i <
V; i++) {
254 for(
int s = 0; s < 4; s++)
255 for(
int c = 0; c < 3; c++) {
256 sFloat a5 = ((s / 2) ? -1.0 : +1.0) * a;
257 tmp[s * 6 + c * 2 + 0] = b* (in[i * 24 + s * 6 + c * 2 + 0] - a5*in[i * 24 + s * 6 + c * 2 + 1]);
258 tmp[s * 6 + c * 2 + 1] = b* (in[i * 24 + s * 6 + c * 2 + 1] + a5*in[i * 24 + s * 6 + c * 2 + 0]);
261 for (
int j=0; j<24; j++) out[i*24+j] =
tmp[j];
272 twistGamma5((
float*)out, (
float*)in, daggerBit, (
float)
kappa, (
float)
mu, flavor,
V, twist);
371 }
else if (!daggerBit) {
411 template <
typename sFloat>
415 sFloat a=0.0, b=0.0, d=0.0;
423 d = 1.0 / (1.0 + a*a - b*b);
425 printf(
"Twist type %d not defined\n", twist);
431 for(
int i = 0; i <
V; i++) {
434 for(
int s = 0; s < 4; s++)
435 for(
int c = 0; c < 3; c++) {
436 sFloat a5 = ((s / 2) ? -1.0 : +1.0) * a;
437 tmp1[s * 6 + c * 2 + 0] = d
438 * (in1[i * 24 + s * 6 + c * 2 + 0] - a5 * in1[i * 24 + s * 6 + c * 2 + 1]
439 + b * in2[i * 24 + s * 6 + c * 2 + 0]);
440 tmp1[s * 6 + c * 2 + 1] = d
441 * (in1[i * 24 + s * 6 + c * 2 + 1] + a5 * in1[i * 24 + s * 6 + c * 2 + 0]
442 + b * in2[i * 24 + s * 6 + c * 2 + 1]);
443 tmp2[s * 6 + c * 2 + 0] = d
444 * (in2[i * 24 + s * 6 + c * 2 + 0] + a5 * in2[i * 24 + s * 6 + c * 2 + 1]
445 + b * in1[i * 24 + s * 6 + c * 2 + 0]);
446 tmp2[s * 6 + c * 2 + 1] = d
447 * (in2[i * 24 + s * 6 + c * 2 + 1] - a5 * in2[i * 24 + s * 6 + c * 2 + 0]
448 + b * in1[i * 24 + s * 6 + c * 2 + 1]);
450 for (
int j=0; j<24; j++) out1[i*24+j] = tmp1[j], out2[i*24+j] = tmp2[j];
468 void tm_ndeg_dslash(
void *res1,
void *res2,
void **gauge,
void *spinorField1,
void *spinorField2,
double kappa,
double mu,
479 ndeg_twist_gamma5(res1, res2, res1, res2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
494 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
497 ndeg_twist_gamma5(outEven1, outEven2, outEven1, outEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
501 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
504 ndeg_twist_gamma5(outEven1, outEven2, outEven1, outEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
509 tmp1, tmp2, inEven1, inEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
512 ndeg_twist_gamma5(tmp1, tmp2, outEven1, outEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
517 tmp1, tmp2, inEven1, inEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
520 ndeg_twist_gamma5(tmp1, tmp2, outEven1, outEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
529 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
535 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
543 ndeg_twist_gamma5(inEven1, inEven2, inEven1, inEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
554 void tm_ndeg_mat(
void *evenOut,
void* oddOut,
void **gauge,
void *evenIn,
void *oddIn,
double kappa,
double mu,
double epsilon,
int daggerBit,
QudaPrecision precision,
QudaGaugeParam &
gauge_param)
557 void *inEven1 = evenIn;
560 void *inOdd1 = oddIn;
563 void *outEven1 = evenOut;
566 void *outOdd1 = oddOut;
583 ndeg_twist_gamma5(tmpEven1, tmpEven2, inEven1, inEven2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
584 ndeg_twist_gamma5(tmpOdd1, tmpOdd2, inOdd1, inOdd2, daggerBit,
kappa,
mu,
epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
QudaGammaBasis gammaBasis
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
QudaFieldOrder fieldOrder
const void ** Ghost() const
void exchangeGhost(QudaParity parity, int nFace, int dagger, const MemoryLocation *pack_destination=nullptr, const MemoryLocation *halo_location=nullptr, bool gdr_send=false, bool gdr_recv=false, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION) const
This is a unified ghost exchange function for doing a complete halo exchange regardless of the type o...
static void * fwdGhostFaceBuffer[QUDA_MAX_DIM]
static void * backGhostFaceBuffer[QUDA_MAX_DIM]
cudaColorSpinorField * tmp
cpuColorSpinorField * spinor
QudaGaugeParam gauge_param
const double projector[10][4][4][2]
enum QudaPrecision_s QudaPrecision
enum QudaTwistFlavorType_s QudaTwistFlavorType
@ QUDA_PARITY_SITE_SUBSET
@ QUDA_DEGRAND_ROSSI_GAMMA_BASIS
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
@ QUDA_GHOST_EXCHANGE_PAD
@ QUDA_MATPC_ODD_ODD_ASYMMETRIC
@ QUDA_MATPC_EVEN_EVEN_ASYMMETRIC
enum QudaMatPCType_s QudaMatPCType
@ QUDA_TWIST_GAMMA5_INVERSE
@ QUDA_TWIST_GAMMA5_DIRECT
@ QUDA_EVEN_ODD_SITE_ORDER
@ QUDA_SPACE_SPIN_COLOR_FIELD_ORDER
@ QUDA_REFERENCE_FIELD_CREATE
enum QudaParity_s QudaParity
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
FloatingPoint< float > Float
__host__ __device__ T sum(const array< T, s > &a)
QudaGhostExchange ghostExchange
QudaSiteSubset siteSubset
void multiplySpinorByDiracProjector(Float *res, int projIdx, Float *spinorIn)
void dslashReference(sFloat *res, gFloat **gaugeFull, sFloat *spinorField, int oddBit, int daggerBit)
void ndeg_twist_gamma5(void *outf1, void *outf2, void *inf1, void *inf2, const int dagger, const double kappa, const double mu, const double epsilon, const int Vf, QudaTwistGamma5Type twist, QudaPrecision precision)
void wil_mat(void *out, void **gauge, void *in, double kappa, int dagger_bit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void twistGamma5(sFloat *out, sFloat *in, const int dagger, const sFloat kappa, const sFloat mu, const QudaTwistFlavorType flavor, const int V, QudaTwistGamma5Type twist)
void ndegTwistGamma5(sFloat *out1, sFloat *out2, sFloat *in1, sFloat *in2, const int dagger, const sFloat kappa, const sFloat mu, const sFloat epsilon, const int V, QudaTwistGamma5Type twist)
void tm_dslash(void *res, void **gaugeFull, void *spinorField, double kappa, double mu, QudaTwistFlavorType flavor, int oddBit, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void tm_mat(void *out, void **gauge, void *in, double kappa, double mu, QudaTwistFlavorType flavor, int dagger_bit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void tm_ndeg_mat(void *evenOut, void *oddOut, void **gauge, void *evenIn, void *oddIn, double kappa, double mu, double epsilon, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void wil_matpc(void *outEven, void **gauge, void *inEven, double kappa, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void wil_dslash(void *out, void **gauge, void *in, int oddBit, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void tm_ndeg_matpc(void *outEven1, void *outEven2, void **gauge, void *inEven1, void *inEven2, double kappa, double mu, double epsilon, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void twist_gamma5(void *out, void *in, int daggerBit, double kappa, double mu, QudaTwistFlavorType flavor, int V, QudaTwistGamma5Type twist, QudaPrecision precision)
void tm_matpc(void *outEven, void **gauge, void *inEven, double kappa, double mu, QudaTwistFlavorType flavor, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
void tm_ndeg_dslash(void *res1, void *res2, void **gauge, void *spinorField1, void *spinorField2, double kappa, double mu, double epsilon, int oddBit, int daggerBit, QudaMatPCType matpc_type, QudaPrecision precision, QudaGaugeParam &gauge_param)