Calculate the Hamiltonian.
171 {
172 const char funcname[] = "Hamilton";
173
174#ifdef USE_GPU
175 double hp;
176 int device=-1;
177 cudaGetDevice(&device);
178
180 hp*=hp;
181#elif defined USE_BLAS
182 double hp = cblas_dnrm2(
kmom, pp, 1);
183 hp*=hp;
184#else
185 double hp=0;
186 for(
unsigned int i = 0; i<
kmom; i++)
187 hp+=pp[i]*pp[i];
188#endif
189 hp*=0.5;
190 double avplaqs, avplaqt;
191 double hg = 0;
192
194
195 alignas(8) double hf = 0; int itercg = 0;
196#ifdef USE_GPU
198#ifdef _DEBUG
199 cudaMallocManaged((
void **)&smallPhi,
kferm2*
sizeof(
Complex),cudaMemAttachGlobal);
200#else
201 cudaMallocAsync((
void **)&smallPhi,
kferm2*
sizeof(
Complex),NULL);
202#endif
203#else
205#endif
207 if(c_sw)
209
210 for(
unsigned short na=0;na<
nf;na++){
211#ifdef USE_GPU
212#if (nproc>1)
213 for(
unsigned short j=0;j<
nc*
ndirac;j++)
215#else
217#endif
218#else
219 for(
unsigned short j=0;j<
nc*
ndirac;j++)
221#endif
223 if(
Congradq(na,res2,X1,smallPhi,ud,ut,clover,iu,
id,gamval,gamval_f,gamin,sigval,sigval_f,sigin,dk,dk_f,\
224 jqq,akappa,c_sw,&itercg))
225 fprintf(stderr,"Trajectory %d\n", traj);
226
227 *ancgh+=itercg;
228#ifdef USE_GPU
229#if (nproc>1)
230 for(
unsigned short j=0;j<
nc*
ndirac;j++)
232#else
234#endif
235#else
236 for(
unsigned short j=0;j<
nc*
ndirac;j++)
238#endif
240#ifdef USE_GPU
242#if (nproc>1)
243 for(
unsigned short j=0;j<
nc*
ndirac;j++){
246 dot+=buff;
247 }
248#else
249 cublasZdotc(
cublas_handle,
kferm2,(cuDoubleComplex *)smallPhi,1,(cuDoubleComplex *) X1,1,(cuDoubleComplex *) &dot);
250#endif
252#elif defined USE_BLAS
254 for(
unsigned short j=0;j<
nc*
ndirac;j++){
257 dot+=buff;
258 }
260#else
261
262
263#pragma omp parallel for simd collapse(2) aligned(smallPhi,X1:AVX)
264 for(
unsigned short j=0;j<
nc*
ndirac;j++)
265 for(
unsigned int i=0;i<
kvol;i++)
267#endif
268 }
269 if(c_sw)
271#ifdef USE_GPU
272#ifdef _DEBUG
273 cudaFree(smallPhi);
274#else
275 cudaFreeAsync(smallPhi,NULL);
276#endif
277#else
278 free(smallPhi);
279#endif
280
281#if(nproc>1)
283#endif
284 *s=hg+hf; *h=(*s)+hp;
285#ifdef _DEBUG
287 printf("hg=%.5e; hf=%.5e; hp=%.5e; h=%.5e\n", hg, hf, hp, *h);
288#endif
289 return 0;
290}
int Average_Plaquette(double *hg, double *avplaqs, double *avplaqt, Complex_f *ut[2], unsigned int *iu, float beta)
Calculates the gauge action using new (how new?) lookup table Follows a routine called qedplaq in som...
void Clover_free(Complex_f *clover[nc])
Free's memory used for clover terms and leaves.
void Clover(Complex_f *clover[2], Complex_f *ut[2], unsigned int *iu, unsigned int *id)
Calculates the clovers in all directions at all sites.
__device__ __forceinline__ T conj(const T &z)
Complex Conjugation.
int Fill_Small_Phi(int na, Complex *smallPhi, Complex *Phi)
Copies necessary (2*4*kvol) elements of Phi into a vector variable.
int Congradq(int na, double res, Complex *X1, Complex *r, Complex *ud[2], Complex_f *ut[2], Complex_f *clover_f[nc], unsigned int *iu, unsigned int *id, Complex gamval[20], Complex_f gamval_f[20], const unsigned short gamin[16], Complex *sigval, Complex_f *sigval_f, unsigned short *sigin, double *dk[2], float *dk_f[2], Complex_f jqq, float akappa, float c_sw, int *itercg)
Matrix Inversion via Conjugate Gradient (up/down flavour partitioning). Solves Implements up/down pa...
int Par_dsum(double *dval)
Performs a reduction on a double dval to get a sum which is then distributed to all ranks.
#define AVX
Alignment of arrays. 64 for AVX-512, 32 for AVX/AVX2. 16 for SSE. Since AVX is standard on modern x86...
#define kmom
sublattice momentum sites
#define kvol
Sublattice volume.
#define Complex
Double precision complex number.
#define nf
Fermion flavours (double it).
#define ndirac
Dirac indices.
cublasHandle_t cublas_handle
Handle for cuBLAS.
#define Complex_f
Single precision complex number.
#define kferm2
sublattice size including Dirac indices
#define kvolHalo
Subvolume + halo size.
cudaStream_t streams[ndirac *ndim *nadj]
An array of concurrent GPU streams to keep it busy.
#define creal(z)
Extract Real Component using C standard notation.