quda-ref/v0.7.0/clover__core_8h_source.html

 // *** CUDA CLOVER ***


 #define CLOVER_SHARED_FLOATS_PER_THREAD 0


 #if ((CUDA_VERSION >= 4010) && (__COMPUTE_CAPABILITY__ >= 200)) // NVVM compiler

 #define VOLATILE

 #else // Open64 compiler

 #define VOLATILE volatile

 #endif

 // input spinor

 #ifdef SPINOR_DOUBLE

 #define spinorFloat double

 #define i00_re I0.x

 #define i00_im I0.y

 #define i01_re I1.x

 #define i01_im I1.y

 #define i02_re I2.x

 #define i02_im I2.y

 #define i10_re I3.x

 #define i10_im I3.y

 #define i11_re I4.x

 #define i11_im I4.y

 #define i12_re I5.x

 #define i12_im I5.y

 #define i20_re I6.x

 #define i20_im I6.y

 #define i21_re I7.x

 #define i21_im I7.y

 #define i22_re I8.x

 #define i22_im I8.y

 #define i30_re I9.x

 #define i30_im I9.y

 #define i31_re I10.x

 #define i31_im I10.y

 #define i32_re I11.x

 #define i32_im I11.y

 #else

 #define spinorFloat float

 #define i00_re I0.x

 #define i00_im I0.y

 #define i01_re I0.z

 #define i01_im I0.w

 #define i02_re I1.x

 #define i02_im I1.y

 #define i10_re I1.z

 #define i10_im I1.w

 #define i11_re I2.x

 #define i11_im I2.y

 #define i12_re I2.z

 #define i12_im I2.w

 #define i20_re I3.x

 #define i20_im I3.y

 #define i21_re I3.z

 #define i21_im I3.w

 #define i22_re I4.x

 #define i22_im I4.y

 #define i30_re I4.z

 #define i30_im I4.w

 #define i31_re I5.x

 #define i31_im I5.y

 #define i32_re I5.z

 #define i32_im I5.w

 #endif // SPINOR_DOUBLE


 // first chiral block of inverted clover term

 #ifdef CLOVER_DOUBLE

 #define c00_00_re C0.x

 #define c01_01_re C0.y

 #define c02_02_re C1.x

 #define c10_10_re C1.y

 #define c11_11_re C2.x

 #define c12_12_re C2.y

 #define c01_00_re C3.x

 #define c01_00_im C3.y

 #define c02_00_re C4.x

 #define c02_00_im C4.y

 #define c10_00_re C5.x

 #define c10_00_im C5.y

 #define c11_00_re C6.x

 #define c11_00_im C6.y

 #define c12_00_re C7.x

 #define c12_00_im C7.y

 #define c02_01_re C8.x

 #define c02_01_im C8.y

 #define c10_01_re C9.x

 #define c10_01_im C9.y

 #define c11_01_re C10.x

 #define c11_01_im C10.y

 #define c12_01_re C11.x

 #define c12_01_im C11.y

 #define c10_02_re C12.x

 #define c10_02_im C12.y

 #define c11_02_re C13.x

 #define c11_02_im C13.y

 #define c12_02_re C14.x

 #define c12_02_im C14.y

 #define c11_10_re C15.x

 #define c11_10_im C15.y

 #define c12_10_re C16.x

 #define c12_10_im C16.y

 #define c12_11_re C17.x

 #define c12_11_im C17.y

 #else

 #define c00_00_re C0.x

 #define c01_01_re C0.y

 #define c02_02_re C0.z

 #define c10_10_re C0.w

 #define c11_11_re C1.x

 #define c12_12_re C1.y

 #define c01_00_re C1.z

 #define c01_00_im C1.w

 #define c02_00_re C2.x

 #define c02_00_im C2.y

 #define c10_00_re C2.z

 #define c10_00_im C2.w

 #define c11_00_re C3.x

 #define c11_00_im C3.y

 #define c12_00_re C3.z

 #define c12_00_im C3.w

 #define c02_01_re C4.x

 #define c02_01_im C4.y

 #define c10_01_re C4.z

 #define c10_01_im C4.w

 #define c11_01_re C5.x

 #define c11_01_im C5.y

 #define c12_01_re C5.z

 #define c12_01_im C5.w

 #define c10_02_re C6.x

 #define c10_02_im C6.y

 #define c11_02_re C6.z

 #define c11_02_im C6.w

 #define c12_02_re C7.x

 #define c12_02_im C7.y

 #define c11_10_re C7.z

 #define c11_10_im C7.w

 #define c12_10_re C8.x

 #define c12_10_im C8.y

 #define c12_11_re C8.z

 #define c12_11_im C8.w

 #endif // CLOVER_DOUBLE


 #define c00_01_re (+c01_00_re)

 #define c00_01_im (-c01_00_im)

 #define c00_02_re (+c02_00_re)

 #define c00_02_im (-c02_00_im)

 #define c01_02_re (+c02_01_re)

 #define c01_02_im (-c02_01_im)

 #define c00_10_re (+c10_00_re)

 #define c00_10_im (-c10_00_im)

 #define c01_10_re (+c10_01_re)

 #define c01_10_im (-c10_01_im)

 #define c02_10_re (+c10_02_re)

 #define c02_10_im (-c10_02_im)

 #define c00_11_re (+c11_00_re)

 #define c00_11_im (-c11_00_im)

 #define c01_11_re (+c11_01_re)

 #define c01_11_im (-c11_01_im)

 #define c02_11_re (+c11_02_re)

 #define c02_11_im (-c11_02_im)

 #define c10_11_re (+c11_10_re)

 #define c10_11_im (-c11_10_im)

 #define c00_12_re (+c12_00_re)

 #define c00_12_im (-c12_00_im)

 #define c01_12_re (+c12_01_re)

 #define c01_12_im (-c12_01_im)

 #define c02_12_re (+c12_02_re)

 #define c02_12_im (-c12_02_im)

 #define c10_12_re (+c12_10_re)

 #define c10_12_im (-c12_10_im)

 #define c11_12_re (+c12_11_re)

 #define c11_12_im (-c12_11_im)


 // second chiral block of inverted clover term (reuses C0,...,C9)

 #define c20_20_re c00_00_re

 #define c21_20_re c01_00_re

 #define c21_20_im c01_00_im

 #define c22_20_re c02_00_re

 #define c22_20_im c02_00_im

 #define c30_20_re c10_00_re

 #define c30_20_im c10_00_im

 #define c31_20_re c11_00_re

 #define c31_20_im c11_00_im

 #define c32_20_re c12_00_re

 #define c32_20_im c12_00_im

 #define c20_21_re c00_01_re

 #define c20_21_im c00_01_im

 #define c21_21_re c01_01_re

 #define c22_21_re c02_01_re

 #define c22_21_im c02_01_im

 #define c30_21_re c10_01_re

 #define c30_21_im c10_01_im

 #define c31_21_re c11_01_re

 #define c31_21_im c11_01_im

 #define c32_21_re c12_01_re

 #define c32_21_im c12_01_im

 #define c20_22_re c00_02_re

 #define c20_22_im c00_02_im

 #define c21_22_re c01_02_re

 #define c21_22_im c01_02_im

 #define c22_22_re c02_02_re

 #define c30_22_re c10_02_re

 #define c30_22_im c10_02_im

 #define c31_22_re c11_02_re

 #define c31_22_im c11_02_im

 #define c32_22_re c12_02_re

 #define c32_22_im c12_02_im

 #define c20_30_re c00_10_re

 #define c20_30_im c00_10_im

 #define c21_30_re c01_10_re

 #define c21_30_im c01_10_im

 #define c22_30_re c02_10_re

 #define c22_30_im c02_10_im

 #define c30_30_re c10_10_re

 #define c31_30_re c11_10_re

 #define c31_30_im c11_10_im

 #define c32_30_re c12_10_re

 #define c32_30_im c12_10_im

 #define c20_31_re c00_11_re

 #define c20_31_im c00_11_im

 #define c21_31_re c01_11_re

 #define c21_31_im c01_11_im

 #define c22_31_re c02_11_re

 #define c22_31_im c02_11_im

 #define c30_31_re c10_11_re

 #define c30_31_im c10_11_im

 #define c31_31_re c11_11_re

 #define c32_31_re c12_11_re

 #define c32_31_im c12_11_im

 #define c20_32_re c00_12_re

 #define c20_32_im c00_12_im

 #define c21_32_re c01_12_re

 #define c21_32_im c01_12_im

 #define c22_32_re c02_12_re

 #define c22_32_im c02_12_im

 #define c30_32_re c10_12_re

 #define c30_32_im c10_12_im

 #define c31_32_re c11_12_re

 #define c31_32_im c11_12_im

 #define c32_32_re c12_12_re


 // output spinor

 VOLATILE spinorFloat o00_re;

 VOLATILE spinorFloat o00_im;

 VOLATILE spinorFloat o01_re;

 VOLATILE spinorFloat o01_im;

 VOLATILE spinorFloat o02_re;

 VOLATILE spinorFloat o02_im;

 VOLATILE spinorFloat o10_re;

 VOLATILE spinorFloat o10_im;

 VOLATILE spinorFloat o11_re;

 VOLATILE spinorFloat o11_im;

 VOLATILE spinorFloat o12_re;

 VOLATILE spinorFloat o12_im;

 VOLATILE spinorFloat o20_re;

 VOLATILE spinorFloat o20_im;

 VOLATILE spinorFloat o21_re;

 VOLATILE spinorFloat o21_im;

 VOLATILE spinorFloat o22_re;

 VOLATILE spinorFloat o22_im;

 VOLATILE spinorFloat o30_re;

 VOLATILE spinorFloat o30_im;

 VOLATILE spinorFloat o31_re;

 VOLATILE spinorFloat o31_im;

 VOLATILE spinorFloat o32_re;

 VOLATILE spinorFloat o32_im;


 #include "read_clover.h"

 #include "io_spinor.h"


 int sid = blockIdx.x*blockDim.x + threadIdx.x;

 if (sid >= param.threads) return;


 // read spinor from device memory

 READ_SPINOR(SPINORTEX, param.sp_stride, sid, sid);

 {

   // change to chiral basis

   {

     spinorFloat a00_re = -i10_re - i30_re;

     spinorFloat a00_im = -i10_im - i30_im;

     spinorFloat a10_re =  i00_re + i20_re;

     spinorFloat a10_im =  i00_im + i20_im;

     spinorFloat a20_re = -i10_re + i30_re;

     spinorFloat a20_im = -i10_im + i30_im;

     spinorFloat a30_re =  i00_re - i20_re;

     spinorFloat a30_im =  i00_im - i20_im;


     o00_re = a00_re;  o00_im = a00_im;

     o10_re = a10_re;  o10_im = a10_im;

     o20_re = a20_re;  o20_im = a20_im;

     o30_re = a30_re;  o30_im = a30_im;

   }


   {

     spinorFloat a01_re = -i11_re - i31_re;

     spinorFloat a01_im = -i11_im - i31_im;

     spinorFloat a11_re =  i01_re + i21_re;

     spinorFloat a11_im =  i01_im + i21_im;

     spinorFloat a21_re = -i11_re + i31_re;

     spinorFloat a21_im = -i11_im + i31_im;

     spinorFloat a31_re =  i01_re - i21_re;

     spinorFloat a31_im =  i01_im - i21_im;


     o01_re = a01_re;  o01_im = a01_im;

     o11_re = a11_re;  o11_im = a11_im;

     o21_re = a21_re;  o21_im = a21_im;

     o31_re = a31_re;  o31_im = a31_im;

   }


   {

     spinorFloat a02_re = -i12_re - i32_re;

     spinorFloat a02_im = -i12_im - i32_im;

     spinorFloat a12_re =  i02_re + i22_re;

     spinorFloat a12_im =  i02_im + i22_im;

     spinorFloat a22_re = -i12_re + i32_re;

     spinorFloat a22_im = -i12_im + i32_im;

     spinorFloat a32_re =  i02_re - i22_re;

     spinorFloat a32_im =  i02_im - i22_im;


     o02_re = a02_re;  o02_im = a02_im;

     o12_re = a12_re;  o12_im = a12_im;

     o22_re = a22_re;  o22_im = a22_im;

     o32_re = a32_re;  o32_im = a32_im;

   }


   // apply first chiral block

   {

     READ_CLOVER(CLOVERTEX, 0)


     spinorFloat a00_re = 0; spinorFloat a00_im = 0;

     spinorFloat a01_re = 0; spinorFloat a01_im = 0;

     spinorFloat a02_re = 0; spinorFloat a02_im = 0;

     spinorFloat a10_re = 0; spinorFloat a10_im = 0;

     spinorFloat a11_re = 0; spinorFloat a11_im = 0;

     spinorFloat a12_re = 0; spinorFloat a12_im = 0;


     a00_re += c00_00_re * o00_re;

     a00_im += c00_00_re * o00_im;

     a00_re += c00_01_re * o01_re;

     a00_re -= c00_01_im * o01_im;

     a00_im += c00_01_re * o01_im;

     a00_im += c00_01_im * o01_re;

     a00_re += c00_02_re * o02_re;

     a00_re -= c00_02_im * o02_im;

     a00_im += c00_02_re * o02_im;

     a00_im += c00_02_im * o02_re;

     a00_re += c00_10_re * o10_re;

     a00_re -= c00_10_im * o10_im;

     a00_im += c00_10_re * o10_im;

     a00_im += c00_10_im * o10_re;

     a00_re += c00_11_re * o11_re;

     a00_re -= c00_11_im * o11_im;

     a00_im += c00_11_re * o11_im;

     a00_im += c00_11_im * o11_re;

     a00_re += c00_12_re * o12_re;

     a00_re -= c00_12_im * o12_im;

     a00_im += c00_12_re * o12_im;

     a00_im += c00_12_im * o12_re;


     a01_re += c01_00_re * o00_re;

     a01_re -= c01_00_im * o00_im;

     a01_im += c01_00_re * o00_im;

     a01_im += c01_00_im * o00_re;

     a01_re += c01_01_re * o01_re;

     a01_im += c01_01_re * o01_im;

     a01_re += c01_02_re * o02_re;

     a01_re -= c01_02_im * o02_im;

     a01_im += c01_02_re * o02_im;

     a01_im += c01_02_im * o02_re;

     a01_re += c01_10_re * o10_re;

     a01_re -= c01_10_im * o10_im;

     a01_im += c01_10_re * o10_im;

     a01_im += c01_10_im * o10_re;

     a01_re += c01_11_re * o11_re;

     a01_re -= c01_11_im * o11_im;

     a01_im += c01_11_re * o11_im;

     a01_im += c01_11_im * o11_re;

     a01_re += c01_12_re * o12_re;

     a01_re -= c01_12_im * o12_im;

     a01_im += c01_12_re * o12_im;

     a01_im += c01_12_im * o12_re;


     a02_re += c02_00_re * o00_re;

     a02_re -= c02_00_im * o00_im;

     a02_im += c02_00_re * o00_im;

     a02_im += c02_00_im * o00_re;

     a02_re += c02_01_re * o01_re;

     a02_re -= c02_01_im * o01_im;

     a02_im += c02_01_re * o01_im;

     a02_im += c02_01_im * o01_re;

     a02_re += c02_02_re * o02_re;

     a02_im += c02_02_re * o02_im;

     a02_re += c02_10_re * o10_re;

     a02_re -= c02_10_im * o10_im;

     a02_im += c02_10_re * o10_im;

     a02_im += c02_10_im * o10_re;

     a02_re += c02_11_re * o11_re;

     a02_re -= c02_11_im * o11_im;

     a02_im += c02_11_re * o11_im;

     a02_im += c02_11_im * o11_re;

     a02_re += c02_12_re * o12_re;

     a02_re -= c02_12_im * o12_im;

     a02_im += c02_12_re * o12_im;

     a02_im += c02_12_im * o12_re;


     a10_re += c10_00_re * o00_re;

     a10_re -= c10_00_im * o00_im;

     a10_im += c10_00_re * o00_im;

     a10_im += c10_00_im * o00_re;

     a10_re += c10_01_re * o01_re;

     a10_re -= c10_01_im * o01_im;

     a10_im += c10_01_re * o01_im;

     a10_im += c10_01_im * o01_re;

     a10_re += c10_02_re * o02_re;

     a10_re -= c10_02_im * o02_im;

     a10_im += c10_02_re * o02_im;

     a10_im += c10_02_im * o02_re;

     a10_re += c10_10_re * o10_re;

     a10_im += c10_10_re * o10_im;

     a10_re += c10_11_re * o11_re;

     a10_re -= c10_11_im * o11_im;

     a10_im += c10_11_re * o11_im;

     a10_im += c10_11_im * o11_re;

     a10_re += c10_12_re * o12_re;

     a10_re -= c10_12_im * o12_im;

     a10_im += c10_12_re * o12_im;

     a10_im += c10_12_im * o12_re;


     a11_re += c11_00_re * o00_re;

     a11_re -= c11_00_im * o00_im;

     a11_im += c11_00_re * o00_im;

     a11_im += c11_00_im * o00_re;

     a11_re += c11_01_re * o01_re;

     a11_re -= c11_01_im * o01_im;

     a11_im += c11_01_re * o01_im;

     a11_im += c11_01_im * o01_re;

     a11_re += c11_02_re * o02_re;

     a11_re -= c11_02_im * o02_im;

     a11_im += c11_02_re * o02_im;

     a11_im += c11_02_im * o02_re;

     a11_re += c11_10_re * o10_re;

     a11_re -= c11_10_im * o10_im;

     a11_im += c11_10_re * o10_im;

     a11_im += c11_10_im * o10_re;

     a11_re += c11_11_re * o11_re;

     a11_im += c11_11_re * o11_im;

     a11_re += c11_12_re * o12_re;

     a11_re -= c11_12_im * o12_im;

     a11_im += c11_12_re * o12_im;

     a11_im += c11_12_im * o12_re;


     a12_re += c12_00_re * o00_re;

     a12_re -= c12_00_im * o00_im;

     a12_im += c12_00_re * o00_im;

     a12_im += c12_00_im * o00_re;

     a12_re += c12_01_re * o01_re;

     a12_re -= c12_01_im * o01_im;

     a12_im += c12_01_re * o01_im;

     a12_im += c12_01_im * o01_re;

     a12_re += c12_02_re * o02_re;

     a12_re -= c12_02_im * o02_im;

     a12_im += c12_02_re * o02_im;

     a12_im += c12_02_im * o02_re;

     a12_re += c12_10_re * o10_re;

     a12_re -= c12_10_im * o10_im;

     a12_im += c12_10_re * o10_im;

     a12_im += c12_10_im * o10_re;

     a12_re += c12_11_re * o11_re;

     a12_re -= c12_11_im * o11_im;

     a12_im += c12_11_re * o11_im;

     a12_im += c12_11_im * o11_re;

     a12_re += c12_12_re * o12_re;

     a12_im += c12_12_re * o12_im;


     o00_re = a00_re;  o00_im = a00_im;

     o01_re = a01_re;  o01_im = a01_im;

     o02_re = a02_re;  o02_im = a02_im;

     o10_re = a10_re;  o10_im = a10_im;

     o11_re = a11_re;  o11_im = a11_im;

     o12_re = a12_re;  o12_im = a12_im;


   }


   // apply second chiral block

   {

     READ_CLOVER(CLOVERTEX, 1)


     spinorFloat a20_re = 0; spinorFloat a20_im = 0;

     spinorFloat a21_re = 0; spinorFloat a21_im = 0;

     spinorFloat a22_re = 0; spinorFloat a22_im = 0;

     spinorFloat a30_re = 0; spinorFloat a30_im = 0;

     spinorFloat a31_re = 0; spinorFloat a31_im = 0;

     spinorFloat a32_re = 0; spinorFloat a32_im = 0;


     a20_re += c20_20_re * o20_re;

     a20_im += c20_20_re * o20_im;

     a20_re += c20_21_re * o21_re;

     a20_re -= c20_21_im * o21_im;

     a20_im += c20_21_re * o21_im;

     a20_im += c20_21_im * o21_re;

     a20_re += c20_22_re * o22_re;

     a20_re -= c20_22_im * o22_im;

     a20_im += c20_22_re * o22_im;

     a20_im += c20_22_im * o22_re;

     a20_re += c20_30_re * o30_re;

     a20_re -= c20_30_im * o30_im;

     a20_im += c20_30_re * o30_im;

     a20_im += c20_30_im * o30_re;

     a20_re += c20_31_re * o31_re;

     a20_re -= c20_31_im * o31_im;

     a20_im += c20_31_re * o31_im;

     a20_im += c20_31_im * o31_re;

     a20_re += c20_32_re * o32_re;

     a20_re -= c20_32_im * o32_im;

     a20_im += c20_32_re * o32_im;

     a20_im += c20_32_im * o32_re;


     a21_re += c21_20_re * o20_re;

     a21_re -= c21_20_im * o20_im;

     a21_im += c21_20_re * o20_im;

     a21_im += c21_20_im * o20_re;

     a21_re += c21_21_re * o21_re;

     a21_im += c21_21_re * o21_im;

     a21_re += c21_22_re * o22_re;

     a21_re -= c21_22_im * o22_im;

     a21_im += c21_22_re * o22_im;

     a21_im += c21_22_im * o22_re;

     a21_re += c21_30_re * o30_re;

     a21_re -= c21_30_im * o30_im;

     a21_im += c21_30_re * o30_im;

     a21_im += c21_30_im * o30_re;

     a21_re += c21_31_re * o31_re;

     a21_re -= c21_31_im * o31_im;

     a21_im += c21_31_re * o31_im;

     a21_im += c21_31_im * o31_re;

     a21_re += c21_32_re * o32_re;

     a21_re -= c21_32_im * o32_im;

     a21_im += c21_32_re * o32_im;

     a21_im += c21_32_im * o32_re;


     a22_re += c22_20_re * o20_re;

     a22_re -= c22_20_im * o20_im;

     a22_im += c22_20_re * o20_im;

     a22_im += c22_20_im * o20_re;

     a22_re += c22_21_re * o21_re;

     a22_re -= c22_21_im * o21_im;

     a22_im += c22_21_re * o21_im;

     a22_im += c22_21_im * o21_re;

     a22_re += c22_22_re * o22_re;

     a22_im += c22_22_re * o22_im;

     a22_re += c22_30_re * o30_re;

     a22_re -= c22_30_im * o30_im;

     a22_im += c22_30_re * o30_im;

     a22_im += c22_30_im * o30_re;

     a22_re += c22_31_re * o31_re;

     a22_re -= c22_31_im * o31_im;

     a22_im += c22_31_re * o31_im;

     a22_im += c22_31_im * o31_re;

     a22_re += c22_32_re * o32_re;

     a22_re -= c22_32_im * o32_im;

     a22_im += c22_32_re * o32_im;

     a22_im += c22_32_im * o32_re;


     a30_re += c30_20_re * o20_re;

     a30_re -= c30_20_im * o20_im;

     a30_im += c30_20_re * o20_im;

     a30_im += c30_20_im * o20_re;

     a30_re += c30_21_re * o21_re;

     a30_re -= c30_21_im * o21_im;

     a30_im += c30_21_re * o21_im;

     a30_im += c30_21_im * o21_re;

     a30_re += c30_22_re * o22_re;

     a30_re -= c30_22_im * o22_im;

     a30_im += c30_22_re * o22_im;

     a30_im += c30_22_im * o22_re;

     a30_re += c30_30_re * o30_re;

     a30_im += c30_30_re * o30_im;

     a30_re += c30_31_re * o31_re;

     a30_re -= c30_31_im * o31_im;

     a30_im += c30_31_re * o31_im;

     a30_im += c30_31_im * o31_re;

     a30_re += c30_32_re * o32_re;

     a30_re -= c30_32_im * o32_im;

     a30_im += c30_32_re * o32_im;

     a30_im += c30_32_im * o32_re;


     a31_re += c31_20_re * o20_re;

     a31_re -= c31_20_im * o20_im;

     a31_im += c31_20_re * o20_im;

     a31_im += c31_20_im * o20_re;

     a31_re += c31_21_re * o21_re;

     a31_re -= c31_21_im * o21_im;

     a31_im += c31_21_re * o21_im;

     a31_im += c31_21_im * o21_re;

     a31_re += c31_22_re * o22_re;

     a31_re -= c31_22_im * o22_im;

     a31_im += c31_22_re * o22_im;

     a31_im += c31_22_im * o22_re;

     a31_re += c31_30_re * o30_re;

     a31_re -= c31_30_im * o30_im;

     a31_im += c31_30_re * o30_im;

     a31_im += c31_30_im * o30_re;

     a31_re += c31_31_re * o31_re;

     a31_im += c31_31_re * o31_im;

     a31_re += c31_32_re * o32_re;

     a31_re -= c31_32_im * o32_im;

     a31_im += c31_32_re * o32_im;

     a31_im += c31_32_im * o32_re;


     a32_re += c32_20_re * o20_re;

     a32_re -= c32_20_im * o20_im;

     a32_im += c32_20_re * o20_im;

     a32_im += c32_20_im * o20_re;

     a32_re += c32_21_re * o21_re;

     a32_re -= c32_21_im * o21_im;

     a32_im += c32_21_re * o21_im;

     a32_im += c32_21_im * o21_re;

     a32_re += c32_22_re * o22_re;

     a32_re -= c32_22_im * o22_im;

     a32_im += c32_22_re * o22_im;

     a32_im += c32_22_im * o22_re;

     a32_re += c32_30_re * o30_re;

     a32_re -= c32_30_im * o30_im;

     a32_im += c32_30_re * o30_im;

     a32_im += c32_30_im * o30_re;

     a32_re += c32_31_re * o31_re;

     a32_re -= c32_31_im * o31_im;

     a32_im += c32_31_re * o31_im;

     a32_im += c32_31_im * o31_re;

     a32_re += c32_32_re * o32_re;

     a32_im += c32_32_re * o32_im;


     o20_re = a20_re;  o20_im = a20_im;

     o21_re = a21_re;  o21_im = a21_im;

     o22_re = a22_re;  o22_im = a22_im;

     o30_re = a30_re;  o30_im = a30_im;

     o31_re = a31_re;  o31_im = a31_im;

     o32_re = a32_re;  o32_im = a32_im;


   }


   // change back from chiral basis

   // (note: required factor of 1/2 is included in clover term normalization)

   {

     spinorFloat a00_re =  o10_re + o30_re;

     spinorFloat a00_im =  o10_im + o30_im;

     spinorFloat a10_re = -o00_re - o20_re;

     spinorFloat a10_im = -o00_im - o20_im;

     spinorFloat a20_re =  o10_re - o30_re;

     spinorFloat a20_im =  o10_im - o30_im;

     spinorFloat a30_re = -o00_re + o20_re;

     spinorFloat a30_im = -o00_im + o20_im;


     o00_re = a00_re;  o00_im = a00_im;

     o10_re = a10_re;  o10_im = a10_im;

     o20_re = a20_re;  o20_im = a20_im;

     o30_re = a30_re;  o30_im = a30_im;

   }


   {

     spinorFloat a01_re =  o11_re + o31_re;

     spinorFloat a01_im =  o11_im + o31_im;

     spinorFloat a11_re = -o01_re - o21_re;

     spinorFloat a11_im = -o01_im - o21_im;

     spinorFloat a21_re =  o11_re - o31_re;

     spinorFloat a21_im =  o11_im - o31_im;

     spinorFloat a31_re = -o01_re + o21_re;

     spinorFloat a31_im = -o01_im + o21_im;


     o01_re = a01_re;  o01_im = a01_im;

     o11_re = a11_re;  o11_im = a11_im;

     o21_re = a21_re;  o21_im = a21_im;

     o31_re = a31_re;  o31_im = a31_im;

   }


   {

     spinorFloat a02_re =  o12_re + o32_re;

     spinorFloat a02_im =  o12_im + o32_im;

     spinorFloat a12_re = -o02_re - o22_re;

     spinorFloat a12_im = -o02_im - o22_im;

     spinorFloat a22_re =  o12_re - o32_re;

     spinorFloat a22_im =  o12_im - o32_im;

     spinorFloat a32_re = -o02_re + o22_re;

     spinorFloat a32_im = -o02_im + o22_im;


     o02_re = a02_re;  o02_im = a02_im;

     o12_re = a12_re;  o12_im = a12_im;

     o22_re = a22_re;  o22_im = a22_im;

     o32_re = a32_re;  o32_im = a32_im;

   }


 #ifdef DSLASH_XPAY


   READ_ACCUM(ACCUMTEX, param.sp_stride)


   o00_re = a*o00_re+acc00_re;

   o00_im = a*o00_im+acc00_im;

   o01_re = a*o01_re+acc01_re;

   o01_im = a*o01_im+acc01_im;

   o02_re = a*o02_re+acc02_re;

   o02_im = a*o02_im+acc02_im;

   o10_re = a*o10_re+acc10_re;

   o10_im = a*o10_im+acc10_im;

   o11_re = a*o11_re+acc11_re;

   o11_im = a*o11_im+acc11_im;

   o12_re = a*o12_re+acc12_re;

   o12_im = a*o12_im+acc12_im;

   o20_re = a*o20_re+acc20_re;

   o20_im = a*o20_im+acc20_im;

   o21_re = a*o21_re+acc21_re;

   o21_im = a*o21_im+acc21_im;

   o22_re = a*o22_re+acc22_re;

   o22_im = a*o22_im+acc22_im;

   o30_re = a*o30_re+acc30_re;

   o30_im = a*o30_im+acc30_im;

   o31_re = a*o31_re+acc31_re;

   o31_im = a*o31_im+acc31_im;

   o32_re = a*o32_re+acc32_re;

   o32_im = a*o32_im+acc32_im;

 #endif // DSLASH_XPAY

 }


 // write spinor field back to device memory

 WRITE_SPINOR(param.sp_stride);


 // undefine to prevent warning when precision is changed

 #undef spinorFloat

 #undef i00_re

 #undef i00_im

 #undef i01_re

 #undef i01_im

 #undef i02_re

 #undef i02_im

 #undef i10_re

 #undef i10_im

 #undef i11_re

 #undef i11_im

 #undef i12_re

 #undef i12_im

 #undef i20_re

 #undef i20_im

 #undef i21_re

 #undef i21_im

 #undef i22_re

 #undef i22_im

 #undef i30_re

 #undef i30_im

 #undef i31_re

 #undef i31_im

 #undef i32_re

 #undef i32_im


 #undef c00_00_re

 #undef c01_01_re

 #undef c02_02_re

 #undef c10_10_re

 #undef c11_11_re

 #undef c12_12_re

 #undef c01_00_re

 #undef c01_00_im

 #undef c02_00_re

 #undef c02_00_im

 #undef c10_00_re

 #undef c10_00_im

 #undef c11_00_re

 #undef c11_00_im

 #undef c12_00_re

 #undef c12_00_im

 #undef c02_01_re

 #undef c02_01_im

 #undef c10_01_re

 #undef c10_01_im

 #undef c11_01_re

 #undef c11_01_im

 #undef c12_01_re

 #undef c12_01_im

 #undef c10_02_re

 #undef c10_02_im

 #undef c11_02_re

 #undef c11_02_im

 #undef c12_02_re

 #undef c12_02_im

 #undef c11_10_re

 #undef c11_10_im

 #undef c12_10_re

 #undef c12_10_im

 #undef c12_11_re

 #undef c12_11_im


 #undef VOLATILE

