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++){
279 //dag is just to do with the output layout and if it has a halo
280 if(dag)
281 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
282 else
283 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
284 }
285 }
286#endif
287 return;
288}
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:716
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.

Also \(\sigma_{\mu\nu}F_{\mu\nu}=\sigma_{\nu\mu}F_{\nu\mu}\) so we double it to take account of that But then we multiply by \(-\frac{1}{2}\) so the \(2\) disappears

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++){
619 //dag is just to do with the output layout and if it has a halo
620 if(dag)
621 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
622 else
623 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
624 }
625 }
626 return;
627 }

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 336 of file clover.c.

336 {
337#ifdef USE_GPU
338 cuByClover_f(phi,r,clover,sigval,akappa,sigin,dag);
339#else
340#pragma omp parallel for simd
341 for(unsigned int i=0;i<kvol;i++){
342 //Prefetched r and Phi array
343 Complex_f phi_s[ngorkov][nc];
344#pragma unroll
345 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
346 for(unsigned short c=0; c<nc; c++){
347 phi_s[igorkov][c]=0;
348 }
349 Complex_f r_s[nc];
350 Complex_f clov_s[nc];
351#pragma unroll
352 for(unsigned short clov=0;clov<6;clov++){
353 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
354 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++){
355 //Mod 4 done bitwise. In general n mod 2^m = n & (2^m-1)
356 const unsigned short idirac = igorkov&3;
357 const unsigned short sind = (igorkov<4) ? sigin[clov*ndirac+idirac] : sigin[clov*ndirac+idirac]+4;
358#pragma unroll
359 for(unsigned short c=0; c<nc; c++)
360 r_s[c]= r[i+kvolHalo*(sind*nc+c)];
362 phi_s[igorkov][0]+=sigval[clov*ndirac+idirac]*(crealf(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
363 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
364 phi_s[igorkov][1]+=sigval[clov*ndirac+idirac]*(conj(clov_s[1])*r_s[0]-crealf(clov_s[0])*r_s[1]);
365 }
366 }
367#pragma unroll
368 for(unsigned short igorkov=0; igorkov<ngorkov; igorkov++)
369 for(unsigned short c=0; c<nc; c++){
372 //dag is just to do with the output layout and if it has a halo
373 if(dag)
374 phi[i+kvol*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
375 else
376 phi[i+kvolHalo*(nc*igorkov+c)]+=akappa*phi_s[igorkov][c];
377 }
378 }
379#endif
380 return;
381}
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:722
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 716 of file cuclover.cu.

716 {
717 Kernels::ByClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
718}
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 722 of file cuclover.cu.

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

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 719 of file cuclover.cu.

719 {
720 Kernels::HbyClover<<<dimGrid,dimBlock>>>(phi,r,clover[0],clover[1],sigval,akappa,sigin,dag);
721}
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:644

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 725 of file cuclover.cu.

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

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 289 of file clover.c.

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

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.

\(\sigma_{\mu\nu}F_{\mu\nu}=\sigma_{\nu\mu}F_{\nu\mu}\) so we double it to take account of that But then we multiply by \(-\frac{1}{2}\) so the \(2\) disappears

Definition at line 644 of file cuclover.cu.

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

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 382 of file clover.c.

382 {
383 const char funcname[] = "HbyClover_f";
384#ifdef USE_GPU
385 cuHbyClover_f(phi,r,clover,sigval,akappa,sigin,dag);
386#else
387#pragma omp parallel for simd
388 for(unsigned int i=0;i<kvol;i++){
389 //Prefetched r and Phi array
390 Complex_f phi_s[ndirac*nc];
391#pragma unroll
392 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
393 for(unsigned short c=0; c<nc; c++){
394 phi_s[idirac+c]=0;
395 }
396 Complex_f r_s[nc]; Complex_f clov_s[nc];
397#pragma unroll
398 for(unsigned short clov=0;clov<6;clov++){
399 clov_s[0]=clover[0][clov*kvol+i]; clov_s[1]=clover[1][clov*kvol+i];
400 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc){
401 const unsigned short sind = sigin[clov*ndirac+(idirac>>1)] << (nc-1);
402#pragma unroll
403 for(unsigned short c=0; c<nc; c++){
404 r_s[c]= r[i+kvolHalo*(sind+c)];
405 }
407 const Complex_f sig=sigval[clov*ndirac+(idirac>>1)];
408 phi_s[idirac+0]+=sig*(crealf(clov_s[0])*r_s[0]+clov_s[1]*r_s[1]);
409 //Clover is in the Lie Algebra, not Lie group. So signs are correct here.
410 phi_s[idirac+1]+=sig*(conj(clov_s[1])*r_s[0]-crealf(clov_s[0])*r_s[1]);
411 }
412 }
413#pragma unroll
414 for(unsigned short idirac=0; idirac<ndirac*nc; idirac+=nc)
415 for(unsigned short c=0; c<nc; c++)
418 //dag is just to do with the output layout and if it has a halo
419 if(dag)
420 phi[i+kvol*(c+idirac)]+=akappa*phi_s[idirac+c];
421 else
422 phi[i+kvolHalo*(c+idirac)]+=akappa*phi_s[idirac+c];
423 }
424#endif
425 return;
426}
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:725

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: