su2hmc
Loading...
Searching...
No Matches
bosonic.c
Go to the documentation of this file.
1
11#include <su2hmc.h>
12
13
14int Average_Plaquette(double *hg, double *avplaqs, double *avplaqt, Complex_f *ut[2], unsigned int *iu, float beta){
15 const char funcname[] = "Average_Plaquette";
16 /*There was a halo exchange here but moved it outside
17 The FORTRAN code used several consecutive loops to get the plaquette
18 Instead we'll just make the arrays variables and do everything in one loop
19 Should work since in the FORTRAN Sigma11[i] only depends on i components for example
20 Since the \nu loop doesn't get called for \mu=0 we'll start at \mu=1
21 */
22#ifdef USE_GPU
23 __managed__ double hgs = 0; __managed__ double hgt = 0;
24 cuAverage_Plaquette(&hgs, &hgt, ut[0], ut[1], iu,dimGrid,dimBlock);
25#else
26 double hgs = 0; double hgt = 0;
27 for(int mu=1;mu<ndim;mu++)
28 for(int nu=0;nu<mu;nu++)
29 //Don't merge into a single loop. Makes vectorisation easier?
30 //Or merge into a single loop and dispense with the a arrays?
31#pragma omp parallel for simd reduction(+:hgs,hgt)
32 for(int i=0;i<kvol;i++){
33 Complex_f Sigma[2];
34 SU2plaq(ut,Sigma,iu,i,mu,nu);
35 switch(mu){
36 //Time component
37 case(ndim-1): hgt -= creal(Sigma[0]);
38 break;
39 //Space component
40 default: hgs -= creal(Sigma[0]);
41 break;
42 }
43 }
44#endif
45#if(nproc>1)
46 Par_dsum(&hgs); Par_dsum(&hgt);
47#endif
48 *avplaqs=-hgs/(3.0*gvol); *avplaqt=-hgt/(gvol*3.0);
49 *hg=(hgs+hgt)*beta;
50#ifdef _DEBUG
51 if(!rank)
52 printf("hgs=%e hgt=%e hg=%e\n", hgs, hgt, *hg);
53#endif
54 return 0;
55}
56#ifndef USE_GPU
57#pragma omp declare simd
58inline int SU2plaq(Complex_f *ut[2], Complex_f Sigma[2], unsigned int *iu, int i, int mu, int nu){
59 const char funcname[] = "SU2plaq";
60 int uidm = iu[mu*kvol+i];
61
62 Sigma[0]=ut[0][i+kvolHalo*mu]*ut[0][uidm+kvolHalo*nu]-ut[1][i+kvolHalo*mu]*conj(ut[1][uidm+kvolHalo*nu]);
63 Sigma[1]=ut[0][i+kvolHalo*mu]*ut[1][uidm+kvolHalo*nu]+ut[1][i+kvolHalo*mu]*conj(ut[0][uidm+kvolHalo*nu]);
64
65 int uidn = iu[nu*kvol+i];
66 Complex_f a11=Sigma[0]*conj(ut[0][uidn+kvolHalo*mu])+Sigma[1]*conj(ut[1][uidn+kvolHalo*mu]);
67 Complex_f a12=-Sigma[0]*ut[1][uidn+kvolHalo*mu]+Sigma[1]*ut[0][uidn+kvolHalo*mu];
68
69 Sigma[0]=a11*conj(ut[0][i+kvolHalo*nu])+a12*conj(ut[1][i+kvolHalo*nu]);
70 Sigma[1]=-a11*ut[1][i+kvolHalo*nu]+a12*ut[0][i+kvolHalo*nu];
71 return 0;
72}
73#endif
74double Polyakov(Complex_f *ut[2]){
75 const char funcname[] = "Polyakov";
76 double poly = 0;
77 Complex_f *Sigma[2];
78#ifdef USE_GPU
79 cuPolyakov(Sigma,ut,dimGrid,dimBlock);
80#else
81 Sigma[0] = (Complex_f *)aligned_alloc(AVX,kvol3*sizeof(Complex_f));
82 Sigma[1] = (Complex_f *)aligned_alloc(AVX,kvol3*sizeof(Complex_f));
83
84 //Extract the time component from each site and save in corresponding Sigma
85 memcpy(Sigma[0],ut[0]+3*kvolHalo,kvol3*sizeof(Complex_f));
86 memcpy(Sigma[1],ut[1]+3*kvolHalo,kvol3*sizeof(Complex_f));
87 /* Some Fortran commentary
88 Changed this routine.
89 ut[0] and ut[1] now defined as normal ie (kvol+halo,4).
90 Copy of Sigma[0] and Sigma[1] is changed so that it copies
91 in blocks of ksizet.
92 Variable indexu also used to select correct element of ut[0] and ut[1]
93 in loop 10 below.
94
95 Change the order of multiplication so that it can
96 be done in parallel. Start at t=1 and go up to t=T:
97 previously started at t+T and looped back to 1, 2, ... T-1
98 Buffers
99 There is a dependency. Can only parallelise the inner loop
100 */
101#pragma unroll
102 for(int it=1;it<ksizet;it++)
103#pragma omp parallel for simd
104 for(int i=0;i<kvol3;i++){
105 //Seems a bit more efficient to increment indexu instead of reassigning
106 //it every single loop
107 int indexu=it*kvol3+i;
108 Complex_f a11=Sigma[0][i]*ut[0][indexu+kvol*3]-Sigma[1][i]*conj(ut[1][indexu+kvol*3]);
109 //Instead of having to store a second buffer just assign it directly
110 Sigma[1][i]=Sigma[0][i]*ut[1][indexu+kvol*3]+Sigma[1][i]*conj(ut[0][indexu+kvol*3]);
111 Sigma[0][i]=a11;
112 }
113#endif
114
115 //Multiply this partial loop with the contributions of the other cores in the
116 //Time-like dimension
117 //
118 //Par_tmul does nothing if there is only a single processor in the time direction. So we only compile
119 //its call if it is required
120#if (npt>1)
121#ifdef __NVCC_
122#error Par_tmul is not yet implimented in CUDA as Sigma[1] is device only memory
123#endif
124#ifdef _DEBUG
125 printf("Multiplying with MPI\n");
126#endif
127 Par_tmul(Sigma[0], Sigma[1]);
128 //end of #if(npt>1)
129#endif
130 /*Now all cores have the value for the complete Polyakov line at all spacial sites
131 We need to globally sum over spacial processors but not across time as these
132 are duplicates. So we zero the value for all but t=0
133 This is (according to the FORTRAN code) a bit of a hack
134 I will expand on this hack and completely avoid any work
135 for this case rather than calculating everything just to set it to zero
136 */
137 if(!pcoord[3+rank*ndim])
138#pragma omp parallel for simd reduction(+:poly)
139 for(unsigned int i=0;i<kvol3;i++)
140 poly+=creal(Sigma[0][i]);
141#ifdef USE_GPU
142 cudaFree(Sigma[0]);
143#else
144 free(Sigma[0]);
145 free(Sigma[1]);
146#endif
147
148#if(nproc>1)
149 Par_dsum(&poly);
150#endif
151 poly/=gvol3;
152 return poly;
153}
double Polyakov(Complex_f *ut[2])
Calculate the Polyakov loop (no prizes for guessing that one...).
Definition bosonic.c:74
int SU2plaq(Complex_f *ut[2], Complex_f Sigma[2], unsigned int *iu, int i, int mu, int nu)
Calculates the plaquette at site i in the direction.
void cuPolyakov(Complex_f *Sigma[2], Complex_f *ut[2], dim3 dimGrid, dim3 dimBlock)
Calculate the Polyakov loop (no prizes for guessing that one...).
Definition cubosonic.cu:140
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...
Definition bosonic.c:14
void cuAverage_Plaquette(double *hgs, double *hgt, Complex_f *u11t, Complex_f *u12t, unsigned int *iu, dim3 dimGrid, dim3 dimBlock)
Calculates the gauge action using new (how new?) lookup table Follows a routine called qedplaq in som...
Definition cubosonic.cu:120
__device__ __forceinline__ T conj(const T &z)
Complex Conjugation.
Definition cusu2hmc.cu:33
int Par_dsum(double *dval)
Performs a reduction on a double dval to get a sum which is then distributed to all ranks.
int rank
The MPI rank.
Definition par_mpi.c:20
int * pcoord
The processor grid.
Definition par_mpi.c:17
#define AVX
Alignment of arrays. 64 for AVX-512, 32 for AVX/AVX2. 16 for SSE. Since AVX is standard on modern x86...
Definition sizes.h:279
#define gvol3
Lattice spatial volume.
Definition sizes.h:100
#define ksizet
Sublattice t extent.
Definition sizes.h:158
#define kvol
Sublattice volume.
Definition sizes.h:163
#define gvol
Lattice volume.
Definition sizes.h:98
#define kvol3
Sublattice spatial volume.
Definition sizes.h:165
#define Complex_f
Single precision complex number.
Definition sizes.h:62
dim3 dimGrid
Default grid size. First component is normally nt. Second and third depend whatever is needed to get ...
Definition cusu2hmc.cu:27
#define ndim
Dimensions.
Definition sizes.h:188
#define kvolHalo
Subvolume + halo size.
Definition sizes.h:234
dim3 dimBlock
Default block size. Usually 128.
Definition cusu2hmc.cu:25
Function declarations for most of the routines.
#define creal(z)
Extract Real Component using C standard notation.