i22_im
#define i22_im
Definition: clover_core.h:57

c31_22_im
#define c31_22_im
Definition: clover_core.h:205

c21_22_re
#define c21_22_re
Definition: clover_core.h:199

acc02_re
#define acc02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:96

acc21_im
#define acc21_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:107

o11_re
VOLATILE spinorFloat o11_re
Definition: clover_core.h:251

c22_30_re
#define c22_30_re
Definition: clover_core.h:212

c22_21_im
#define c22_21_im
Definition: clover_core.h:190

a22_re
#define a22_re
Definition: llfat_core.h:131

c31_32_im
#define c31_32_im
Definition: clover_core.h:239

o31_re
VOLATILE spinorFloat o31_re
Definition: clover_core.h:263

c32_31_re
#define c32_31_re
Definition: clover_core.h:228

c31_30_im
#define c31_30_im
Definition: clover_core.h:216

i11_im
#define i11_im
Definition: clover_core.h:49

c00_02_re
#define c00_02_re
Definition: clover_core.h:145

c11_10_im
#define c11_10_im
Definition: clover_core.h:136

c11_00_re
#define c11_00_re
Definition: clover_core.h:117

acc22_im
#define acc22_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:109

c20_30_im
#define c20_30_im
Definition: clover_core.h:209

