22 {{1,0}, {0,0}, {0,0}, {0,-1}},
23 {{0,0}, {1,0}, {0,-1}, {0,0}},
24 {{0,0}, {0,1}, {1,0}, {0,0}},
25 {{0,1}, {0,0}, {0,0}, {1,0}}
28 {{1,0}, {0,0}, {0,0}, {0,1}},
29 {{0,0}, {1,0}, {0,1}, {0,0}},
30 {{0,0}, {0,-1}, {1,0}, {0,0}},
31 {{0,-1}, {0,0}, {0,0}, {1,0}}
34 {{1,0}, {0,0}, {0,0}, {1,0}},
35 {{0,0}, {1,0}, {-1,0}, {0,0}},
36 {{0,0}, {-1,0}, {1,0}, {0,0}},
37 {{1,0}, {0,0}, {0,0}, {1,0}}
40 {{1,0}, {0,0}, {0,0}, {-1,0}},
41 {{0,0}, {1,0}, {1,0}, {0,0}},
42 {{0,0}, {1,0}, {1,0}, {0,0}},
43 {{-1,0}, {0,0}, {0,0}, {1,0}}
46 {{1,0}, {0,0}, {0,-1}, {0,0}},
47 {{0,0}, {1,0}, {0,0}, {0,1}},
48 {{0,1}, {0,0}, {1,0}, {0,0}},
49 {{0,0}, {0,-1}, {0,0}, {1,0}}
52 {{1,0}, {0,0}, {0,1}, {0,0}},
53 {{0,0}, {1,0}, {0,0}, {0,-1}},
54 {{0,-1}, {0,0}, {1,0}, {0,0}},
55 {{0,0}, {0,1}, {0,0}, {1,0}}
58 {{1,0}, {0,0}, {-1,0}, {0,0}},
59 {{0,0}, {1,0}, {0,0}, {-1,0}},
60 {{-1,0}, {0,0}, {1,0}, {0,0}},
61 {{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}},
66 {{1,0}, {0,0}, {1,0}, {0,0}},
67 {{0,0}, {1,0}, {0,0}, {1,0}}
73 template <
typename Float>
75 for (
int i=0;
i<4*3*2;
i++) res[
i] = 0.0;
77 for (
int s = 0;
s < 4;
s++) {
78 for (
int t = 0;
t < 4;
t++) {
82 for (
int m = 0; m < 3; m++) {
83 Float spinorRe = spinorIn[
t*(3*2) + m*(2) + 0];
84 Float spinorIm = spinorIn[
t*(3*2) + m*(2) + 1];
85 res[
s*(3*2) + m*(2) + 0] += projRe*spinorRe - projIm*spinorIm;
86 res[
s*(3*2) + m*(2) + 1] += projRe*spinorIm + projIm*spinorRe;
105 template <
typename sFloat,
typename gFloat>
106 void dslashReference(sFloat *res, gFloat **gaugeFull, sFloat *spinorField,
int oddBit,
int daggerBit) {
109 gFloat *gaugeEven[4], *gaugeOdd[4];
110 for (
int dir = 0; dir < 4; dir++) {
111 gaugeEven[dir] = gaugeFull[dir];
115 for (
int i = 0;
i <
Vh;
i++) {
116 for (
int dir = 0; dir < 8; dir++) {
117 gFloat *gauge =
gaugeLink(
i, dir, oddBit, gaugeEven, gaugeOdd, 1);
120 sFloat projectedSpinor[4*3*2], gaugedSpinor[4*3*2];
121 int projIdx = 2*(dir/2)+(dir+daggerBit)%2;
124 for (
int s = 0;
s < 4;
s++) {
125 if (dir % 2 == 0)
su3Mul(&gaugedSpinor[
s*(3*2)], gauge, &projectedSpinor[
s*(3*2)]);
126 else su3Tmul(&gaugedSpinor[
s*(3*2)], gauge, &projectedSpinor[
s*(3*2)]);
129 sum(&res[
i*(4*3*2)], &res[
i*(4*3*2)], gaugedSpinor, 4*3*2);
136 template <
typename sFloat,
typename gFloat>
137 void dslashReference(sFloat *res, gFloat **gaugeFull, gFloat **ghostGauge, sFloat *spinorField,
138 sFloat **fwdSpinor, 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)]);
166 sum(&res[
i*(4*3*2)], &res[
i*(4*3*2)], gaugedSpinor, 4*3*2);
188 void **ghostGauge = (
void**)cpu.
Ghost();
213 else errorQuda(
"ERROR: full parity not supported in function %s", __FUNCTION__);
216 inField.exchangeGhost(otherParity, nFace, daggerBit);
218 void** fwd_nbr_spinor = inField.fwdGhostFaceBuffer;
219 void** back_nbr_spinor = inField.backGhostFaceBuffer;
223 (
double**)fwd_nbr_spinor, (
double**)back_nbr_spinor, oddBit, daggerBit);
226 (
float**)fwd_nbr_spinor, (
float**)back_nbr_spinor, oddBit, daggerBit);
234 template <
typename sFloat>
244 b = 1.0 / (1.0 +
a*
a);
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];
373 }
else if (!daggerBit) {
413 template <
typename sFloat>
417 sFloat
a=0.0,
b=0.0,
d=0.0;
420 b = -2.0 *
kappa * epsilon;
424 b = 2.0 *
kappa * epsilon;
425 d = 1.0 / (1.0 +
a*
a -
b*
b);
433 for(
int i = 0;
i <
V;
i++) {
436 for(
int s = 0;
s < 4;
s++)
437 for(
int c = 0;
c < 3;
c++) {
438 sFloat a5 = ((
s / 2) ? -1.0 : +1.0) *
a;
439 tmp1[
s * 6 +
c * 2 + 0] =
d* (in1[
i * 24 +
s * 6 +
c * 2 + 0] - a5*in1[
i * 24 +
s * 6 +
c * 2 + 1] +
b*in2[
i * 24 +
s * 6 +
c * 2 + 0]);
440 tmp1[
s * 6 +
c * 2 + 1] =
d* (in1[
i * 24 +
s * 6 +
c * 2 + 1] + a5*in1[
i * 24 +
s * 6 +
c * 2 + 0] +
b*in2[
i * 24 +
s * 6 +
c * 2 + 1]);
441 tmp2[
s * 6 +
c * 2 + 0] =
d* (in2[
i * 24 +
s * 6 +
c * 2 + 0] + a5*in2[
i * 24 +
s * 6 +
c * 2 + 1] +
b*in1[
i * 24 +
s * 6 +
c * 2 + 0]);
442 tmp2[
s * 6 +
c * 2 + 1] =
d* (in2[
i * 24 +
s * 6 +
c * 2 + 1] - a5*in2[
i * 24 +
s * 6 +
c * 2 + 0] +
b*in1[
i * 24 +
s * 6 +
c * 2 + 1]);
444 for (
int j=0; j<24; j++) out1[
i*24+j] =
tmp1[j], out2[
i*24+j] =
tmp2[j];
462 void tm_ndeg_dslash(
void *res1,
void *res2,
void **gauge,
void *spinorField1,
void *spinorField2,
double kappa,
double mu,
466 ndeg_twist_gamma5(spinorField1, spinorField2, spinorField1, spinorField2, daggerBit,
kappa, -
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
472 ndeg_twist_gamma5(res1, res2, res1, res2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
477 void tm_ndeg_matpc(
void *outEven1,
void *outEven2,
void **gauge,
void *inEven1,
void *inEven2,
double kappa,
double mu,
double epsilon,
487 ndeg_twist_gamma5(
tmp1,
tmp2,
tmp1,
tmp2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
490 ndeg_twist_gamma5(outEven1, outEven2, outEven1, outEven2, 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, inEven1, inEven2, daggerBit,
kappa, -
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
504 ndeg_twist_gamma5(
tmp1,
tmp2, outEven1, outEven2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
508 ndeg_twist_gamma5(
tmp1,
tmp2, inEven1, inEven2, daggerBit,
kappa, -
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
511 ndeg_twist_gamma5(
tmp1,
tmp2, outEven1, outEven2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
520 ndeg_twist_gamma5(
tmp1,
tmp2,
tmp1,
tmp2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
526 ndeg_twist_gamma5(
tmp1,
tmp2,
tmp1,
tmp2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
534 ndeg_twist_gamma5(inEven1, inEven2, inEven1, inEven2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
545 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)
548 void *inEven1 = evenIn;
551 void *inOdd1 = oddIn;
554 void *outEven1 = evenOut;
557 void *outOdd1 = oddOut;
574 ndeg_twist_gamma5(tmpEven1, tmpEven2, inEven1, inEven2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
575 ndeg_twist_gamma5(tmpOdd1, tmpOdd2, inOdd1, inOdd2, daggerBit,
kappa,
mu, epsilon,
Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
QudaGhostExchange ghostExchange
void xpay(ColorSpinorField &x, const double &a, ColorSpinorField &y)
enum QudaPrecision_s QudaPrecision
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_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)
cudaColorSpinorField * tmp
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)
QudaGaugeParam gauge_param
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 twist_gamma5(void *out, void *in, int daggerBit, double kappa, double mu, QudaTwistFlavorType flavor, int V, QudaTwistGamma5Type twist, QudaPrecision precision)
QudaSiteSubset siteSubset
void exit(int) __attribute__((noreturn))
static Float * spinorNeighbor(int i, int dir, int oddBit, Float *spinorField, int neighbor_distance)
void * malloc(size_t __size) __attribute__((__warn_unused_result__)) __attribute__((alloc_size(1)))
int printf(const char *,...) __attribute__((__format__(__printf__
VOLATILE spinorFloat kappa
QudaFieldOrder fieldOrder
__host__ __device__ void sum(double &a, double &b)
enum QudaMatPCType_s QudaMatPCType
QudaGammaBasis gammaBasis
void multiplySpinorByDiracProjector(Float *res, int projIdx, Float *spinorIn)
static const double projector[8][4][4][2]
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)
const void ** Ghost() const
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)
enum QudaParity_s QudaParity
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)
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 dslashReference(sFloat *res, gFloat **gaugeFull, sFloat *spinorField, int oddBit, int daggerBit)
static Float * gaugeLink(int i, int dir, int oddBit, Float **gaugeEven, Float **gaugeOdd, int nbr_distance)
void wil_mat(void *out, void **gauge, void *in, double kappa, int dagger_bit, QudaPrecision precision, QudaGaugeParam &gauge_param)
cpuColorSpinorField * out
void wil_matpc(void *outEven, void **gauge, void *inEven, double kappa, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
static void su3Mul(sFloat *res, gFloat *mat, sFloat *vec)
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
void wil_dslash(void *out, void **gauge, void *in, int oddBit, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
static void su3Tmul(sFloat *res, gFloat *mat, sFloat *vec)
static __inline__ size_t size_t d
cpuColorSpinorField * spinor
void twistGamma5(sFloat *out, sFloat *in, const int dagger, const sFloat kappa, const sFloat mu, const QudaTwistFlavorType flavor, const int V, QudaTwistGamma5Type twist)
enum QudaTwistFlavorType_s QudaTwistFlavorType