20 static const double projector[8][4][4][2] = {
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>
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);
118 sFloat *
spinor = spinorNeighbor(i, dir, oddBit, spinorField, 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];
148 ghostGaugeOdd[dir] = ghostGauge[dir] + (
faceVolume[dir]/2)*gaugeSiteSize;
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);
180 dslashReference((
double*)out, (
double**)gauge, (
double*)in, oddBit, daggerBit);
182 dslashReference((
float*)out, (
float**)gauge, (
float*)in, oddBit, daggerBit);
187 void **ghostGauge = (
void**)cpu.
Ghost();
196 for (
int d=0; d<4; d++) csParam.
x[d] =
Z[d];
212 else errorQuda(
"ERROR: full parity not supported in function %s", __FUNCTION__);
215 FaceBuffer faceBuf(
Z, 4, mySpinorSiteSize, nFace, precision);
218 void** fwd_nbr_spinor = inField.fwdGhostFaceBuffer;
219 void** back_nbr_spinor = inField.backGhostFaceBuffer;
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>
240 a = 2.0 * kappa * mu * flavor;
243 a = -2.0 * kappa * mu * flavor;
244 b = 1.0 / (1.0 + a*a);
246 printf(
"Twist type %d not defined\n", twist);
250 if (dagger) a *= -1.0;
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];
270 twistGamma5((
double*)out, (
double*)in, daggerBit, kappa, mu, flavor, V, twist);
272 twistGamma5((
float*)out, (
float*)in, daggerBit, (
float)kappa, (
float)mu, flavor, V, twist);
285 wil_dslash(res, gaugeFull, spinorField, oddBit, daggerBit, precision, gauge_param);
291 twist_gamma5(spinorField, spinorField, daggerBit, kappa, mu, flavor,
304 wil_dslash(outOdd, gauge, inEven, 1, dagger_bit, precision, gauge_param);
305 wil_dslash(outEven, gauge, inOdd, 0, dagger_bit, precision, gauge_param);
309 else xpay((
float*)in, -(
float)kappa, (
float*)out,
V*spinorSiteSize);
322 wil_dslash(outOdd, gauge, inEven, 1, dagger_bit, precision, gauge_param);
323 wil_dslash(outEven, gauge, inOdd, 0, dagger_bit, precision, gauge_param);
330 else xpay((
float*)tmp, -(
float)kappa, (
float*)out,
V*spinorSiteSize);
345 wil_dslash(tmp, gauge, inEven, 1, daggerBit, precision, gauge_param);
346 wil_dslash(outEven, gauge, tmp, 0, daggerBit, precision, gauge_param);
348 wil_dslash(tmp, gauge, inEven, 0, daggerBit, precision, gauge_param);
349 wil_dslash(outEven, gauge, tmp, 1, daggerBit, precision, gauge_param);
353 double kappa2 = -kappa*
kappa;
355 else xpay((
float*)inEven, (
float)kappa2, (
float*)outEven, Vh*spinorSiteSize);
367 wil_dslash(tmp, gauge, inEven, 1, daggerBit, precision, gauge_param);
369 wil_dslash(outEven, gauge, tmp, 0, daggerBit, precision, gauge_param);
372 wil_dslash(tmp, gauge, inEven, 0, daggerBit, precision, gauge_param);
374 wil_dslash(outEven, gauge, tmp, 1, daggerBit, precision, gauge_param);
376 }
else if (!daggerBit) {
378 wil_dslash(tmp, gauge, inEven, 1, daggerBit, precision, gauge_param);
380 wil_dslash(outEven, gauge, tmp, 0, daggerBit, precision, gauge_param);
383 wil_dslash(tmp, gauge, inEven, 0, daggerBit, precision, gauge_param);
385 wil_dslash(outEven, gauge, tmp, 1, daggerBit, precision, gauge_param);
391 wil_dslash(tmp, gauge, inEven, 1, daggerBit, precision, gauge_param);
393 wil_dslash(outEven, gauge, tmp, 0, daggerBit, precision, gauge_param);
397 wil_dslash(tmp, gauge, inEven, 0, daggerBit, precision, gauge_param);
399 wil_dslash(outEven, gauge, tmp, 1, daggerBit, precision, gauge_param);
404 double kappa2 = -kappa*
kappa;
407 else xpay((
float*)inEven, (
float)kappa2, (
float*)outEven, Vh*spinorSiteSize);
410 else xpay((
float*)tmp, (
float)kappa2, (
float*)outEven, Vh*spinorSiteSize);
418 template <
typename sFloat>
422 sFloat a=0.0, b=0.0, d=0.0;
424 a = 2.0 * kappa *
mu;
425 b = -2.0 * kappa * epsilon;
428 a = -2.0 * kappa *
mu;
429 b = 2.0 * kappa * epsilon;
430 d = 1.0 / (1.0 + a*a - b*b);
432 printf(
"Twist type %d not defined\n", twist);
436 if (dagger) a *= -1.0;
438 for(
int i = 0; i <
V; i++) {
441 for(
int s = 0;
s < 4;
s++)
442 for(
int c = 0; c < 3; c++) {
443 sFloat a5 = ((
s / 2) ? -1.0 : +1.0) * a;
444 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]);
445 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]);
446 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]);
447 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]);
449 for (
int j=0; j<24; j++) out1[i*24+j] = tmp1[j], out2[i*24+j] = tmp2[j];
459 ndegTwistGamma5((
double*)outf1, (
double*)outf2, (
double*)inf1, (
double*)inf2, dagger, kappa, mu, epsilon, Vf, twist);
463 ndegTwistGamma5((
float*)outf1, (
float*)outf2, (
float*)inf1, (
float*)inf2, dagger, (
float)kappa, (
float)mu, (
float)epsilon, Vf, twist);
467 void tm_ndeg_dslash(
void *res1,
void *res2,
void **gauge,
void *spinorField1,
void *spinorField2,
double kappa,
double mu,
471 ndeg_twist_gamma5(spinorField1, spinorField2, spinorField1, spinorField2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
473 wil_dslash(res1, gauge, spinorField1, oddBit, daggerBit, precision, gauge_param);
474 wil_dslash(res2, gauge, spinorField2, oddBit, daggerBit, precision, gauge_param);
477 ndeg_twist_gamma5(res1, res2, res1, res2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
482 void tm_ndeg_matpc(
void *outEven1,
void *outEven2,
void **gauge,
void *inEven1,
void *inEven2,
double kappa,
double mu,
double epsilon,
490 wil_dslash(tmp1, gauge, inEven1, 1, daggerBit, precision, gauge_param);
491 wil_dslash(tmp2, gauge, inEven2, 1, daggerBit, precision, gauge_param);
492 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
493 wil_dslash(outEven1, gauge, tmp1, 0, daggerBit, precision, gauge_param);
494 wil_dslash(outEven2, gauge, tmp2, 0, daggerBit, precision, gauge_param);
495 ndeg_twist_gamma5(outEven1, outEven2, outEven1, outEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
497 wil_dslash(tmp1, gauge, inEven1, 0, daggerBit, precision, gauge_param);
498 wil_dslash(tmp2, gauge, inEven2, 0, daggerBit, precision, gauge_param);
499 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
500 wil_dslash(outEven1, gauge, tmp1, 1, daggerBit, precision, gauge_param);
501 wil_dslash(outEven2, gauge, tmp2, 1, daggerBit, precision, gauge_param);
502 ndeg_twist_gamma5(outEven1, outEven2, outEven1, outEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
506 ndeg_twist_gamma5(tmp1, tmp2, inEven1, inEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
507 wil_dslash(outEven1, gauge, tmp1, 1, daggerBit, precision, gauge_param);
508 wil_dslash(outEven2, gauge, tmp2, 1, daggerBit, precision, gauge_param);
509 ndeg_twist_gamma5(tmp1, tmp2, outEven1, outEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
510 wil_dslash(outEven1, gauge, tmp1, 0, daggerBit, precision, gauge_param);
511 wil_dslash(outEven2, gauge, tmp2, 0, daggerBit, precision, gauge_param);
513 ndeg_twist_gamma5(tmp1, tmp2, inEven1, inEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
514 wil_dslash(outEven1, gauge, tmp1, 0, daggerBit, precision, gauge_param);
515 wil_dslash(outEven2, gauge, tmp2, 0, daggerBit, precision, gauge_param);
516 ndeg_twist_gamma5(tmp1, tmp2, outEven1, outEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
517 wil_dslash(outEven1, gauge, tmp1, 1, daggerBit, precision, gauge_param);
518 wil_dslash(outEven2, gauge, tmp2, 1, daggerBit, precision, gauge_param);
523 wil_dslash(tmp1, gauge, inEven1, 1, daggerBit, precision, gauge_param);
524 wil_dslash(tmp2, gauge, inEven2, 1, daggerBit, precision, gauge_param);
525 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
526 wil_dslash(outEven1, gauge, tmp1, 0, daggerBit, precision, gauge_param);
527 wil_dslash(outEven2, gauge, tmp2, 0, daggerBit, precision, gauge_param);
529 wil_dslash(tmp1, gauge, inEven1, 0, daggerBit, precision, gauge_param);
530 wil_dslash(tmp2, gauge, inEven2, 0, daggerBit, precision, gauge_param);
531 ndeg_twist_gamma5(tmp1, tmp2, tmp1, tmp2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_INVERSE, precision);
532 wil_dslash(outEven1, gauge, tmp1, 1, daggerBit, precision, gauge_param);
533 wil_dslash(outEven2, gauge, tmp2, 1, daggerBit, precision, gauge_param);
537 double kappa2 = -kappa*
kappa;
539 ndeg_twist_gamma5(inEven1, inEven2, inEven1, inEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
544 xpay((
double*)inEven2, kappa2, (
double*)outEven2, Vh*spinorSiteSize);
548 xpay((
float*)inEven2, (
float)kappa2, (
float*)outEven2, Vh*spinorSiteSize);
556 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)
559 void *inEven1 = evenIn;
562 void *inOdd1 = oddIn;
565 void *outEven1 = evenOut;
568 void *outOdd1 = oddOut;
578 wil_dslash(outOdd1, gauge, inEven1, 1, daggerBit, precision, gauge_param);
579 wil_dslash(outOdd2, gauge, inEven2, 1, daggerBit, precision, gauge_param);
581 wil_dslash(outEven1, gauge, inOdd1, 0, daggerBit, precision, gauge_param);
582 wil_dslash(outEven2, gauge, inOdd2, 0, daggerBit, precision, gauge_param);
585 ndeg_twist_gamma5(tmpEven1, tmpEven2, inEven1, inEven2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
586 ndeg_twist_gamma5(tmpOdd1, tmpOdd2, inOdd1, inOdd2, daggerBit, kappa, mu, epsilon, Vh,
QUDA_TWIST_GAMMA5_DIRECT, precision);
590 xpay((
double*)tmpOdd2, -kappa, (
double*)outOdd2, Vh*spinorSiteSize);
592 xpay((
double*)tmpEven1, -kappa, (
double*)outEven1, Vh*spinorSiteSize);
593 xpay((
double*)tmpEven2, -kappa, (
double*)outEven2, Vh*spinorSiteSize);
597 xpay((
float*)tmpOdd2, (
float)(-kappa), (
float*)outOdd2, Vh*spinorSiteSize);
599 xpay((
float*)tmpEven1, (
float)(-kappa), (
float*)outEven1, Vh*spinorSiteSize);
600 xpay((
float*)tmpEven2, (
float)(-kappa), (
float*)outEven2, Vh*spinorSiteSize);
QudaGaugeParam gauge_param
const void ** Ghost() const
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)
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
cudaColorSpinorField * tmp1
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)
cpuColorSpinorField * spinor
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
cudaColorSpinorField * tmp2
cudaColorSpinorField * tmp
VOLATILE spinorFloat kappa
QudaFieldOrder fieldOrder
FloatingPoint< float > Float
void exchangeCpuSpinor(quda::cpuColorSpinorField &in, int parity, int dagger)
enum QudaMatPCType_s QudaMatPCType
QudaGammaBasis gammaBasis
void multiplySpinorByDiracProjector(Float *res, int projIdx, Float *spinorIn)
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_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)
void wil_mat(void *out, void **gauge, void *in, double kappa, int dagger_bit, QudaPrecision precision, QudaGaugeParam &gauge_param)
cpuColorSpinorField * out
const double projector[10][4][4][2]
void wil_matpc(void *outEven, void **gauge, void *inEven, double kappa, QudaMatPCType matpc_type, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
void wil_dslash(void *out, void **gauge, void *in, int oddBit, int daggerBit, 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)
enum QudaTwistFlavorType_s QudaTwistFlavorType