c00_00_re
#define c00_00_re
Definition: clover_core.h:105

c12_10_re
#define c12_10_re
Definition: clover_core.h:137

c00_10_im
#define c00_10_im
Definition: clover_core.h:150

c21_30_re
#define c21_30_re
Definition: clover_core.h:210

o21_re
VOLATILE spinorFloat o21_re
Definition: clover_core.h:257

acc00_im
#define acc00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:93

c21_30_im
#define c21_30_im
Definition: clover_core.h:211

o20_im
VOLATILE spinorFloat o20_im
Definition: clover_core.h:256

i00_im
#define i00_im
Definition: clover_core.h:41

CLOVERTEX
#define CLOVERTEX
Definition: clover_def.h:101

c31_30_re
#define c31_30_re
Definition: clover_core.h:215

c10_00_re
#define c10_00_re
Definition: clover_core.h:115

o12_re
VOLATILE spinorFloat o12_re
Definition: clover_core.h:253

READ_SPINOR
READ_SPINOR(SPINORTEX, param.sp_stride, sid, sid)

a02_im
#define a02_im
Definition: llfat_core.h:120

c12_02_re
#define c12_02_re
Definition: clover_core.h:133

c22_32_im
#define c22_32_im
Definition: clover_core.h:235

c10_11_re
#define c10_11_re
Definition: clover_core.h:161

