su2hmc
Loading...
Searching...
No Matches
Clover Multiplication routines
Clover | Fermion matrix products
Collaboration diagram for Clover Multiplication routines:

Functions

void ByClover (Complex *phi, Complex *r, Complex *clover[2], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
void ByClover_f (Complex_f *phi, Complex_f *r, Complex_f *clover[2], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
void HbyClover (Complex *phi, Complex *r, Complex *clover[2], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
void HbyClover_f (Complex_f *phi, Complex_f *r, Complex_f *clover[2], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
void cuByClover (Complex *phi, Complex *r, Complex *clover[nc], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
 CUDA wrapper for ByClover.
void cuHbyClover (Complex *phi, Complex *r, Complex *clover[nc], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
 CUDA wrapper for HbyClover.
void cuByClover_f (Complex_f *phi, Complex_f *r, Complex_f *clover[nc], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
 CUDA wrapper for ByClover_f.
void cuHbyClover_f (Complex_f *phi, Complex_f *r, Complex_f *clover[nc], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
 CUDA wrapper for HbyClover_f.
template<typename T>
__global__ void Kernels::ByClover (complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
template<typename T>
__global__ void Kernels::HbyClover (complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
 Clover analogue of the Hdslash operation. The H in front is for half, as we only act on the fermions of flavour 1.

Detailed Description

Function Documentation

◆ ByClover() [1/2]

void ByClover ( Complex * phi,
Complex * r,
Complex * clover[2],
Complex * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered output has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 243 of file clover.c.

243 {
244#ifdef USE_GPU
245 cuByClover(phi,r,clover,sigval,akappa,sigin,dag);
246#else
247#pragma omp parallel for simd
248 for(unsigned int i=0;i<kvol;i++){
249 //Prefetched r and Phi array
250 Complex phi_s[ngorkov][nc];
251#pragma unroll
252 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
253 for(unsigned short c=0; c<nc; c++){
254 phi_s[igorkov][c]=0;
255 }
256 Complex r_s[nc];
257 Complex clov_s[nc];
258#pragma unroll
259 for(unsigned short clov=0;clov<6;clov++){
260 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
261 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++){
262 //Mod 4 done bitwise. In general n mod 2^m = n & (2^m-1)
263 const unsigned short idirac = igorkov&3;
264 const unsigned short sind = (igorkov<4) ? sigin[clov*ndirac+idirac] : sigin[clov*ndirac+idirac]+4;
265#pragma unroll
266 for(unsigned short c=0; c<nc; c++)
267 r_s[c]= r[i+kvolHalo*(sind*nc+c)];
269 phi_s[igorkov][0]+=sigval[clov*ndirac+idirac]*(creal(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
270 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
271 phi_s[igorkov][1]+=sigval[clov*ndirac+idirac]*(conj(clov_s[1])*r_s[0]-creal(clov_s[0])*r_s[1]);
272 }
273 }
274#pragma unroll
275 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
276 for(unsigned short c=0; c<nc; c++){
277 //dag is just to do with the output layout and if it has a halo
278 if(dag)
279 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
280 else
281 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
282 }
283 }
284#endif
285 return;
286}
cuByClover
void cuByClover(Complex *phi, Complex *r, Complex *clover[nc], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
CUDA wrapper for ByClover.
Definition cuclover.cu:712
conj
__device__ __forceinline__ T conj(const T &z)
Complex Conjugation.
Definition cusu2hmc.cu:33
nc
#define nc
Colours.
Definition sizes.h:182
ngorkov
#define ngorkov
Gor'kov indices.
Definition sizes.h:190
kvol
#define kvol
Sublattice volume.
Definition sizes.h:163
Complex
#define Complex
Double precision complex number.
Definition sizes.h:64
ndirac
#define ndirac
Dirac indices.
Definition sizes.h:186
kvolHalo
#define kvolHalo
Subvolume + halo size.
Definition sizes.h:234
creal
#define creal(z)
Extract Real Component using C standard notation.
Definition thrust_complex.h:26

References Complex, conj(), creal, cuByClover(), kvol, kvolHalo, nc, ndirac, and ngorkov.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ ByClover() [2/2]

template<typename T>
__global__ void Kernels::ByClover ( complex< T > * phi,
complex< T > * r,
complex< T > * clover1,
complex< T > * clover2,
complex< T > * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.

Parameters
[out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]clover1,clover2Array of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered output has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Note that \(\sigma_{\mu\nu}\) was scaled by \(\frac{c_\text{SW}}{2}\) when we defined it.

Definition at line 581 of file cuclover.cu.

581 {
582 const unsigned int gsize = gridDim.x*gridDim.y*gridDim.z;
583 const unsigned int bsize = blockDim.x*blockDim.y*blockDim.z;
584 const unsigned int blockId = blockIdx.x+ blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
585 const unsigned int bthreadId= (threadIdx.z * blockDim.y+ threadIdx.y)* blockDim.x+ threadIdx.x;
586 const unsigned int gthreadId= blockId * bsize+bthreadId;
587
588 for(unsigned int i=gthreadId;i<kvol;i+=bsize*gsize){
589 //Prefetched r and Phi array
590 complex<T> phi_s[ngorkov][nc];
591#pragma unroll
592 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
593 for(unsigned short c=0; c<nc; c++){
594 phi_s[igorkov][c]=0;
595 }
596 complex<T> r_s[nc];
597 complex<T> clov_s[nc];
598#pragma unroll
599 for(unsigned short clov=0;clov<6;clov++){
600 clov_s[0]=clover1[clov*kvol+i]; clov_s[1]=clover2[clov*kvol+i];
601 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++){
602 //Mod 4 done bitwise. In general n mod 2^m = n & (2^m-1)
603 const unsigned short idirac = igorkov&3;
604 const unsigned short sind = (igorkov<4) ? sigin[clov*ndirac+idirac] : sigin[clov*ndirac+idirac]+4;
605#pragma unroll
606 for(unsigned short c=0; c<nc; c++)
607 r_s[c]= r[i+kvolHalo*(sind*nc+c)];
609 phi_s[igorkov][0]+=sigval[clov*ndirac+idirac]*(creal(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
610 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
611 phi_s[igorkov][1]+=sigval[clov*ndirac+idirac]*(conj(clov_s[1])*r_s[0]-creal(clov_s[0])*r_s[1]);
612 }
613 }
614#pragma unroll
615 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
616 for(unsigned short c=0; c<nc; c++){
617 //dag is just to do with the output layout and if it has a halo
618 if(dag)
619 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
620 else
621 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
622 }
623 }
624 return;
625 }

References conj(), creal, kvol, kvolHalo, nc, ndirac, and ngorkov.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ ByClover_f()

void ByClover_f ( Complex_f * phi,
Complex_f * r,
Complex_f * clover[2],
Complex_f * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered output has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 332 of file clover.c.

332 {
333#ifdef USE_GPU
334 cuByClover_f(phi,r,clover,sigval,akappa,sigin,dag);
335#else
336#pragma omp parallel for simd
337 for(unsigned int i=0;i<kvol;i++){
338 //Prefetched r and Phi array
339 Complex_f phi_s[ngorkov][nc];
340#pragma unroll
341 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
342 for(unsigned short c=0; c<nc; c++){
343 phi_s[igorkov][c]=0;
344 }
345 Complex_f r_s[nc];
346 Complex_f clov_s[nc];
347#pragma unroll
348 for(unsigned short clov=0;clov<6;clov++){
349 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
350 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++){
351 //Mod 4 done bitwise. In general n mod 2^m = n & (2^m-1)
352 const unsigned short idirac = igorkov&3;
353 const unsigned short sind = (igorkov<4) ? sigin[clov*ndirac+idirac] : sigin[clov*ndirac+idirac]+4;
354#pragma unroll
355 for(unsigned short c=0; c<nc; c++)
356 r_s[c]= r[i+kvolHalo*(sind*nc+c)];
358 phi_s[igorkov][0]+=sigval[clov*ndirac+idirac]*(crealf(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
359 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
360 phi_s[igorkov][1]+=sigval[clov*ndirac+idirac]*(conj(clov_s[1])*r_s[0]-crealf(clov_s[0])*r_s[1]);
361 }
362 }
363#pragma unroll
364 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
365 for(unsigned short c=0; c<nc; c++){
366 //dag is just to do with the output layout and if it has a halo
367 if(dag)
368 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
369 else
370 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
371 }
372 }
373#endif
374 return;
375}
cuByClover_f
void cuByClover_f(Complex_f *phi, Complex_f *r, Complex_f *clover[nc], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
CUDA wrapper for ByClover_f.
Definition cuclover.cu:718
Complex_f
#define Complex_f
Single precision complex number.
Definition sizes.h:62

References Complex_f, conj(), cuByClover_f(), kvol, kvolHalo, nc, ndirac, and ngorkov.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ cuByClover()

void cuByClover ( Complex * phi,
Complex * r,
Complex * clover[nc],
Complex * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

CUDA wrapper for ByClover.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 712 of file cuclover.cu.

712 {
713 Kernels::ByClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
714}
Kernels::ByClover
__global__ void ByClover(complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dsla...
Definition cuclover.cu:581

References Kernels::ByClover(), Complex, and nc.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ cuByClover_f()

void cuByClover_f ( Complex_f * phi,
Complex_f * r,
Complex_f * clover[nc],
Complex_f * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

CUDA wrapper for ByClover_f.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 718 of file cuclover.cu.

718 {
719 Kernels::ByClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
720}

References Kernels::ByClover(), Complex_f, and nc.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ cuHbyClover()

void cuHbyClover ( Complex * phi,
Complex * r,
Complex * clover[nc],
Complex * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

CUDA wrapper for HbyClover.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 715 of file cuclover.cu.

715 {
716 Kernels::HbyClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
717}
Kernels::HbyClover
__global__ void HbyClover(complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
Clover analogue of the Hdslash operation. The H in front is for half, as we only act on the fermions ...
Definition cuclover.cu:642

References Complex, Kernels::HbyClover(), and nc.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ cuHbyClover_f()

void cuHbyClover_f ( Complex_f * phi,
Complex_f * r,
Complex_f * clover[nc],
Complex_f * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

CUDA wrapper for HbyClover_f.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 721 of file cuclover.cu.

721 {
722 Kernels::HbyClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
723}

References Complex_f, Kernels::HbyClover(), and nc.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ HbyClover() [1/2]

void HbyClover ( Complex * phi,
Complex * r,
Complex * clover[2],
Complex * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered output has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 287 of file clover.c.

287 {
288 const char funcname[] = "HbyClover";
289#ifdef USE_GPU
290 cuHbyClover(phi,r,clover,sigval,akappa,sigin,dag);
291#else
292#pragma omp parallel for simd
293 for(unsigned int i=0;i<kvol;i++){
294 //Prefetched r and Phi array
295 Complex phi_s[ndirac*nc];
296#pragma unroll
297 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
298 for(unsigned short c=0; c<nc; c++){
299 phi_s[idirac+c]=0;
300 }
301 Complex r_s[nc]; Complex clov_s[nc];
302#pragma unroll
303 for(unsigned short clov=0;clov<6;clov++){
304 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
305 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc){
306 const unsigned short sind = sigin[clov*ndirac+(idirac>>1)] << (nc-1);
307#pragma unroll
308 for(unsigned short c=0; c<nc; c++){
309 r_s[c]= r[i+kvolHalo*(sind+c)];
310 }
312 const Complex sig=sigval[clov*ndirac+(idirac>>1)];
313 //creal just an optimisation. Compiler can't optimise out the zero imag.
314 phi_s[idirac+0]+=sig*(creal(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
315 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
316 phi_s[idirac+1]+=sig*(conj(clov_s[1])*r_s[0]-creal(clov_s[0])*r_s[1]);
317 }
318 }
319#pragma unroll
320 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
321 for(unsigned short c=0; c<nc; c++)
322 //dag is just to do with the output layout and if it has a halo
323 if(dag)
324 phi[i+kvol*(c+idirac)]+=akappa*phi_s[idirac+c];
325 else
326 phi[i+kvolHalo*(c+idirac)]+=akappa*phi_s[idirac+c];
327 }
328#endif
329 return;
330}
cuHbyClover
void cuHbyClover(Complex *phi, Complex *r, Complex *clover[nc], Complex *sigval, const float akappa, unsigned short *sigin, bool dag)
CUDA wrapper for HbyClover.
Definition cuclover.cu:715

References Complex, conj(), creal, cuHbyClover(), kvol, kvolHalo, nc, and ndirac.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ HbyClover() [2/2]

template<typename T>
__global__ void Kernels::HbyClover ( complex< T > * phi,
complex< T > * r,
complex< T > * clover1,
complex< T > * clover2,
complex< T > * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Hdslash operation. The H in front is for half, as we only act on the fermions of flavour 1.

Parameters
[out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]clover1,clover2Array of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered has no MPI halo, but undaggered does.
Postcondition
Result added to phi.

Note that \(\sigma_{\mu\nu}\) was scaled by \(\frac{c_\text{SW}}{2}\) when we defined it.

Definition at line 642 of file cuclover.cu.

642 {
643 const unsigned int gsize = gridDim.x*gridDim.y*gridDim.z;
644 const unsigned int bsize = blockDim.x*blockDim.y*blockDim.z;
645 const unsigned int blockId = blockIdx.x+ blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
646 const unsigned int bthreadId= (threadIdx.z * blockDim.y+ threadIdx.y)* blockDim.x+ threadIdx.x;
647 const unsigned int gthreadId= blockId * bsize+bthreadId;
648
649 for(unsigned int i=gthreadId;i<kvol;i+=bsize*gsize){
650 //Prefetched r and Phi array
651 complex<T> phi_s[ndirac*nc];
652#pragma unroll
653 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
654 for(unsigned short c=0; c<nc; c++){
655 phi_s[idirac+c]=0;
656 }
657 complex<T> r_s[nc]; complex<T> clov_s[nc];
658#pragma unroll
659 for(unsigned short clov=0;clov<6;clov++){
660 clov_s[0]=clover1[clov*kvol+i]; clov_s[1]=clover2[clov*kvol+i];
661 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc){
662 const unsigned short sind = sigin[clov*ndirac+(idirac>>1)] << (nc-1);
663#pragma unroll
664 for(unsigned short c=0; c<nc; c++){
665 r_s[c]= r[i+kvolHalo*(sind+c)];
666 }
668 const complex<T> sig=sigval[clov*ndirac+(idirac>>1)];
669 phi_s[idirac+0]+=sig*(creal(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
670 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
671 phi_s[idirac+1]+=sig*(conj(clov_s[1])*r_s[0]-creal(clov_s[0])*r_s[1]);
672 }
673 }
674#pragma unroll
675 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
676 for(unsigned short c=0; c<nc; c++)
677 //dag is just to do with the output layout and if it has a halo
678 if(dag)
679 phi[i+kvol*(c+idirac)]+=akappa*phi_s[idirac+c];
680 else
681 phi[i+kvolHalo*(c+idirac)]+=akappa*phi_s[idirac+c];
682 }
683 return;
684 }

References conj(), creal, kvol, kvolHalo, nc, and ndirac.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ HbyClover_f()

void HbyClover_f ( Complex_f * phi,
Complex_f * r,
Complex_f * clover[2],
Complex_f * sigval,
const float akappa,
unsigned short * sigin,
bool dag )

Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.

Parameters
[in,out]phiFinal pseudofermion field. This is almost always multiplied by Dslash before calling this function
[in]rPseudofermion field before multiplication. The thing we want to multiply by the clover
[in]cloverArray of clovers
[in]sigval\( \sigma_{\mu\nu}\) entries scaled by \( c_{sw}\)
[in]akappaHopping Parameter
[in]siginWhat element of the spinor is multiplied by row idirac each sigma matrix?
[in]dagDaggered output has no MPI halo, but undaggered does.
Postcondition
Result added to phi

Definition at line 376 of file clover.c.

376 {
377 const char funcname[] = "HbyClover_f";
378#ifdef USE_GPU
379 cuHbyClover_f(phi,r,clover,sigval,akappa,sigin,dag);
380#else
381#pragma omp parallel for simd
382 for(unsigned int i=0;i<kvol;i++){
383 //Prefetched r and Phi array
384 Complex_f phi_s[ndirac*nc];
385#pragma unroll
386 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
387 for(unsigned short c=0; c<nc; c++){
388 phi_s[idirac+c]=0;
389 }
390 Complex_f r_s[nc]; Complex_f clov_s[nc];
391#pragma unroll
392 for(unsigned short clov=0;clov<6;clov++){
393 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
394 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc){
395 const unsigned short sind = sigin[clov*ndirac+(idirac>>1)] << (nc-1);
396#pragma unroll
397 for(unsigned short c=0; c<nc; c++){
398 r_s[c]= r[i+kvolHalo*(sind+c)];
399 }
401 const Complex_f sig=sigval[clov*ndirac+(idirac>>1)];
402 phi_s[idirac+0]+=sig*(crealf(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
403 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
404 phi_s[idirac+1]+=sig*(conj(clov_s[1])*r_s[0]-crealf(clov_s[0])*r_s[1]);
405 }
406 }
407#pragma unroll
408 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
409 for(unsigned short c=0; c<nc; c++)
410 //dag is just to do with the output layout and if it has a halo
411 if(dag)
412 phi[i+kvol*(c+idirac)]+=akappa*phi_s[idirac+c];
413 else
414 phi[i+kvolHalo*(c+idirac)]+=akappa*phi_s[idirac+c];
415 }
416#endif
417 return;
418}
cuHbyClover_f
void cuHbyClover_f(Complex_f *phi, Complex_f *r, Complex_f *clover[nc], Complex_f *sigval, const float akappa, unsigned short *sigin, bool dag)
CUDA wrapper for HbyClover_f.
Definition cuclover.cu:721

References Complex_f, conj(), cuHbyClover_f(), kvol, kvolHalo, nc, and ndirac.

Here is the call graph for this function:
Here is the caller graph for this function: