mirror of
https://github.com/paboyle/Grid.git
synced 2025-06-15 14:27:06 +01:00
Compare commits
72 Commits
gauge-grou
...
feature/ca
Author | SHA1 | Date | |
---|---|---|---|
59282f25ec | |||
b0bd173899 | |||
135808dcfa | |||
7f7d06d963 | |||
2bf3b4d576 | |||
f34d34bd17 | |||
e32d5141b4 | |||
6d5277f2d7 | |||
14d82777e0 | |||
2a4e739513 | |||
8079dc2a14 | |||
6ceb556684 | |||
76cde73705 | |||
cc094366a9 | |||
41a575ff9b | |||
12ef413065 | |||
829a328451 | |||
402523c62e | |||
d7bef70b5c | |||
2ad1811642 | |||
a65a497bae | |||
b27b12828e | |||
fe9edf8526 | |||
44204c7e06 | |||
33b3789598 | |||
195ab2888d | |||
85f750d753 | |||
a4ce6e42c7 | |||
5398b7e7e3 | |||
fd13a3f2be | |||
c144b32368 | |||
ba7e371b90 | |||
99e7a5d18a | |||
f824d99059 | |||
749b8022a4 | |||
7e0057d2c4 | |||
cfe9e870d3 | |||
e9c4f06cbf | |||
1f9688417a | |||
16c2a99965 | |||
cda915a345 | |||
7c16189e16 | |||
ecbfccea43 | |||
a8eda8f6da | |||
9b1a0653cf | |||
7cb1ff7395 | |||
ab6ea29913 | |||
b5c81a02b6 | |||
d899ee80fc | |||
4016e705fc | |||
2f4e85e5d6 | |||
8ed0b57b09 | |||
7e130076d6 | |||
6efdad6f21 | |||
a822c48565 | |||
014fb76e88 | |||
30e5311b43 | |||
11ee8a1061 | |||
770680669d | |||
0cdfc5cf22 | |||
428b8ba907 | |||
54c6b1376d | |||
f3f11b586f | |||
8083e3f7e8 | |||
364793154b | |||
3e2ae1e9af | |||
d38ae2fd18 | |||
030e7754e4 | |||
3b7fce1e76 | |||
4d15417f93 | |||
ab3c855f65 | |||
92e2c517d8 |
@ -358,7 +358,7 @@ public:
|
||||
autoView( in_v , in, AcceleratorRead);
|
||||
autoView( out_v , out, AcceleratorWrite);
|
||||
autoView( Stencil_v , Stencil, AcceleratorRead);
|
||||
auto& geom_v = geom;
|
||||
int npoint = geom.npoint;
|
||||
typedef LatticeView<Cobj> Aview;
|
||||
|
||||
Vector<Aview> AcceleratorViewContainer;
|
||||
@ -380,7 +380,7 @@ public:
|
||||
int ptype;
|
||||
StencilEntry *SE;
|
||||
|
||||
for(int point=0;point<geom_v.npoint;point++){
|
||||
for(int point=0;point<npoint;point++){
|
||||
|
||||
SE=Stencil_v.GetEntry(ptype,point,ss);
|
||||
|
||||
@ -424,7 +424,7 @@ public:
|
||||
autoView( in_v , in, AcceleratorRead);
|
||||
autoView( out_v , out, AcceleratorWrite);
|
||||
autoView( Stencil_v , Stencil, AcceleratorRead);
|
||||
auto& geom_v = geom;
|
||||
int npoint = geom.npoint;
|
||||
typedef LatticeView<Cobj> Aview;
|
||||
|
||||
Vector<Aview> AcceleratorViewContainer;
|
||||
@ -454,7 +454,7 @@ public:
|
||||
int ptype;
|
||||
StencilEntry *SE;
|
||||
|
||||
for(int p=0;p<geom_v.npoint;p++){
|
||||
for(int p=0;p<npoint;p++){
|
||||
int point = points_p[p];
|
||||
|
||||
SE=Stencil_v.GetEntry(ptype,point,ss);
|
||||
|
@ -52,6 +52,7 @@ public:
|
||||
virtual void AdjOp (const Field &in, Field &out) = 0; // Abstract base
|
||||
virtual void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2)=0;
|
||||
virtual void HermOp(const Field &in, Field &out)=0;
|
||||
virtual ~LinearOperatorBase(){};
|
||||
};
|
||||
|
||||
|
||||
@ -507,7 +508,7 @@ class SchurStaggeredOperator : public SchurOperatorBase<Field> {
|
||||
virtual void MpcDag (const Field &in, Field &out){
|
||||
Mpc(in,out);
|
||||
}
|
||||
virtual void MpcDagMpc(const Field &in, Field &out,RealD &ni,RealD &no) {
|
||||
virtual void MpcDagMpc(const Field &in, Field &out) {
|
||||
assert(0);// Never need with staggered
|
||||
}
|
||||
};
|
||||
@ -585,6 +586,7 @@ class HermOpOperatorFunction : public OperatorFunction<Field> {
|
||||
template<typename Field>
|
||||
class PlainHermOp : public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
LinearOperatorBase<Field> &_Linop;
|
||||
|
||||
PlainHermOp(LinearOperatorBase<Field>& linop) : _Linop(linop)
|
||||
@ -598,6 +600,7 @@ public:
|
||||
template<typename Field>
|
||||
class FunctionHermOp : public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
OperatorFunction<Field> & _poly;
|
||||
LinearOperatorBase<Field> &_Linop;
|
||||
|
||||
|
@ -30,13 +30,19 @@ Author: Azusa Yamaguchi <ayamaguc@staffmail.ed.ac.uk>
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
template<class Field> using Preconditioner = LinearFunction<Field> ;
|
||||
|
||||
/*
|
||||
template<class Field> class Preconditioner : public LinearFunction<Field> {
|
||||
using LinearFunction<Field>::operator();
|
||||
virtual void operator()(const Field &src, Field & psi)=0;
|
||||
};
|
||||
*/
|
||||
|
||||
template<class Field> class TrivialPrecon : public Preconditioner<Field> {
|
||||
public:
|
||||
void operator()(const Field &src, Field & psi){
|
||||
using Preconditioner<Field>::operator();
|
||||
virtual void operator()(const Field &src, Field & psi){
|
||||
psi = src;
|
||||
}
|
||||
TrivialPrecon(void){};
|
||||
|
@ -48,6 +48,7 @@ public:
|
||||
virtual void Mdiag (const Field &in, Field &out)=0;
|
||||
virtual void Mdir (const Field &in, Field &out,int dir, int disp)=0;
|
||||
virtual void MdirAll (const Field &in, std::vector<Field> &out)=0;
|
||||
virtual ~SparseMatrixBase() {};
|
||||
};
|
||||
|
||||
/////////////////////////////////////////////////////////////////////////////////////////////
|
||||
@ -72,7 +73,7 @@ public:
|
||||
virtual void MeooeDag (const Field &in, Field &out)=0;
|
||||
virtual void MooeeDag (const Field &in, Field &out)=0;
|
||||
virtual void MooeeInvDag (const Field &in, Field &out)=0;
|
||||
|
||||
virtual ~CheckerBoardedSparseMatrixBase() {};
|
||||
};
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
|
@ -37,6 +37,7 @@ template<class FieldD, class FieldF, typename std::enable_if< getPrecision<Field
|
||||
class MixedPrecisionBiCGSTAB : public LinearFunction<FieldD>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<FieldD>::operator();
|
||||
RealD Tolerance;
|
||||
RealD InnerTolerance; // Initial tolerance for inner CG. Defaults to Tolerance but can be changed
|
||||
Integer MaxInnerIterations;
|
||||
|
@ -36,6 +36,7 @@ NAMESPACE_BEGIN(Grid);
|
||||
typename std::enable_if< getPrecision<FieldF>::value == 1, int>::type = 0>
|
||||
class MixedPrecisionConjugateGradient : public LinearFunction<FieldD> {
|
||||
public:
|
||||
using LinearFunction<FieldD>::operator();
|
||||
RealD Tolerance;
|
||||
RealD InnerTolerance; //Initial tolerance for inner CG. Defaults to Tolerance but can be changed
|
||||
Integer MaxInnerIterations;
|
||||
|
@ -33,16 +33,19 @@ namespace Grid {
|
||||
template<class Field>
|
||||
class ZeroGuesser: public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
virtual void operator()(const Field &src, Field &guess) { guess = Zero(); };
|
||||
};
|
||||
template<class Field>
|
||||
class DoNothingGuesser: public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
virtual void operator()(const Field &src, Field &guess) { };
|
||||
};
|
||||
template<class Field>
|
||||
class SourceGuesser: public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
virtual void operator()(const Field &src, Field &guess) { guess = src; };
|
||||
};
|
||||
|
||||
@ -57,6 +60,7 @@ private:
|
||||
const unsigned int N;
|
||||
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
|
||||
DeflatedGuesser(const std::vector<Field> & _evec,const std::vector<RealD> & _eval)
|
||||
: DeflatedGuesser(_evec, _eval, _evec.size())
|
||||
@ -87,6 +91,7 @@ private:
|
||||
const std::vector<RealD> &eval_coarse;
|
||||
public:
|
||||
|
||||
using LinearFunction<FineField>::operator();
|
||||
LocalCoherenceDeflatedGuesser(const std::vector<FineField> &_subspace,
|
||||
const std::vector<CoarseField> &_evec_coarse,
|
||||
const std::vector<RealD> &_eval_coarse)
|
||||
|
@ -67,6 +67,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis>
|
||||
class ProjectedHermOp : public LinearFunction<Lattice<iVector<CComplex,nbasis > > > {
|
||||
public:
|
||||
using LinearFunction<Lattice<iVector<CComplex,nbasis > > >::operator();
|
||||
typedef iVector<CComplex,nbasis > CoarseSiteVector;
|
||||
typedef Lattice<CoarseSiteVector> CoarseField;
|
||||
typedef Lattice<CComplex> CoarseScalar; // used for inner products on fine field
|
||||
@ -97,6 +98,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis>
|
||||
class ProjectedFunctionHermOp : public LinearFunction<Lattice<iVector<CComplex,nbasis > > > {
|
||||
public:
|
||||
using LinearFunction<Lattice<iVector<CComplex,nbasis > > >::operator();
|
||||
typedef iVector<CComplex,nbasis > CoarseSiteVector;
|
||||
typedef Lattice<CoarseSiteVector> CoarseField;
|
||||
typedef Lattice<CComplex> CoarseScalar; // used for inner products on fine field
|
||||
|
@ -43,7 +43,7 @@ NAMESPACE_BEGIN(Grid);
|
||||
template<class Field>
|
||||
class PrecGeneralisedConjugateResidual : public LinearFunction<Field> {
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
RealD Tolerance;
|
||||
Integer MaxIterations;
|
||||
int verbose;
|
||||
|
@ -43,7 +43,7 @@ NAMESPACE_BEGIN(Grid);
|
||||
template<class Field>
|
||||
class PrecGeneralisedConjugateResidualNonHermitian : public LinearFunction<Field> {
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
RealD Tolerance;
|
||||
Integer MaxIterations;
|
||||
int verbose;
|
||||
@ -119,7 +119,8 @@ public:
|
||||
RealD GCRnStep(const Field &src, Field &psi,RealD rsq){
|
||||
|
||||
RealD cp;
|
||||
ComplexD a, b, zAz;
|
||||
ComplexD a, b;
|
||||
// ComplexD zAz;
|
||||
RealD zAAz;
|
||||
ComplexD rq;
|
||||
|
||||
@ -146,7 +147,7 @@ public:
|
||||
//////////////////////////////////
|
||||
MatTimer.Start();
|
||||
Linop.Op(psi,Az);
|
||||
zAz = innerProduct(Az,psi);
|
||||
// zAz = innerProduct(Az,psi);
|
||||
zAAz= norm2(Az);
|
||||
MatTimer.Stop();
|
||||
|
||||
@ -170,7 +171,7 @@ public:
|
||||
|
||||
LinalgTimer.Start();
|
||||
|
||||
zAz = innerProduct(Az,psi);
|
||||
// zAz = innerProduct(Az,psi);
|
||||
zAAz= norm2(Az);
|
||||
|
||||
//p[0],q[0],qq[0]
|
||||
@ -212,7 +213,7 @@ public:
|
||||
MatTimer.Start();
|
||||
Linop.Op(z,Az);
|
||||
MatTimer.Stop();
|
||||
zAz = innerProduct(Az,psi);
|
||||
// zAz = innerProduct(Az,psi);
|
||||
zAAz= norm2(Az);
|
||||
|
||||
LinalgTimer.Start();
|
||||
|
@ -9,14 +9,30 @@ NAMESPACE_BEGIN(Grid);
|
||||
#define AccSmall (3)
|
||||
#define Shared (4)
|
||||
#define SharedSmall (5)
|
||||
#undef GRID_MM_VERBOSE
|
||||
uint64_t total_shared;
|
||||
uint64_t total_device;
|
||||
uint64_t total_host;;
|
||||
void MemoryManager::PrintBytes(void)
|
||||
{
|
||||
std::cout << " MemoryManager : "<<total_shared<<" shared bytes "<<std::endl;
|
||||
std::cout << " MemoryManager : "<<total_device<<" accelerator bytes "<<std::endl;
|
||||
std::cout << " MemoryManager : "<<total_host <<" cpu bytes "<<std::endl;
|
||||
std::cout << " MemoryManager : ------------------------------------ "<<std::endl;
|
||||
std::cout << " MemoryManager : PrintBytes "<<std::endl;
|
||||
std::cout << " MemoryManager : ------------------------------------ "<<std::endl;
|
||||
std::cout << " MemoryManager : "<<(total_shared>>20)<<" shared Mbytes "<<std::endl;
|
||||
std::cout << " MemoryManager : "<<(total_device>>20)<<" accelerator Mbytes "<<std::endl;
|
||||
std::cout << " MemoryManager : "<<(total_host>>20) <<" cpu Mbytes "<<std::endl;
|
||||
uint64_t cacheBytes;
|
||||
cacheBytes = CacheBytes[Cpu];
|
||||
std::cout << " MemoryManager : "<<(cacheBytes>>20) <<" cpu cache Mbytes "<<std::endl;
|
||||
cacheBytes = CacheBytes[Acc];
|
||||
std::cout << " MemoryManager : "<<(cacheBytes>>20) <<" acc cache Mbytes "<<std::endl;
|
||||
cacheBytes = CacheBytes[Shared];
|
||||
std::cout << " MemoryManager : "<<(cacheBytes>>20) <<" shared cache Mbytes "<<std::endl;
|
||||
|
||||
#ifdef GRID_CUDA
|
||||
cuda_mem();
|
||||
#endif
|
||||
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
@ -24,86 +40,114 @@ void MemoryManager::PrintBytes(void)
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
MemoryManager::AllocationCacheEntry MemoryManager::Entries[MemoryManager::NallocType][MemoryManager::NallocCacheMax];
|
||||
int MemoryManager::Victim[MemoryManager::NallocType];
|
||||
int MemoryManager::Ncache[MemoryManager::NallocType] = { 8, 32, 8, 32, 8, 32 };
|
||||
|
||||
int MemoryManager::Ncache[MemoryManager::NallocType] = { 2, 8, 2, 8, 2, 8 };
|
||||
uint64_t MemoryManager::CacheBytes[MemoryManager::NallocType];
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Actual allocation and deallocation utils
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
void *MemoryManager::AcceleratorAllocate(size_t bytes)
|
||||
{
|
||||
total_device+=bytes;
|
||||
void *ptr = (void *) Lookup(bytes,Acc);
|
||||
if ( ptr == (void *) NULL ) {
|
||||
ptr = (void *) acceleratorAllocDevice(bytes);
|
||||
total_device+=bytes;
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"AcceleratorAllocate "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
return ptr;
|
||||
}
|
||||
void MemoryManager::AcceleratorFree (void *ptr,size_t bytes)
|
||||
{
|
||||
total_device-=bytes;
|
||||
void *__freeme = Insert(ptr,bytes,Acc);
|
||||
if ( __freeme ) {
|
||||
acceleratorFreeDevice(__freeme);
|
||||
total_device-=bytes;
|
||||
// PrintBytes();
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"AcceleratorFree "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
}
|
||||
void *MemoryManager::SharedAllocate(size_t bytes)
|
||||
{
|
||||
total_shared+=bytes;
|
||||
void *ptr = (void *) Lookup(bytes,Shared);
|
||||
if ( ptr == (void *) NULL ) {
|
||||
ptr = (void *) acceleratorAllocShared(bytes);
|
||||
total_shared+=bytes;
|
||||
// std::cout <<"AcceleratorAllocate: allocated Shared pointer "<<std::hex<<ptr<<std::dec<<std::endl;
|
||||
// PrintBytes();
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"SharedAllocate "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
return ptr;
|
||||
}
|
||||
void MemoryManager::SharedFree (void *ptr,size_t bytes)
|
||||
{
|
||||
total_shared-=bytes;
|
||||
void *__freeme = Insert(ptr,bytes,Shared);
|
||||
if ( __freeme ) {
|
||||
acceleratorFreeShared(__freeme);
|
||||
total_shared-=bytes;
|
||||
// PrintBytes();
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"SharedFree "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
}
|
||||
#ifdef GRID_UVM
|
||||
void *MemoryManager::CpuAllocate(size_t bytes)
|
||||
{
|
||||
total_host+=bytes;
|
||||
void *ptr = (void *) Lookup(bytes,Cpu);
|
||||
if ( ptr == (void *) NULL ) {
|
||||
ptr = (void *) acceleratorAllocShared(bytes);
|
||||
total_host+=bytes;
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"CpuAllocate "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
return ptr;
|
||||
}
|
||||
void MemoryManager::CpuFree (void *_ptr,size_t bytes)
|
||||
{
|
||||
total_host-=bytes;
|
||||
NotifyDeletion(_ptr);
|
||||
void *__freeme = Insert(_ptr,bytes,Cpu);
|
||||
if ( __freeme ) {
|
||||
acceleratorFreeShared(__freeme);
|
||||
total_host-=bytes;
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"CpuFree "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
}
|
||||
#else
|
||||
void *MemoryManager::CpuAllocate(size_t bytes)
|
||||
{
|
||||
total_host+=bytes;
|
||||
void *ptr = (void *) Lookup(bytes,Cpu);
|
||||
if ( ptr == (void *) NULL ) {
|
||||
ptr = (void *) acceleratorAllocCpu(bytes);
|
||||
total_host+=bytes;
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"CpuAllocate "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
return ptr;
|
||||
}
|
||||
void MemoryManager::CpuFree (void *_ptr,size_t bytes)
|
||||
{
|
||||
total_host-=bytes;
|
||||
NotifyDeletion(_ptr);
|
||||
void *__freeme = Insert(_ptr,bytes,Cpu);
|
||||
if ( __freeme ) {
|
||||
acceleratorFreeCpu(__freeme);
|
||||
total_host-=bytes;
|
||||
}
|
||||
#ifdef GRID_MM_VERBOSE
|
||||
std::cout <<"CpuFree "<<std::endl;
|
||||
PrintBytes();
|
||||
#endif
|
||||
}
|
||||
#endif
|
||||
|
||||
@ -115,7 +159,6 @@ void MemoryManager::Init(void)
|
||||
|
||||
char * str;
|
||||
int Nc;
|
||||
int NcS;
|
||||
|
||||
str= getenv("GRID_ALLOC_NCACHE_LARGE");
|
||||
if ( str ) {
|
||||
@ -181,13 +224,13 @@ void *MemoryManager::Insert(void *ptr,size_t bytes,int type)
|
||||
#ifdef ALLOCATION_CACHE
|
||||
bool small = (bytes < GRID_ALLOC_SMALL_LIMIT);
|
||||
int cache = type + small;
|
||||
return Insert(ptr,bytes,Entries[cache],Ncache[cache],Victim[cache]);
|
||||
return Insert(ptr,bytes,Entries[cache],Ncache[cache],Victim[cache],CacheBytes[cache]);
|
||||
#else
|
||||
return ptr;
|
||||
#endif
|
||||
}
|
||||
|
||||
void *MemoryManager::Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries,int ncache,int &victim)
|
||||
void *MemoryManager::Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries,int ncache,int &victim, uint64_t &cacheBytes)
|
||||
{
|
||||
assert(ncache>0);
|
||||
#ifdef GRID_OMP
|
||||
@ -211,6 +254,7 @@ void *MemoryManager::Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries
|
||||
|
||||
if ( entries[v].valid ) {
|
||||
ret = entries[v].address;
|
||||
cacheBytes -= entries[v].bytes;
|
||||
entries[v].valid = 0;
|
||||
entries[v].address = NULL;
|
||||
entries[v].bytes = 0;
|
||||
@ -219,6 +263,7 @@ void *MemoryManager::Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries
|
||||
entries[v].address=ptr;
|
||||
entries[v].bytes =bytes;
|
||||
entries[v].valid =1;
|
||||
cacheBytes += bytes;
|
||||
|
||||
return ret;
|
||||
}
|
||||
@ -228,13 +273,13 @@ void *MemoryManager::Lookup(size_t bytes,int type)
|
||||
#ifdef ALLOCATION_CACHE
|
||||
bool small = (bytes < GRID_ALLOC_SMALL_LIMIT);
|
||||
int cache = type+small;
|
||||
return Lookup(bytes,Entries[cache],Ncache[cache]);
|
||||
return Lookup(bytes,Entries[cache],Ncache[cache],CacheBytes[cache]);
|
||||
#else
|
||||
return NULL;
|
||||
#endif
|
||||
}
|
||||
|
||||
void *MemoryManager::Lookup(size_t bytes,AllocationCacheEntry *entries,int ncache)
|
||||
void *MemoryManager::Lookup(size_t bytes,AllocationCacheEntry *entries,int ncache,uint64_t & cacheBytes)
|
||||
{
|
||||
assert(ncache>0);
|
||||
#ifdef GRID_OMP
|
||||
@ -243,6 +288,7 @@ void *MemoryManager::Lookup(size_t bytes,AllocationCacheEntry *entries,int ncach
|
||||
for(int e=0;e<ncache;e++){
|
||||
if ( entries[e].valid && ( entries[e].bytes == bytes ) ) {
|
||||
entries[e].valid = 0;
|
||||
cacheBytes -= entries[e].bytes;
|
||||
return entries[e].address;
|
||||
}
|
||||
}
|
||||
|
@ -82,14 +82,15 @@ private:
|
||||
static AllocationCacheEntry Entries[NallocType][NallocCacheMax];
|
||||
static int Victim[NallocType];
|
||||
static int Ncache[NallocType];
|
||||
static uint64_t CacheBytes[NallocType];
|
||||
|
||||
/////////////////////////////////////////////////
|
||||
// Free pool
|
||||
/////////////////////////////////////////////////
|
||||
static void *Insert(void *ptr,size_t bytes,int type) ;
|
||||
static void *Lookup(size_t bytes,int type) ;
|
||||
static void *Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries,int ncache,int &victim) ;
|
||||
static void *Lookup(size_t bytes,AllocationCacheEntry *entries,int ncache) ;
|
||||
static void *Insert(void *ptr,size_t bytes,AllocationCacheEntry *entries,int ncache,int &victim,uint64_t &cbytes) ;
|
||||
static void *Lookup(size_t bytes,AllocationCacheEntry *entries,int ncache,uint64_t &cbytes) ;
|
||||
|
||||
static void PrintBytes(void);
|
||||
public:
|
||||
@ -113,6 +114,11 @@ private:
|
||||
static uint64_t HostToDeviceXfer;
|
||||
static uint64_t DeviceToHostXfer;
|
||||
|
||||
static uint64_t DeviceAccesses;
|
||||
static uint64_t HostAccesses;
|
||||
static uint64_t DeviceAccessBytes;
|
||||
static uint64_t HostAccessBytes;
|
||||
|
||||
private:
|
||||
#ifndef GRID_UVM
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
@ -151,6 +157,7 @@ private:
|
||||
|
||||
// static void LRUupdate(AcceleratorViewEntry &AccCache);
|
||||
static void LRUinsert(AcceleratorViewEntry &AccCache);
|
||||
static void LRUinsertback(AcceleratorViewEntry &AccCache);
|
||||
static void LRUremove(AcceleratorViewEntry &AccCache);
|
||||
|
||||
// manage entries in the table
|
||||
@ -169,6 +176,7 @@ private:
|
||||
|
||||
public:
|
||||
static void Print(void);
|
||||
static void PrintState( void* CpuPtr);
|
||||
static int isOpen (void* CpuPtr);
|
||||
static void ViewClose(void* CpuPtr,ViewMode mode);
|
||||
static void *ViewOpen (void* CpuPtr,size_t bytes,ViewMode mode,ViewAdvise hint);
|
||||
|
@ -3,7 +3,7 @@
|
||||
|
||||
#warning "Using explicit device memory copies"
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
//define dprintf(...) printf ( __VA_ARGS__ ); fflush(stdout);
|
||||
//#define dprintf(...) printf ( __VA_ARGS__ ); fflush(stdout);
|
||||
#define dprintf(...)
|
||||
|
||||
|
||||
@ -23,6 +23,11 @@ uint64_t MemoryManager::HostToDeviceBytes;
|
||||
uint64_t MemoryManager::DeviceToHostBytes;
|
||||
uint64_t MemoryManager::HostToDeviceXfer;
|
||||
uint64_t MemoryManager::DeviceToHostXfer;
|
||||
uint64_t MemoryManager::DeviceAccesses;
|
||||
uint64_t MemoryManager::HostAccesses;
|
||||
uint64_t MemoryManager::DeviceAccessBytes;
|
||||
uint64_t MemoryManager::HostAccessBytes;
|
||||
|
||||
|
||||
////////////////////////////////////
|
||||
// Priority ordering for unlocked entries
|
||||
@ -86,6 +91,14 @@ void MemoryManager::LRUinsert(AcceleratorViewEntry &AccCache)
|
||||
AccCache.LRU_valid = 1;
|
||||
DeviceLRUBytes+=AccCache.bytes;
|
||||
}
|
||||
void MemoryManager::LRUinsertback(AcceleratorViewEntry &AccCache)
|
||||
{
|
||||
assert(AccCache.LRU_valid==0);
|
||||
LRU.push_back(AccCache.CpuPtr);
|
||||
AccCache.LRU_entry = --LRU.end();
|
||||
AccCache.LRU_valid = 1;
|
||||
DeviceLRUBytes+=AccCache.bytes;
|
||||
}
|
||||
void MemoryManager::LRUremove(AcceleratorViewEntry &AccCache)
|
||||
{
|
||||
assert(AccCache.LRU_valid==1);
|
||||
@ -129,6 +142,7 @@ void MemoryManager::Evict(AcceleratorViewEntry &AccCache)
|
||||
dprintf("MemoryManager: Evict(%llx) %llx\n",(uint64_t)AccCache.CpuPtr,(uint64_t)AccCache.AccPtr);
|
||||
assert(AccCache.accLock==0);
|
||||
assert(AccCache.cpuLock==0);
|
||||
|
||||
if(AccCache.state==AccDirty) {
|
||||
Flush(AccCache);
|
||||
}
|
||||
@ -231,6 +245,9 @@ uint64_t MemoryManager::AcceleratorViewOpen(uint64_t CpuPtr,size_t bytes,ViewMod
|
||||
EntryCreate(CpuPtr,bytes,mode,hint);
|
||||
}
|
||||
|
||||
DeviceAccesses++;
|
||||
DeviceAccessBytes+=bytes;
|
||||
|
||||
auto AccCacheIterator = EntryLookup(CpuPtr);
|
||||
auto & AccCache = AccCacheIterator->second;
|
||||
if (!AccCache.AccPtr) {
|
||||
@ -349,6 +366,10 @@ void MemoryManager::CpuViewClose(uint64_t CpuPtr)
|
||||
assert(AccCache.accLock==0);
|
||||
|
||||
AccCache.cpuLock--;
|
||||
|
||||
if(AccCache.cpuLock==0) {
|
||||
LRUinsertback(AccCache);
|
||||
}
|
||||
}
|
||||
/*
|
||||
* Action State StateNext Flush Clone
|
||||
@ -371,6 +392,9 @@ uint64_t MemoryManager::CpuViewOpen(uint64_t CpuPtr,size_t bytes,ViewMode mode,V
|
||||
EntryCreate(CpuPtr,bytes,mode,transient);
|
||||
}
|
||||
|
||||
HostAccesses++;
|
||||
HostAccessBytes+=bytes;
|
||||
|
||||
auto AccCacheIterator = EntryLookup(CpuPtr);
|
||||
auto & AccCache = AccCacheIterator->second;
|
||||
|
||||
@ -416,6 +440,12 @@ uint64_t MemoryManager::CpuViewOpen(uint64_t CpuPtr,size_t bytes,ViewMode mode,V
|
||||
|
||||
AccCache.transient= transient? EvictNext : 0;
|
||||
|
||||
// If view is opened on host remove from LRU
|
||||
// Host close says evict next from device
|
||||
if(AccCache.LRU_valid==1){
|
||||
LRUremove(AccCache);
|
||||
}
|
||||
|
||||
return AccCache.CpuPtr;
|
||||
}
|
||||
void MemoryManager::NotifyDeletion(void *_ptr)
|
||||
@ -429,6 +459,7 @@ void MemoryManager::NotifyDeletion(void *_ptr)
|
||||
}
|
||||
void MemoryManager::Print(void)
|
||||
{
|
||||
PrintBytes();
|
||||
std::cout << GridLogDebug << "--------------------------------------------" << std::endl;
|
||||
std::cout << GridLogDebug << "Memory Manager " << std::endl;
|
||||
std::cout << GridLogDebug << "--------------------------------------------" << std::endl;
|
||||
@ -473,6 +504,32 @@ int MemoryManager::isOpen (void* _CpuPtr)
|
||||
}
|
||||
}
|
||||
|
||||
void MemoryManager::PrintState(void* _CpuPtr)
|
||||
{
|
||||
uint64_t CpuPtr = (uint64_t)_CpuPtr;
|
||||
|
||||
if ( EntryPresent(CpuPtr) ){
|
||||
auto AccCacheIterator = EntryLookup(CpuPtr);
|
||||
auto & AccCache = AccCacheIterator->second;
|
||||
std::string str;
|
||||
if ( AccCache.state==Empty ) str = std::string("Empty");
|
||||
if ( AccCache.state==CpuDirty ) str = std::string("CpuDirty");
|
||||
if ( AccCache.state==AccDirty ) str = std::string("AccDirty");
|
||||
if ( AccCache.state==Consistent)str = std::string("Consistent");
|
||||
if ( AccCache.state==EvictNext) str = std::string("EvictNext");
|
||||
|
||||
std::cout << GridLogMessage << "CpuAddr\t\tAccAddr\t\tState\t\tcpuLock\taccLock\tLRU_valid "<<std::endl;
|
||||
std::cout << GridLogMessage << "0x"<<std::hex<<AccCache.CpuPtr<<std::dec
|
||||
<< "\t0x"<<std::hex<<AccCache.AccPtr<<std::dec<<"\t" <<str
|
||||
<< "\t" << AccCache.cpuLock
|
||||
<< "\t" << AccCache.accLock
|
||||
<< "\t" << AccCache.LRU_valid<<std::endl;
|
||||
|
||||
} else {
|
||||
std::cout << GridLogMessage << "No Entry in AccCache table." << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
|
||||
#endif
|
||||
|
@ -12,10 +12,18 @@ uint64_t MemoryManager::HostToDeviceBytes;
|
||||
uint64_t MemoryManager::DeviceToHostBytes;
|
||||
uint64_t MemoryManager::HostToDeviceXfer;
|
||||
uint64_t MemoryManager::DeviceToHostXfer;
|
||||
uint64_t MemoryManager::DeviceAccesses;
|
||||
uint64_t MemoryManager::HostAccesses;
|
||||
uint64_t MemoryManager::DeviceAccessBytes;
|
||||
uint64_t MemoryManager::HostAccessBytes;
|
||||
|
||||
void MemoryManager::ViewClose(void* AccPtr,ViewMode mode){};
|
||||
void *MemoryManager::ViewOpen(void* CpuPtr,size_t bytes,ViewMode mode,ViewAdvise hint){ return CpuPtr; };
|
||||
int MemoryManager::isOpen (void* CpuPtr) { return 0;}
|
||||
void MemoryManager::PrintState(void* CpuPtr)
|
||||
{
|
||||
std::cout << GridLogMessage << "Host<->Device memory movement not currently managed by Grid." << std::endl;
|
||||
};
|
||||
void MemoryManager::Print(void){};
|
||||
void MemoryManager::NotifyDeletion(void *ptr){};
|
||||
|
||||
|
@ -388,8 +388,8 @@ double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsReques
|
||||
// TODO : make a OMP loop on CPU, call threaded bcopy
|
||||
void *shm = (void *) this->ShmBufferTranslate(dest,recv);
|
||||
assert(shm!=NULL);
|
||||
// std::cout <<"acceleratorCopyDeviceToDeviceAsynch"<< std::endl;
|
||||
acceleratorCopyDeviceToDeviceAsynch(xmit,shm,bytes);
|
||||
acceleratorCopySynchronise(); // MPI prob slower
|
||||
}
|
||||
|
||||
if ( CommunicatorPolicy == CommunicatorPolicySequential ) {
|
||||
@ -400,6 +400,9 @@ double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsReques
|
||||
}
|
||||
void CartesianCommunicator::StencilSendToRecvFromComplete(std::vector<CommsRequest_t> &list,int dir)
|
||||
{
|
||||
// std::cout << "Copy Synchronised\n"<<std::endl;
|
||||
acceleratorCopySynchronise();
|
||||
|
||||
int nreq=list.size();
|
||||
|
||||
if (nreq==0) return;
|
||||
|
@ -88,6 +88,13 @@ public:
|
||||
LatticeView<vobj> accessor(*( (LatticeAccelerator<vobj> *) this),mode);
|
||||
accessor.ViewClose();
|
||||
}
|
||||
|
||||
// Helper function to print the state of this object in the AccCache
|
||||
void PrintCacheState(void)
|
||||
{
|
||||
MemoryManager::PrintState(this->_odata);
|
||||
}
|
||||
|
||||
/////////////////////////////////////////////////////////////////////////////////
|
||||
// Return a view object that may be dereferenced in site loops.
|
||||
// The view is trivially copy constructible and may be copied to an accelerator device
|
||||
|
@ -42,7 +42,6 @@ void getNumBlocksAndThreads(const Iterator n, const size_t sizeofsobj, Iterator
|
||||
std::cout << GridLogDebug << "\twarpSize = " << warpSize << std::endl;
|
||||
std::cout << GridLogDebug << "\tsharedMemPerBlock = " << sharedMemPerBlock << std::endl;
|
||||
std::cout << GridLogDebug << "\tmaxThreadsPerBlock = " << maxThreadsPerBlock << std::endl;
|
||||
std::cout << GridLogDebug << "\tmaxThreadsPerBlock = " << warpSize << std::endl;
|
||||
std::cout << GridLogDebug << "\tmultiProcessorCount = " << multiProcessorCount << std::endl;
|
||||
|
||||
if (warpSize != WARP_SIZE) {
|
||||
@ -52,6 +51,10 @@ void getNumBlocksAndThreads(const Iterator n, const size_t sizeofsobj, Iterator
|
||||
|
||||
// let the number of threads in a block be a multiple of 2, starting from warpSize
|
||||
threads = warpSize;
|
||||
if ( threads*sizeofsobj > sharedMemPerBlock ) {
|
||||
std::cout << GridLogError << "The object is too large for the shared memory." << std::endl;
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
while( 2*threads*sizeofsobj < sharedMemPerBlock && 2*threads <= maxThreadsPerBlock ) threads *= 2;
|
||||
// keep all the streaming multiprocessors busy
|
||||
blocks = nextPow2(multiProcessorCount);
|
||||
|
@ -85,6 +85,76 @@ template<class vobj> inline void setCheckerboard(Lattice<vobj> &full,const Latti
|
||||
});
|
||||
}
|
||||
|
||||
template<class vobj> inline void acceleratorPickCheckerboard(int cb,Lattice<vobj> &half,const Lattice<vobj> &full, int checker_dim_half=0)
|
||||
{
|
||||
half.Checkerboard() = cb;
|
||||
autoView(half_v, half, AcceleratorWrite);
|
||||
autoView(full_v, full, AcceleratorRead);
|
||||
Coordinate rdim_full = full.Grid()->_rdimensions;
|
||||
Coordinate rdim_half = half.Grid()->_rdimensions;
|
||||
unsigned long ndim_half = half.Grid()->_ndimension;
|
||||
Coordinate checker_dim_mask_half = half.Grid()->_checker_dim_mask;
|
||||
Coordinate ostride_half = half.Grid()->_ostride;
|
||||
accelerator_for(ss, full.Grid()->oSites(),full.Grid()->Nsimd(),{
|
||||
|
||||
Coordinate coor;
|
||||
int cbos;
|
||||
int linear=0;
|
||||
|
||||
Lexicographic::CoorFromIndex(coor,ss,rdim_full);
|
||||
assert(coor.size()==ndim_half);
|
||||
|
||||
for(int d=0;d<ndim_half;d++){
|
||||
if(checker_dim_mask_half[d]) linear += coor[d];
|
||||
}
|
||||
cbos = (linear&0x1);
|
||||
|
||||
if (cbos==cb) {
|
||||
int ssh=0;
|
||||
for(int d=0;d<ndim_half;d++) {
|
||||
if (d == checker_dim_half) ssh += ostride_half[d] * ((coor[d] / 2) % rdim_half[d]);
|
||||
else ssh += ostride_half[d] * (coor[d] % rdim_half[d]);
|
||||
}
|
||||
coalescedWrite(half_v[ssh],full_v(ss));
|
||||
}
|
||||
});
|
||||
}
|
||||
template<class vobj> inline void acceleratorSetCheckerboard(Lattice<vobj> &full,const Lattice<vobj> &half, int checker_dim_half=0)
|
||||
{
|
||||
int cb = half.Checkerboard();
|
||||
autoView(half_v , half, AcceleratorRead);
|
||||
autoView(full_v , full, AcceleratorWrite);
|
||||
Coordinate rdim_full = full.Grid()->_rdimensions;
|
||||
Coordinate rdim_half = half.Grid()->_rdimensions;
|
||||
unsigned long ndim_half = half.Grid()->_ndimension;
|
||||
Coordinate checker_dim_mask_half = half.Grid()->_checker_dim_mask;
|
||||
Coordinate ostride_half = half.Grid()->_ostride;
|
||||
accelerator_for(ss,full.Grid()->oSites(),full.Grid()->Nsimd(),{
|
||||
|
||||
Coordinate coor;
|
||||
int cbos;
|
||||
int linear=0;
|
||||
|
||||
Lexicographic::CoorFromIndex(coor,ss,rdim_full);
|
||||
assert(coor.size()==ndim_half);
|
||||
|
||||
for(int d=0;d<ndim_half;d++){
|
||||
if(checker_dim_mask_half[d]) linear += coor[d];
|
||||
}
|
||||
cbos = (linear&0x1);
|
||||
|
||||
if (cbos==cb) {
|
||||
int ssh=0;
|
||||
for(int d=0;d<ndim_half;d++){
|
||||
if (d == checker_dim_half) ssh += ostride_half[d] * ((coor[d] / 2) % rdim_half[d]);
|
||||
else ssh += ostride_half[d] * (coor[d] % rdim_half[d]);
|
||||
}
|
||||
coalescedWrite(full_v[ss],half_v(ssh));
|
||||
}
|
||||
|
||||
});
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// Flexible Type Conversion for internal promotion to double as well as graceful
|
||||
// treatment of scalar-compatible types
|
||||
|
@ -576,6 +576,8 @@ class ScidacReader : public GridLimeReader {
|
||||
std::string rec_name(ILDG_BINARY_DATA);
|
||||
while ( limeReaderNextRecord(LimeR) == LIME_SUCCESS ) {
|
||||
if ( !strncmp(limeReaderType(LimeR), rec_name.c_str(),strlen(rec_name.c_str()) ) ) {
|
||||
// in principle should do the line below, but that breaks backard compatibility with old data
|
||||
// skipPastObjectRecord(std::string(GRID_FIELD_NORM));
|
||||
skipPastObjectRecord(std::string(SCIDAC_CHECKSUM));
|
||||
return;
|
||||
}
|
||||
|
@ -828,6 +828,7 @@ void CayleyFermion5D<Impl>::SeqConservedCurrent(PropagatorField &q_in,
|
||||
|
||||
#if (!defined(GRID_HIP))
|
||||
int tshift = (mu == Nd-1) ? 1 : 0;
|
||||
unsigned int LLt = GridDefaultLatt()[Tp];
|
||||
////////////////////////////////////////////////
|
||||
// GENERAL CAYLEY CASE
|
||||
////////////////////////////////////////////////
|
||||
@ -880,7 +881,7 @@ void CayleyFermion5D<Impl>::SeqConservedCurrent(PropagatorField &q_in,
|
||||
}
|
||||
|
||||
std::vector<RealD> G_s(Ls,1.0);
|
||||
RealD sign = 1; // sign flip for vector/tadpole
|
||||
RealD sign = 1.0; // sign flip for vector/tadpole
|
||||
if ( curr_type == Current::Axial ) {
|
||||
for(int s=0;s<Ls/2;s++){
|
||||
G_s[s] = -1.0;
|
||||
@ -890,7 +891,7 @@ void CayleyFermion5D<Impl>::SeqConservedCurrent(PropagatorField &q_in,
|
||||
auto b=this->_b;
|
||||
auto c=this->_c;
|
||||
if ( b == 1 && c == 0 ) {
|
||||
sign = -1;
|
||||
sign = -1.0;
|
||||
}
|
||||
else {
|
||||
std::cerr << "Error: Tadpole implementation currently unavailable for non-Shamir actions." << std::endl;
|
||||
@ -934,7 +935,13 @@ void CayleyFermion5D<Impl>::SeqConservedCurrent(PropagatorField &q_in,
|
||||
tmp = Cshift(tmp,mu,-1);
|
||||
Impl::multLinkField(Utmp,this->Umu,tmp,mu+Nd); // Adjoint link
|
||||
tmp = -G_s[s]*( Utmp + gmu*Utmp );
|
||||
tmp = where((lcoor>=tmin+tshift),tmp,zz); // Mask the time
|
||||
// Mask the time
|
||||
if (tmax == LLt - 1 && tshift == 1){ // quick fix to include timeslice 0 if tmax + tshift is over the last timeslice
|
||||
unsigned int t0 = 0;
|
||||
tmp = where(((lcoor==t0) || (lcoor>=tmin+tshift)),tmp,zz);
|
||||
} else {
|
||||
tmp = where((lcoor>=tmin+tshift),tmp,zz);
|
||||
}
|
||||
L_Q += where((lcoor<=tmax+tshift),tmp,zz); // Position of current complicated
|
||||
|
||||
InsertSlice(L_Q, q_out, s , 0);
|
||||
|
@ -77,23 +77,23 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
#define REGISTER
|
||||
|
||||
#ifdef GRID_SIMT
|
||||
#define LOAD_CHIMU(ptype) \
|
||||
#define LOAD_CHIMU(Ptype) \
|
||||
{const SiteSpinor & ref (in[offset]); \
|
||||
Chimu_00=coalescedReadPermute<ptype>(ref()(0)(0),perm,lane); \
|
||||
Chimu_01=coalescedReadPermute<ptype>(ref()(0)(1),perm,lane); \
|
||||
Chimu_02=coalescedReadPermute<ptype>(ref()(0)(2),perm,lane); \
|
||||
Chimu_10=coalescedReadPermute<ptype>(ref()(1)(0),perm,lane); \
|
||||
Chimu_11=coalescedReadPermute<ptype>(ref()(1)(1),perm,lane); \
|
||||
Chimu_12=coalescedReadPermute<ptype>(ref()(1)(2),perm,lane); \
|
||||
Chimu_20=coalescedReadPermute<ptype>(ref()(2)(0),perm,lane); \
|
||||
Chimu_21=coalescedReadPermute<ptype>(ref()(2)(1),perm,lane); \
|
||||
Chimu_22=coalescedReadPermute<ptype>(ref()(2)(2),perm,lane); \
|
||||
Chimu_30=coalescedReadPermute<ptype>(ref()(3)(0),perm,lane); \
|
||||
Chimu_31=coalescedReadPermute<ptype>(ref()(3)(1),perm,lane); \
|
||||
Chimu_32=coalescedReadPermute<ptype>(ref()(3)(2),perm,lane); }
|
||||
Chimu_00=coalescedReadPermute<Ptype>(ref()(0)(0),perm,lane); \
|
||||
Chimu_01=coalescedReadPermute<Ptype>(ref()(0)(1),perm,lane); \
|
||||
Chimu_02=coalescedReadPermute<Ptype>(ref()(0)(2),perm,lane); \
|
||||
Chimu_10=coalescedReadPermute<Ptype>(ref()(1)(0),perm,lane); \
|
||||
Chimu_11=coalescedReadPermute<Ptype>(ref()(1)(1),perm,lane); \
|
||||
Chimu_12=coalescedReadPermute<Ptype>(ref()(1)(2),perm,lane); \
|
||||
Chimu_20=coalescedReadPermute<Ptype>(ref()(2)(0),perm,lane); \
|
||||
Chimu_21=coalescedReadPermute<Ptype>(ref()(2)(1),perm,lane); \
|
||||
Chimu_22=coalescedReadPermute<Ptype>(ref()(2)(2),perm,lane); \
|
||||
Chimu_30=coalescedReadPermute<Ptype>(ref()(3)(0),perm,lane); \
|
||||
Chimu_31=coalescedReadPermute<Ptype>(ref()(3)(1),perm,lane); \
|
||||
Chimu_32=coalescedReadPermute<Ptype>(ref()(3)(2),perm,lane); }
|
||||
#define PERMUTE_DIR(dir) ;
|
||||
#else
|
||||
#define LOAD_CHIMU(ptype) \
|
||||
#define LOAD_CHIMU(Ptype) \
|
||||
{const SiteSpinor & ref (in[offset]); \
|
||||
Chimu_00=ref()(0)(0);\
|
||||
Chimu_01=ref()(0)(1);\
|
||||
@ -371,10 +371,11 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
result_32-= UChi_12;
|
||||
|
||||
#define HAND_STENCIL_LEGB(PROJ,PERM,DIR,RECON) \
|
||||
{int ptype; \
|
||||
SE=st.GetEntry(ptype,DIR,ss); \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
auto offset = SE->_offset; \
|
||||
auto local = SE->_is_local; \
|
||||
auto perm = SE->_permute; \
|
||||
if ( local ) { \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
@ -386,14 +387,14 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
} \
|
||||
acceleratorSynchronise(); \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON;
|
||||
RECON; }
|
||||
|
||||
#define HAND_STENCIL_LEG(PROJ,PERM,DIR,RECON) \
|
||||
SE=&st_p[DIR+8*ss]; \
|
||||
ptype=st_perm[DIR]; \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
{ SE=&st_p[DIR+8*ss]; \
|
||||
auto ptype=st_perm[DIR]; \
|
||||
auto offset = SE->_offset; \
|
||||
auto local = SE->_is_local; \
|
||||
auto perm = SE->_permute; \
|
||||
if ( local ) { \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
@ -405,24 +406,25 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
} \
|
||||
acceleratorSynchronise(); \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON;
|
||||
RECON; }
|
||||
|
||||
#define HAND_STENCIL_LEGA(PROJ,PERM,DIR,RECON) \
|
||||
SE=&st_p[DIR+8*ss]; \
|
||||
ptype=st_perm[DIR]; \
|
||||
{ SE=&st_p[DIR+8*ss]; \
|
||||
auto ptype=st_perm[DIR]; \
|
||||
/*SE=st.GetEntry(ptype,DIR,ss);*/ \
|
||||
offset = SE->_offset; \
|
||||
perm = SE->_permute; \
|
||||
auto offset = SE->_offset; \
|
||||
auto perm = SE->_permute; \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON;
|
||||
RECON; }
|
||||
|
||||
#define HAND_STENCIL_LEG_INT(PROJ,PERM,DIR,RECON) \
|
||||
{ int ptype; \
|
||||
SE=st.GetEntry(ptype,DIR,ss); \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
auto offset = SE->_offset; \
|
||||
auto local = SE->_is_local; \
|
||||
auto perm = SE->_permute; \
|
||||
if ( local ) { \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
@ -437,18 +439,19 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON; \
|
||||
} \
|
||||
acceleratorSynchronise();
|
||||
acceleratorSynchronise(); }
|
||||
|
||||
#define HAND_STENCIL_LEG_EXT(PROJ,PERM,DIR,RECON) \
|
||||
{ int ptype; \
|
||||
SE=st.GetEntry(ptype,DIR,ss); \
|
||||
offset = SE->_offset; \
|
||||
auto offset = SE->_offset; \
|
||||
if((!SE->_is_local)&&(!st.same_node[DIR]) ) { \
|
||||
LOAD_CHI; \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON; \
|
||||
nmu++; \
|
||||
} \
|
||||
acceleratorSynchronise();
|
||||
acceleratorSynchronise(); }
|
||||
|
||||
#define HAND_RESULT(ss) \
|
||||
{ \
|
||||
@ -563,7 +566,6 @@ WilsonKernels<Impl>::HandDhopSiteSycl(StencilVector st_perm,StencilEntry *st_p,
|
||||
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
int offset,local,perm, ptype;
|
||||
StencilEntry *SE;
|
||||
HAND_STENCIL_LEG(XM_PROJ,3,Xp,XM_RECON);
|
||||
HAND_STENCIL_LEG(YM_PROJ,2,Yp,YM_RECON_ACCUM);
|
||||
@ -593,9 +595,7 @@ WilsonKernels<Impl>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,Site
|
||||
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
int offset,local,perm, ptype;
|
||||
StencilEntry *SE;
|
||||
|
||||
HAND_STENCIL_LEG(XM_PROJ,3,Xp,XM_RECON);
|
||||
HAND_STENCIL_LEG(YM_PROJ,2,Yp,YM_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(ZM_PROJ,1,Zp,ZM_RECON_ACCUM);
|
||||
@ -623,8 +623,6 @@ void WilsonKernels<Impl>::HandDhopSiteDag(StencilView &st,DoubledGaugeFieldView
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
StencilEntry *SE;
|
||||
int offset,local,perm, ptype;
|
||||
|
||||
HAND_STENCIL_LEG(XP_PROJ,3,Xp,XP_RECON);
|
||||
HAND_STENCIL_LEG(YP_PROJ,2,Yp,YP_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(ZP_PROJ,1,Zp,ZP_RECON_ACCUM);
|
||||
@ -640,8 +638,8 @@ template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
{
|
||||
auto st_p = st._entries_p;
|
||||
auto st_perm = st._permute_type;
|
||||
// auto st_p = st._entries_p;
|
||||
// auto st_perm = st._permute_type;
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -652,7 +650,6 @@ WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,Si
|
||||
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
int offset,local,perm, ptype;
|
||||
StencilEntry *SE;
|
||||
ZERO_RESULT;
|
||||
HAND_STENCIL_LEG_INT(XM_PROJ,3,Xp,XM_RECON_ACCUM);
|
||||
@ -670,8 +667,8 @@ template<class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
{
|
||||
auto st_p = st._entries_p;
|
||||
auto st_perm = st._permute_type;
|
||||
// auto st_p = st._entries_p;
|
||||
// auto st_perm = st._permute_type;
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
typedef decltype( coalescedRead( in[0]()(0)(0) )) Simt;
|
||||
@ -682,7 +679,6 @@ void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldVi
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
StencilEntry *SE;
|
||||
int offset,local,perm, ptype;
|
||||
ZERO_RESULT;
|
||||
HAND_STENCIL_LEG_INT(XP_PROJ,3,Xp,XP_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG_INT(YP_PROJ,2,Yp,YP_RECON_ACCUM);
|
||||
@ -699,8 +695,8 @@ template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
{
|
||||
auto st_p = st._entries_p;
|
||||
auto st_perm = st._permute_type;
|
||||
// auto st_p = st._entries_p;
|
||||
// auto st_perm = st._permute_type;
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -711,7 +707,7 @@ WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,Si
|
||||
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
int offset, ptype;
|
||||
// int offset, ptype;
|
||||
StencilEntry *SE;
|
||||
int nmu=0;
|
||||
ZERO_RESULT;
|
||||
@ -730,8 +726,8 @@ template<class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
{
|
||||
auto st_p = st._entries_p;
|
||||
auto st_perm = st._permute_type;
|
||||
// auto st_p = st._entries_p;
|
||||
// auto st_perm = st._permute_type;
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
typedef decltype( coalescedRead( in[0]()(0)(0) )) Simt;
|
||||
@ -742,7 +738,7 @@ void WilsonKernels<Impl>::HandDhopSiteDagExt(StencilView &st,DoubledGaugeFieldVi
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
StencilEntry *SE;
|
||||
int offset, ptype;
|
||||
// int offset, ptype;
|
||||
int nmu=0;
|
||||
ZERO_RESULT;
|
||||
HAND_STENCIL_LEG_EXT(XP_PROJ,3,Xp,XP_RECON_ACCUM);
|
||||
|
@ -1,61 +1,63 @@
|
||||
Using HMC in Grid version 0.5.1
|
||||
# Using HMC in Grid
|
||||
|
||||
These are the instructions to use the Generalised HMC on Grid version 0.5.1.
|
||||
Disclaimer: GRID is still under active development so any information here can be changed in future releases.
|
||||
These are the instructions to use the Generalised HMC on Grid as of commit `749b802`.
|
||||
Disclaimer: Grid is still under active development so any information here can be changed in future releases.
|
||||
|
||||
|
||||
Command line options
|
||||
===================
|
||||
(relevant file GenericHMCrunner.h)
|
||||
## Command line options
|
||||
|
||||
(relevant file `GenericHMCrunner.h`)
|
||||
The initial configuration can be changed at the command line using
|
||||
--StartType <your choice>
|
||||
valid choices, one among these
|
||||
HotStart, ColdStart, TepidStart, CheckpointStart
|
||||
default: HotStart
|
||||
`--StartingType STARTING_TYPE`, where `STARTING_TYPE` is one of
|
||||
`HotStart`, `ColdStart`, `TepidStart`, and `CheckpointStart`.
|
||||
Default: `--StartingType HotStart`
|
||||
|
||||
example
|
||||
./My_hmc_exec --StartType HotStart
|
||||
Example:
|
||||
```
|
||||
./My_hmc_exec --StartingType HotStart
|
||||
```
|
||||
|
||||
The CheckpointStart option uses the prefix for the configurations and rng seed files defined in your executable and the initial configuration is specified by
|
||||
--StartTrajectory <integer>
|
||||
default: 0
|
||||
The `CheckpointStart` option uses the prefix for the configurations and rng seed files defined in your executable and the initial configuration is specified by
|
||||
`--StartingTrajectory STARTING_TRAJECTORY`, where `STARTING_TRAJECTORY` is an integer.
|
||||
Default: `--StartingTrajectory 0`
|
||||
|
||||
The number of trajectories for a specific run are specified at command line by
|
||||
--Trajectories <integer>
|
||||
default: 1
|
||||
`--Trajectories TRAJECTORIES`, where `TRAJECTORIES` is an integer.
|
||||
Default: `--Trajectories 1`
|
||||
|
||||
The number of thermalization steps (i.e. steps when the Metropolis acceptance check is turned off) is specified by
|
||||
--Thermalizations <integer>
|
||||
default: 10
|
||||
|
||||
`--Thermalizations THERMALIZATIONS`, where `THERMALIZATIONS` is an integer.
|
||||
Default: `--Thermalizations 10`
|
||||
|
||||
Any other parameter is defined in the source for the executable.
|
||||
|
||||
HMC controls
|
||||
===========
|
||||
## HMC controls
|
||||
|
||||
The lines
|
||||
|
||||
```
|
||||
std::vector<int> SerSeed({1, 2, 3, 4, 5});
|
||||
std::vector<int> ParSeed({6, 7, 8, 9, 10});
|
||||
```
|
||||
|
||||
define the seeds for the serial and the parallel RNG.
|
||||
|
||||
The line
|
||||
|
||||
```
|
||||
TheHMC.MDparameters.set(20, 1.0);// MDsteps, traj length
|
||||
```
|
||||
|
||||
declares the number of molecular dynamics steps and the total trajectory length.
|
||||
|
||||
|
||||
Actions
|
||||
======
|
||||
## Actions
|
||||
|
||||
Action names are defined in the file
|
||||
lib/qcd/Actions.h
|
||||
Action names are defined in the directory `Grid/qcd/action`.
|
||||
|
||||
Gauge actions list:
|
||||
Gauge actions list (from `Grid/qcd/action/gauge/Gauge.h`):
|
||||
|
||||
```
|
||||
WilsonGaugeActionR;
|
||||
WilsonGaugeActionF;
|
||||
WilsonGaugeActionD;
|
||||
@ -68,8 +70,9 @@ IwasakiGaugeActionD;
|
||||
SymanzikGaugeActionR;
|
||||
SymanzikGaugeActionF;
|
||||
SymanzikGaugeActionD;
|
||||
```
|
||||
|
||||
|
||||
```
|
||||
ConjugateWilsonGaugeActionR;
|
||||
ConjugateWilsonGaugeActionF;
|
||||
ConjugateWilsonGaugeActionD;
|
||||
@ -82,26 +85,23 @@ ConjugateIwasakiGaugeActionD;
|
||||
ConjugateSymanzikGaugeActionR;
|
||||
ConjugateSymanzikGaugeActionF;
|
||||
ConjugateSymanzikGaugeActionD;
|
||||
```
|
||||
|
||||
Each of these action accepts one single parameter at creation time (beta).
|
||||
Example for creating a Symanzik action with beta=4.0
|
||||
|
||||
```
|
||||
SymanzikGaugeActionR(4.0)
|
||||
```
|
||||
|
||||
Scalar actions list (from `Grid/qcd/action/scalar/Scalar.h`):
|
||||
|
||||
```
|
||||
ScalarActionR;
|
||||
ScalarActionF;
|
||||
ScalarActionD;
|
||||
```
|
||||
|
||||
|
||||
each of these action accept one single parameter at creation time (beta).
|
||||
Example for creating a Symanzik action with beta=4.0
|
||||
|
||||
SymanzikGaugeActionR(4.0)
|
||||
|
||||
The suffixes R,F,D in the action names refer to the Real
|
||||
(the precision is defined at compile time by the --enable-precision flag in the configure),
|
||||
Float and Double, that force the precision of the action to be 32, 64 bit respectively.
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
The suffixes `R`, `F`, `D` in the action names refer to the `Real`
|
||||
(the precision is defined at compile time by the `--enable-precision` flag in the configure),
|
||||
`Float` and `Double`, that force the precision of the action to be 32, 64 bit respectively.
|
||||
|
@ -322,8 +322,8 @@ public:
|
||||
int simd_layout = _grid->_simd_layout[dimension];
|
||||
int comm_dim = _grid->_processors[dimension] >1 ;
|
||||
|
||||
int recv_from_rank;
|
||||
int xmit_to_rank;
|
||||
// int recv_from_rank;
|
||||
// int xmit_to_rank;
|
||||
|
||||
if ( ! comm_dim ) return 1;
|
||||
if ( displacement == 0 ) return 1;
|
||||
|
@ -47,20 +47,20 @@ NAMESPACE_BEGIN(Grid);
|
||||
class TypePair {
|
||||
public:
|
||||
T _internal[2];
|
||||
TypePair<T>& operator=(const Grid::Zero& o) {
|
||||
accelerator TypePair<T>& operator=(const Grid::Zero& o) {
|
||||
_internal[0] = Zero();
|
||||
_internal[1] = Zero();
|
||||
return *this;
|
||||
}
|
||||
|
||||
TypePair<T> operator+(const TypePair<T>& o) const {
|
||||
accelerator TypePair<T> operator+(const TypePair<T>& o) const {
|
||||
TypePair<T> r;
|
||||
r._internal[0] = _internal[0] + o._internal[0];
|
||||
r._internal[1] = _internal[1] + o._internal[1];
|
||||
return r;
|
||||
}
|
||||
|
||||
TypePair<T>& operator+=(const TypePair<T>& o) {
|
||||
accelerator TypePair<T>& operator+=(const TypePair<T>& o) {
|
||||
_internal[0] += o._internal[0];
|
||||
_internal[1] += o._internal[1];
|
||||
return *this;
|
||||
|
@ -74,29 +74,43 @@ void acceleratorInit(void)
|
||||
// GPU_PROP(singleToDoublePrecisionPerfRatio);
|
||||
}
|
||||
}
|
||||
|
||||
MemoryManager::DeviceMaxBytes = (8*totalDeviceMem)/10; // Assume 80% ours
|
||||
#undef GPU_PROP_FMT
|
||||
#undef GPU_PROP
|
||||
|
||||
#ifdef GRID_DEFAULT_GPU
|
||||
int device = 0;
|
||||
// IBM Jsrun makes cuda Device numbering screwy and not match rank
|
||||
if ( world_rank == 0 ) {
|
||||
printf("AcceleratorCudaInit: using default device \n");
|
||||
printf("AcceleratorCudaInit: assume user either uses a) IBM jsrun, or \n");
|
||||
printf("AcceleratorCudaInit: assume user either uses\n");
|
||||
printf("AcceleratorCudaInit: a) IBM jsrun, or \n");
|
||||
printf("AcceleratorCudaInit: b) invokes through a wrapping script to set CUDA_VISIBLE_DEVICES, UCX_NET_DEVICES, and numa binding \n");
|
||||
printf("AcceleratorCudaInit: Configure options --enable-setdevice=no \n");
|
||||
}
|
||||
#else
|
||||
int device = rank;
|
||||
printf("AcceleratorCudaInit: rank %d setting device to node rank %d\n",world_rank,rank);
|
||||
printf("AcceleratorCudaInit: Configure options --enable-setdevice=yes \n");
|
||||
cudaSetDevice(rank);
|
||||
#endif
|
||||
|
||||
cudaSetDevice(device);
|
||||
cudaStreamCreate(©Stream);
|
||||
const int len=64;
|
||||
char busid[len];
|
||||
if( rank == world_rank ) {
|
||||
cudaDeviceGetPCIBusId(busid, len, device);
|
||||
printf("local rank %d device %d bus id: %s\n", rank, device, busid);
|
||||
}
|
||||
|
||||
if ( world_rank == 0 ) printf("AcceleratorCudaInit: ================================================\n");
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef GRID_HIP
|
||||
hipDeviceProp_t *gpu_props;
|
||||
hipStream_t copyStream;
|
||||
void acceleratorInit(void)
|
||||
{
|
||||
int nDevices = 1;
|
||||
@ -154,16 +168,25 @@ void acceleratorInit(void)
|
||||
#ifdef GRID_DEFAULT_GPU
|
||||
if ( world_rank == 0 ) {
|
||||
printf("AcceleratorHipInit: using default device \n");
|
||||
printf("AcceleratorHipInit: assume user either uses a wrapping script to set CUDA_VISIBLE_DEVICES, UCX_NET_DEVICES, and numa binding \n");
|
||||
printf("AcceleratorHipInit: Configure options --enable-summit, --enable-select-gpu=no \n");
|
||||
printf("AcceleratorHipInit: assume user or srun sets ROCR_VISIBLE_DEVICES and numa binding \n");
|
||||
printf("AcceleratorHipInit: Configure options --enable-setdevice=no \n");
|
||||
}
|
||||
int device = 0;
|
||||
#else
|
||||
if ( world_rank == 0 ) {
|
||||
printf("AcceleratorHipInit: rank %d setting device to node rank %d\n",world_rank,rank);
|
||||
printf("AcceleratorHipInit: Configure options --enable-select-gpu=yes \n");
|
||||
printf("AcceleratorHipInit: Configure options --enable-setdevice=yes \n");
|
||||
}
|
||||
hipSetDevice(rank);
|
||||
int device = rank;
|
||||
#endif
|
||||
hipSetDevice(device);
|
||||
hipStreamCreate(©Stream);
|
||||
const int len=64;
|
||||
char busid[len];
|
||||
if( rank == world_rank ) {
|
||||
hipDeviceGetPCIBusId(busid, len, device);
|
||||
printf("local rank %d device %d bus id: %s\n", rank, device, busid);
|
||||
}
|
||||
if ( world_rank == 0 ) printf("AcceleratorHipInit: ================================================\n");
|
||||
}
|
||||
#endif
|
||||
|
@ -95,6 +95,7 @@ void acceleratorInit(void);
|
||||
//////////////////////////////////////////////
|
||||
|
||||
#ifdef GRID_CUDA
|
||||
|
||||
#include <cuda.h>
|
||||
|
||||
#ifdef __CUDA_ARCH__
|
||||
@ -115,6 +116,14 @@ accelerator_inline int acceleratorSIMTlane(int Nsimd) {
|
||||
#endif
|
||||
} // CUDA specific
|
||||
|
||||
inline void cuda_mem(void)
|
||||
{
|
||||
size_t free_t,total_t,used_t;
|
||||
cudaMemGetInfo(&free_t,&total_t);
|
||||
used_t=total_t-free_t;
|
||||
std::cout << " MemoryManager : GPU used "<<used_t<<" free "<<free_t<< " total "<<total_t<<std::endl;
|
||||
}
|
||||
|
||||
#define accelerator_for2dNB( iter1, num1, iter2, num2, nsimd, ... ) \
|
||||
{ \
|
||||
int nt=acceleratorThreads(); \
|
||||
@ -221,6 +230,7 @@ inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes
|
||||
cudaMemcpyAsync(to,from,bytes, cudaMemcpyDeviceToDevice,copyStream);
|
||||
}
|
||||
inline void acceleratorCopySynchronise(void) { cudaStreamSynchronize(copyStream); };
|
||||
|
||||
inline int acceleratorIsCommunicable(void *ptr)
|
||||
{
|
||||
// int uvm=0;
|
||||
@ -297,7 +307,7 @@ inline void acceleratorFreeDevice(void *ptr){free(ptr,*theGridAccelerator);};
|
||||
inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes) {
|
||||
theGridAccelerator->memcpy(to,from,bytes);
|
||||
}
|
||||
inline void acceleratorCopySynchronise(void) { theGridAccelerator->wait(); }
|
||||
inline void acceleratorCopySynchronise(void) { theGridAccelerator->wait(); std::cout<<"acceleratorCopySynchronise() wait "<<std::endl; }
|
||||
inline void acceleratorCopyToDevice(void *from,void *to,size_t bytes) { theGridAccelerator->memcpy(to,from,bytes); theGridAccelerator->wait();}
|
||||
inline void acceleratorCopyFromDevice(void *from,void *to,size_t bytes){ theGridAccelerator->memcpy(to,from,bytes); theGridAccelerator->wait();}
|
||||
inline void acceleratorMemSet(void *base,int value,size_t bytes) { theGridAccelerator->memset(base,value,bytes); theGridAccelerator->wait();}
|
||||
@ -328,6 +338,7 @@ NAMESPACE_BEGIN(Grid);
|
||||
#define accelerator __host__ __device__
|
||||
#define accelerator_inline __host__ __device__ inline
|
||||
|
||||
extern hipStream_t copyStream;
|
||||
/*These routines define mapping from thread grid to loop & vector lane indexing */
|
||||
accelerator_inline int acceleratorSIMTlane(int Nsimd) {
|
||||
#ifdef GRID_SIMT
|
||||
@ -402,10 +413,16 @@ inline void acceleratorFreeShared(void *ptr){ hipFree(ptr);};
|
||||
inline void acceleratorFreeDevice(void *ptr){ hipFree(ptr);};
|
||||
inline void acceleratorCopyToDevice(void *from,void *to,size_t bytes) { hipMemcpy(to,from,bytes, hipMemcpyHostToDevice);}
|
||||
inline void acceleratorCopyFromDevice(void *from,void *to,size_t bytes){ hipMemcpy(to,from,bytes, hipMemcpyDeviceToHost);}
|
||||
inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes) { hipMemcpy(to,from,bytes, hipMemcpyDeviceToDevice);}
|
||||
inline void acceleratorCopySynchronise(void) { }
|
||||
//inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes) { hipMemcpy(to,from,bytes, hipMemcpyDeviceToDevice);}
|
||||
//inline void acceleratorCopySynchronise(void) { }
|
||||
inline void acceleratorMemSet(void *base,int value,size_t bytes) { hipMemset(base,value,bytes);}
|
||||
|
||||
inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes) // Asynch
|
||||
{
|
||||
hipMemcpyAsync(to,from,bytes, hipMemcpyDeviceToDevice,copyStream);
|
||||
}
|
||||
inline void acceleratorCopySynchronise(void) { hipStreamSynchronize(copyStream); };
|
||||
|
||||
#endif
|
||||
|
||||
//////////////////////////////////////////////
|
||||
@ -476,18 +493,12 @@ inline void acceleratorFreeCpu (void *ptr){free(ptr);};
|
||||
///////////////////////////////////////////////////
|
||||
// Synchronise across local threads for divergence resynch
|
||||
///////////////////////////////////////////////////
|
||||
accelerator_inline void acceleratorSynchronise(void)
|
||||
accelerator_inline void acceleratorSynchronise(void) // Only Nvidia needs
|
||||
{
|
||||
#ifdef GRID_SIMT
|
||||
#ifdef GRID_CUDA
|
||||
__syncwarp();
|
||||
#endif
|
||||
#ifdef GRID_SYCL
|
||||
//cl::sycl::detail::workGroupBarrier();
|
||||
#endif
|
||||
#ifdef GRID_HIP
|
||||
__syncthreads();
|
||||
#endif
|
||||
#endif
|
||||
return;
|
||||
}
|
||||
|
@ -88,7 +88,7 @@ public:
|
||||
// Coordinate class, maxdims = 8 for now.
|
||||
////////////////////////////////////////////////////////////////
|
||||
#define GRID_MAX_LATTICE_DIMENSION (8)
|
||||
#define GRID_MAX_SIMD (16)
|
||||
#define GRID_MAX_SIMD (32)
|
||||
|
||||
static constexpr int MaxDims = GRID_MAX_LATTICE_DIMENSION;
|
||||
|
||||
|
@ -137,7 +137,7 @@ int main (int argc, char ** argv)
|
||||
|
||||
Eigen::MatrixXd mean(nVol, 4), stdDev(nVol, 4), rob(nVol, 4);
|
||||
Eigen::VectorXd avMean(4), avStdDev(4), avRob(4);
|
||||
double n = BENCH_IO_NPASS;
|
||||
// double n = BENCH_IO_NPASS;
|
||||
|
||||
stats(mean, stdDev, perf);
|
||||
stats(avMean, avStdDev, avPerf);
|
||||
@ -164,7 +164,7 @@ int main (int argc, char ** argv)
|
||||
mean(volInd(l), gWrite), stdDev(volInd(l), gWrite));
|
||||
}
|
||||
MSG << std::endl;
|
||||
MSG << "Robustness of individual results, in \%. (rob = 100\% - std dev / mean)" << std::endl;
|
||||
MSG << "Robustness of individual results, in %. (rob = 100% - std dev / mean)" << std::endl;
|
||||
MSG << std::endl;
|
||||
grid_printf("%4s %12s %12s %12s %12s\n",
|
||||
"L", "std read", "std write", "Grid read", "Grid write");
|
||||
@ -185,7 +185,7 @@ int main (int argc, char ** argv)
|
||||
avMean(sRead), avStdDev(sRead), avMean(sWrite), avStdDev(sWrite),
|
||||
avMean(gRead), avStdDev(gRead), avMean(gWrite), avStdDev(gWrite));
|
||||
MSG << std::endl;
|
||||
MSG << "Robustness of volume-averaged results, in \%. (rob = 100\% - std dev / mean)" << std::endl;
|
||||
MSG << "Robustness of volume-averaged results, in %. (rob = 100% - std dev / mean)" << std::endl;
|
||||
MSG << std::endl;
|
||||
grid_printf("%12s %12s %12s %12s\n",
|
||||
"std read", "std write", "Grid read", "Grid write");
|
||||
|
@ -142,7 +142,7 @@ public:
|
||||
// bzero((void *)rbuf[d],lat*lat*lat*Ls*sizeof(HalfSpinColourVectorD));
|
||||
}
|
||||
|
||||
int ncomm;
|
||||
// int ncomm;
|
||||
double dbytes;
|
||||
|
||||
for(int dir=0;dir<8;dir++) {
|
||||
@ -290,7 +290,7 @@ public:
|
||||
LatticeSU4 z(&Grid); z=Zero();
|
||||
LatticeSU4 x(&Grid); x=Zero();
|
||||
LatticeSU4 y(&Grid); y=Zero();
|
||||
double a=2.0;
|
||||
// double a=2.0;
|
||||
|
||||
uint64_t Nloop=NLOOP;
|
||||
|
||||
|
@ -72,7 +72,7 @@ int main (int argc, char ** argv)
|
||||
|
||||
std::cout << GridLogMessage << "Number of iterations to average: "<< Nloop << std::endl;
|
||||
std::vector<double> t_time(Nloop);
|
||||
time_statistics timestat;
|
||||
// time_statistics timestat;
|
||||
|
||||
std::cout<<GridLogMessage << "===================================================================================================="<<std::endl;
|
||||
std::cout<<GridLogMessage << "= Benchmarking sequential halo exchange from host memory "<<std::endl;
|
||||
|
@ -126,19 +126,10 @@ int main (int argc, char ** argv)
|
||||
// Naive wilson implementation
|
||||
////////////////////////////////////
|
||||
// replicate across fifth dimension
|
||||
LatticeGaugeFieldF Umu5d(FGrid);
|
||||
std::vector<LatticeColourMatrixF> U(4,FGrid);
|
||||
{
|
||||
autoView( Umu5d_v, Umu5d, CpuWrite);
|
||||
autoView( Umu_v , Umu , CpuRead);
|
||||
for(int ss=0;ss<Umu.Grid()->oSites();ss++){
|
||||
for(int s=0;s<Ls;s++){
|
||||
Umu5d_v[Ls*ss+s] = Umu_v[ss];
|
||||
}
|
||||
}
|
||||
}
|
||||
// LatticeGaugeFieldF Umu5d(FGrid);
|
||||
std::vector<LatticeColourMatrixF> U(4,UGrid);
|
||||
for(int mu=0;mu<Nd;mu++){
|
||||
U[mu] = PeekIndex<LorentzIndex>(Umu5d,mu);
|
||||
U[mu] = PeekIndex<LorentzIndex>(Umu,mu);
|
||||
}
|
||||
std::cout << GridLogMessage << "Setting up Cshift based reference " << std::endl;
|
||||
|
||||
@ -147,10 +138,28 @@ int main (int argc, char ** argv)
|
||||
ref = Zero();
|
||||
for(int mu=0;mu<Nd;mu++){
|
||||
|
||||
tmp = U[mu]*Cshift(src,mu+1,1);
|
||||
tmp = Cshift(src,mu+1,1);
|
||||
{
|
||||
autoView( tmp_v , tmp , CpuWrite);
|
||||
autoView( U_v , U[mu] , CpuRead);
|
||||
for(int ss=0;ss<U[mu].Grid()->oSites();ss++){
|
||||
for(int s=0;s<Ls;s++){
|
||||
tmp_v[Ls*ss+s] = U_v[ss]*tmp_v[Ls*ss+s];
|
||||
}
|
||||
}
|
||||
}
|
||||
ref=ref + tmp - Gamma(Gmu[mu])*tmp;
|
||||
|
||||
tmp =adj(U[mu])*src;
|
||||
{
|
||||
autoView( tmp_v , tmp , CpuWrite);
|
||||
autoView( U_v , U[mu] , CpuRead);
|
||||
autoView( src_v, src , CpuRead);
|
||||
for(int ss=0;ss<U[mu].Grid()->oSites();ss++){
|
||||
for(int s=0;s<Ls;s++){
|
||||
tmp_v[Ls*ss+s] = adj(U_v[ss])*src_v[Ls*ss+s];
|
||||
}
|
||||
}
|
||||
}
|
||||
tmp =Cshift(tmp,mu+1,-1);
|
||||
ref=ref + tmp + Gamma(Gmu[mu])*tmp;
|
||||
}
|
||||
@ -182,7 +191,7 @@ int main (int argc, char ** argv)
|
||||
std::cout << GridLogMessage<< "*****************************************************************" <<std::endl;
|
||||
|
||||
DomainWallFermionF Dw(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5);
|
||||
int ncall =3000;
|
||||
int ncall =300;
|
||||
|
||||
if (1) {
|
||||
FGrid->Barrier();
|
||||
@ -242,16 +251,30 @@ int main (int argc, char ** argv)
|
||||
for(int mu=0;mu<Nd;mu++){
|
||||
|
||||
// ref = src - Gamma(Gamma::Algebra::GammaX)* src ; // 1+gamma_x
|
||||
tmp = U[mu]*Cshift(src,mu+1,1);
|
||||
tmp = Cshift(src,mu+1,1);
|
||||
{
|
||||
autoView( ref_v, ref, CpuWrite);
|
||||
autoView( tmp_v, tmp, CpuRead);
|
||||
for(int i=0;i<ref_v.size();i++){
|
||||
ref_v[i]+= tmp_v[i] + Gamma(Gmu[mu])*tmp_v[i]; ;
|
||||
autoView( U_v , U[mu] , CpuRead);
|
||||
for(int ss=0;ss<U[mu].Grid()->oSites();ss++){
|
||||
for(int s=0;s<Ls;s++){
|
||||
int i=s+Ls*ss;
|
||||
ref_v[i]+= U_v[ss]*(tmp_v[i] + Gamma(Gmu[mu])*tmp_v[i]); ;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
tmp =adj(U[mu])*src;
|
||||
{
|
||||
autoView( tmp_v , tmp , CpuWrite);
|
||||
autoView( U_v , U[mu] , CpuRead);
|
||||
autoView( src_v, src , CpuRead);
|
||||
for(int ss=0;ss<U[mu].Grid()->oSites();ss++){
|
||||
for(int s=0;s<Ls;s++){
|
||||
tmp_v[Ls*ss+s] = adj(U_v[ss])*src_v[Ls*ss+s];
|
||||
}
|
||||
}
|
||||
}
|
||||
// tmp =adj(U[mu])*src;
|
||||
tmp =Cshift(tmp,mu+1,-1);
|
||||
{
|
||||
autoView( ref_v, ref, CpuWrite);
|
||||
|
@ -184,8 +184,10 @@ int main (int argc, char ** argv)
|
||||
|
||||
double bytes=1.0*vol*Nvec*sizeof(Real);
|
||||
double flops=vol*Nvec*2;// mul,add
|
||||
std::cout<<GridLogMessage<<std::setprecision(3) << lat<<"\t\t"<<bytes<<" \t\t"<<bytes/time<<"\t\t"<<flops/time<< "\t\t"<<(stop-start)/1000./1000.<< "\t\t " <<std::endl;
|
||||
|
||||
std::cout<<GridLogMessage<<std::setprecision(3) << lat<<"\t\t"
|
||||
<<bytes<<" \t\t"<<bytes/time<<"\t\t"<<flops/time<< "\t\t"
|
||||
<<(stop-start)/1000./1000.<< "\t\t " <<std::endl;
|
||||
assert(nn==nn);
|
||||
}
|
||||
|
||||
Grid_finalize();
|
||||
|
@ -4,7 +4,7 @@ using namespace Grid;
|
||||
template<class Field>
|
||||
void SimpleConjugateGradient(LinearOperatorBase<Field> &HPDop,const Field &b, Field &x)
|
||||
{
|
||||
RealD cp, c, alpha, d, beta, ssq, qq;
|
||||
RealD cp, c, alpha, d, beta, ssq;
|
||||
RealD Tolerance=1.0e-10;
|
||||
int MaxIterations=10000;
|
||||
|
||||
|
539
examples/Example_wall_wall_3pt.cc
Normal file
539
examples/Example_wall_wall_3pt.cc
Normal file
@ -0,0 +1,539 @@
|
||||
/*
|
||||
* Warning: This code illustrative only: not well tested, and not meant for production use
|
||||
* without regression / tests being applied
|
||||
*/
|
||||
|
||||
#include <Grid/Grid.h>
|
||||
|
||||
using namespace std;
|
||||
using namespace Grid;
|
||||
typedef SpinColourMatrix Propagator;
|
||||
typedef SpinColourVector Fermion;
|
||||
typedef PeriodicGimplR GimplR;
|
||||
|
||||
template<class Gimpl,class Field> class CovariantLaplacianCshift : public SparseMatrixBase<Field>
|
||||
{
|
||||
public:
|
||||
INHERIT_GIMPL_TYPES(Gimpl);
|
||||
|
||||
GridBase *grid;
|
||||
GaugeField U;
|
||||
|
||||
CovariantLaplacianCshift(GaugeField &_U) :
|
||||
grid(_U.Grid()),
|
||||
U(_U) { };
|
||||
|
||||
virtual GridBase *Grid(void) { return grid; };
|
||||
|
||||
virtual void M (const Field &in, Field &out)
|
||||
{
|
||||
out=Zero();
|
||||
for(int mu=0;mu<Nd-1;mu++) {
|
||||
GaugeLinkField Umu = PeekIndex<LorentzIndex>(U, mu); // NB: Inefficent
|
||||
out = out - Gimpl::CovShiftForward(Umu,mu,in);
|
||||
out = out - Gimpl::CovShiftBackward(Umu,mu,in);
|
||||
out = out + 2.0*in;
|
||||
}
|
||||
};
|
||||
virtual void Mdag (const Field &in, Field &out) { M(in,out);}; // Laplacian is hermitian
|
||||
virtual void Mdiag (const Field &in, Field &out) {assert(0);}; // Unimplemented need only for multigrid
|
||||
virtual void Mdir (const Field &in, Field &out,int dir, int disp){assert(0);}; // Unimplemented need only for multigrid
|
||||
virtual void MdirAll (const Field &in, std::vector<Field> &out) {assert(0);}; // Unimplemented need only for multigrid
|
||||
};
|
||||
|
||||
void MakePhase(Coordinate mom,LatticeComplex &phase)
|
||||
{
|
||||
GridBase *grid = phase.Grid();
|
||||
auto latt_size = grid->GlobalDimensions();
|
||||
ComplexD ci(0.0,1.0);
|
||||
phase=Zero();
|
||||
|
||||
LatticeComplex coor(phase.Grid());
|
||||
for(int mu=0;mu<Nd;mu++){
|
||||
RealD TwoPiL = M_PI * 2.0/ latt_size[mu];
|
||||
LatticeCoordinate(coor,mu);
|
||||
phase = phase + (TwoPiL * mom[mu]) * coor;
|
||||
}
|
||||
phase = exp(phase*ci);
|
||||
}
|
||||
void LinkSmear(int nstep, RealD rho,LatticeGaugeField &Uin,LatticeGaugeField &Usmr)
|
||||
{
|
||||
Smear_Stout<GimplR> Stout(rho);
|
||||
LatticeGaugeField Utmp(Uin.Grid());
|
||||
Utmp = Uin;
|
||||
for(int i=0;i<nstep;i++){
|
||||
Stout.smear(Usmr,Utmp);
|
||||
Utmp = Usmr;
|
||||
}
|
||||
}
|
||||
void PointSource(Coordinate &coor,LatticePropagator &source)
|
||||
{
|
||||
// Coordinate coor({0,0,0,0});
|
||||
source=Zero();
|
||||
SpinColourMatrix kronecker; kronecker=1.0;
|
||||
pokeSite(kronecker,source,coor);
|
||||
}
|
||||
void GFWallSource(int tslice,LatticePropagator &source)
|
||||
{
|
||||
GridBase *grid = source.Grid();
|
||||
LatticeComplex one(grid); one = ComplexD(1.0,0.0);
|
||||
LatticeComplex zz(grid); zz=Zero();
|
||||
LatticeInteger t(grid);
|
||||
LatticeCoordinate(t,Tdir);
|
||||
one = where(t==Integer(tslice), one, zz);
|
||||
source = 1.0;
|
||||
source = source * one;
|
||||
}
|
||||
|
||||
void Z2WallSource(GridParallelRNG &RNG,int tslice,LatticePropagator &source)
|
||||
{
|
||||
GridBase *grid = source.Grid();
|
||||
LatticeComplex noise(grid);
|
||||
LatticeComplex zz(grid); zz=Zero();
|
||||
LatticeInteger t(grid);
|
||||
|
||||
RealD nrm=1.0/sqrt(2);
|
||||
bernoulli(RNG, noise); // 0,1 50:50
|
||||
|
||||
noise = (2.*noise - Complex(1,1))*nrm;
|
||||
|
||||
LatticeCoordinate(t,Tdir);
|
||||
noise = where(t==Integer(tslice), noise, zz);
|
||||
|
||||
source = 1.0;
|
||||
source = source*noise;
|
||||
std::cout << " Z2 wall " << norm2(source) << std::endl;
|
||||
}
|
||||
void GaugeFix(LatticeGaugeField &U,LatticeGaugeField &Ufix)
|
||||
{
|
||||
Real alpha=0.05;
|
||||
|
||||
Real plaq=WilsonLoops<GimplR>::avgPlaquette(U);
|
||||
|
||||
std::cout << " Initial plaquette "<<plaq << std::endl;
|
||||
|
||||
LatticeColourMatrix xform(U.Grid());
|
||||
Ufix = U;
|
||||
int orthog=Nd-1;
|
||||
FourierAcceleratedGaugeFixer<GimplR>::SteepestDescentGaugeFix(Ufix,xform,alpha,100000,1.0e-14, 1.0e-14,true,orthog);
|
||||
|
||||
plaq=WilsonLoops<GimplR>::avgPlaquette(Ufix);
|
||||
|
||||
std::cout << " Final plaquette "<<plaq << std::endl;
|
||||
}
|
||||
template<class Field>
|
||||
void GaussianSmear(LatticeGaugeField &U,Field &unsmeared,Field &smeared)
|
||||
{
|
||||
typedef CovariantLaplacianCshift <GimplR,Field> Laplacian_t;
|
||||
Laplacian_t Laplacian(U);
|
||||
|
||||
Integer Iterations = 40;
|
||||
Real width = 2.0;
|
||||
Real coeff = (width*width) / Real(4*Iterations);
|
||||
|
||||
Field tmp(U.Grid());
|
||||
smeared=unsmeared;
|
||||
// chi = (1-p^2/2N)^N kronecker
|
||||
for(int n = 0; n < Iterations; ++n) {
|
||||
Laplacian.M(smeared,tmp);
|
||||
smeared = smeared - coeff*tmp;
|
||||
std::cout << " smear iter " << n<<" " <<norm2(smeared)<<std::endl;
|
||||
}
|
||||
}
|
||||
void GaussianSource(Coordinate &site,LatticeGaugeField &U,LatticePropagator &source)
|
||||
{
|
||||
LatticePropagator tmp(source.Grid());
|
||||
PointSource(site,source);
|
||||
std::cout << " GaussianSource Kronecker "<< norm2(source)<<std::endl;
|
||||
tmp = source;
|
||||
GaussianSmear(U,tmp,source);
|
||||
std::cout << " GaussianSource Smeared "<< norm2(source)<<std::endl;
|
||||
}
|
||||
void GaussianWallSource(GridParallelRNG &RNG,int tslice,LatticeGaugeField &U,LatticePropagator &source)
|
||||
{
|
||||
Z2WallSource(RNG,tslice,source);
|
||||
auto tmp = source;
|
||||
GaussianSmear(U,tmp,source);
|
||||
}
|
||||
void SequentialSource(int tslice,Coordinate &mom,LatticePropagator &spectator,LatticePropagator &source)
|
||||
{
|
||||
assert(mom.size()==Nd);
|
||||
assert(mom[Tdir] == 0);
|
||||
|
||||
GridBase * grid = spectator.Grid();
|
||||
|
||||
LatticeInteger ts(grid);
|
||||
LatticeCoordinate(ts,Tdir);
|
||||
source = Zero();
|
||||
source = where(ts==Integer(tslice),spectator,source); // Stick in a slice of the spectator, zero everywhere else
|
||||
|
||||
LatticeComplex phase(grid);
|
||||
MakePhase(mom,phase);
|
||||
|
||||
source = source *phase;
|
||||
}
|
||||
template<class Action>
|
||||
void Solve(Action &D,LatticePropagator &source,LatticePropagator &propagator)
|
||||
{
|
||||
GridBase *UGrid = D.GaugeGrid();
|
||||
GridBase *FGrid = D.FermionGrid();
|
||||
|
||||
LatticeFermion src4 (UGrid);
|
||||
LatticeFermion src5 (FGrid);
|
||||
LatticeFermion result5(FGrid);
|
||||
LatticeFermion result4(UGrid);
|
||||
|
||||
ConjugateGradient<LatticeFermion> CG(1.0e-12,100000);
|
||||
SchurRedBlackDiagTwoSolve<LatticeFermion> schur(CG);
|
||||
ZeroGuesser<LatticeFermion> ZG; // Could be a DeflatedGuesser if have eigenvectors
|
||||
for(int s=0;s<Nd;s++){
|
||||
for(int c=0;c<Nc;c++){
|
||||
PropToFerm<Action>(src4,source,s,c);
|
||||
|
||||
D.ImportPhysicalFermionSource(src4,src5);
|
||||
|
||||
result5=Zero();
|
||||
schur(D,src5,result5,ZG);
|
||||
std::cout<<GridLogMessage
|
||||
<<"spin "<<s<<" color "<<c
|
||||
<<" norm2(src5d) " <<norm2(src5)
|
||||
<<" norm2(result5d) "<<norm2(result5)<<std::endl;
|
||||
|
||||
D.ExportPhysicalFermionSolution(result5,result4);
|
||||
|
||||
FermToProp<Action>(propagator,result4,s,c);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
class MesonFile: Serializable {
|
||||
public:
|
||||
GRID_SERIALIZABLE_CLASS_MEMBERS(MesonFile, std::vector<std::vector<Complex> >, data);
|
||||
};
|
||||
|
||||
void MesonTrace(std::string file,LatticePropagator &q1,LatticePropagator &q2,LatticeComplex &phase)
|
||||
{
|
||||
const int nchannel=4;
|
||||
Gamma::Algebra Gammas[nchannel][2] = {
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::Gamma5},
|
||||
{Gamma::Algebra::GammaTGamma5,Gamma::Algebra::GammaTGamma5},
|
||||
{Gamma::Algebra::GammaTGamma5,Gamma::Algebra::Gamma5},
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaTGamma5}
|
||||
};
|
||||
|
||||
Gamma G5(Gamma::Algebra::Gamma5);
|
||||
|
||||
LatticeComplex meson_CF(q1.Grid());
|
||||
MesonFile MF;
|
||||
|
||||
for(int ch=0;ch<nchannel;ch++){
|
||||
|
||||
Gamma Gsrc(Gammas[ch][0]);
|
||||
Gamma Gsnk(Gammas[ch][1]);
|
||||
|
||||
meson_CF = trace(G5*adj(q1)*G5*Gsnk*q2*adj(Gsrc));
|
||||
|
||||
std::vector<TComplex> meson_T;
|
||||
sliceSum(meson_CF,meson_T, Tdir);
|
||||
|
||||
int nt=meson_T.size();
|
||||
|
||||
std::vector<Complex> corr(nt);
|
||||
for(int t=0;t<nt;t++){
|
||||
corr[t] = TensorRemove(meson_T[t]); // Yes this is ugly, not figured a work around
|
||||
std::cout << " channel "<<ch<<" t "<<t<<" " <<corr[t]<<std::endl;
|
||||
}
|
||||
MF.data.push_back(corr);
|
||||
}
|
||||
|
||||
{
|
||||
XmlWriter WR(file);
|
||||
write(WR,"MesonFile",MF);
|
||||
}
|
||||
}
|
||||
|
||||
void Meson3pt(std::string file,LatticePropagator &q1,LatticePropagator &q2,LatticeComplex &phase)
|
||||
{
|
||||
const int nchannel=4;
|
||||
Gamma::Algebra Gammas[nchannel][2] = {
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaX},
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaY},
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaZ},
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaT}
|
||||
};
|
||||
|
||||
Gamma G5(Gamma::Algebra::Gamma5);
|
||||
|
||||
LatticeComplex meson_CF(q1.Grid());
|
||||
MesonFile MF;
|
||||
|
||||
for(int ch=0;ch<nchannel;ch++){
|
||||
|
||||
Gamma Gsrc(Gammas[ch][0]);
|
||||
Gamma Gsnk(Gammas[ch][1]);
|
||||
|
||||
meson_CF = trace(G5*adj(q1)*G5*Gsnk*q2*adj(Gsrc));
|
||||
|
||||
std::vector<TComplex> meson_T;
|
||||
sliceSum(meson_CF,meson_T, Tdir);
|
||||
|
||||
int nt=meson_T.size();
|
||||
|
||||
std::vector<Complex> corr(nt);
|
||||
for(int t=0;t<nt;t++){
|
||||
corr[t] = TensorRemove(meson_T[t]); // Yes this is ugly, not figured a work around
|
||||
std::cout << " channel "<<ch<<" t "<<t<<" " <<corr[t]<<std::endl;
|
||||
}
|
||||
MF.data.push_back(corr);
|
||||
}
|
||||
|
||||
{
|
||||
XmlWriter WR(file);
|
||||
write(WR,"MesonFile",MF);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void WallSinkMesonTrace(std::string file,std::vector<Propagator> &q1,std::vector<Propagator> &q2)
|
||||
{
|
||||
const int nchannel=4;
|
||||
Gamma::Algebra Gammas[nchannel][2] = {
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::Gamma5},
|
||||
{Gamma::Algebra::GammaTGamma5,Gamma::Algebra::GammaTGamma5},
|
||||
{Gamma::Algebra::GammaTGamma5,Gamma::Algebra::Gamma5},
|
||||
{Gamma::Algebra::Gamma5 ,Gamma::Algebra::GammaTGamma5}
|
||||
};
|
||||
|
||||
Gamma G5(Gamma::Algebra::Gamma5);
|
||||
int nt=q1.size();
|
||||
std::vector<Complex> meson_CF(nt);
|
||||
MesonFile MF;
|
||||
|
||||
for(int ch=0;ch<nchannel;ch++){
|
||||
|
||||
Gamma Gsrc(Gammas[ch][0]);
|
||||
Gamma Gsnk(Gammas[ch][1]);
|
||||
|
||||
std::vector<Complex> corr(nt);
|
||||
for(int t=0;t<nt;t++){
|
||||
meson_CF[t] = trace(G5*adj(q1[t])*G5*Gsnk*q2[t]*adj(Gsrc));
|
||||
corr[t] = TensorRemove(meson_CF[t]); // Yes this is ugly, not figured a work around
|
||||
std::cout << " channel "<<ch<<" t "<<t<<" " <<corr[t]<<std::endl;
|
||||
}
|
||||
MF.data.push_back(corr);
|
||||
}
|
||||
|
||||
{
|
||||
XmlWriter WR(file);
|
||||
write(WR,"MesonFile",MF);
|
||||
}
|
||||
}
|
||||
int make_idx(int p, int m,int nmom)
|
||||
{
|
||||
if (m==0) return p;
|
||||
assert(p==0);
|
||||
return nmom + m - 1;
|
||||
}
|
||||
|
||||
int main (int argc, char ** argv)
|
||||
{
|
||||
Grid_init(&argc,&argv);
|
||||
|
||||
// Double precision grids
|
||||
auto latt = GridDefaultLatt();
|
||||
GridCartesian * UGrid = SpaceTimeGrid::makeFourDimGrid(GridDefaultLatt(),
|
||||
GridDefaultSimd(Nd,vComplex::Nsimd()),
|
||||
GridDefaultMpi());
|
||||
GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
|
||||
|
||||
|
||||
LatticeGaugeField Umu(UGrid);
|
||||
LatticeGaugeField Utmp(UGrid);
|
||||
LatticeGaugeField Usmr(UGrid);
|
||||
std::string config;
|
||||
if( argc > 1 && argv[1][0] != '-' )
|
||||
{
|
||||
std::cout<<GridLogMessage <<"Loading configuration from "<<argv[1]<<std::endl;
|
||||
FieldMetaData header;
|
||||
NerscIO::readConfiguration(Umu, header, argv[1]);
|
||||
config=argv[1];
|
||||
}
|
||||
else
|
||||
{
|
||||
std::cout<<GridLogMessage <<"Using hot configuration"<<std::endl;
|
||||
SU<Nc>::ColdConfiguration(Umu);
|
||||
config="ColdConfig";
|
||||
}
|
||||
// GaugeFix(Umu,Utmp);
|
||||
// Umu=Utmp;
|
||||
|
||||
int nsmr=3;
|
||||
RealD rho=0.1;
|
||||
LinkSmear(nsmr,rho,Umu,Usmr);
|
||||
|
||||
|
||||
std::vector<int> smeared_link({ 0,0,1} );
|
||||
std::vector<RealD> masses({ 0.004,0.02477,0.447} ); // u/d, s, c ??
|
||||
std::vector<RealD> M5s ({ 1.8,1.8,1.0} );
|
||||
std::vector<RealD> bs ({ 1.0,1.0,1.5} ); // DDM
|
||||
std::vector<RealD> cs ({ 0.0,0.0,0.5} ); // DDM
|
||||
std::vector<int> Ls_s ({ 16,16,12} );
|
||||
std::vector<GridCartesian *> FGrids;
|
||||
std::vector<GridRedBlackCartesian *> FrbGrids;
|
||||
|
||||
std::vector<Coordinate> momenta;
|
||||
momenta.push_back(Coordinate({0,0,0,0}));
|
||||
momenta.push_back(Coordinate({1,0,0,0}));
|
||||
momenta.push_back(Coordinate({2,0,0,0}));
|
||||
|
||||
int nmass = masses.size();
|
||||
int nmom = momenta.size();
|
||||
|
||||
std::vector<MobiusFermionR *> FermActs;
|
||||
|
||||
std::cout<<GridLogMessage <<"======================"<<std::endl;
|
||||
std::cout<<GridLogMessage <<"MobiusFermion action as Scaled Shamir kernel"<<std::endl;
|
||||
std::cout<<GridLogMessage <<"======================"<<std::endl;
|
||||
|
||||
std::vector<Complex> boundary = {1,1,1,-1};
|
||||
typedef MobiusFermionR FermionAction;
|
||||
FermionAction::ImplParams Params(boundary);
|
||||
|
||||
for(int m=0;m<masses.size();m++) {
|
||||
|
||||
RealD mass = masses[m];
|
||||
RealD M5 = M5s[m];
|
||||
RealD b = bs[m];
|
||||
RealD c = cs[m];
|
||||
int Ls = Ls_s[m];
|
||||
|
||||
if ( smeared_link[m] ) Utmp = Usmr;
|
||||
else Utmp = Umu;
|
||||
|
||||
FGrids.push_back(SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid));
|
||||
FrbGrids.push_back(SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid));
|
||||
|
||||
FermActs.push_back(new MobiusFermionR(Utmp,*FGrids[m],*FrbGrids[m],*UGrid,*UrbGrid,mass,M5,b,c,Params));
|
||||
}
|
||||
|
||||
LatticePropagator z2wall_source(UGrid);
|
||||
LatticePropagator gfwall_source(UGrid);
|
||||
LatticePropagator phased_prop(UGrid);
|
||||
|
||||
int tslice = 0;
|
||||
int tseq=(tslice+16)%latt[Nd-1];
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// RNG seeded for Z2 wall
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// You can manage seeds however you like.
|
||||
// Recommend SeedUniqueString.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
GridParallelRNG RNG4(UGrid); RNG4.SeedUniqueString("Study2-Source_Z2_p_0_0_0_t_0-880");
|
||||
Z2WallSource (RNG4,tslice,z2wall_source);
|
||||
GFWallSource (tslice,gfwall_source);
|
||||
|
||||
std::vector<LatticeComplex> phase(nmom,UGrid);
|
||||
for(int m=0;m<nmom;m++){
|
||||
MakePhase(momenta[m],phase[m]);
|
||||
}
|
||||
|
||||
std::vector<LatticePropagator> Z2Props (nmom+nmass-1,UGrid);
|
||||
std::vector<LatticePropagator> GFProps (nmom+nmass-1,UGrid);
|
||||
for(int p=0;p<nmom;p++) {
|
||||
int m=0;
|
||||
int idx = make_idx(p,m,nmom);
|
||||
phased_prop = z2wall_source * phase[p];
|
||||
Solve(*FermActs[m],phased_prop ,Z2Props[idx]);
|
||||
|
||||
phased_prop = gfwall_source * phase[p];
|
||||
Solve(*FermActs[m],phased_prop ,GFProps[idx]);
|
||||
}
|
||||
for(int m=1;m<nmass;m++) {
|
||||
int p=0;
|
||||
int idx = make_idx(p,m,nmom);
|
||||
phased_prop = z2wall_source;
|
||||
Solve(*FermActs[m],phased_prop ,Z2Props[idx]);
|
||||
|
||||
phased_prop = gfwall_source;
|
||||
Solve(*FermActs[m],phased_prop ,GFProps[idx]);
|
||||
}
|
||||
|
||||
std::vector<std::vector<Propagator> > wsnk_z2Props(nmom+nmass-1);
|
||||
std::vector<std::vector<Propagator> > wsnk_gfProps(nmom+nmass-1);
|
||||
|
||||
// Non-zero kaon and point and D two point
|
||||
// WW stick momentum on m1 (lighter)
|
||||
// zero momentum on m2
|
||||
for(int m1=0;m1<nmass;m1++) {
|
||||
for(int m2=m1;m2<nmass;m2++) {
|
||||
int pmax = (m1==0)? nmom:1;
|
||||
for(int p=0;p<pmax;p++){
|
||||
|
||||
std::stringstream ssg,ssz;
|
||||
std::stringstream wssg,wssz;
|
||||
|
||||
int idx1 = make_idx(p,m1,nmom);
|
||||
int idx2 = make_idx(0,m2,nmom);
|
||||
|
||||
/// Point sinks
|
||||
ssg<<config<<"_p"<<p<< "_m" << m1 << "_m"<< m2 << "_p_gf_meson.xml";
|
||||
ssz<<config<<"_p"<<p<< "_m" << m1 << "_m"<< m2 << "_p_z2_meson.xml";
|
||||
MesonTrace(ssz.str(),Z2Props[idx1],Z2Props[idx2],phase[p]); // Q1 is conjugated
|
||||
MesonTrace(ssg.str(),GFProps[idx1],GFProps[idx2],phase[p]);
|
||||
|
||||
/// Wall sinks
|
||||
wssg<<config<<"_p"<<p<< "_m" << m1 << "_m"<< m2 << "_w_gf_meson.xml";
|
||||
wssz<<config<<"_p"<<p<< "_m" << m1 << "_m"<< m2 << "_w_z2_meson.xml";
|
||||
|
||||
phased_prop = GFProps[m2] * phase[p];
|
||||
sliceSum(phased_prop,wsnk_gfProps[m1],Tdir);
|
||||
sliceSum(GFProps[m1],wsnk_gfProps[m2],Tdir);
|
||||
WallSinkMesonTrace(wssg.str(),wsnk_gfProps[m1],wsnk_gfProps[m2]);
|
||||
|
||||
phased_prop = Z2Props[m2] * phase[p];
|
||||
sliceSum(phased_prop,wsnk_gfProps[m1],Tdir);
|
||||
sliceSum(Z2Props[m1],wsnk_gfProps[m2],Tdir);
|
||||
WallSinkMesonTrace(wssz.str(),wsnk_z2Props[m1],wsnk_z2Props[m2]);
|
||||
}
|
||||
}}
|
||||
|
||||
|
||||
/////////////////////////////////////
|
||||
// Sequential solves
|
||||
/////////////////////////////////////
|
||||
LatticePropagator seq_wsnk_z2src(UGrid);
|
||||
LatticePropagator seq_wsnk_gfsrc(UGrid);
|
||||
LatticePropagator seq_psnk_z2src(UGrid);
|
||||
LatticePropagator seq_psnk_gfsrc(UGrid);
|
||||
LatticePropagator source(UGrid);
|
||||
for(int m=0;m<nmass-1;m++){
|
||||
int spect_idx = make_idx(0,m,nmom);
|
||||
int charm=nmass-1;
|
||||
|
||||
SequentialSource(tseq,momenta[0],GFProps[spect_idx],source);
|
||||
Solve(*FermActs[charm],source,seq_psnk_gfsrc);
|
||||
|
||||
SequentialSource(tseq,momenta[0],Z2Props[spect_idx],source);
|
||||
Solve(*FermActs[charm],source,seq_psnk_z2src);
|
||||
|
||||
// Todo need wall sequential solve
|
||||
for(int p=0;p<nmom;p++){
|
||||
int active_idx = make_idx(p,0,nmom);
|
||||
std::stringstream seq_3pt_p_z2;
|
||||
std::stringstream seq_3pt_p_gf;
|
||||
std::stringstream seq_3pt_w_z2;
|
||||
std::stringstream seq_3pt_w_gf;
|
||||
seq_3pt_p_z2 <<config<<"_3pt_p"<<p<< "_m" << m << "_p_z2_meson.xml";
|
||||
seq_3pt_p_gf <<config<<"_3pt_p"<<p<< "_m" << m << "_p_gf_meson.xml";
|
||||
seq_3pt_w_z2 <<config<<"_3pt_p"<<p<< "_m" << m << "_w_z2_meson.xml";
|
||||
seq_3pt_w_gf <<config<<"_3pt_p"<<p<< "_m" << m << "_w_gf_meson.xml";
|
||||
Meson3pt(seq_3pt_p_gf.str(),GFProps[active_idx],seq_psnk_gfsrc,phase[p]);
|
||||
Meson3pt(seq_3pt_p_z2.str(),Z2Props[active_idx],seq_psnk_z2src,phase[p]);
|
||||
}
|
||||
}
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
|
||||
|
||||
|
@ -9,6 +9,7 @@ using namespace std;
|
||||
using namespace Grid;
|
||||
typedef SpinColourMatrix Propagator;
|
||||
typedef SpinColourVector Fermion;
|
||||
typedef PeriodicGimplR GimplR;
|
||||
|
||||
template<class Gimpl,class Field> class CovariantLaplacianCshift : public SparseMatrixBase<Field>
|
||||
{
|
||||
@ -55,6 +56,16 @@ void MakePhase(Coordinate mom,LatticeComplex &phase)
|
||||
}
|
||||
phase = exp(phase*ci);
|
||||
}
|
||||
void LinkSmear(int nstep, RealD rho,LatticeGaugeField &Uin,LatticeGaugeField &Usmr)
|
||||
{
|
||||
Smear_Stout<GimplR> Stout(rho);
|
||||
LatticeGaugeField Utmp(Uin.Grid());
|
||||
Utmp = Uin;
|
||||
for(int i=0;i<nstep;i++){
|
||||
Stout.smear(Usmr,Utmp);
|
||||
Utmp = Usmr;
|
||||
}
|
||||
}
|
||||
void PointSource(Coordinate &coor,LatticePropagator &source)
|
||||
{
|
||||
// Coordinate coor({0,0,0,0});
|
||||
@ -97,23 +108,23 @@ void GaugeFix(LatticeGaugeField &U,LatticeGaugeField &Ufix)
|
||||
{
|
||||
Real alpha=0.05;
|
||||
|
||||
Real plaq=WilsonLoops<PeriodicGimplR>::avgPlaquette(U);
|
||||
Real plaq=WilsonLoops<GimplR>::avgPlaquette(U);
|
||||
|
||||
std::cout << " Initial plaquette "<<plaq << std::endl;
|
||||
|
||||
LatticeColourMatrix xform(U.Grid());
|
||||
Ufix = U;
|
||||
int orthog=Nd-1;
|
||||
FourierAcceleratedGaugeFixer<PeriodicGimplR>::SteepestDescentGaugeFix(Ufix,xform,alpha,10000,1.0e-12, 1.0e-12,true,orthog);
|
||||
FourierAcceleratedGaugeFixer<GimplR>::SteepestDescentGaugeFix(Ufix,xform,alpha,100000,1.0e-14, 1.0e-14,true,orthog);
|
||||
|
||||
plaq=WilsonLoops<PeriodicGimplR>::avgPlaquette(Ufix);
|
||||
plaq=WilsonLoops<GimplR>::avgPlaquette(Ufix);
|
||||
|
||||
std::cout << " Final plaquette "<<plaq << std::endl;
|
||||
}
|
||||
template<class Field>
|
||||
void GaussianSmear(LatticeGaugeField &U,Field &unsmeared,Field &smeared)
|
||||
{
|
||||
typedef CovariantLaplacianCshift <PeriodicGimplR,Field> Laplacian_t;
|
||||
typedef CovariantLaplacianCshift <GimplR,Field> Laplacian_t;
|
||||
Laplacian_t Laplacian(U);
|
||||
|
||||
Integer Iterations = 40;
|
||||
@ -167,19 +178,21 @@ void Solve(Action &D,LatticePropagator &source,LatticePropagator &propagator)
|
||||
GridBase *UGrid = D.GaugeGrid();
|
||||
GridBase *FGrid = D.FermionGrid();
|
||||
|
||||
LatticeFermion src4 (UGrid);
|
||||
LatticeFermion src4 (UGrid); src4 = Zero();
|
||||
LatticeFermion src5 (FGrid);
|
||||
LatticeFermion result5(FGrid);
|
||||
LatticeFermion result4(UGrid);
|
||||
|
||||
ConjugateGradient<LatticeFermion> CG(1.0e-8,100000);
|
||||
SchurRedBlackDiagMooeeSolve<LatticeFermion> schur(CG);
|
||||
ConjugateGradient<LatticeFermion> CG(1.0e-12,100000);
|
||||
SchurRedBlackDiagTwoSolve<LatticeFermion> schur(CG);
|
||||
ZeroGuesser<LatticeFermion> ZG; // Could be a DeflatedGuesser if have eigenvectors
|
||||
std::cout<<GridLogMessage<< " source4 "<<norm2(source)<<std::endl;
|
||||
for(int s=0;s<Nd;s++){
|
||||
for(int c=0;c<Nc;c++){
|
||||
PropToFerm<Action>(src4,source,s,c);
|
||||
|
||||
std::cout<<GridLogMessage<< s<<c<<" src4 "<<norm2(src4)<<std::endl;
|
||||
D.ImportPhysicalFermionSource(src4,src5);
|
||||
std::cout<<GridLogMessage<< s<<c<<" src5 "<<norm2(src5)<<std::endl;
|
||||
|
||||
result5=Zero();
|
||||
schur(D,src5,result5,ZG);
|
||||
@ -287,15 +300,10 @@ int main (int argc, char ** argv)
|
||||
GridDefaultMpi());
|
||||
GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// You can manage seeds however you like.
|
||||
// Recommend SeedUniqueString.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
std::vector<int> seeds4({1,2,3,4});
|
||||
GridParallelRNG RNG4(UGrid); RNG4.SeedFixedIntegers(seeds4);
|
||||
|
||||
LatticeGaugeField Umu(UGrid);
|
||||
LatticeGaugeField Ufixed(UGrid);
|
||||
LatticeGaugeField Utmp(UGrid);
|
||||
LatticeGaugeField Usmr(UGrid);
|
||||
std::string config;
|
||||
if( argc > 1 && argv[1][0] != '-' )
|
||||
{
|
||||
@ -308,13 +316,20 @@ int main (int argc, char ** argv)
|
||||
{
|
||||
std::cout<<GridLogMessage <<"Using hot configuration"<<std::endl;
|
||||
SU<Nc>::ColdConfiguration(Umu);
|
||||
// SU<Nc>::HotConfiguration(RNG4,Umu);
|
||||
config="HotConfig";
|
||||
config="ColdConfig";
|
||||
}
|
||||
GaugeFix(Umu,Ufixed);
|
||||
Umu=Ufixed;
|
||||
// GaugeFix(Umu,Utmp);
|
||||
// Umu=Utmp;
|
||||
|
||||
int nsmr=3;
|
||||
RealD rho=0.1;
|
||||
RealD plaq_gf =WilsonLoops<GimplR>::avgPlaquette(Umu);
|
||||
LinkSmear(nsmr,rho,Umu,Usmr);
|
||||
RealD plaq_smr=WilsonLoops<GimplR>::avgPlaquette(Usmr);
|
||||
std::cout << GridLogMessage << " GF Plaquette " <<plaq_gf<<std::endl;
|
||||
std::cout << GridLogMessage << " SM Plaquette " <<plaq_smr<<std::endl;
|
||||
|
||||
std::vector<int> smeared_link({ 0,0,1} );
|
||||
std::vector<RealD> masses({ 0.004,0.02477,0.447} ); // u/d, s, c ??
|
||||
std::vector<RealD> M5s ({ 1.8,1.8,1.0} );
|
||||
std::vector<RealD> bs ({ 1.0,1.0,1.5} ); // DDM
|
||||
@ -330,6 +345,9 @@ int main (int argc, char ** argv)
|
||||
std::cout<<GridLogMessage <<"======================"<<std::endl;
|
||||
std::cout<<GridLogMessage <<"MobiusFermion action as Scaled Shamir kernel"<<std::endl;
|
||||
std::cout<<GridLogMessage <<"======================"<<std::endl;
|
||||
std::vector<Complex> boundary = {1,1,1,-1};
|
||||
typedef MobiusFermionR FermionAction;
|
||||
FermionAction::ImplParams Params(boundary);
|
||||
|
||||
for(int m=0;m<masses.size();m++) {
|
||||
|
||||
@ -339,31 +357,41 @@ int main (int argc, char ** argv)
|
||||
RealD c = cs[m];
|
||||
int Ls = Ls_s[m];
|
||||
|
||||
if ( smeared_link[m] ) Utmp = Usmr;
|
||||
else Utmp = Umu;
|
||||
|
||||
FGrids.push_back(SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid));
|
||||
FrbGrids.push_back(SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid));
|
||||
|
||||
FermActs.push_back(new MobiusFermionR(Umu,*FGrids[m],*FrbGrids[m],*UGrid,*UrbGrid,mass,M5,b,c));
|
||||
FermActs.push_back(new MobiusFermionR(Utmp,*FGrids[m],*FrbGrids[m],*UGrid,*UrbGrid,mass,M5,b,c,Params));
|
||||
}
|
||||
|
||||
LatticePropagator point_source(UGrid);
|
||||
LatticePropagator z2wall_source(UGrid);
|
||||
LatticePropagator gfwall_source(UGrid);
|
||||
|
||||
Coordinate Origin({0,0,0,0});
|
||||
PointSource (Origin,point_source);
|
||||
Z2WallSource (RNG4,0,z2wall_source);
|
||||
GFWallSource (0,gfwall_source);
|
||||
int tslice = 0;
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// RNG seeded for Z2 wall
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// You can manage seeds however you like.
|
||||
// Recommend SeedUniqueString.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
GridParallelRNG RNG4(UGrid); RNG4.SeedUniqueString("Study2-Source_Z2_p_0_0_0_t_0-880");
|
||||
Z2WallSource (RNG4,tslice,z2wall_source);
|
||||
GFWallSource (tslice,gfwall_source);
|
||||
|
||||
std::vector<LatticePropagator> PointProps(nmass,UGrid);
|
||||
std::vector<LatticePropagator> GaussProps(nmass,UGrid);
|
||||
std::vector<LatticePropagator> Z2Props (nmass,UGrid);
|
||||
std::vector<LatticePropagator> GFProps (nmass,UGrid);
|
||||
|
||||
for(int m=0;m<nmass;m++) {
|
||||
|
||||
std::cout << GridLogMessage << " Mass " <<m << " z2wall source "<<norm2(z2wall_source)<<std::endl;
|
||||
Solve(*FermActs[m],z2wall_source ,Z2Props[m]);
|
||||
std::cout << GridLogMessage << " Mass " <<m << " gfwall source "<<norm2(gfwall_source)<<std::endl;
|
||||
Solve(*FermActs[m],gfwall_source ,GFProps[m]);
|
||||
|
||||
std::cout << GridLogMessage << " Mass " <<m << " z2wall source "<<norm2(z2wall_source)<< " " << norm2(gfwall_source)<<std::endl;
|
||||
|
||||
}
|
||||
|
||||
LatticeComplex phase(UGrid);
|
||||
@ -383,14 +411,15 @@ int main (int argc, char ** argv)
|
||||
std::stringstream wssg,wssz;
|
||||
|
||||
/// Point sinks
|
||||
ssg<<config<< "_m" << m1 << "_m"<< m2 << "p_gf_meson.xml";
|
||||
ssz<<config<< "_m" << m1 << "_m"<< m2 << "p_z2_meson.xml";
|
||||
ssg<<config<< "_m" << m1 << "_m"<< m2 << "_p_gf_meson.xml";
|
||||
ssz<<config<< "_m" << m1 << "_m"<< m2 << "_p_z2_meson.xml";
|
||||
|
||||
MesonTrace(ssz.str(),Z2Props[m1],Z2Props[m2],phase);
|
||||
MesonTrace(ssg.str(),GFProps[m1],GFProps[m2],phase);
|
||||
|
||||
/// Wall sinks
|
||||
wssg<<config<< "_m" << m1 << "_m"<< m2 << "w_gf_meson.xml";
|
||||
wssz<<config<< "_m" << m1 << "_m"<< m2 << "w_z2_meson.xml";
|
||||
wssg<<config<< "_m" << m1 << "_m"<< m2 << "_w_gf_meson.xml";
|
||||
wssz<<config<< "_m" << m1 << "_m"<< m2 << "_w_z2_meson.xml";
|
||||
|
||||
WallSinkMesonTrace(wssg.str(),wsnk_gfProps[m1],wsnk_gfProps[m2]);
|
||||
WallSinkMesonTrace(wssz.str(),wsnk_z2Props[m1],wsnk_z2Props[m2]);
|
||||
|
26
systems/Spock/comms.slurm
Normal file
26
systems/Spock/comms.slurm
Normal file
@ -0,0 +1,26 @@
|
||||
#!/bin/bash
|
||||
# Begin LSF Directives
|
||||
#SBATCH -A LGT104
|
||||
#SBATCH -t 01:00:00
|
||||
##SBATCH -U openmpThu
|
||||
#SBATCH -p ecp
|
||||
#SBATCH -J comms
|
||||
#SBATCH -o comms.%J
|
||||
#SBATCH -e comms.%J
|
||||
#SBATCH -N 1
|
||||
#SBATCH -n 2
|
||||
|
||||
DIR=.
|
||||
module list
|
||||
export MPIR_CVAR_GPU_EAGER_DEVICE_MEM=0
|
||||
export MPICH_GPU_SUPPORT_ENABLED=1
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=XPMEM
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=CMA
|
||||
export MPICH_SMP_SINGLE_COPY_MODE=NONE
|
||||
export OMP_NUM_THREADS=8
|
||||
|
||||
AT=8
|
||||
echo MPICH_SMP_SINGLE_COPY_MODE $MPICH_SMP_SINGLE_COPY_MODE
|
||||
PARAMS=" --accelerator-threads ${AT} --grid 64.64.32.32 --mpi 2.1.1.1 "
|
||||
srun -n2 --label -c$OMP_NUM_THREADS --gpus-per-task=1 ./mpiwrapper.sh ./benchmarks/Benchmark_comms_host_device $PARAMS
|
||||
|
12
systems/Spock/config-command
Normal file
12
systems/Spock/config-command
Normal file
@ -0,0 +1,12 @@
|
||||
../../configure --enable-comms=mpi-auto \
|
||||
--enable-unified=no \
|
||||
--enable-shm=nvlink \
|
||||
--enable-accelerator=hip \
|
||||
--enable-gen-simd-width=64 \
|
||||
--enable-simd=GPU \
|
||||
--disable-fermion-reps \
|
||||
--disable-gparity \
|
||||
CXX=hipcc MPICXX=mpicxx \
|
||||
CXXFLAGS="-fPIC -I/opt/rocm-4.3.0/include/ -std=c++14 -I${MPICH_DIR}/include " \
|
||||
--prefix=/ccs/home/chulwoo/Grid \
|
||||
LDFLAGS=" -L${MPICH_DIR}/lib -lmpi -L${CRAY_MPICH_ROOTDIR}/gtl/lib -lmpi_gtl_hsa "
|
26
systems/Spock/dwf.slurm
Normal file
26
systems/Spock/dwf.slurm
Normal file
@ -0,0 +1,26 @@
|
||||
#!/bin/bash
|
||||
# Begin LSF Directives
|
||||
#SBATCH -A LGT104
|
||||
#SBATCH -t 01:00:00
|
||||
##SBATCH -U openmpThu
|
||||
#SBATCH -p ecp
|
||||
#SBATCH -J DWF
|
||||
#SBATCH -o DWF.%J
|
||||
#SBATCH -e DWF.%J
|
||||
#SBATCH -N 1
|
||||
#SBATCH -n 1
|
||||
|
||||
DIR=.
|
||||
module list
|
||||
export MPIR_CVAR_GPU_EAGER_DEVICE_MEM=0
|
||||
export MPICH_GPU_SUPPORT_ENABLED=1
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=XPMEM
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=NONE
|
||||
export MPICH_SMP_SINGLE_COPY_MODE=CMA
|
||||
export OMP_NUM_THREADS=8
|
||||
|
||||
AT=8
|
||||
echo MPICH_SMP_SINGLE_COPY_MODE $MPICH_SMP_SINGLE_COPY_MODE
|
||||
PARAMS=" --accelerator-threads ${AT} --grid 32.32.32.32 --mpi 1.1.1.1 --comms-overlap"
|
||||
srun -n1 --label -c$OMP_NUM_THREADS --gpus-per-task=1 ./mpiwrapper.sh ./benchmarks/Benchmark_dwf_fp32 $PARAMS
|
||||
|
26
systems/Spock/dwf4.slurm
Normal file
26
systems/Spock/dwf4.slurm
Normal file
@ -0,0 +1,26 @@
|
||||
#!/bin/bash
|
||||
# Begin LSF Directives
|
||||
#SBATCH -A LGT104
|
||||
#SBATCH -t 01:00:00
|
||||
##SBATCH -U openmpThu
|
||||
#SBATCH -p ecp
|
||||
#SBATCH -J DWF
|
||||
#SBATCH -o DWF.%J
|
||||
#SBATCH -e DWF.%J
|
||||
#SBATCH -N 1
|
||||
#SBATCH -n 4
|
||||
|
||||
DIR=.
|
||||
module list
|
||||
export MPIR_CVAR_GPU_EAGER_DEVICE_MEM=0
|
||||
export MPICH_GPU_SUPPORT_ENABLED=1
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=XPMEM
|
||||
export MPICH_SMP_SINGLE_COPY_MODE=NONE
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=CMA
|
||||
export OMP_NUM_THREADS=8
|
||||
|
||||
AT=8
|
||||
echo MPICH_SMP_SINGLE_COPY_MODE $MPICH_SMP_SINGLE_COPY_MODE
|
||||
PARAMS=" --accelerator-threads ${AT} --grid 32.32.64.64 --mpi 1.1.2.2 --comms-overlap --shm 2048 --shm-mpi 0"
|
||||
srun -n4 --label -c$OMP_NUM_THREADS --gpus-per-task=1 ./mpiwrapper.sh ./benchmarks/Benchmark_dwf_fp32 $PARAMS
|
||||
|
26
systems/Spock/dwf8.slurm
Normal file
26
systems/Spock/dwf8.slurm
Normal file
@ -0,0 +1,26 @@
|
||||
#!/bin/bash
|
||||
# Begin LSF Directives
|
||||
#SBATCH -A LGT104
|
||||
#SBATCH -t 01:00:00
|
||||
##SBATCH -U openmpThu
|
||||
#SBATCH -p ecp
|
||||
#SBATCH -J DWF
|
||||
#SBATCH -o DWF.%J
|
||||
#SBATCH -e DWF.%J
|
||||
#SBATCH -N 2
|
||||
#SBATCH -n 8
|
||||
|
||||
DIR=.
|
||||
module list
|
||||
export MPIR_CVAR_GPU_EAGER_DEVICE_MEM=0
|
||||
export MPICH_GPU_SUPPORT_ENABLED=1
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=XPMEM
|
||||
export MPICH_SMP_SINGLE_COPY_MODE=NONE
|
||||
#export MPICH_SMP_SINGLE_COPY_MODE=CMA
|
||||
export OMP_NUM_THREADS=8
|
||||
|
||||
AT=8
|
||||
echo MPICH_SMP_SINGLE_COPY_MODE $MPICH_SMP_SINGLE_COPY_MODE
|
||||
PARAMS=" --accelerator-threads ${AT} --grid 32.64.64.64 --mpi 1.2.2.2 --comms-overlap --shm 2048 --shm-mpi 0"
|
||||
srun -n8 --label -c$OMP_NUM_THREADS --gpus-per-task=1 ./mpiwrapper.sh ./benchmarks/Benchmark_dwf_fp32 $PARAMS
|
||||
|
12
systems/Spock/mpiwrapper.sh
Executable file
12
systems/Spock/mpiwrapper.sh
Executable file
@ -0,0 +1,12 @@
|
||||
#!/bin/bash
|
||||
|
||||
lrank=$SLURM_LOCALID
|
||||
|
||||
export ROCR_VISIBLE_DEVICES=$SLURM_LOCALID
|
||||
|
||||
echo "`hostname` - $lrank device=$ROCR_VISIBLE_DEVICES binding=$BINDING"
|
||||
|
||||
$*
|
||||
|
||||
|
||||
|
5
systems/Spock/sourceme.sh
Normal file
5
systems/Spock/sourceme.sh
Normal file
@ -0,0 +1,5 @@
|
||||
module load PrgEnv-gnu
|
||||
module load rocm/4.3.0
|
||||
module load gmp
|
||||
module load cray-fftw
|
||||
module load craype-accel-amd-gfx908
|
179
systems/Summit/comms.4node
Normal file
179
systems/Summit/comms.4node
Normal file
@ -0,0 +1,179 @@
|
||||
OPENMPI detected
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device Number : 0
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device identifier: Tesla V100-SXM2-16GB
|
||||
AcceleratorCudaInit[0]: totalGlobalMem: 16911433728
|
||||
AcceleratorCudaInit[0]: managedMemory: 1
|
||||
AcceleratorCudaInit[0]: isMultiGpuBoard: 0
|
||||
AcceleratorCudaInit[0]: warpSize: 32
|
||||
AcceleratorCudaInit[0]: pciBusID: 4
|
||||
AcceleratorCudaInit[0]: pciDeviceID: 0
|
||||
AcceleratorCudaInit[0]: maxGridSize (2147483647,65535,65535)
|
||||
AcceleratorCudaInit: rank 0 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 0 device 0 bus id: 0004:04:00.0
|
||||
AcceleratorCudaInit: ================================================
|
||||
SharedMemoryMpi: World communicator of size 24
|
||||
SharedMemoryMpi: Node communicator of size 6
|
||||
0SharedMemoryMpi: SharedMemoryMPI.cc acceleratorAllocDevice 1073741824bytes at 0x200060000000 for comms buffers
|
||||
Setting up IPC
|
||||
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|_ | | | | | | | | | | | | _|__
|
||||
__|_ _|__
|
||||
__|_ GGGG RRRR III DDDD _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G GG RRRR I D D _|__
|
||||
__|_ G G R R I D D _|__
|
||||
__|_ GGGG R R III DDDD _|__
|
||||
__|_ _|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
| | | | | | | | | | | | | |
|
||||
|
||||
|
||||
Copyright (C) 2015 Peter Boyle, Azusa Yamaguchi, Guido Cossu, Antonin Portelli and other authors
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
Current Grid git commit hash=7cb1ff7395a5833ded6526c43891bd07a0436290: (HEAD -> develop, origin/develop, origin/HEAD) clean
|
||||
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : MPI is initialised and logging filters activated
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : Requested 1073741824 byte stencil comms buffers
|
||||
AcceleratorCudaInit: rank 1 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 1 device 1 bus id: 0004:05:00.0
|
||||
AcceleratorCudaInit: rank 2 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 2 device 2 bus id: 0004:06:00.0
|
||||
AcceleratorCudaInit: rank 5 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 5 device 5 bus id: 0035:05:00.0
|
||||
AcceleratorCudaInit: rank 4 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 4 device 4 bus id: 0035:04:00.0
|
||||
AcceleratorCudaInit: rank 3 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 3 device 3 bus id: 0035:03:00.0
|
||||
Grid : Message : MemoryManager Cache 13529146982 bytes
|
||||
Grid : Message : MemoryManager::Init() setting up
|
||||
Grid : Message : MemoryManager::Init() cache pool for recent allocations: SMALL 8 LARGE 2
|
||||
Grid : Message : MemoryManager::Init() Non unified: Caching accelerator data in dedicated memory
|
||||
Grid : Message : MemoryManager::Init() Using cudaMalloc
|
||||
Grid : Message : 2.137929 s : Grid is setup to use 6 threads
|
||||
Grid : Message : 2.137941 s : Number of iterations to average: 250
|
||||
Grid : Message : 2.137950 s : ====================================================================================================
|
||||
Grid : Message : 2.137958 s : = Benchmarking sequential halo exchange from host memory
|
||||
Grid : Message : 2.137966 s : ====================================================================================================
|
||||
Grid : Message : 2.137974 s : L Ls bytes MB/s uni MB/s bidi
|
||||
AcceleratorCudaInit: rank 22 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 10 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 15 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 21 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 20 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 7 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 9 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 11 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 8 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 6 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 19 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 23 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 18 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 12 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 16 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 13 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 14 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 17 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
Grid : Message : 2.604949 s : 8 8 393216 89973.9 179947.8
|
||||
Grid : Message : 2.668249 s : 8 8 393216 18650.3 37300.5
|
||||
Grid : Message : 2.732288 s : 8 8 393216 18428.5 36857.1
|
||||
Grid : Message : 2.753565 s : 8 8 393216 55497.2 110994.4
|
||||
Grid : Message : 2.808960 s : 12 8 1327104 100181.5 200363.0
|
||||
Grid : Message : 3.226900 s : 12 8 1327104 20600.5 41201.0
|
||||
Grid : Message : 3.167459 s : 12 8 1327104 24104.6 48209.2
|
||||
Grid : Message : 3.227660 s : 12 8 1327104 66156.7 132313.5
|
||||
Grid : Message : 3.413570 s : 16 8 3145728 56174.4 112348.8
|
||||
Grid : Message : 3.802697 s : 16 8 3145728 24255.9 48511.7
|
||||
Grid : Message : 4.190498 s : 16 8 3145728 24336.7 48673.4
|
||||
Grid : Message : 4.385171 s : 16 8 3145728 48484.1 96968.2
|
||||
Grid : Message : 4.805284 s : 20 8 6144000 46380.5 92761.1
|
||||
Grid : Message : 5.562975 s : 20 8 6144000 24328.5 48656.9
|
||||
Grid : Message : 6.322562 s : 20 8 6144000 24266.7 48533.4
|
||||
Grid : Message : 6.773598 s : 20 8 6144000 40868.5 81736.9
|
||||
Grid : Message : 7.600999 s : 24 8 10616832 40198.3 80396.6
|
||||
Grid : Message : 8.912917 s : 24 8 10616832 24279.5 48559.1
|
||||
Grid : Message : 10.220961 s : 24 8 10616832 24350.2 48700.4
|
||||
Grid : Message : 11.728250 s : 24 8 10616832 37390.9 74781.8
|
||||
Grid : Message : 12.497258 s : 28 8 16859136 36792.2 73584.5
|
||||
Grid : Message : 14.585387 s : 28 8 16859136 24222.2 48444.3
|
||||
Grid : Message : 16.664783 s : 28 8 16859136 24323.4 48646.8
|
||||
Grid : Message : 17.955238 s : 28 8 16859136 39194.7 78389.4
|
||||
Grid : Message : 20.136479 s : 32 8 25165824 35718.3 71436.5
|
||||
Grid : Message : 23.241958 s : 32 8 25165824 24311.4 48622.9
|
||||
Grid : Message : 26.344810 s : 32 8 25165824 24331.9 48663.7
|
||||
Grid : Message : 28.384420 s : 32 8 25165824 37016.3 74032.7
|
||||
Grid : Message : 28.388879 s : ====================================================================================================
|
||||
Grid : Message : 28.388894 s : = Benchmarking sequential halo exchange from GPU memory
|
||||
Grid : Message : 28.388909 s : ====================================================================================================
|
||||
Grid : Message : 28.388924 s : L Ls bytes MB/s uni MB/s bidi
|
||||
Grid : Message : 28.553993 s : 8 8 393216 8272.4 16544.7
|
||||
Grid : Message : 28.679592 s : 8 8 393216 9395.4 18790.8
|
||||
Grid : Message : 28.811112 s : 8 8 393216 8971.0 17942.0
|
||||
Grid : Message : 28.843770 s : 8 8 393216 36145.6 72291.2
|
||||
Grid : Message : 28.981754 s : 12 8 1327104 49591.6 99183.2
|
||||
Grid : Message : 29.299764 s : 12 8 1327104 12520.8 25041.7
|
||||
Grid : Message : 29.620288 s : 12 8 1327104 12422.2 24844.4
|
||||
Grid : Message : 29.657645 s : 12 8 1327104 106637.5 213275.1
|
||||
Grid : Message : 29.952933 s : 16 8 3145728 43939.2 87878.5
|
||||
Grid : Message : 30.585411 s : 16 8 3145728 14922.1 29844.2
|
||||
Grid : Message : 31.219781 s : 16 8 3145728 14877.2 29754.4
|
||||
Grid : Message : 31.285017 s : 16 8 3145728 144724.3 289448.7
|
||||
Grid : Message : 31.706443 s : 20 8 6144000 54676.2 109352.4
|
||||
Grid : Message : 32.739205 s : 20 8 6144000 17848.0 35696.1
|
||||
Grid : Message : 33.771852 s : 20 8 6144000 17849.9 35699.7
|
||||
Grid : Message : 33.871981 s : 20 8 6144000 184141.4 368282.8
|
||||
Grid : Message : 34.536808 s : 24 8 10616832 55784.3 111568.6
|
||||
Grid : Message : 36.275648 s : 24 8 10616832 18317.6 36635.3
|
||||
Grid : Message : 37.997181 s : 24 8 10616832 18501.7 37003.4
|
||||
Grid : Message : 38.140442 s : 24 8 10616832 222383.9 444767.9
|
||||
Grid : Message : 39.177222 s : 28 8 16859136 56609.7 113219.4
|
||||
Grid : Message : 41.874755 s : 28 8 16859136 18749.9 37499.8
|
||||
Grid : Message : 44.529381 s : 28 8 16859136 19052.9 38105.8
|
||||
Grid : Message : 44.742192 s : 28 8 16859136 237717.1 475434.2
|
||||
Grid : Message : 46.184000 s : 32 8 25165824 57091.2 114182.4
|
||||
Grid : Message : 50.734740 s : 32 8 25165824 19411.0 38821.9
|
||||
Grid : Message : 53.931228 s : 32 8 25165824 19570.6 39141.2
|
||||
Grid : Message : 54.238467 s : 32 8 25165824 245765.6 491531.2
|
||||
Grid : Message : 54.268664 s : ====================================================================================================
|
||||
Grid : Message : 54.268680 s : = All done; Bye Bye
|
||||
Grid : Message : 54.268691 s : ====================================================================================================
|
14
systems/Summit/config-command
Normal file
14
systems/Summit/config-command
Normal file
@ -0,0 +1,14 @@
|
||||
../../configure --enable-comms=mpi \
|
||||
--enable-simd=GPU \
|
||||
--enable-gen-simd-width=32 \
|
||||
--enable-unified=no \
|
||||
--enable-shm=nvlink \
|
||||
--disable-gparity \
|
||||
--enable-setdevice \
|
||||
--disable-fermion-reps \
|
||||
--enable-accelerator=cuda \
|
||||
--prefix /ccs/home/paboyle/prefix \
|
||||
CXX=nvcc \
|
||||
LDFLAGS=-L/ccs/home/paboyle/prefix/lib/ \
|
||||
CXXFLAGS="-ccbin mpicxx -gencode arch=compute_70,code=sm_70 -I/ccs/home/paboyle/prefix/include/ -std=c++14"
|
||||
|
206
systems/Summit/dwf.24.4node
Normal file
206
systems/Summit/dwf.24.4node
Normal file
@ -0,0 +1,206 @@
|
||||
OPENMPI detected
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device Number : 0
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device identifier: Tesla V100-SXM2-16GB
|
||||
AcceleratorCudaInit[0]: totalGlobalMem: 16911433728
|
||||
AcceleratorCudaInit[0]: managedMemory: 1
|
||||
AcceleratorCudaInit[0]: isMultiGpuBoard: 0
|
||||
AcceleratorCudaInit[0]: warpSize: 32
|
||||
AcceleratorCudaInit[0]: pciBusID: 4
|
||||
AcceleratorCudaInit[0]: pciDeviceID: 0
|
||||
AcceleratorCudaInit[0]: maxGridSize (2147483647,65535,65535)
|
||||
AcceleratorCudaInit: rank 0 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 0 device 0 bus id: 0004:04:00.0
|
||||
AcceleratorCudaInit: ================================================
|
||||
SharedMemoryMpi: World communicator of size 24
|
||||
SharedMemoryMpi: Node communicator of size 6
|
||||
0SharedMemoryMpi: SharedMemoryMPI.cc acceleratorAllocDevice 2147483648bytes at 0x200080000000 for comms buffers
|
||||
AcceleratorCudaInit: rank 3 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 3 device 3 bus id: 0035:03:00.0
|
||||
AcceleratorCudaInit: rank 5 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 5 device 5 bus id: 0035:05:00.0
|
||||
Setting up IPC
|
||||
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|_ | | | | | | | | | | | | _|__
|
||||
__|_ _|__
|
||||
__|_ GGGG RRRR III DDDD _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G GG RRRR I D D _|__
|
||||
__|_ G G R R I D D _|__
|
||||
__|_ GGGG R R III DDDD _|__
|
||||
__|_ _|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
| | | | | | | | | | | | | |
|
||||
|
||||
|
||||
Copyright (C) 2015 Peter Boyle, Azusa Yamaguchi, Guido Cossu, Antonin Portelli and other authors
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
AcceleratorCudaInit: rank 4 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 4 device 4 bus id: 0035:04:00.0
|
||||
AcceleratorCudaInit: rank 1 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 1 device 1 bus id: 0004:05:00.0
|
||||
AcceleratorCudaInit: rank 2 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 2 device 2 bus id: 0004:06:00.0
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
Current Grid git commit hash=7cb1ff7395a5833ded6526c43891bd07a0436290: (HEAD -> develop, origin/develop, origin/HEAD) clean
|
||||
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : MPI is initialised and logging filters activated
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : Requested 2147483648 byte stencil comms buffers
|
||||
Grid : Message : MemoryManager Cache 8388608000 bytes
|
||||
Grid : Message : MemoryManager::Init() setting up
|
||||
Grid : Message : MemoryManager::Init() cache pool for recent allocations: SMALL 8 LARGE 2
|
||||
Grid : Message : MemoryManager::Init() Non unified: Caching accelerator data in dedicated memory
|
||||
Grid : Message : MemoryManager::Init() Using cudaMalloc
|
||||
Grid : Message : 1.731905 s : Grid Layout
|
||||
Grid : Message : 1.731915 s : Global lattice size : 48 48 48 72
|
||||
Grid : Message : 1.731928 s : OpenMP threads : 6
|
||||
Grid : Message : 1.731938 s : MPI tasks : 2 2 2 3
|
||||
AcceleratorCudaInit: rank 9 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 23 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 22 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 21 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 18 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 6 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 7 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 10 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 8 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 11 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 20 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 19 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 13 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 12 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 14 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 16 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 15 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 17 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
Grid : Message : 2.683494 s : Making s innermost grids
|
||||
Grid : Message : 2.780034 s : Initialising 4d RNG
|
||||
Grid : Message : 2.833099 s : Intialising parallel RNG with unique string 'The 4D RNG'
|
||||
Grid : Message : 2.833121 s : Seed SHA256: 49db4542db694e3b1a74bf2592a8c1b83bfebbe18401693c2609a4c3af1
|
||||
Grid : Message : 2.916841 s : Initialising 5d RNG
|
||||
Grid : Message : 3.762880 s : Intialising parallel RNG with unique string 'The 5D RNG'
|
||||
Grid : Message : 3.762902 s : Seed SHA256: b6316f2fac44ce14111f93e0296389330b077bfd0a7b359f781c58589f8a
|
||||
Grid : Message : 5.264345 s : Initialised RNGs
|
||||
Grid : Message : 6.489904 s : Drawing gauge field
|
||||
Grid : Message : 6.729262 s : Random gauge initialised
|
||||
Grid : Message : 7.781273 s : Setting up Cshift based reference
|
||||
Grid : Message : 8.725313 s : *****************************************************************
|
||||
Grid : Message : 8.725332 s : * Kernel options --dslash-generic, --dslash-unroll, --dslash-asm
|
||||
Grid : Message : 8.725342 s : *****************************************************************
|
||||
Grid : Message : 8.725352 s : *****************************************************************
|
||||
Grid : Message : 8.725362 s : * Benchmarking DomainWallFermionR::Dhop
|
||||
Grid : Message : 8.725372 s : * Vectorising space-time by 4
|
||||
Grid : Message : 8.725383 s : * VComplexF size is 32 B
|
||||
Grid : Message : 8.725395 s : * SINGLE precision
|
||||
Grid : Message : 8.725405 s : * Using Overlapped Comms/Compute
|
||||
Grid : Message : 8.725415 s : * Using GENERIC Nc WilsonKernels
|
||||
Grid : Message : 8.725425 s : *****************************************************************
|
||||
Grid : Message : 9.465229 s : Called warmup
|
||||
Grid : Message : 58.646066 s : Called Dw 3000 times in 4.91764e+07 us
|
||||
Grid : Message : 58.646121 s : mflop/s = 1.02592e+07
|
||||
Grid : Message : 58.646134 s : mflop/s per rank = 427468
|
||||
Grid : Message : 58.646145 s : mflop/s per node = 2.56481e+06
|
||||
Grid : Message : 58.646156 s : RF GiB/s (base 2) = 20846.5
|
||||
Grid : Message : 58.646166 s : mem GiB/s (base 2) = 13029.1
|
||||
Grid : Message : 58.648008 s : norm diff 1.04778e-13
|
||||
Grid : Message : 58.734885 s : #### Dhop calls report
|
||||
Grid : Message : 58.734897 s : WilsonFermion5D Number of DhopEO Calls : 6002
|
||||
Grid : Message : 58.734909 s : WilsonFermion5D TotalTime /Calls : 8217.71 us
|
||||
Grid : Message : 58.734922 s : WilsonFermion5D CommTime /Calls : 7109.5 us
|
||||
Grid : Message : 58.734933 s : WilsonFermion5D FaceTime /Calls : 446.623 us
|
||||
Grid : Message : 58.734943 s : WilsonFermion5D ComputeTime1/Calls : 18.0558 us
|
||||
Grid : Message : 58.734953 s : WilsonFermion5D ComputeTime2/Calls : 731.097 us
|
||||
Grid : Message : 58.734979 s : Average mflops/s per call : 4.8157e+09
|
||||
Grid : Message : 58.734989 s : Average mflops/s per call per rank : 2.00654e+08
|
||||
Grid : Message : 58.734999 s : Average mflops/s per call per node : 1.20393e+09
|
||||
Grid : Message : 58.735008 s : Average mflops/s per call (full) : 1.04183e+07
|
||||
Grid : Message : 58.735017 s : Average mflops/s per call per rank (full): 434094
|
||||
Grid : Message : 58.735026 s : Average mflops/s per call per node (full): 2.60456e+06
|
||||
Grid : Message : 58.735035 s : WilsonFermion5D Stencil
|
||||
Grid : Message : 58.735043 s : WilsonFermion5D StencilEven
|
||||
Grid : Message : 58.735051 s : WilsonFermion5D StencilOdd
|
||||
Grid : Message : 58.735059 s : WilsonFermion5D Stencil Reporti()
|
||||
Grid : Message : 58.735067 s : WilsonFermion5D StencilEven Reporti()
|
||||
Grid : Message : 58.735075 s : WilsonFermion5D StencilOdd Reporti()
|
||||
Grid : Message : 64.934380 s : Compare to naive wilson implementation Dag to verify correctness
|
||||
Grid : Message : 64.934740 s : Called DwDag
|
||||
Grid : Message : 64.934870 s : norm dag result 12.0422
|
||||
Grid : Message : 64.120756 s : norm dag ref 12.0422
|
||||
Grid : Message : 64.149389 s : norm dag diff 7.6644e-14
|
||||
Grid : Message : 64.317786 s : Calling Deo and Doe and //assert Deo+Doe == Dunprec
|
||||
Grid : Message : 64.465331 s : src_e0.499995
|
||||
Grid : Message : 64.524653 s : src_o0.500005
|
||||
Grid : Message : 64.558706 s : *********************************************************
|
||||
Grid : Message : 64.558717 s : * Benchmarking DomainWallFermionF::DhopEO
|
||||
Grid : Message : 64.558727 s : * Vectorising space-time by 4
|
||||
Grid : Message : 64.558737 s : * SINGLE precision
|
||||
Grid : Message : 64.558745 s : * Using Overlapped Comms/Compute
|
||||
Grid : Message : 64.558753 s : * Using GENERIC Nc WilsonKernels
|
||||
Grid : Message : 64.558761 s : *********************************************************
|
||||
Grid : Message : 92.702145 s : Deo mflop/s = 8.97692e+06
|
||||
Grid : Message : 92.702185 s : Deo mflop/s per rank 374038
|
||||
Grid : Message : 92.702198 s : Deo mflop/s per node 2.24423e+06
|
||||
Grid : Message : 92.702209 s : #### Dhop calls report
|
||||
Grid : Message : 92.702223 s : WilsonFermion5D Number of DhopEO Calls : 3001
|
||||
Grid : Message : 92.702240 s : WilsonFermion5D TotalTime /Calls : 9377.88 us
|
||||
Grid : Message : 92.702257 s : WilsonFermion5D CommTime /Calls : 8221.84 us
|
||||
Grid : Message : 92.702277 s : WilsonFermion5D FaceTime /Calls : 543.548 us
|
||||
Grid : Message : 92.702301 s : WilsonFermion5D ComputeTime1/Calls : 20.936 us
|
||||
Grid : Message : 92.702322 s : WilsonFermion5D ComputeTime2/Calls : 732.33 us
|
||||
Grid : Message : 92.702376 s : Average mflops/s per call : 4.13001e+09
|
||||
Grid : Message : 92.702387 s : Average mflops/s per call per rank : 1.72084e+08
|
||||
Grid : Message : 92.702397 s : Average mflops/s per call per node : 1.0325e+09
|
||||
Grid : Message : 92.702407 s : Average mflops/s per call (full) : 9.12937e+06
|
||||
Grid : Message : 92.702416 s : Average mflops/s per call per rank (full): 380391
|
||||
Grid : Message : 92.702426 s : Average mflops/s per call per node (full): 2.28234e+06
|
||||
Grid : Message : 92.702435 s : WilsonFermion5D Stencil
|
||||
Grid : Message : 92.702443 s : WilsonFermion5D StencilEven
|
||||
Grid : Message : 92.702451 s : WilsonFermion5D StencilOdd
|
||||
Grid : Message : 92.702459 s : WilsonFermion5D Stencil Reporti()
|
||||
Grid : Message : 92.702467 s : WilsonFermion5D StencilEven Reporti()
|
||||
Grid : Message : 92.702475 s : WilsonFermion5D StencilOdd Reporti()
|
||||
Grid : Message : 92.772983 s : r_e6.02121
|
||||
Grid : Message : 92.786384 s : r_o6.02102
|
||||
Grid : Message : 92.799622 s : res12.0422
|
||||
Grid : Message : 93.860500 s : norm diff 0
|
||||
Grid : Message : 93.162026 s : norm diff even 0
|
||||
Grid : Message : 93.197529 s : norm diff odd 0
|
206
systems/Summit/dwf.32.4node
Normal file
206
systems/Summit/dwf.32.4node
Normal file
@ -0,0 +1,206 @@
|
||||
OPENMPI detected
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device Number : 0
|
||||
AcceleratorCudaInit[0]: ========================
|
||||
AcceleratorCudaInit[0]: Device identifier: Tesla V100-SXM2-16GB
|
||||
AcceleratorCudaInit[0]: totalGlobalMem: 16911433728
|
||||
AcceleratorCudaInit[0]: managedMemory: 1
|
||||
AcceleratorCudaInit[0]: isMultiGpuBoard: 0
|
||||
AcceleratorCudaInit[0]: warpSize: 32
|
||||
AcceleratorCudaInit[0]: pciBusID: 4
|
||||
AcceleratorCudaInit[0]: pciDeviceID: 0
|
||||
AcceleratorCudaInit[0]: maxGridSize (2147483647,65535,65535)
|
||||
AcceleratorCudaInit: rank 0 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 0 device 0 bus id: 0004:04:00.0
|
||||
AcceleratorCudaInit: ================================================
|
||||
SharedMemoryMpi: World communicator of size 24
|
||||
SharedMemoryMpi: Node communicator of size 6
|
||||
0SharedMemoryMpi: SharedMemoryMPI.cc acceleratorAllocDevice 2147483648bytes at 0x200080000000 for comms buffers
|
||||
Setting up IPC
|
||||
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|_ | | | | | | | | | | | | _|__
|
||||
__|_ _|__
|
||||
__|_ GGGG RRRR III DDDD _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G R R I D D _|__
|
||||
__|_ G GG RRRR I D D _|__
|
||||
__|_ G G R R I D D _|__
|
||||
__|_ GGGG R R III DDDD _|__
|
||||
__|_ _|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
__|__|__|__|__|__|__|__|__|__|__|__|__|__|__
|
||||
| | | | | | | | | | | | | |
|
||||
|
||||
|
||||
Copyright (C) 2015 Peter Boyle, Azusa Yamaguchi, Guido Cossu, Antonin Portelli and other authors
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
AcceleratorCudaInit: rank 2 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 2 device 2 bus id: 0004:06:00.0
|
||||
AcceleratorCudaInit: rank 1 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 1 device 1 bus id: 0004:05:00.0
|
||||
AcceleratorCudaInit: rank 4 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 4 device 4 bus id: 0035:04:00.0
|
||||
AcceleratorCudaInit: rank 3 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 3 device 3 bus id: 0035:03:00.0
|
||||
AcceleratorCudaInit: rank 5 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
local rank 5 device 5 bus id: 0035:05:00.0
|
||||
GNU General Public License for more details.
|
||||
Current Grid git commit hash=7cb1ff7395a5833ded6526c43891bd07a0436290: (HEAD -> develop, origin/develop, origin/HEAD) clean
|
||||
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : MPI is initialised and logging filters activated
|
||||
Grid : Message : ================================================
|
||||
Grid : Message : Requested 2147483648 byte stencil comms buffers
|
||||
Grid : Message : MemoryManager Cache 8388608000 bytes
|
||||
Grid : Message : MemoryManager::Init() setting up
|
||||
Grid : Message : MemoryManager::Init() cache pool for recent allocations: SMALL 8 LARGE 2
|
||||
Grid : Message : MemoryManager::Init() Non unified: Caching accelerator data in dedicated memory
|
||||
Grid : Message : MemoryManager::Init() Using cudaMalloc
|
||||
Grid : Message : 1.544984 s : Grid Layout
|
||||
Grid : Message : 1.544992 s : Global lattice size : 64 64 64 96
|
||||
Grid : Message : 1.545003 s : OpenMP threads : 6
|
||||
Grid : Message : 1.545011 s : MPI tasks : 2 2 2 3
|
||||
AcceleratorCudaInit: rank 8 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 6 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 11 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 16 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 17 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 13 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 12 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 21 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 23 setting device to node rank 5
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 22 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 19 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 18 setting device to node rank 0
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 7 setting device to node rank 1
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 10 setting device to node rank 4
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 9 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 14 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 15 setting device to node rank 3
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
AcceleratorCudaInit: rank 20 setting device to node rank 2
|
||||
AcceleratorCudaInit: Configure options --enable-setdevice=yes
|
||||
Grid : Message : 2.994920 s : Making s innermost grids
|
||||
Grid : Message : 2.232502 s : Initialising 4d RNG
|
||||
Grid : Message : 2.397047 s : Intialising parallel RNG with unique string 'The 4D RNG'
|
||||
Grid : Message : 2.397069 s : Seed SHA256: 49db4542db694e3b1a74bf2592a8c1b83bfebbe18401693c2609a4c3af1
|
||||
Grid : Message : 2.653140 s : Initialising 5d RNG
|
||||
Grid : Message : 5.285347 s : Intialising parallel RNG with unique string 'The 5D RNG'
|
||||
Grid : Message : 5.285369 s : Seed SHA256: b6316f2fac44ce14111f93e0296389330b077bfd0a7b359f781c58589f8a
|
||||
Grid : Message : 9.994738 s : Initialised RNGs
|
||||
Grid : Message : 13.153426 s : Drawing gauge field
|
||||
Grid : Message : 13.825697 s : Random gauge initialised
|
||||
Grid : Message : 18.537657 s : Setting up Cshift based reference
|
||||
Grid : Message : 22.296755 s : *****************************************************************
|
||||
Grid : Message : 22.296781 s : * Kernel options --dslash-generic, --dslash-unroll, --dslash-asm
|
||||
Grid : Message : 22.296791 s : *****************************************************************
|
||||
Grid : Message : 22.296800 s : *****************************************************************
|
||||
Grid : Message : 22.296809 s : * Benchmarking DomainWallFermionR::Dhop
|
||||
Grid : Message : 22.296818 s : * Vectorising space-time by 4
|
||||
Grid : Message : 22.296828 s : * VComplexF size is 32 B
|
||||
Grid : Message : 22.296838 s : * SINGLE precision
|
||||
Grid : Message : 22.296847 s : * Using Overlapped Comms/Compute
|
||||
Grid : Message : 22.296855 s : * Using GENERIC Nc WilsonKernels
|
||||
Grid : Message : 22.296863 s : *****************************************************************
|
||||
Grid : Message : 24.746452 s : Called warmup
|
||||
Grid : Message : 137.525756 s : Called Dw 3000 times in 1.12779e+08 us
|
||||
Grid : Message : 137.525818 s : mflop/s = 1.41383e+07
|
||||
Grid : Message : 137.525831 s : mflop/s per rank = 589097
|
||||
Grid : Message : 137.525843 s : mflop/s per node = 3.53458e+06
|
||||
Grid : Message : 137.525854 s : RF GiB/s (base 2) = 28728.7
|
||||
Grid : Message : 137.525864 s : mem GiB/s (base 2) = 17955.5
|
||||
Grid : Message : 137.693645 s : norm diff 1.04885e-13
|
||||
Grid : Message : 137.965585 s : #### Dhop calls report
|
||||
Grid : Message : 137.965598 s : WilsonFermion5D Number of DhopEO Calls : 6002
|
||||
Grid : Message : 137.965612 s : WilsonFermion5D TotalTime /Calls : 18899.7 us
|
||||
Grid : Message : 137.965624 s : WilsonFermion5D CommTime /Calls : 16041.4 us
|
||||
Grid : Message : 137.965634 s : WilsonFermion5D FaceTime /Calls : 859.705 us
|
||||
Grid : Message : 137.965644 s : WilsonFermion5D ComputeTime1/Calls : 70.5881 us
|
||||
Grid : Message : 137.965654 s : WilsonFermion5D ComputeTime2/Calls : 2094.8 us
|
||||
Grid : Message : 137.965682 s : Average mflops/s per call : 3.87638e+09
|
||||
Grid : Message : 137.965692 s : Average mflops/s per call per rank : 1.61516e+08
|
||||
Grid : Message : 137.965702 s : Average mflops/s per call per node : 9.69095e+08
|
||||
Grid : Message : 137.965712 s : Average mflops/s per call (full) : 1.43168e+07
|
||||
Grid : Message : 137.965721 s : Average mflops/s per call per rank (full): 596533
|
||||
Grid : Message : 137.965730 s : Average mflops/s per call per node (full): 3.5792e+06
|
||||
Grid : Message : 137.965740 s : WilsonFermion5D Stencil
|
||||
Grid : Message : 137.965748 s : WilsonFermion5D StencilEven
|
||||
Grid : Message : 137.965756 s : WilsonFermion5D StencilOdd
|
||||
Grid : Message : 137.965764 s : WilsonFermion5D Stencil Reporti()
|
||||
Grid : Message : 137.965772 s : WilsonFermion5D StencilEven Reporti()
|
||||
Grid : Message : 137.965780 s : WilsonFermion5D StencilOdd Reporti()
|
||||
Grid : Message : 156.554605 s : Compare to naive wilson implementation Dag to verify correctness
|
||||
Grid : Message : 156.554632 s : Called DwDag
|
||||
Grid : Message : 156.554642 s : norm dag result 12.0421
|
||||
Grid : Message : 156.639265 s : norm dag ref 12.0421
|
||||
Grid : Message : 156.888281 s : norm dag diff 7.62057e-14
|
||||
Grid : Message : 157.609797 s : Calling Deo and Doe and //assert Deo+Doe == Dunprec
|
||||
Grid : Message : 158.208630 s : src_e0.499996
|
||||
Grid : Message : 158.162447 s : src_o0.500004
|
||||
Grid : Message : 158.267780 s : *********************************************************
|
||||
Grid : Message : 158.267791 s : * Benchmarking DomainWallFermionF::DhopEO
|
||||
Grid : Message : 158.267801 s : * Vectorising space-time by 4
|
||||
Grid : Message : 158.267811 s : * SINGLE precision
|
||||
Grid : Message : 158.267820 s : * Using Overlapped Comms/Compute
|
||||
Grid : Message : 158.267828 s : * Using GENERIC Nc WilsonKernels
|
||||
Grid : Message : 158.267836 s : *********************************************************
|
||||
Grid : Message : 216.487829 s : Deo mflop/s = 1.37283e+07
|
||||
Grid : Message : 216.487869 s : Deo mflop/s per rank 572011
|
||||
Grid : Message : 216.487881 s : Deo mflop/s per node 3.43206e+06
|
||||
Grid : Message : 216.487893 s : #### Dhop calls report
|
||||
Grid : Message : 216.487903 s : WilsonFermion5D Number of DhopEO Calls : 3001
|
||||
Grid : Message : 216.487913 s : WilsonFermion5D TotalTime /Calls : 19399.6 us
|
||||
Grid : Message : 216.487923 s : WilsonFermion5D CommTime /Calls : 16475.4 us
|
||||
Grid : Message : 216.487933 s : WilsonFermion5D FaceTime /Calls : 972.393 us
|
||||
Grid : Message : 216.487943 s : WilsonFermion5D ComputeTime1/Calls : 49.8474 us
|
||||
Grid : Message : 216.487953 s : WilsonFermion5D ComputeTime2/Calls : 2089.93 us
|
||||
Grid : Message : 216.488001 s : Average mflops/s per call : 5.39682e+09
|
||||
Grid : Message : 216.488011 s : Average mflops/s per call per rank : 2.24867e+08
|
||||
Grid : Message : 216.488020 s : Average mflops/s per call per node : 1.3492e+09
|
||||
Grid : Message : 216.488030 s : Average mflops/s per call (full) : 1.39479e+07
|
||||
Grid : Message : 216.488039 s : Average mflops/s per call per rank (full): 581162
|
||||
Grid : Message : 216.488048 s : Average mflops/s per call per node (full): 3.48697e+06
|
||||
Grid : Message : 216.488057 s : WilsonFermion5D Stencil
|
||||
Grid : Message : 216.488065 s : WilsonFermion5D StencilEven
|
||||
Grid : Message : 216.488073 s : WilsonFermion5D StencilOdd
|
||||
Grid : Message : 216.488081 s : WilsonFermion5D Stencil Reporti()
|
||||
Grid : Message : 216.488089 s : WilsonFermion5D StencilEven Reporti()
|
||||
Grid : Message : 216.488097 s : WilsonFermion5D StencilOdd Reporti()
|
||||
Grid : Message : 217.384495 s : r_e6.02113
|
||||
Grid : Message : 217.426121 s : r_o6.02096
|
||||
Grid : Message : 217.472636 s : res12.0421
|
||||
Grid : Message : 218.200068 s : norm diff 0
|
||||
Grid : Message : 218.645673 s : norm diff even 0
|
||||
Grid : Message : 218.816561 s : norm diff odd 0
|
25
systems/Summit/dwf16.lsf
Normal file
25
systems/Summit/dwf16.lsf
Normal file
@ -0,0 +1,25 @@
|
||||
#!/bin/bash
|
||||
#BSUB -P LGT104
|
||||
#BSUB -W 2:00
|
||||
#BSUB -nnodes 16
|
||||
#BSUB -J DWF
|
||||
|
||||
export OMP_NUM_THREADS=6
|
||||
export PAMI_IBV_ADAPTER_AFFINITY=1
|
||||
export PAMI_ENABLE_STRIPING=1
|
||||
export OPT="--comms-concurrent --comms-overlap "
|
||||
|
||||
APP="./benchmarks/Benchmark_comms_host_device --mpi 4.4.4.3 "
|
||||
jsrun --nrs 16 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > comms.16node.log
|
||||
|
||||
APP="./benchmarks/Benchmark_dwf_fp32 --grid 96.96.96.72 --mpi 4.4.4.3 --shm 2048 --shm-force-mpi 1 --device-mem 8000 --shm-force-mpi 1 $OPT "
|
||||
jsrun --nrs 16 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > dwf.16node.24.log
|
||||
|
||||
APP="./benchmarks/Benchmark_dwf_fp32 --grid 128.128.128.96 --mpi 4.4.4.3 --shm 2048 --shm-force-mpi 1 --device-mem 8000 --shm-force-mpi 1 $OPT "
|
||||
jsrun --nrs 16 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > dwf.16node.32.log
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
25
systems/Summit/dwf4.lsf
Normal file
25
systems/Summit/dwf4.lsf
Normal file
@ -0,0 +1,25 @@
|
||||
#!/bin/bash
|
||||
#BSUB -P LGT104
|
||||
#BSUB -W 2:00
|
||||
#BSUB -nnodes 4
|
||||
#BSUB -J DWF
|
||||
|
||||
export OMP_NUM_THREADS=6
|
||||
export PAMI_IBV_ADAPTER_AFFINITY=1
|
||||
export PAMI_ENABLE_STRIPING=1
|
||||
export OPT="--comms-concurrent --comms-overlap "
|
||||
#export GRID_ALLOC_NCACHE_LARGE=1
|
||||
export APP="./benchmarks/Benchmark_comms_host_device --mpi 2.2.2.3 "
|
||||
jsrun --nrs 4 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > comms.4node
|
||||
|
||||
APP="./benchmarks/Benchmark_dwf_fp32 --grid 48.48.48.72 --mpi 2.2.2.3 --shm 2048 --shm-force-mpi 1 --device-mem 8000 --shm-force-mpi 1 $OPT "
|
||||
jsrun --nrs 4 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > dwf.24.4node
|
||||
|
||||
APP="./benchmarks/Benchmark_dwf_fp32 --grid 64.64.64.96 --mpi 2.2.2.3 --shm 2048 --shm-force-mpi 1 --device-mem 8000 --shm-force-mpi 1 $OPT "
|
||||
jsrun --nrs 4 -a6 -g6 -c42 -dpacked -b packed:7 --latency_priority gpu-cpu --smpiargs=-gpu $APP > dwf.32.4node
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
8
systems/Summit/sourceme-cuda10.sh
Normal file
8
systems/Summit/sourceme-cuda10.sh
Normal file
@ -0,0 +1,8 @@
|
||||
export UCX_GDR_COPY_RCACHE=no
|
||||
export UCX_MEMTYPE_CACHE=n
|
||||
export UCX_RNDV_SCHEME=put_zcopy
|
||||
module load gcc/7.5.0
|
||||
module load cuda/10.2.89
|
||||
#cuda/11.4.0
|
||||
export LD_LIBRARY_PATH=/ccs/home/paboyle/prefix/lib/:$LD_LIBRARY_PATH
|
||||
|
@ -5,7 +5,7 @@
|
||||
--enable-gen-simd-width=64 \
|
||||
--enable-accelerator=cuda \
|
||||
--with-lime=/mnt/lustre/tursafs1/home/tc002/tc002/dc-boyl1/spack/spack/opt/spack/linux-rhel8-zen/gcc-8.4.1/c-lime-2-3-9-e6wxqrid6rqmd45z7n32dxkvkykpvyez \
|
||||
--disable-accelerator-cshift \
|
||||
--enable-accelerator-cshift \
|
||||
--disable-unified \
|
||||
CXX=nvcc \
|
||||
LDFLAGS="-cudart shared " \
|
||||
|
@ -1,2 +1,6 @@
|
||||
spack load c-lime
|
||||
module load cuda/11.4.1 openmpi/4.1.1 ucx/1.10.1
|
||||
module load cuda/11.4.1 openmpi/4.1.1-cuda11.4.1 ucx/1.12.0-cuda11.4.1
|
||||
#module load cuda/11.4.1 openmpi/4.1.1 ucx/1.10.1
|
||||
export PREFIX=/home/tc002/tc002/shared/env/prefix/
|
||||
export LD_LIBRARY_PATH=$PREFIX/lib/:$LD_LIBRARY_PATH
|
||||
unset SBATCH_EXPORT
|
||||
|
||||
|
@ -235,7 +235,6 @@ void TestWhat(What & Ddwf,
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<What,LatticeFermion> HermOpEO(Ddwf);
|
||||
HermOpEO.MpcDagMpc(chi_e,dchi_e);
|
||||
|
@ -215,7 +215,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd , chi_o, chi);
|
||||
pickCheckerboard(Even, phi_e, phi);
|
||||
pickCheckerboard(Odd , phi_o, phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<DomainWallEOFAFermionR,LatticeFermion> HermOpEO(Ddwf);
|
||||
HermOpEO.MpcDagMpc(chi_e, dchi_e);
|
||||
|
@ -212,8 +212,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
|
||||
SchurDiagMooeeOperator<DomainWallFermionR,LatticeFermion> HermOpEO(Ddwf);
|
||||
HermOpEO.MpcDagMpc(chi_e,dchi_e);
|
||||
|
@ -181,8 +181,8 @@ void checkAdj(const Gamma::Algebra a)
|
||||
|
||||
void checkProject(GridSerialRNG &rng)
|
||||
{
|
||||
SpinVector rv, recon, full;
|
||||
HalfSpinVector hsp, hsm;
|
||||
SpinVector rv, recon;
|
||||
HalfSpinVector hsm;
|
||||
|
||||
random(rng, rv);
|
||||
|
||||
|
@ -198,7 +198,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<GparityWilsonFermionR,FermionField> HermOpEO(Dw);
|
||||
HermOpEO.MpcDagMpc(chi_e,dchi_e);
|
||||
|
@ -364,14 +364,12 @@ int main(int argc, char **argv) {
|
||||
|
||||
{ // Peek-ology and Poke-ology, with a little app-ology
|
||||
Complex c;
|
||||
ColourMatrix c_m;
|
||||
SpinMatrix s_m;
|
||||
SpinColourMatrix sc_m;
|
||||
ColourMatrix c_m = Zero();
|
||||
SpinMatrix s_m = Zero();
|
||||
SpinColourMatrix sc_m = Zero();
|
||||
|
||||
s_m = TensorIndexRecursion<ColourIndex>::traceIndex(
|
||||
sc_m); // Map to traceColour
|
||||
c_m = TensorIndexRecursion<SpinIndex>::traceIndex(
|
||||
sc_m); // map to traceSpin
|
||||
s_m = TensorIndexRecursion<ColourIndex>::traceIndex(sc_m); // Map to traceColour
|
||||
c_m = TensorIndexRecursion<SpinIndex>::traceIndex(sc_m); // map to traceSpin
|
||||
|
||||
c = TensorIndexRecursion<SpinIndex>::traceIndex(s_m);
|
||||
c = TensorIndexRecursion<ColourIndex>::traceIndex(c_m);
|
||||
|
@ -217,7 +217,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd , chi_o, chi);
|
||||
pickCheckerboard(Even, phi_e, phi);
|
||||
pickCheckerboard(Odd , phi_o, phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<MobiusEOFAFermionR,LatticeFermion> HermOpEO(Ddwf);
|
||||
HermOpEO.MpcDagMpc(chi_e, dchi_e);
|
||||
|
@ -262,7 +262,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
|
||||
SchurDiagMooeeOperator<MobiusFermionR,LatticeFermion> HermOpEO(Ddwf);
|
||||
|
@ -144,7 +144,7 @@ int main (int argc, char ** argv)
|
||||
Ds.Dhop(src,result,0);
|
||||
}
|
||||
double t1=usecond();
|
||||
double t2;
|
||||
|
||||
double flops=(16*(3*(6+8+8)) + 15*3*2)*volume*ncall; // == 66*16 + == 1146
|
||||
|
||||
std::cout<<GridLogMessage << "Called Ds"<<std::endl;
|
||||
|
@ -162,7 +162,6 @@ int main (int argc, char ** argv)
|
||||
}
|
||||
double t1=usecond();
|
||||
|
||||
double t2;
|
||||
double flops=(16*(3*(6+8+8)) + 15*3*2)*volume*ncall; // == 66*16 + == 1146
|
||||
|
||||
std::cout<<GridLogMessage << "Called Ds"<<std::endl;
|
||||
|
@ -30,7 +30,6 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
|
||||
using namespace std;
|
||||
using namespace Grid;
|
||||
;
|
||||
|
||||
int main (int argc, char ** argv)
|
||||
{
|
||||
@ -135,7 +134,6 @@ int main (int argc, char ** argv)
|
||||
Ds.Dhop(src,result,0);
|
||||
}
|
||||
double t1=usecond();
|
||||
double t2;
|
||||
double flops=(16*(3*(6+8+8)) + 15*3*2)*volume*ncall; // == 66*16 + == 1146
|
||||
|
||||
std::cout<<GridLogMessage << "Called Ds"<<std::endl;
|
||||
|
@ -204,7 +204,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<WilsonFermionR,LatticeFermion> HermOpEO(Dw);
|
||||
HermOpEO.MpcDagMpc(chi_e,dchi_e);
|
||||
|
@ -205,7 +205,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
SchurDiagMooeeOperator<WilsonTMFermionR,LatticeFermion> HermOpEO(Dw);
|
||||
HermOpEO.MpcDagMpc(chi_e,dchi_e);
|
||||
|
@ -276,7 +276,6 @@ int main (int argc, char ** argv)
|
||||
pickCheckerboard(Odd ,chi_o,chi);
|
||||
pickCheckerboard(Even,phi_e,phi);
|
||||
pickCheckerboard(Odd ,phi_o,phi);
|
||||
RealD t1,t2;
|
||||
|
||||
|
||||
SchurDiagMooeeOperator<ZMobiusFermionR,LatticeFermion> HermOpEO(Ddwf);
|
||||
|
@ -57,7 +57,6 @@ int main (int argc, char ** argv)
|
||||
SU<Nc>::HotConfiguration(pRNG,U);
|
||||
|
||||
double beta = 1.0;
|
||||
double c1 = -0.331;
|
||||
|
||||
IwasakiGaugeActionR Action(beta);
|
||||
// PlaqPlusRectangleActionR Action(beta,c1);
|
||||
|
@ -40,6 +40,7 @@ using namespace Grid;
|
||||
template<class Fobj,class CComplex,int nbasis>
|
||||
class ProjectedHermOp : public LinearFunction<Lattice<iVector<CComplex,nbasis > > > {
|
||||
public:
|
||||
using LinearFunction<Lattice<iVector<CComplex,nbasis > > >::operator();
|
||||
typedef iVector<CComplex,nbasis > CoarseSiteVector;
|
||||
typedef Lattice<CoarseSiteVector> CoarseField;
|
||||
typedef Lattice<CComplex> CoarseScalar; // used for inner products on fine field
|
||||
@ -67,6 +68,8 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis>
|
||||
class ProjectedFunctionHermOp : public LinearFunction<Lattice<iVector<CComplex,nbasis > > > {
|
||||
public:
|
||||
using LinearFunction<Lattice<iVector<CComplex,nbasis > > >::operator ();
|
||||
|
||||
typedef iVector<CComplex,nbasis > CoarseSiteVector;
|
||||
typedef Lattice<CoarseSiteVector> CoarseField;
|
||||
typedef Lattice<CComplex> CoarseScalar; // used for inner products on fine field
|
||||
|
@ -55,6 +55,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -78,6 +79,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -108,6 +110,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, class CoarseSolver>
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
|
@ -57,7 +57,7 @@ private:
|
||||
CheckerBoardedSparseMatrixBase<Field> & _Matrix;
|
||||
SchurRedBlackBase<Field> & _Solver;
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
/////////////////////////////////////////////////////
|
||||
@ -75,6 +75,7 @@ public:
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -98,6 +99,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -128,6 +130,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, class CoarseSolver>
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
|
@ -55,6 +55,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -78,6 +79,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -109,6 +111,8 @@ template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, clas
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
typedef typename Aggregation<Fobj,CComplex,nbasis>::CoarseVector CoarseVector;
|
||||
|
@ -56,6 +56,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -79,6 +80,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -108,6 +110,7 @@ public:
|
||||
template<class Field,class Matrix> class RedBlackSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
RealD tol;
|
||||
@ -134,6 +137,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, class CoarseSolver>
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
@ -241,7 +245,7 @@ int main (int argc, char ** argv)
|
||||
Grid_init(&argc,&argv);
|
||||
|
||||
const int Ls=16;
|
||||
const int rLs=8;
|
||||
// const int rLs=8;
|
||||
|
||||
GridCartesian * UGrid = SpaceTimeGrid::makeFourDimGrid(GridDefaultLatt(), GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
|
||||
GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
|
||||
@ -388,7 +392,7 @@ int main (int argc, char ** argv)
|
||||
// RedBlackSmoother<LatticeFermion,DomainWallFermionR> FineRBSmoother(0.00,0.001,100,Ddwf);
|
||||
|
||||
// Wrap the 2nd level solver in a MultiGrid preconditioner acting on the fine space
|
||||
ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
// ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
TwoLevelMG TwoLevelPrecon(Aggregates, LDOp,
|
||||
HermIndefOp,Ddwf,
|
||||
FineSmoother,
|
||||
|
@ -57,7 +57,7 @@ private:
|
||||
CheckerBoardedSparseMatrixBase<Field> & _Matrix;
|
||||
SchurRedBlackBase<Field> & _Solver;
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
/////////////////////////////////////////////////////
|
||||
@ -75,6 +75,7 @@ public:
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -98,6 +99,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -128,6 +130,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, class CoarseSolver>
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
|
@ -55,6 +55,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -78,6 +79,7 @@ public:
|
||||
template<class Field,class Matrix> class MirsSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & SmootherMatrix;
|
||||
FineOperator & SmootherOperator;
|
||||
@ -108,6 +110,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class Matrix, class Guesser, class CoarseSolver>
|
||||
class MultiGridPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
|
@ -57,6 +57,7 @@ private:
|
||||
OperatorFunction<Field> & _Solver;
|
||||
LinearFunction<Field> & _Guess;
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
@ -118,6 +119,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -174,6 +176,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class CoarseSolver>
|
||||
class HDCRPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef CoarsenedMatrix<Fobj,CComplex,nbasis> CoarseOperator;
|
||||
|
@ -456,8 +456,8 @@ public:
|
||||
|
||||
siteVector *CBp=Stencil.CommBuf();
|
||||
|
||||
int ptype;
|
||||
int nb2=nbasis/2;
|
||||
// int ptype;
|
||||
// int nb2=nbasis/2;
|
||||
|
||||
autoView(in_v , in, AcceleratorRead);
|
||||
autoView(st, Stencil, AcceleratorRead);
|
||||
@ -471,7 +471,7 @@ public:
|
||||
typedef decltype(coalescedRead(in_v[0])) calcVector;
|
||||
typedef decltype(coalescedRead(in_v[0](0))) calcComplex;
|
||||
int sU = sF/Ls;
|
||||
int s = sF%Ls;
|
||||
// int s = sF%Ls;
|
||||
|
||||
calcComplex res = Zero();
|
||||
calcVector nbr;
|
||||
@ -517,14 +517,14 @@ public:
|
||||
autoView(st, Stencil, AcceleratorRead);
|
||||
siteVector *CBp=Stencil.CommBuf();
|
||||
|
||||
int ptype;
|
||||
int nb2=nbasis/2;
|
||||
// int ptype;
|
||||
// int nb2=nbasis/2;
|
||||
accelerator_for2d(sF, Coarse5D->oSites(), b, nbasis, Nsimd, {
|
||||
|
||||
typedef decltype(coalescedRead(in_v[0])) calcVector;
|
||||
typedef decltype(coalescedRead(in_v[0](0))) calcComplex;
|
||||
int sU = sF/Ls;
|
||||
int s = sF%Ls;
|
||||
// int s = sF%Ls;
|
||||
|
||||
calcComplex res = Zero();
|
||||
|
||||
@ -650,7 +650,7 @@ private:
|
||||
OperatorFunction<Field> & _Solver;
|
||||
LinearFunction<Field> & _Guess;
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
/////////////////////////////////////////////////////
|
||||
@ -712,6 +712,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -735,6 +736,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class CoarseSolver>
|
||||
class MGPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef typename Aggregation<Fobj,CComplex,nbasis>::CoarseVector CoarseVector;
|
||||
@ -831,6 +833,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class CoarseSolver>
|
||||
class HDCRPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef typename Aggregation<Fobj,CComplex,nbasis>::CoarseVector CoarseVector;
|
||||
@ -1174,18 +1177,18 @@ int main (int argc, char ** argv)
|
||||
PlainHermOp<CoarseCoarseVector> IRLOpL2 (IRLHermOpL2);
|
||||
ImplicitlyRestartedLanczos<CoarseCoarseVector> IRLL2(IRLOpChebyL2,IRLOpL2,cNstop,cNk,cNm,1.0e-3,20);
|
||||
|
||||
int cNconv;
|
||||
cNm=0;
|
||||
std::vector<RealD> eval2(cNm);
|
||||
std::vector<CoarseCoarseVector> evec2(cNm,CoarseCoarse5d);
|
||||
cc_src=1.0;
|
||||
// int cNconv;
|
||||
// IRLL2.calc(eval2,evec2,cc_src,cNconv);
|
||||
|
||||
ConjugateGradient<CoarseCoarseVector> CoarseCoarseCG(0.02,10000);
|
||||
DeflatedGuesser<CoarseCoarseVector> DeflCoarseCoarseGuesser(evec2,eval2);
|
||||
NormalEquations<CoarseCoarseVector> DeflCoarseCoarseCGNE(cc_Dwf,CoarseCoarseCG,DeflCoarseCoarseGuesser);
|
||||
|
||||
ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
// ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
ZeroGuesser<CoarseCoarseVector> CoarseCoarseZeroGuesser;
|
||||
|
||||
std::cout<<GridLogMessage << "**************************************************"<< std::endl;
|
||||
|
@ -456,8 +456,8 @@ public:
|
||||
|
||||
siteVector *CBp=Stencil.CommBuf();
|
||||
|
||||
int ptype;
|
||||
int nb2=nbasis/2;
|
||||
//int ptype;
|
||||
// int nb2=nbasis/2;
|
||||
|
||||
autoView(in_v , in, AcceleratorRead);
|
||||
autoView(st, Stencil, AcceleratorRead);
|
||||
@ -471,7 +471,7 @@ public:
|
||||
typedef decltype(coalescedRead(in_v[0])) calcVector;
|
||||
typedef decltype(coalescedRead(in_v[0](0))) calcComplex;
|
||||
int sU = sF/Ls;
|
||||
int s = sF%Ls;
|
||||
// int s = sF%Ls;
|
||||
|
||||
calcComplex res = Zero();
|
||||
calcVector nbr;
|
||||
@ -517,14 +517,14 @@ public:
|
||||
autoView(st, Stencil, AcceleratorRead);
|
||||
siteVector *CBp=Stencil.CommBuf();
|
||||
|
||||
int ptype;
|
||||
int nb2=nbasis/2;
|
||||
// int ptype;
|
||||
// int nb2=nbasis/2;
|
||||
accelerator_for2d(sF, Coarse5D->oSites(), b, nbasis, Nsimd, {
|
||||
|
||||
typedef decltype(coalescedRead(in_v[0])) calcVector;
|
||||
typedef decltype(coalescedRead(in_v[0](0))) calcComplex;
|
||||
int sU = sF/Ls;
|
||||
int s = sF%Ls;
|
||||
// int s = sF%Ls;
|
||||
|
||||
calcComplex res = Zero();
|
||||
|
||||
@ -648,7 +648,7 @@ private:
|
||||
CheckerBoardedSparseMatrixBase<Field> & _Matrix;
|
||||
SchurRedBlackBase<Field> & _Solver;
|
||||
public:
|
||||
|
||||
using LinearFunction<Field>::operator();
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
/////////////////////////////////////////////////////
|
||||
@ -669,6 +669,7 @@ private:
|
||||
OperatorFunction<Field> & _Solver;
|
||||
LinearFunction<Field> & _Guess;
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
|
||||
/////////////////////////////////////////////////////
|
||||
// Wrap the usual normal equations trick
|
||||
@ -731,6 +732,7 @@ RealD InverseApproximation(RealD x){
|
||||
template<class Field,class Matrix> class ChebyshevSmoother : public LinearFunction<Field>
|
||||
{
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
typedef LinearOperatorBase<Field> FineOperator;
|
||||
Matrix & _SmootherMatrix;
|
||||
FineOperator & _SmootherOperator;
|
||||
@ -754,6 +756,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class CoarseSolver>
|
||||
class MGPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef typename Aggregation<Fobj,CComplex,nbasis>::CoarseVector CoarseVector;
|
||||
@ -850,6 +853,7 @@ public:
|
||||
template<class Fobj,class CComplex,int nbasis, class CoarseSolver>
|
||||
class HDCRPreconditioner : public LinearFunction< Lattice<Fobj> > {
|
||||
public:
|
||||
using LinearFunction<Lattice<Fobj> >::operator();
|
||||
|
||||
typedef Aggregation<Fobj,CComplex,nbasis> Aggregates;
|
||||
typedef typename Aggregation<Fobj,CComplex,nbasis>::CoarseVector CoarseVector;
|
||||
@ -1194,11 +1198,11 @@ int main (int argc, char ** argv)
|
||||
PlainHermOp<CoarseCoarseVector> IRLOpL2 (IRLHermOpL2);
|
||||
ImplicitlyRestartedLanczos<CoarseCoarseVector> IRLL2(IRLOpChebyL2,IRLOpL2,cNstop,cNk,cNm,1.0e-3,20);
|
||||
|
||||
int cNconv;
|
||||
cNm=0;
|
||||
std::vector<RealD> eval2(cNm);
|
||||
std::vector<CoarseCoarseVector> evec2(cNm,CoarseCoarse5d);
|
||||
cc_src=1.0;
|
||||
// int cNconv;
|
||||
// IRLL2.calc(eval2,evec2,cc_src,cNconv);
|
||||
|
||||
std::vector<RealD> tols ({0.005,0.001});
|
||||
@ -1218,10 +1222,10 @@ int main (int argc, char ** argv)
|
||||
for(auto c_hi : c_his ) {
|
||||
for(auto f_lo : f_los ) {
|
||||
for(auto f_hi : f_his ) {
|
||||
ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
ZeroGuesser<CoarseCoarseVector> CoarseCoarseZeroGuesser;
|
||||
// ZeroGuesser<CoarseVector> CoarseZeroGuesser;
|
||||
// ZeroGuesser<CoarseCoarseVector> CoarseCoarseZeroGuesser;
|
||||
ConjugateGradient<CoarseCoarseVector> CoarseCoarseCG(tol,10000);
|
||||
ZeroGuesser<CoarseCoarseVector> CoarseCoarseGuesser;
|
||||
// ZeroGuesser<CoarseCoarseVector> CoarseCoarseGuesser;
|
||||
SchurRedBlackDiagMooeeSolve<CoarseCoarseVector> CoarseCoarseRBCG(CoarseCoarseCG);
|
||||
SchurSolverWrapper<CoarseCoarseVector> CoarseCoarseSolver(cc_Dwf,CoarseCoarseRBCG);
|
||||
|
||||
|
@ -143,6 +143,7 @@ public:
|
||||
|
||||
template<class Field> class MultiGridPreconditionerBase : public LinearFunction<Field> {
|
||||
public:
|
||||
using LinearFunction<Field>::operator();
|
||||
virtual ~MultiGridPreconditionerBase() = default;
|
||||
virtual void setup() = 0;
|
||||
virtual void operator()(Field const &in, Field &out) = 0;
|
||||
@ -156,6 +157,7 @@ public:
|
||||
/////////////////////////////////////////////
|
||||
// Type Definitions
|
||||
/////////////////////////////////////////////
|
||||
using MultiGridPreconditionerBase<Lattice<Fobj>>::operator();
|
||||
|
||||
// clang-format off
|
||||
typedef Aggregation<Fobj, CComplex, nBasis> Aggregates;
|
||||
@ -568,6 +570,7 @@ public:
|
||||
/////////////////////////////////////////////
|
||||
// Type Definitions
|
||||
/////////////////////////////////////////////
|
||||
using MultiGridPreconditionerBase<Lattice<Fobj>>::operator();
|
||||
|
||||
typedef Matrix FineDiracMatrix;
|
||||
typedef Lattice<Fobj> FineVector;
|
||||
|
@ -56,7 +56,6 @@ int main (int argc, char ** argv)
|
||||
QuasiMinimalResidual<LatticeFermion> QMR(1.0e-8,10000);
|
||||
|
||||
RealD mass=0.0;
|
||||
RealD M5=1.8;
|
||||
WilsonFermionR Dw(Umu,*Grid,*rbGrid,mass);
|
||||
|
||||
NonHermitianLinearOperator<WilsonFermionR,LatticeFermion> NonHermOp(Dw);
|
||||
|
Reference in New Issue
Block a user