o00_re
VOLATILE spinorFloat o00_re
Definition: clover_core.h:243

c31_22_re
#define c31_22_re
Definition: clover_core.h:204

c12_01_re
#define c12_01_re
Definition: clover_core.h:127

c30_31_re
#define c30_31_re
Definition: clover_core.h:225

i10_re
#define i10_re
Definition: clover_core.h:46

VOLATILE
#define VOLATILE
Definition: clover_core.h:9

o32_im
VOLATILE spinorFloat o32_im
Definition: clover_core.h:266

c32_31_im
#define c32_31_im
Definition: clover_core.h:229

a22_im
#define a22_im
Definition: llfat_core.h:132

c02_02_re
#define c02_02_re
Definition: clover_core.h:107

acc31_im
#define acc31_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:113

c30_22_im
#define c30_22_im
Definition: clover_core.h:203

c31_21_im
#define c31_21_im
Definition: clover_core.h:194

c21_31_im
#define c21_31_im
Definition: clover_core.h:222

c02_00_re
#define c02_00_re
Definition: clover_core.h:113

o01_re
VOLATILE spinorFloat o01_re
Definition: clover_core.h:245

acc02_im
#define acc02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:97

c11_00_im
#define c11_00_im
Definition: clover_core.h:118

c32_32_re
#define c32_32_re
Definition: clover_core.h:240

