QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
fermion_force_reference.cpp
Go to the documentation of this file.
1 #include <stdio.h>
2 #include <stdlib.h>
3 #include <math.h>
4 #include <string.h>
5 
6 #include "quda.h"
7 #include "test_util.h"
8 #include "misc.h"
10 
11 extern int Z[4];
12 extern int V;
13 extern int Vh;
14 
15 
16 #define CADD(a,b,c) { (c).real = (a).real + (b).real; \
17  (c).imag = (a).imag + (b).imag; }
18 #define CMUL(a,b,c) { (c).real = (a).real*(b).real - (a).imag*(b).imag; \
19  (c).imag = (a).real*(b).imag + (a).imag*(b).real; }
20 #define CSUM(a,b) { (a).real += (b).real; (a).imag += (b).imag; }
21 
22 /* c = a* * b */
23 #define CMULJ_(a,b,c) { (c).real = (a).real*(b).real + (a).imag*(b).imag; \
24  (c).imag = (a).real*(b).imag - (a).imag*(b).real; }
25 
26 /* c = a * b* */
27 #define CMUL_J(a,b,c) { (c).real = (a).real*(b).real + (a).imag*(b).imag; \
28  (c).imag = (a).imag*(b).real - (a).real*(b).imag; }
29 
30 #define CONJG(a,b) { (b).real = (a).real; (b).imag = -(a).imag; }
31 
32 typedef struct {
33  float real;
34  float imag;
35 } fcomplex;
36 
37 /* specific for double complex */
38 typedef struct {
39  double real;
40  double imag;
41 } dcomplex;
42 
43 typedef struct { fcomplex e[3][3]; } fsu3_matrix;
44 typedef struct { fcomplex c[3]; } fsu3_vector;
45 typedef struct { dcomplex e[3][3]; } dsu3_matrix;
46 typedef struct { dcomplex c[3]; } dsu3_vector;
47 
48 typedef struct {
49  fcomplex m01,m02,m12;
50  float m00im,m11im,m22im;
51  float space;
53 
54 typedef struct {
55  dcomplex m01,m02,m12;
56  double m00im,m11im,m22im;
57  double space;
59 
60 typedef struct { fsu3_vector h[2]; } fhalf_wilson_vector;
61 typedef struct { dsu3_vector h[2]; } dhalf_wilson_vector;
62 
63 
64 template<typename su3_matrix>
65 static void
66 su3_adjoint( su3_matrix *a, su3_matrix *b )
67 {
68  int i,j;
69  for(i=0;i<3;i++)for(j=0;j<3;j++){
70  CONJG( a->e[j][i], b->e[i][j] );
71  }
72 }
73 
74 
75 template<typename su3_matrix, typename anti_hermitmat>
76 static void
77 make_anti_hermitian( su3_matrix *m3, anti_hermitmat *ah3 )
78 {
79 
80  typeof(ah3->m00im) temp =
81  (m3->e[0][0].imag + m3->e[1][1].imag + m3->e[2][2].imag)*0.33333333333333333;
82  ah3->m00im = m3->e[0][0].imag - temp;
83  ah3->m11im = m3->e[1][1].imag - temp;
84  ah3->m22im = m3->e[2][2].imag - temp;
85  ah3->m01.real = (m3->e[0][1].real - m3->e[1][0].real)*0.5;
86  ah3->m02.real = (m3->e[0][2].real - m3->e[2][0].real)*0.5;
87  ah3->m12.real = (m3->e[1][2].real - m3->e[2][1].real)*0.5;
88  ah3->m01.imag = (m3->e[0][1].imag + m3->e[1][0].imag)*0.5;
89  ah3->m02.imag = (m3->e[0][2].imag + m3->e[2][0].imag)*0.5;
90  ah3->m12.imag = (m3->e[1][2].imag + m3->e[2][1].imag)*0.5;
91 
92 }
93 
94 template <typename anti_hermitmat, typename su3_matrix>
95 static void
96 uncompress_anti_hermitian(anti_hermitmat *mat_antihermit,
97  su3_matrix *mat_su3 )
98 {
99  typeof(mat_antihermit->m00im) temp1;
100  mat_su3->e[0][0].imag=mat_antihermit->m00im;
101  mat_su3->e[0][0].real=0.;
102  mat_su3->e[1][1].imag=mat_antihermit->m11im;
103  mat_su3->e[1][1].real=0.;
104  mat_su3->e[2][2].imag=mat_antihermit->m22im;
105  mat_su3->e[2][2].real=0.;
106  mat_su3->e[0][1].imag=mat_antihermit->m01.imag;
107  temp1=mat_antihermit->m01.real;
108  mat_su3->e[0][1].real=temp1;
109  mat_su3->e[1][0].real= -temp1;
110  mat_su3->e[1][0].imag=mat_antihermit->m01.imag;
111  mat_su3->e[0][2].imag=mat_antihermit->m02.imag;
112  temp1=mat_antihermit->m02.real;
113  mat_su3->e[0][2].real=temp1;
114  mat_su3->e[2][0].real= -temp1;
115  mat_su3->e[2][0].imag=mat_antihermit->m02.imag;
116  mat_su3->e[1][2].imag=mat_antihermit->m12.imag;
117  temp1=mat_antihermit->m12.real;
118  mat_su3->e[1][2].real=temp1;
119  mat_su3->e[2][1].real= -temp1;
120  mat_su3->e[2][1].imag=mat_antihermit->m12.imag;
121 }
122 
123 template <typename su3_matrix, typename Float>
124 static void
125 scalar_mult_sub_su3_matrix(su3_matrix *a,su3_matrix *b, Float s, su3_matrix *c)
126 {
127  int i,j;
128  for(i=0;i<3;i++){
129  for(j=0;j<3;j++){
130  c->e[i][j].real = a->e[i][j].real - s*b->e[i][j].real;
131  c->e[i][j].imag = a->e[i][j].imag - s*b->e[i][j].imag;
132  }
133  }
134 }
135 
136 template <typename su3_matrix, typename Float>
137 static void
138 scalar_mult_add_su3_matrix(su3_matrix *a,su3_matrix *b, Float s, su3_matrix *c)
139 {
140  int i,j;
141  for(i=0;i<3;i++){
142  for(j=0;j<3;j++){
143  c->e[i][j].real = a->e[i][j].real + s*b->e[i][j].real;
144  c->e[i][j].imag = a->e[i][j].imag + s*b->e[i][j].imag;
145  }
146  }
147 }
148 
149 
150 template<typename su3_matrix, typename su3_vector>
151 static void
152 mult_su3_mat_vec( su3_matrix *a, su3_vector *b, su3_vector *c )
153 {
154  int i,j;
155  typeof(a->e[0][0]) x,y;
156  for(i=0;i<3;i++){
157  x.real=x.imag=0.0;
158  for(j=0;j<3;j++){
159  CMUL( a->e[i][j] , b->c[j] , y );
160  CSUM( x , y );
161  }
162  c->c[i] = x;
163  }
164 }
165 template<typename su3_matrix, typename su3_vector>
166 static void
167 mult_adj_su3_mat_vec( su3_matrix *a, su3_vector *b, su3_vector *c )
168 {
169  int i,j;
170  typeof(a->e[0][0]) x,y,z;
171  for(i=0;i<3;i++){
172  x.real=x.imag=0.0;
173  for(j=0;j<3;j++){
174  CONJG( a->e[j][i], z );
175  CMUL( z , b->c[j], y );
176  CSUM( x , y );
177  }
178  c->c[i] = x;
179  }
180 }
181 
182 template<typename su3_vector, typename su3_matrix>
183 static void
184 su3_projector( su3_vector *a, su3_vector *b, su3_matrix *c )
185 {
186  int i,j;
187  for(i=0;i<3;i++)for(j=0;j<3;j++){
188  CMUL_J( a->c[i], b->c[j], c->e[i][j] );
189  }
190 }
191 
192 template<typename su3_vector, typename Real>
193 static void
194 scalar_mult_add_su3_vector(su3_vector *a, su3_vector *b, Real s,
195  su3_vector *c)
196 {
197  int i;
198  for(i=0;i<3;i++){
199  c->c[i].real = a->c[i].real + s*b->c[i].real;
200  c->c[i].imag = a->c[i].imag + s*b->c[i].imag;
201  }
202 }
203 
204 
205 
206 template < typename su3_matrix>
207 static void
208 print_su3_matrix(su3_matrix *a)
209 {
210  int i, j;
211  for(i=0;i < 3; i++){
212  for(j=0;j < 3;j++){
213  printf("(%f %f)\t", a->e[i][j].real, a->e[i][j].imag);
214  }
215  printf("\n");
216  }
217 
218 }
219 
220 
221 
222 
223 template <typename su3_matrix, typename anti_hermitmat, typename Float>
224 static void
225 update_mom(anti_hermitmat* momentum, int dir, su3_matrix* sitelink,
226  su3_matrix* staple, Float eb3)
227 {
228  int i;
229  for(i=0;i <V; i++){
230  su3_matrix tmat1;
231  su3_matrix tmat2;
232  su3_matrix tmat3;
233 
234  su3_matrix* lnk = sitelink + 4*i+dir;
235  su3_matrix* stp = staple + i;
236  anti_hermitmat* mom = momentum + 4*i+dir;
237 
238  mult_su3_na(lnk, stp, &tmat1);
239  uncompress_anti_hermitian(mom, &tmat2);
240 
241  scalar_mult_sub_su3_matrix(&tmat2, &tmat1, eb3, &tmat3);
242  make_anti_hermitian(&tmat3, mom);
243 
244  }
245 
246 }
247 
248 
249 template<typename half_wilson_vector, typename su3_matrix>
250 static void
251 u_shift_hw(half_wilson_vector *src, half_wilson_vector *dest, int dir, su3_matrix* sitelink )
252 {
253  int i ;
254  int dx[4];
255 
256  dx[3]=dx[2]=dx[1]=dx[0]=0;
257 
258  if(GOES_FORWARDS(dir)){
259  dx[dir]=1;
260  for(i=0;i < V; i++){
261  int nbr_idx = neighborIndexFullLattice(i, dx[3], dx[2], dx[1], dx[0]);
262  half_wilson_vector* hw = src + nbr_idx;
263  su3_matrix* link = sitelink + i*4 + dir;
264  mult_su3_mat_vec(link, &hw->h[0], &dest[i].h[0]);
265  mult_su3_mat_vec(link, &hw->h[1], &dest[i].h[1]);
266  }
267  }else{
268  dx[OPP_DIR(dir)]=-1;
269  for(i=0;i < V; i++){
270  int nbr_idx = neighborIndexFullLattice(i, dx[3], dx[2], dx[1], dx[0]);
271  half_wilson_vector* hw = src + nbr_idx;
272  su3_matrix* link = sitelink + nbr_idx*4 + OPP_DIR(dir);
273  mult_adj_su3_mat_vec(link, &hw->h[0], &dest[i].h[0]);
274  mult_adj_su3_mat_vec(link, &hw->h[1], &dest[i].h[1]);
275  }
276  }
277 
278 }
279 
280 
281 template <class su3_matrix, typename half_wilson_vector,
282  typename anti_hermitmat, typename Real>
283 static void
284 add_3f_force_to_mom(half_wilson_vector *back, half_wilson_vector *forw,
285  int dir, Real coeff[2], anti_hermitmat* momentum)
286 {
287  Real my_coeff[2] ;
288  Real tmp_coeff[2] ;
289  int mydir;
290  int i;
291 
292  if(GOES_BACKWARDS(dir)){
293  mydir = OPP_DIR(dir);
294  my_coeff[0] = -coeff[0];
295  my_coeff[1] = -coeff[1];
296  }else{
297  mydir = dir;
298  my_coeff[0] = coeff[0];
299  my_coeff[1] = coeff[1];
300  }
301 
302  for(i=0;i < V;i++){
303  if (i < Vh){
304  tmp_coeff[0] = my_coeff[0];
305  tmp_coeff[1] = my_coeff[1];
306  }else{
307  tmp_coeff[0] = -my_coeff[0];
308  tmp_coeff[1] = -my_coeff[1] ;
309  }
310 
311  su3_matrix tmat;
312  su3_matrix mom_matrix;
313  anti_hermitmat* mom = momentum+ 4* i + mydir;
314  uncompress_anti_hermitian(mom, &mom_matrix);
315  su3_projector( &back[i].h[0], &forw[i].h[0], &tmat);
316  scalar_mult_add_su3_matrix(&mom_matrix, &tmat, tmp_coeff[0], &mom_matrix );
317  su3_projector( &back[i].h[1], &forw[i].h[1], &tmat);
318  scalar_mult_add_su3_matrix(&mom_matrix, &tmat, tmp_coeff[1], &mom_matrix );
319  make_anti_hermitian(&mom_matrix, mom);
320  }
321 }
322 
323 template<class su3_matrix, typename Real, typename half_wilson_vector, typename anti_hermitmat>
324 static void
325 side_link_3f_force(int mu, int nu, Real coeff[2], half_wilson_vector *Path ,
326  half_wilson_vector *Path_nu, half_wilson_vector *Path_mu,
327  half_wilson_vector *Path_numu, anti_hermitmat* mom)
328 {
329 
330  Real m_coeff[2] ;
331 
332  m_coeff[0] = -coeff[0] ;
333  m_coeff[1] = -coeff[1] ;
334 
335  if(GOES_FORWARDS(mu)){
336  if(GOES_FORWARDS(nu)){
337  add_3f_force_to_mom<su3_matrix>(Path_numu, Path, mu, coeff, mom) ;
338  }else{
339  add_3f_force_to_mom<su3_matrix>(Path,Path_numu,OPP_DIR(mu),m_coeff, mom);
340  }
341  }
342  else{
343  if(GOES_FORWARDS(nu))
344  add_3f_force_to_mom<su3_matrix>(Path_nu, Path_mu, mu, m_coeff, mom);
345  else
346  add_3f_force_to_mom<su3_matrix>(Path_mu, Path_nu, OPP_DIR(mu), coeff, mom) ;
347  }
348 }
349 
350 
351 
352 #define Pmu tempvec[0]
353 #define Pnumu tempvec[1]
354 #define Prhonumu tempvec[2]
355 #define P7 tempvec[3]
356 #define P7rho tempvec[4]
357 #define P7rhonu tempvec[5]
358 #define P5 tempvec[6]
359 #define P3 tempvec[7]
360 #define P5nu tempvec[3]
361 #define P3mu tempvec[3]
362 #define Popmu tempvec[4]
363 #define Pmumumu tempvec[4]
364 
365 
366 
367 template <typename Real, typename half_wilson_vector, typename anti_hermitmat,
368  typename su3_matrix>
369 static void
370 do_fermion_force_reference(Real eps, Real weight1, Real weight2,
371  half_wilson_vector* temp_x, Real* act_path_coeff,
372  su3_matrix* sitelink, anti_hermitmat* mom)
373 {
374  int i;
375  int mu, nu, rho, sig;
376  Real coeff[2];
377  Real OneLink[2], Lepage[2], Naik[2], FiveSt[2], ThreeSt[2], SevenSt[2] ;
378  Real mNaik[2], mLepage[2], mFiveSt[2], mThreeSt[2], mSevenSt[2];
379  half_wilson_vector *tempvec[8] ;
380  int sites_on_node = V;
381  int DirectLinks[8] ;
382 
383  Real ferm_epsilon;
384  ferm_epsilon = 2.0*weight1*eps;
385  OneLink[0] = act_path_coeff[0]*ferm_epsilon ;
386  Naik[0] = act_path_coeff[1]*ferm_epsilon ; mNaik[0] = -Naik[0];
387  ThreeSt[0] = act_path_coeff[2]*ferm_epsilon ; mThreeSt[0] = -ThreeSt[0];
388  FiveSt[0] = act_path_coeff[3]*ferm_epsilon ; mFiveSt[0] = -FiveSt[0];
389  SevenSt[0] = act_path_coeff[4]*ferm_epsilon ; mSevenSt[0] = -SevenSt[0];
390  Lepage[0] = act_path_coeff[5]*ferm_epsilon ; mLepage[0] = -Lepage[0];
391 
392  ferm_epsilon = 2.0*weight2*eps;
393  OneLink[1] = act_path_coeff[0]*ferm_epsilon ;
394  Naik[1] = act_path_coeff[1]*ferm_epsilon ; mNaik[1] = -Naik[1];
395  ThreeSt[1] = act_path_coeff[2]*ferm_epsilon ; mThreeSt[1] = -ThreeSt[1];
396  FiveSt[1] = act_path_coeff[3]*ferm_epsilon ; mFiveSt[1] = -FiveSt[1];
397  SevenSt[1] = act_path_coeff[4]*ferm_epsilon ; mSevenSt[1] = -SevenSt[1];
398  Lepage[1] = act_path_coeff[5]*ferm_epsilon ; mLepage[1] = -Lepage[1];
399 
400  for(mu=0;mu<8;mu++){
401  DirectLinks[mu] = 0 ;
402  }
403 
404  for(mu=0;mu<8;mu++){
405  tempvec[mu] = (half_wilson_vector *)malloc( sites_on_node*sizeof(half_wilson_vector) );
406  }
407 
408  for(sig=0; sig < 8; sig++){
409  for(mu = 0; mu < 8; mu++){
410  if ( (mu == sig) || (mu == OPP_DIR(sig))){
411  continue;
412  }
413 
414  // 3 link path
415  u_shift_hw(temp_x, Pmu, OPP_DIR(mu), sitelink);
416  u_shift_hw(Pmu, P3, sig, sitelink);
417  if (GOES_FORWARDS(sig)){
418  add_3f_force_to_mom<su3_matrix>(P3, Pmu, sig, mThreeSt, mom);
419  }
420  for(nu=0; nu < 8; nu++){
421  if (nu == sig || nu == OPP_DIR(sig)
422  || nu == mu || nu == OPP_DIR(mu)){
423  continue;
424  }
425 
426  // 5 link path
427  u_shift_hw(Pmu, Pnumu, OPP_DIR(nu), sitelink);
428  u_shift_hw(Pnumu, P5, sig, sitelink);
429  if (GOES_FORWARDS(sig)){
430  add_3f_force_to_mom<su3_matrix>(P5, Pnumu, sig, FiveSt, mom);
431  }
432 
433 
434  for(rho =0; rho < 8; rho++){
435  if (rho == sig || rho == OPP_DIR(sig)
436  || rho == mu || rho == OPP_DIR(mu)
437  || rho == nu || rho == OPP_DIR(nu)){
438  continue;
439  }
440 
441  //7 link path
442  u_shift_hw(Pnumu, Prhonumu, OPP_DIR(rho), sitelink);
443  u_shift_hw(Prhonumu, P7,sig, sitelink);
444  if(GOES_FORWARDS(sig)){
445  add_3f_force_to_mom<su3_matrix>(P7, Prhonumu, sig, mSevenSt, mom) ;
446  }
447  u_shift_hw(P7, P7rho, rho, sitelink);
448  side_link_3f_force<su3_matrix>(rho,sig,SevenSt, Pnumu, P7, Prhonumu, P7rho, mom);
449  if(FiveSt[0] != 0)coeff[0] = SevenSt[0]/FiveSt[0] ; else coeff[0] = 0;
450  if(FiveSt[1] != 0)coeff[1] = SevenSt[1]/FiveSt[1] ; else coeff[1] = 0;
451  for(i=0;i < V; i++){
452  scalar_mult_add_su3_vector(&(P5[i].h[0]),&(P7rho[i].h[0]), coeff[0],
453  &(P5[i].h[0]));
454  scalar_mult_add_su3_vector(&(P5[i].h[1]),&(P7rho[i].h[1]), coeff[1],
455  &(P5[i].h[1]));
456  }
457 
458  }//rho
459 
460  u_shift_hw(P5,P5nu, nu, sitelink);
461  side_link_3f_force<su3_matrix>(nu,sig,mFiveSt,Pmu,P5, Pnumu,P5nu, mom);
462  if(ThreeSt[0] != 0)coeff[0] = FiveSt[0]/ThreeSt[0] ; else coeff[0] = 0;
463  if(ThreeSt[1] != 0)coeff[1] = FiveSt[1]/ThreeSt[1] ; else coeff[1] = 0;
464  for(i=0; i < V; i++){
465  scalar_mult_add_su3_vector(&(P3[i].h[0]),&(P5nu[i].h[0]),coeff[0],&(P3[i].h[0]));
466  scalar_mult_add_su3_vector(&(P3[i].h[1]),&(P5nu[i].h[1]),coeff[1],&(P3[i].h[1]));
467  }
468 
469  }//nu
470 
471  //Lepage term
472  u_shift_hw(Pmu, Pnumu, OPP_DIR(mu), sitelink);
473  u_shift_hw(Pnumu, P5, sig, sitelink);
474  if(GOES_FORWARDS(sig)){
475  add_3f_force_to_mom<su3_matrix>(P5, Pnumu, sig, Lepage, mom);
476  }
477 
478  u_shift_hw(P5,P5nu, mu, sitelink);
479  side_link_3f_force<su3_matrix>(mu, sig, mLepage, Pmu, P5, Pnumu, P5nu, mom);
480  if(ThreeSt[0] != 0) coeff[0] = Lepage[0]/ThreeSt[0] ; else coeff[0] = 0;
481  if(ThreeSt[1] != 0) coeff[1] = Lepage[1]/ThreeSt[1] ; else coeff[1] = 0;
482  for(i=0;i < V;i++){
483  scalar_mult_add_su3_vector(&(P3[i].h[0]),&(P5nu[i].h[0]),coeff[0],&(P3[i].h[0]));
484  scalar_mult_add_su3_vector(&(P3[i].h[1]),&(P5nu[i].h[1]),coeff[1],&(P3[i].h[1]));
485  }
486 
487  if(GOES_FORWARDS(mu)) {
488  u_shift_hw(P3,P3mu, mu, sitelink);
489  }
490  side_link_3f_force<su3_matrix>(mu, sig, ThreeSt, temp_x, P3, Pmu, P3mu, mom);
491 
492  //One link term and Naik term
493  if( (!DirectLinks[mu]) ){
494  DirectLinks[mu]=1 ;
495  if(GOES_BACKWARDS(mu)){
496  //one link term
497  add_3f_force_to_mom<su3_matrix>(Pmu, temp_x, OPP_DIR(mu), OneLink, mom);
498  u_shift_hw(temp_x, Popmu, mu, sitelink);
499  add_3f_force_to_mom<su3_matrix>(Pnumu, Popmu, OPP_DIR(mu), mNaik, mom);
500  u_shift_hw(Pnumu, Pmumumu, OPP_DIR(mu), sitelink);
501  add_3f_force_to_mom<su3_matrix>(Pmumumu, temp_x, OPP_DIR(mu), Naik, mom);
502  }else{
503  u_shift_hw(temp_x, Popmu, mu, sitelink);
504  add_3f_force_to_mom<su3_matrix>(Popmu, Pnumu, mu, Naik, mom);
505  }
506  }//if
507  }//mu
508  }//sig
509 
510  for(mu=0;mu<8;mu++){
511  free(tempvec[mu]);
512  }
513 
514 }
515 
516 #undef Pmu
517 #undef Pnumu
518 #undef Prhonumu
519 #undef P7
520 #undef P7rho
521 #undef P7rhonu
522 #undef P5
523 #undef P3
524 #undef P5nu
525 #undef P3mu
526 #undef Popmu
527 #undef Pmumumu
528 
529 
530 
531 void
532 fermion_force_reference(float eps, float weight1, float weight2,
533  void* act_path_coeff, void* temp_x, void* sitelink, void* mom)
534 {
535  do_fermion_force_reference((float)eps, (float)weight1, (float)weight2,
536  (fhalf_wilson_vector*) temp_x, (float*)act_path_coeff,
537  (fsu3_matrix*)sitelink, (fanti_hermitmat*) mom);
538 }
539 
540 
541 
542 void
543 fermion_force_reference(double eps, double weight1, double weight2,
544  void* act_path_coeff, void* temp_x, void* sitelink, void* mom)
545 {
546  do_fermion_force_reference((double)eps, (double)weight1, (double)weight2,
547  (dhalf_wilson_vector*) temp_x, (double*)act_path_coeff,
548  (dsu3_matrix*)sitelink, (danti_hermitmat*) mom);
549 }
550 
551 
#define P7
int V
Definition: test_util.cpp:29
int y[4]
void print_su3_matrix(su3_matrix *a)
#define CSUM(a, b)
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
void make_anti_hermitian(su3_matrix *m3, anti_hermitmat *ah3)
#define CMUL(a, b, c)
int neighborIndexFullLattice(int i, int dx4, int dx3, int dx2, int dx1)
Definition: test_util.cpp:521
#define CMUL_J(a, b, c)
void fermion_force_reference(float eps, float weight1, float weight2, void *act_path_coeff, void *temp_x, void *sitelink, void *mom)
int Vh
#define OPP_DIR(dir)
Definition: force_common.h:16
#define Pmumumu
void su3_adjoint(su3_matrix *a, su3_matrix *b)
FloatingPoint< float > Float
Definition: gtest.h:7350
for(int s=0;s< param.Ls;s++)
#define P5nu
#define P3mu
__constant__ double coeff
#define P5
#define CONJG(a, b)
#define P7rho
int x[4]
#define Pnumu
#define P3
int dx[4]
int Z[4]
Definition: test_util.cpp:28
#define Prhonumu
#define GOES_BACKWARDS(dir)
Definition: force_common.h:18
Main header file for the QUDA library.
#define Popmu
#define Pmu
#define GOES_FORWARDS(dir)
Definition: force_common.h:17
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int sig
VOLATILE spinorFloat * s