sid
int sid
Definition: clover_core.h:271

o30_re
VOLATILE spinorFloat o30_re
Definition: clover_core.h:261

c30_32_im
#define c30_32_im
Definition: clover_core.h:237

c00_10_re
#define c00_10_re
Definition: clover_core.h:149

c22_31_re
#define c22_31_re
Definition: clover_core.h:223

c11_02_re
#define c11_02_re
Definition: clover_core.h:131

acc32_re
#define acc32_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:114

acc00_re
#define acc00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:92

c11_02_im
#define c11_02_im
Definition: clover_core.h:132

a01_re
#define a01_re
Definition: llfat_core.h:117

c01_01_re
#define c01_01_re
Definition: clover_core.h:106

spinorFloat
#define spinorFloat
Definition: clover_core.h:39

c10_00_im
#define c10_00_im
Definition: clover_core.h:116

c12_00_re
#define c12_00_re
Definition: clover_core.h:119

i01_re
#define i01_re
Definition: clover_core.h:42

c20_20_re
#define c20_20_re
Definition: clover_core.h:175

c21_32_re
#define c21_32_re
Definition: clover_core.h:232

c02_11_im
#define c02_11_im
Definition: clover_core.h:160

c02_10_re
#define c02_10_re
Definition: clover_core.h:153

acc30_re
#define acc30_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:110

acc01_re
#define acc01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:94

a02_re
#define a02_re
Definition: llfat_core.h:119

a20_re
#define a20_re
Definition: llfat_core.h:127

c10_10_re
#define c10_10_re
Definition: clover_core.h:108

a12_im
#define a12_im
Definition: llfat_core.h:126

c21_20_im
#define c21_20_im
Definition: clover_core.h:177

acc21_re
#define acc21_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:106

c10_11_im
#define c10_11_im
Definition: clover_core.h:162

a20_im
#define a20_im
Definition: llfat_core.h:128

acc31_re
#define acc31_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:112

c12_02_im
#define c12_02_im
Definition: clover_core.h:134

c01_12_im
#define c01_12_im
Definition: clover_core.h:166

c12_12_re
#define c12_12_re
Definition: clover_core.h:110

c22_30_im
#define c22_30_im
Definition: clover_core.h:213

param
QudaGaugeParam param
Definition: pack_test.cpp:17

c12_11_re
#define c12_11_re
Definition: clover_core.h:139

i12_re
#define i12_re
Definition: clover_core.h:50

o10_re
VOLATILE spinorFloat o10_re
Definition: clover_core.h:249

c31_20_re
#define c31_20_re
Definition: clover_core.h:182

c10_02_im
#define c10_02_im
Definition: clover_core.h:130

c01_02_im
#define c01_02_im
Definition: clover_core.h:148

c22_20_im
#define c22_20_im
Definition: clover_core.h:179

i20_im
#define i20_im
Definition: clover_core.h:53

c30_31_im
#define c30_31_im
Definition: clover_core.h:226

o20_re
VOLATILE spinorFloat o20_re
Definition: clover_core.h:255

c01_10_im
#define c01_10_im
Definition: clover_core.h:152

acc12_re
#define acc12_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:102

i11_re
#define i11_re
Definition: clover_core.h:48

o31_im
VOLATILE spinorFloat o31_im
Definition: clover_core.h:264

c30_22_re
#define c30_22_re
Definition: clover_core.h:202

i02_im
#define i02_im
Definition: clover_core.h:45

c00_11_re
#define c00_11_re
Definition: clover_core.h:155

c31_32_re
#define c31_32_re
Definition: clover_core.h:238

i02_re
#define i02_re
Definition: clover_core.h:44

c01_10_re
#define c01_10_re
Definition: clover_core.h:151

o11_im
VOLATILE spinorFloat o11_im
Definition: clover_core.h:252

c32_22_re
#define c32_22_re
Definition: clover_core.h:206

a01_im
#define a01_im
Definition: llfat_core.h:118

c11_10_re
#define c11_10_re
Definition: clover_core.h:135

o30_im
VOLATILE spinorFloat o30_im
Definition: clover_core.h:262

a12_re
#define a12_re
Definition: llfat_core.h:125

c02_12_re
#define c02_12_re
Definition: clover_core.h:167

a11_re
#define a11_re
Definition: llfat_core.h:123

c02_10_im
#define c02_10_im
Definition: clover_core.h:154

acc20_im
#define acc20_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:105

c10_01_im
#define c10_01_im
Definition: clover_core.h:124

o01_im
VOLATILE spinorFloat o01_im
Definition: clover_core.h:246

o02_im
VOLATILE spinorFloat o02_im
Definition: clover_core.h:248

c21_21_re
#define c21_21_re
Definition: clover_core.h:188

c30_21_im
#define c30_21_im
Definition: clover_core.h:192

c22_22_re
#define c22_22_re
Definition: clover_core.h:201

c30_20_im
#define c30_20_im
Definition: clover_core.h:181

c10_12_im
#define c10_12_im
Definition: clover_core.h:170

c11_01_re
#define c11_01_re
Definition: clover_core.h:125

c20_30_re
#define c20_30_re
Definition: clover_core.h:208

acc20_re
#define acc20_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:104

o12_im
VOLATILE spinorFloat o12_im
Definition: clover_core.h:254

c10_02_re
#define c10_02_re
Definition: clover_core.h:129

i22_re
#define i22_re
Definition: clover_core.h:56

c00_01_re
#define c00_01_re
Definition: clover_core.h:143

c01_11_re
#define c01_11_re
Definition: clover_core.h:157

c21_32_im
#define c21_32_im
Definition: clover_core.h:233

c32_21_im
#define c32_21_im
Definition: clover_core.h:196

c01_11_im
#define c01_11_im
Definition: clover_core.h:158

i32_im
#define i32_im
Definition: clover_core.h:63

c21_20_re
#define c21_20_re
Definition: clover_core.h:176

c01_00_im
#define c01_00_im
Definition: clover_core.h:112

io_spinor.h

SPINORTEX
#define SPINORTEX
Definition: clover_def.h:40

c12_00_im
#define c12_00_im
Definition: clover_core.h:120

i30_re
#define i30_re
Definition: clover_core.h:58

c21_22_im
#define c21_22_im
Definition: clover_core.h:200

acc11_re
#define acc11_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:100

i12_im
#define i12_im
Definition: clover_core.h:51

c30_30_re
#define c30_30_re
Definition: clover_core.h:214

c02_12_im
#define c02_12_im
Definition: clover_core.h:168

i21_re
#define i21_re
Definition: clover_core.h:54

acc01_im
#define acc01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:95

c01_00_re
#define c01_00_re
Definition: clover_core.h:111

c02_11_re
#define c02_11_re
Definition: clover_core.h:159

c22_32_re
#define c22_32_re
Definition: clover_core.h:234

c02_01_re
#define c02_01_re
Definition: clover_core.h:121

c30_21_re
#define c30_21_re
Definition: clover_core.h:191

c22_21_re
#define c22_21_re
Definition: clover_core.h:189

c10_01_re
#define c10_01_re
Definition: clover_core.h:123

c32_21_re
#define c32_21_re
Definition: clover_core.h:195

o22_re
VOLATILE spinorFloat o22_re
Definition: clover_core.h:259

o00_im
VOLATILE spinorFloat o00_im
Definition: clover_core.h:244

a00_re
#define a00_re
Definition: llfat_core.h:115

i32_re
#define i32_re
Definition: clover_core.h:62

c12_01_im
#define c12_01_im
Definition: clover_core.h:128

o02_re
VOLATILE spinorFloat o02_re
Definition: clover_core.h:247

c01_12_re
#define c01_12_re
Definition: clover_core.h:165

c20_32_im
#define c20_32_im
Definition: clover_core.h:231

i21_im
#define i21_im
Definition: clover_core.h:55

c11_11_re
#define c11_11_re
Definition: clover_core.h:109

c30_20_re
#define c30_20_re
Definition: clover_core.h:180

c11_12_im
#define c11_12_im
Definition: clover_core.h:172

c32_30_im
#define c32_30_im
Definition: clover_core.h:218

c32_22_im
#define c32_22_im
Definition: clover_core.h:207

c00_02_im
#define c00_02_im
Definition: clover_core.h:146

a11_im
#define a11_im
Definition: llfat_core.h:124

i00_re
#define i00_re
Definition: clover_core.h:40

acc12_im
#define acc12_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:103

c31_31_re
#define c31_31_re
Definition: clover_core.h:227

a10_re
#define a10_re
Definition: llfat_core.h:121

c32_30_re
#define c32_30_re
Definition: clover_core.h:217

c11_12_re
#define c11_12_re
Definition: clover_core.h:171

c10_12_re
#define c10_12_re
Definition: clover_core.h:169

c20_31_re
#define c20_31_re
Definition: clover_core.h:219

acc10_re
#define acc10_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:98

c32_20_re
#define c32_20_re
Definition: clover_core.h:184

c00_11_im
#define c00_11_im
Definition: clover_core.h:156

c11_01_im
#define c11_01_im
Definition: clover_core.h:126

o10_im
VOLATILE spinorFloat o10_im
Definition: clover_core.h:250

c31_21_re
#define c31_21_re
Definition: clover_core.h:193

c12_10_im
#define c12_10_im
Definition: clover_core.h:138

acc22_re
#define acc22_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:108

c00_01_im
#define c00_01_im
Definition: clover_core.h:144

acc30_im
#define acc30_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:111

c30_32_re
#define c30_32_re
Definition: clover_core.h:236

o32_re
VOLATILE spinorFloat o32_re
Definition: clover_core.h:265

c00_12_im
#define c00_12_im
Definition: clover_core.h:164

c02_01_im
#define c02_01_im
Definition: clover_core.h:122

read_clover.h

acc10_im
#define acc10_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:99

a10_im
#define a10_im
Definition: llfat_core.h:122

acc32_im
#define acc32_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:115

a21_re
#define a21_re
Definition: llfat_core.h:129

i01_im
#define i01_im
Definition: clover_core.h:43

WRITE_SPINOR
#define WRITE_SPINOR
Definition: clover_def.h:48

c22_20_re
#define c22_20_re
Definition: clover_core.h:178

c20_21_im
#define c20_21_im
Definition: clover_core.h:187

i31_re
#define i31_re
Definition: clover_core.h:60

c21_31_re
#define c21_31_re
Definition: clover_core.h:221

c20_32_re
#define c20_32_re
Definition: clover_core.h:230

i30_im
#define i30_im
Definition: clover_core.h:59

i10_im
#define i10_im
Definition: clover_core.h:47

c32_20_im
#define c32_20_im
Definition: clover_core.h:185

c22_31_im
#define c22_31_im
Definition: clover_core.h:224

c02_00_im
#define c02_00_im
Definition: clover_core.h:114

i20_re
#define i20_re
Definition: clover_core.h:52

c31_20_im
#define c31_20_im
Definition: clover_core.h:183

a21_im
#define a21_im
Definition: llfat_core.h:130

c20_31_im
#define c20_31_im
Definition: clover_core.h:220

READ_CLOVER
#define READ_CLOVER
Definition: clover_def.h:103

c01_02_re
#define c01_02_re
Definition: clover_core.h:147

c00_12_re
#define c00_12_re
Definition: clover_core.h:163

c20_22_re
#define c20_22_re
Definition: clover_core.h:197

c12_11_im
#define c12_11_im
Definition: clover_core.h:140

c20_22_im
#define c20_22_im
Definition: clover_core.h:198

o21_im
VOLATILE spinorFloat o21_im
Definition: clover_core.h:258

c20_21_re
#define c20_21_re
Definition: clover_core.h:186

a00_im
#define a00_im
Definition: llfat_core.h:116

acc11_im
#define acc11_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:101

o22_im
VOLATILE spinorFloat o22_im
Definition: clover_core.h:260

i31_im
#define i31_im
Definition: clover_core.h:61