Compare commits

...
4 Commits
Author SHA1 Message Date
Peter Boyle 87baaf1d37 Updated workaround 2026-09-29 14:58:18 -04:00
Peter Boyle 45c1a525b9 Useful script and keep the mixed precision job 2026-09-29 13:47:38 -04:00
Peter Boyle 3a3a20b8c9 Updates 2026-09-29 13:46:09 -04:00
Peter Boyle 3bd5e8883f Synch multigrid rewrite to repository and
broad unicode elimination effort
2026-09-29 12:36:34 -04:00
82 changed files with 3587 additions and 864 deletions

No files matched your search

+2 -2
View File
@@ -106,7 +106,7 @@ Tests and benchmarks that need optional fermion representations are guarded by `
### GPU acceleration and the view/memory-manager discipline
### Multigrid (`Grid/algorithms/multigrid/`)
Aggregation-based algebraic multigrid for Wilson-type fermions. Key files: `GeneralCoarsenedMatrixMultiRHSV2.h` (the multi-RHS coarse operator, "V2": coarsening by Fourier probing, batched-GEMM apply on a D+1 grid with rhs innermost and unvectorised), `MultiRHSBlockProject.h` (in `deflation/`; the batched-GEMM transfer operators), `PVdagMMultiGrid.h` (three-level non-Hermitian PVdagM chain with dense bottom) and `HDCGMultiGrid.h` (two-level Hermitian HDCG chain) with their `*Params.h` (XML-serialisable parameters), `MrhsMultiGrid.h` (mrhs V-cycle, preconditioner interface, fp64/fp32 seam), `Smoothers.h`, `MultiGridIO.h`, `Aggregates.h` (near-null vector construction), `Geometry.h` (coarse stencils: max-norm-1 boxes of Manhattan radius 1, 2 or 4), `CoarsenedMatrix.h` (the original nearest-neighbour Wilson multigrid). `deprecated/` holds the V1 coarse operators and the Aggregation-based single-RHS ADEF-2, kept only for the pre-2026 drivers. `MultiGrid.h` is the top-level include.
Aggregation-based algebraic multigrid for Wilson-type fermions. Key files: `GeneralCoarsenedMatrixMultiRHS.h` (`MultiGeneralCoarsenedOperator`, the coarse operator for any number of right-hand sides: coarsening by Fourier probing, batched-GEMM apply on a D+1 grid with rhs innermost and unvectorised), `MultiRHSBlockProject.h` (in `deflation/`; the batched-GEMM transfer operators), `PVdagMMultiGrid.h` (three-level non-Hermitian PVdagM chain with dense bottom) and `HDCGMultiGrid.h` (two-level Hermitian HDCG chain) with their `*Params.h` (XML-serialisable parameters), `MrhsMultiGrid.h` (mrhs V-cycle, preconditioner interface, fp64/fp32 seam), `Smoothers.h`, `MultiGridIO.h`, `Aggregates.h` (near-null vector construction), `Geometry.h` (coarse stencils: max-norm-1 boxes of Manhattan radius 1, 2 or 4), `CoarsenedMatrix.h` (the original nearest-neighbour Wilson multigrid). `deprecated/` holds the superseded coarse operators and the Aggregation-based single-RHS ADEF-2, with class names prefixed `Deprecated` (`DeprecatedGeneralCoarsenedMatrix`, `DeprecatedMultiGeneralCoarsenedMatrix`, `DeprecatedTwoLevelADEF2`); nothing in the library uses them, only older drivers in `examples/` and `tests/debug/`. `MultiGrid.h` is the top-level include.
### GPU acceleration
@@ -174,7 +174,7 @@ Every programme is wrapped in `Grid_init(&argc, &argv)` / `Grid_finalize()` (`Gr
- Template structure: most classes are templated on `<_FImpl>` (fermion impl) or `<Gimpl>` (gauge impl), which encode the representation and precision. Instantiation is controlled by `--enable-fermion-instantiations`.
- **Tensor indices are positional, not labelled.** The `Grid/tensors/` arithmetic recurses structurally over the `iScalar`/`iVector`/`iMatrix` nest: each level defines only the {scalar,vector,matrix}² products at its own level, with element types resolved by automatic type deduction, so every colour/spin/lorentz combination composes from ~200 lines (versus the pre-C++11 QDP++/PETE approach of machine-generating every case). An index's meaning derives entirely from its nesting depth counted from the outside; `iScalar` is the identity/broadcast case at every level. Never insert or remove a nesting level casually — the multiplication tables contract by position.
- **Multigrid coarsening deepens the tensor nest by one level.** A coarse site vector is `iVector<CComplex,nbasis>`, and `innerProduct` on it returns `iScalar<CComplex>` — one level deeper than the fine block scalar. So the block-inner-product scalar type gains one `iScalar` wrapper per MG level (fine: `vTComplex`; level 2: `iScalar<vTComplex>`; see `examples/Example_pvdagm_3level.cc`). When calling `blockInnerProduct`/`blockZAXPY`/`blockOrthogonalise` on coarse fields, the coarse scalar type must match `decltype(innerProduct(siteVector(),siteVector()))` exactly; a wrong depth fails to compile (no viable `operator=` deep in the instantiation chain) rather than mis-contracting.
- **Grids are borrowed, never owned.** `conformable` is pointer identity, so every object that interoperates must hold the *same* `GridCartesian *`; a class that minted its own grid internally could never conform with anything else. Ownership is therefore not available, and lifetime is managed by scope discipline instead of reference counting: whoever creates a grid retains it beyond every object it handed a reference to. Anything *derived* from a grid inherits this — `~PaddedCell` dereferences its `unpadded_grid`, so a `PaddedCell` cannot even be **destroyed** after its parent grid, only used. Where a consumer must let go early, it offers an explicit hand-back (`MultiGeneralCoarsenedOperatorV2::ReleaseGrid()`) to be called *before* the grid is destroyed.
- **Grids are borrowed, never owned.** `conformable` is pointer identity, so every object that interoperates must hold the *same* `GridCartesian *`; a class that minted its own grid internally could never conform with anything else. Ownership is therefore not available, and lifetime is managed by scope discipline instead of reference counting: whoever creates a grid retains it beyond every object it handed a reference to. Anything *derived* from a grid inherits this — `~PaddedCell` dereferences its `unpadded_grid`, so a `PaddedCell` cannot even be **destroyed** after its parent grid, only used. Where a consumer must let go early, it offers an explicit hand-back (`MultiGeneralCoarsenedOperator::ReleaseGrid()`) to be called *before* the grid is destroyed.
- The `RealD`/`RealF`/`ComplexD`/`ComplexF` typedefs are used everywhere; avoid raw `double`/`float`.
- Use `GRID_ASSERT(cond)` (defined in `Grid/GridStd.h`), not bare `assert` — it prints a Grid-formatted message and aborts cleanly under MPI.
+3
View File
@@ -667,6 +667,9 @@ public:
(*this)(in[i], out[i]);
}
}
// The identity, if a derived class says so. A caller can then use its input
// where it would otherwise have materialised the output, and skip the copy.
virtual int isTrivial(void) { return 0; };
virtual ~LinearFunction(){};
};
+1
View File
@@ -45,6 +45,7 @@ public:
virtual void operator()(const Field &src, Field & psi){
psi = src;
}
virtual int isTrivial(void) { return 1; };
TrivialPrecon(void){};
};
@@ -1,6 +1,6 @@
/*************************************************************************************
Grid physics library, www.github.com/paboyle/Grid
Grid physics library, www.github.com/paboyle/Grid
Source file: ./lib/algorithms/iterative/PrecGeneralisedConjugateResidual.h
@@ -28,213 +28,48 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
/* END LEGAL */
#ifndef GRID_PREC_GCR_H
#define GRID_PREC_GCR_H
#include <Grid/algorithms/iterative/PrecGeneralisedConjugateResidualNonHermitian.h>
///////////////////////////////////////////////////////////////////////////////////////////////////////
//VPGCR Abe and Zhang, 2005.
//INTERNATIONAL JOURNAL OF NUMERICAL ANALYSIS AND MODELING
//Computing and Information Volume 2, Number 2, Pages 147-161
//NB. Likely not original reference since they are focussing on a preconditioner variant.
// but VPGCR was nicely written up in their paper
///////////////////////////////////////////////////////////////////////////////////////////////////////
NAMESPACE_BEGIN(Grid);
#define GCRLogLevel std::cout << GridLogMessage <<std::string(level,'\t')<< " Level "<<level<<" "
///////////////////////////////////////////////////////////////////////////////////////////////////////
// Hermitian spelling of PrecGeneralisedConjugateResidualNonHermitian.
//
// The algorithm is the same; only the operator entry point differs. The solver
// drives Linop.Op(), while a caller of this class means "invert HermOp", so the
// operator is wrapped in HermOpAdaptor. The wrap is what makes the two classes
// one: for an operator whose Op() already IS its HermOp() (HermitianLinearOperator,
// Gamma5R5HermitianLinearOperator, HermOpAdaptor) it changes nothing, and for one
// where they differ (MdagMLinearOperator: Op()=M, HermOp()=MdagM) it preserves the
// Hermitian meaning the caller asked for.
//
// A non-zero initial guess is respected. Callers that guarantee x0 = 0 can call
// SetZeroGuess(1) to skip the first r0 = src - A x0 apply.
///////////////////////////////////////////////////////////////////////////////////////////////////////
template<class Field>
class PrecGeneralisedConjugateResidual : public LinearFunction<Field> {
public:
// Declaration order is initialisation order: the adaptor is complete before
// the solver binds its reference to it.
HermOpAdaptor<Field> HermLinop;
PrecGeneralisedConjugateResidualNonHermitian<Field> GCR;
public:
using LinearFunction<Field>::operator();
RealD Tolerance;
Integer MaxIterations;
int verbose;
int mmax;
int nstep;
int steps;
int level;
GridStopWatch PrecTimer;
GridStopWatch MatTimer;
GridStopWatch LinalgTimer;
LinearFunction<Field> &Preconditioner;
LinearOperatorBase<Field> &Linop;
PrecGeneralisedConjugateResidual(RealD tol,Integer maxit,LinearOperatorBase<Field> &Linop,
LinearFunction<Field> &Prec,int mmax,int nstep)
: HermLinop(Linop),
GCR(tol,maxit,HermLinop,Prec,mmax,nstep) {};
void Level(int lv) { level=lv; };
void operator() (const Field &src, Field &psi) { GCR(src,psi); };
PrecGeneralisedConjugateResidual(RealD tol,Integer maxit,LinearOperatorBase<Field> &_Linop,LinearFunction<Field> &Prec,int _mmax,int _nstep) :
Tolerance(tol),
MaxIterations(maxit),
Linop(_Linop),
Preconditioner(Prec),
mmax(_mmax),
nstep(_nstep)
{
level=1;
verbose=1;
};
void operator() (const Field &src, Field &psi){
psi=Zero();
RealD cp, ssq,rsq;
ssq=norm2(src);
rsq=Tolerance*Tolerance*ssq;
Field r(src.Grid());
PrecTimer.Reset();
MatTimer.Reset();
LinalgTimer.Reset();
GridStopWatch SolverTimer;
SolverTimer.Start();
steps=0;
for(int k=0;k<MaxIterations;k++){
cp=GCRnStep(src,psi,rsq);
GCRLogLevel <<"PGCR("<<mmax<<","<<nstep<<") "<< steps <<" steps cp = "<<cp<<" target "<<rsq <<std::endl;
if(cp<rsq) {
SolverTimer.Stop();
Linop.HermOp(psi,r);
axpy(r,-1.0,src,r);
RealD tr = norm2(r);
GCRLogLevel<<"PGCR: Converged on iteration " <<steps
<< " computed residual "<<sqrt(cp/ssq)
<< " true residual " <<sqrt(tr/ssq)
<< " target " <<Tolerance <<std::endl;
GCRLogLevel<<"PGCR Time elapsed: Total "<< SolverTimer.Elapsed() <<std::endl;
/*
GCRLogLevel<<"PGCR Time elapsed: Precon "<< PrecTimer.Elapsed() <<std::endl;
GCRLogLevel<<"PGCR Time elapsed: Matrix "<< MatTimer.Elapsed() <<std::endl;
GCRLogLevel<<"PGCR Time elapsed: Linalg "<< LinalgTimer.Elapsed() <<std::endl;
*/
return;
}
}
GCRLogLevel<<"Variable Preconditioned GCR did not converge"<<std::endl;
// GRID_ASSERT(0);
}
RealD GCRnStep(const Field &src, Field &psi,RealD rsq){
RealD cp;
RealD a, b;
RealD zAz, zAAz;
RealD rq;
GridBase *grid = src.Grid();
Field r(grid);
Field z(grid);
Field tmp(grid);
Field ttmp(grid);
Field Az(grid);
////////////////////////////////
// history for flexible orthog
////////////////////////////////
std::vector<Field> q(mmax,grid);
std::vector<Field> p(mmax,grid);
std::vector<RealD> qq(mmax);
GCRLogLevel<< "PGCR nStep("<<nstep<<")"<<std::endl;
//////////////////////////////////
// initial guess x0 is taken as nonzero.
// r0=src-A x0 = src
//////////////////////////////////
MatTimer.Start();
Linop.HermOpAndNorm(psi,Az,zAz,zAAz);
MatTimer.Stop();
LinalgTimer.Start();
r=src-Az;
LinalgTimer.Stop();
GCRLogLevel<< "PGCR true residual r = src - A psi "<<norm2(r) <<std::endl;
/////////////////////
// p = Prec(r)
/////////////////////
PrecTimer.Start();
Preconditioner(r,z);
PrecTimer.Stop();
MatTimer.Start();
Linop.HermOpAndNorm(z,Az,zAz,zAAz);
MatTimer.Stop();
LinalgTimer.Start();
//p[0],q[0],qq[0]
p[0]= z;
q[0]= Az;
qq[0]= zAAz;
cp =norm2(r);
LinalgTimer.Stop();
for(int k=0;k<nstep;k++){
steps++;
int kp = k+1;
int peri_k = k %mmax;
int peri_kp= kp%mmax;
LinalgTimer.Start();
rq= real(innerProduct(r,q[peri_k])); // what if rAr not real?
a = rq/qq[peri_k];
axpy(psi,a,p[peri_k],psi);
cp = axpy_norm(r,-a,q[peri_k],r);
LinalgTimer.Stop();
GCRLogLevel<< "PGCR step["<<steps<<"] resid " << cp << " target " <<rsq<<std::endl;
if((k==nstep-1)||(cp<rsq)){
return cp;
}
PrecTimer.Start();
Preconditioner(r,z);// solve Az = r
PrecTimer.Stop();
MatTimer.Start();
Linop.HermOpAndNorm(z,Az,zAz,zAAz);
MatTimer.Stop();
LinalgTimer.Start();
q[peri_kp]=Az;
p[peri_kp]=z;
int northog = ((kp)>(mmax-1))?(mmax-1):(kp); // if more than mmax done, we orthog all mmax history.
for(int back=0;back<northog;back++){
int peri_back=(k-back)%mmax; GRID_ASSERT((k-back)>=0);
b=-real(innerProduct(q[peri_back],Az))/qq[peri_back];
p[peri_kp]=p[peri_kp]+b*p[peri_back];
q[peri_kp]=q[peri_kp]+b*q[peri_back];
}
qq[peri_kp]=norm2(q[peri_kp]); // could use axpy_norm
LinalgTimer.Stop();
}
GRID_ASSERT(0); // never reached
return cp;
}
void Name(std::string n) { GCR.Name(n); };
void Level(int n) { GCR.Level(n); };
void SetZeroGuess(int z) { GCR.SetZeroGuess(z); };
int Steps(void) const { return GCR.Steps(); };
void LogCoefficients(int l) { GCR.LogCoefficients(l); };
void SetCoefficientRecorder(GCRCoefficients *r){ GCR.SetCoefficientRecorder(r); };
void ReleaseHistory(void) { GCR.ReleaseHistory(); };
};
NAMESPACE_END(Grid);
#undef GCRLogLevel
#endif
@@ -43,8 +43,51 @@ NAMESPACE_BEGIN(Grid);
template<class Field>
class PrecGeneralisedConjugateResidualNonHermitian : public LinearFunction<Field> {
public:
public:
using LinearFunction<Field>::operator();
PrecGeneralisedConjugateResidualNonHermitian(RealD tol,Integer maxit,LinearOperatorBase<Field> &_Linop,LinearFunction<Field> &Prec,int _mmax,int _nstep) :
Tolerance(tol),
MaxIterations(maxit),
mmax(_mmax),
nstep(_nstep),
Preconditioner(Prec),
Linop(_Linop)
{
Level(1);
verbose=1;
trivial_prec = Preconditioner.isTrivial();
};
// The name is also the trace range's, so a profile separates the four
// instances (Fouter, Fsmoother, Couter, Csmoother) that otherwise nest
// indistinguishably as one shared range.
void Name(std::string _name) { name = _name; trace_step = name + " PGCR_step"; };
void Level(int n) { Name("Level " + std::to_string(n)); level = n; }
void SetZeroGuess(int z) { ZeroGuess = z; };
// Steps taken by the last solve.
int Steps(void) const { return steps; };
// Coefficient logging: one line per step with the step length a_k and the
// orthogonalisation coefficients b_{k,j}. These are the data from which a
// FIXED polynomial smoother can be harvested: if they are stable from call
// to call, the adaptive GCR can be replaced by a stationary p(A) with the
// same applies and no reductions. Off by default; boss rank prints.
void LogCoefficients(int l) { LogCoeffs = l; };
// Optional recorder of the per-step coefficients (means over calls), for
// replay by GCRReplaySmoother (Smoothers.h). Records only; no effect on
// the iteration.
void SetCoefficientRecorder(GCRCoefficients *r) { Recorder = r; if(r) r->mmax = mmax; };
// Free the persistent history (e.g. when this solver is replaced by a
// replayed polynomial): re-made on the next call if ever needed again.
void ReleaseHistory(void) { q.clear(); p.clear(); qq.clear(); hist_grid = nullptr; };
private:
RealD Tolerance;
RealD SSQ;
Integer MaxIterations;
@@ -65,45 +108,14 @@ public:
std::vector<Field> p;
std::vector<RealD> qq;
int FirstCycle = 0;
int LogCoeffs = 0;
int trivial_prec; // Preconditioner is the identity: p == r, copies fold away
GCRCoefficients *Recorder = nullptr;
LinearFunction<Field> &Preconditioner;
LinearOperatorBase<Field> &Linop;
// The name is also the trace range's, so a profile separates the four
// instances (Fouter, Fsmoother, Couter, Csmoother) that otherwise nest
// indistinguishably as one shared range.
void Name(std::string _name) { name = _name; trace_step = name + " PGCR_step"; };
void Level(int n) { Name("Level " + std::to_string(n)); level = n; }
void SetZeroGuess(int z) { ZeroGuess = z; };
// Coefficient logging: one line per step with the step length a_k and the
// orthogonalisation coefficients b_{k,j}. These are the data from which a
// FIXED polynomial smoother can be harvested: if they are stable from call
// to call, the adaptive GCR can be replaced by a stationary p(A) with the
// same applies and no reductions. Off by default; boss rank prints.
int LogCoeffs = 0;
void LogCoefficients(int l) { LogCoeffs = l; };
// Optional recorder of the per-step coefficients (means over calls), for
// replay by GCRReplaySmoother (Smoothers.h). Records only; no effect on
// the iteration.
GCRCoefficients *Recorder = nullptr;
void SetCoefficientRecorder(GCRCoefficients *r) { Recorder = r; if(r) r->mmax = mmax; };
// Free the persistent history (e.g. when this solver is replaced by a
// replayed polynomial): re-made on the next call if ever needed again.
void ReleaseHistory(void) { q.clear(); p.clear(); qq.clear(); hist_grid = nullptr; };
PrecGeneralisedConjugateResidualNonHermitian(RealD tol,Integer maxit,LinearOperatorBase<Field> &_Linop,LinearFunction<Field> &Prec,int _mmax,int _nstep) :
Tolerance(tol),
MaxIterations(maxit),
Linop(_Linop),
Preconditioner(Prec),
mmax(_mmax),
nstep(_nstep)
{
Level(1);
verbose=1;
};
public:
void operator() (const Field &src, Field &psi){
@@ -151,6 +163,7 @@ public:
// GRID_ASSERT(0);
}
private:
RealD GCRnStep(const Field &src, Field &psi,RealD rsq){
RealD cp;
@@ -189,18 +202,20 @@ public:
// contract (enforced here), so r0 = src exactly; skip the apply.
// Restart cycles (psi!=0) always do the full computation.
//////////////////////////////////
if (ZeroGuess && FirstCycle) {
psi = Zero();
LinalgTimer.Start();
r = src;
LinalgTimer.Stop();
} else {
// The residual is READ through rp and only materialised in r when it is
// first written, at the k=0 update. On a zero-guess first cycle r0 IS src,
// so no copy is made; a restart cycle computes r and rp points at it.
const Field *rp = &src;
GRID_ASSERT(nstep>=1); // psi is first written at k=0; see zero_start below
int zero_start = (ZeroGuess && FirstCycle);
if ( !zero_start ) {
MatTimer.Start();
Linop.Op(psi,Az);
MatTimer.Stop();
LinalgTimer.Start();
r=src-Az;
LinalgTimer.Stop();
rp = &r;
}
FirstCycle=0;
@@ -208,9 +223,11 @@ public:
// p = Prec(r)
/////////////////////
// p[0] = Prec(r), q[0] = A p[0], written straight into the history slots
// p[0] = Prec(r), q[0] = A p[0], written straight into the history slots.
// This copy is forced even for a trivial preconditioner: p[0] is history and
// must survive while the residual evolves.
PrecTimer.Start();
Preconditioner(r,p[0]);
Preconditioner(*rp,p[0]);
PrecTimer.Stop();
MatTimer.Start();
@@ -221,7 +238,7 @@ public:
qq[0]= norm2(q[0]);
cp =norm2(r);
cp =norm2(*rp);
LinalgTimer.Stop();
GCRLogLevel<< "PGCR true residual "<< sqrt(cp/SSQ) <<std::endl;
@@ -234,13 +251,32 @@ public:
int peri_kp= kp%mmax;
LinalgTimer.Start();
rq= innerProduct(q[peri_k],r); // what if rAr not real?
rq= innerProduct(q[peri_k],*rp); // what if rAr not real?
a = rq/qq[peri_k];
if ( Recorder ) Recorder->RecordA(k,a);
axpy(psi,a,p[peri_k],psi);
// On a zero-guess first step psi is still unwritten, and the update would
// be a*p added to zero: assign it instead, so psi is never zeroed.
if ( zero_start && (k==0) ) psi = a*p[peri_k];
else axpy(psi,a,p[peri_k],psi);
cp = axpy_norm(r,-a,q[peri_k],r);
#ifdef GRID_GCR_RESIDUAL_RECURRENCE
// |r - a q|^2 = |r|^2 - |<q,r>|^2/|q|^2 with a = <q,r>/|q|^2, so the new
// residual norm follows from scalars already in hand and the reduction in
// axpy_norm is not needed. cp then DRIFTS rather than being measured; the
// true residual is recomputed at restart and at convergence, which bounds
// the exposure. The drift only matters for an instance driven to a tight
// tolerance -- the outer solver -- and there it is the same computed-vs-true
// residual gap that CG lives with; the smoother and coarse instances run at
// fixed work or loose tolerance and are indifferent to it.
// Explicit re/im: ComplexD is thrust::complex under HIP.
axpy(r,-a,q[peri_k],*rp);
cp = cp - (real(rq)*real(rq)+imag(rq)*imag(rq))/qq[peri_k];
if ( cp < 0.0 ) cp = 0.0; // drift can undershoot; keeps the log finite
#else
cp = axpy_norm(r,-a,q[peri_k],*rp);
#endif
rp = &r; // the residual now lives in r; src is untouched from here on
LinalgTimer.Stop();
if ( LogCoeffs ) {
GCRLogLevel<<"coeff["<<k<<"] a=("<<real(a)<<","<<imag(a)<<")"
@@ -254,13 +290,21 @@ public:
}
// New direction written straight into its history slot: p = Prec(r), q = A p.
PrecTimer.Start();
Preconditioner(r,p[peri_kp]);
PrecTimer.Stop();
MatTimer.Start();
Linop.Op(p[peri_kp],q[peri_kp]);
MatTimer.Stop();
// For a trivial preconditioner p[kp] would only be a copy of the residual,
// so the operator is applied to the residual directly and the copy folds
// into the orthogonalisation pass below.
if ( trivial_prec ) {
MatTimer.Start();
Linop.Op(*rp,q[peri_kp]);
MatTimer.Stop();
} else {
PrecTimer.Start();
Preconditioner(*rp,p[peri_kp]);
PrecTimer.Stop();
MatTimer.Start();
Linop.Op(p[peri_kp],q[peri_kp]);
MatTimer.Stop();
}
LinalgTimer.Start();
@@ -289,16 +333,19 @@ public:
if ( LogCoeffs ) bs<<" b["<<back<<"]="<<bcoef[back];
}
if ( Recorder && northog ) Recorder->RecordB(k,bcoef);
// axpyMultiNorm returns |q|^2 from the update pass, so qq is complete here:
// a separate norm2 would be a second reduction and a second GlobalSum.
if ( northog ) {
axpyMulti(p[peri_kp],bcoef,pwin);
qq[peri_kp]=axpyMultiNorm(q[peri_kp],bcoef,qwin);
if ( trivial_prec ) axpyMultiFrom(p[peri_kp],*rp,bcoef,pwin);
else axpyMulti(p[peri_kp],bcoef,pwin);
qq[peri_kp]=axpyMultiNorm(q[peri_kp],bcoef,qwin);
} else {
qq[peri_kp]=norm2(q[peri_kp]);
if ( trivial_prec ) p[peri_kp] = *rp; // no pass to fold the copy into
qq[peri_kp]=norm2(q[peri_kp]);
}
if ( LogCoeffs && northog ) {
GCRLogLevel<<"coeff["<<k<<"]"<<bs.str()<<std::endl;
}
qq[peri_kp]=norm2(q[peri_kp]); // could use axpy_norm
LinalgTimer.Stop();
}
GRID_ASSERT(0); // never reached
+1 -2
View File
@@ -39,8 +39,7 @@ typedef ComplexD DenseInverseScalar;
// BlockCyclicLayout: the index arithmetic of a 2D block-cyclic distribution
// of an N x N matrix over a Pr x Pc logical process grid with block size nb.
//
// This is stage 1 of the 2D distributed dense inverse
// (documentation/DistributedDenseInverse2D.tex). It is deliberately
// This is stage 1 of the 2D distributed dense inverse. It is deliberately
// COMMUNICATOR-FREE: every mapping is a static pure function of
// (N, nb, Pr, Pc), so the whole layout is exhaustively unit-testable on one
// rank with no MPI in the loop (Test_blockcyclic). A thin instance layer
@@ -273,11 +273,15 @@ public:
static void RowsToCyclic(GridBase *grid, const std::vector<int64_t> &rowStart,
DenseInverseScalar *rows1d, int64_t myrows, BlockCyclicMatrix &A)
{ Redistribute(+1, grid, rowStart, rows1d, myrows, A); }
{
Redistribute(+1, grid, rowStart, rows1d, myrows, A);
}
static void CyclicToRows(GridBase *grid, const std::vector<int64_t> &rowStart,
BlockCyclicMatrix &A, DenseInverseScalar *rows1d, int64_t myrows)
{ Redistribute(-1, grid, rowStart, rows1d, myrows, A); }
{
Redistribute(-1, grid, rowStart, rows1d, myrows, A);
}
};
NAMESPACE_END(Grid);
@@ -215,7 +215,11 @@ public:
// SendToRecvFrom with root (symmetric byte count: the reverse direction
// carries a same-size dummy -- a leaf-local cost, accepted for simplicity).
///////////////////////////////////////////////////////////////////////////
static int64_t FirstBlock(int64_t b0, int p, int Pg){ int64_t r = ((b0 % Pg) <= p) ? b0 - (b0 % Pg) + p : b0 - (b0 % Pg) + Pg + p; return r; }
static int64_t FirstBlock(int64_t b0, int p, int Pg)
{
int64_t r = ((b0 % Pg) <= p) ? b0 - (b0 % Pg) + p : b0 - (b0 % Pg) + Pg + p;
return r;
}
void BigLeaf(BlockCyclicMatrix &A, int64_t b0, int64_t b1)
{
GRID_TRACE("SchurBigLeaf");
+48 -13
View File
@@ -38,7 +38,7 @@ NAMESPACE_BEGIN(Grid);
//////////////////////////////////////////////////////////////////////////////////////
// DenseCoarseMatrix: a coarsened operator treated as a DENSE matrix -- explicit,
// row-distributed A^{-1} of a GeneralCoarsenedMatrix.
// row-distributed A^{-1} of a coarsened operator.
//
// - Stencil -> dense DIRECT IMPORT. The coarse operator IS the dense matrix
// unrolled: Dense[(s,a),(s+shift_p,b)] += A[p][s]_{a,b}. Rows of my sites are
@@ -176,6 +176,7 @@ public:
{
double t0 = usecond();
{
CheckMrhsSlices(Op); // the batched apply keeps its slices apart
ImportDense(Op); // slab <- my rows of A (LOCAL, no comms)
ImportCertificate(Op); // dense apply == Op.M, before inversion
InvertDense(Op); // slab <- my rows of A^{-1}
@@ -274,32 +275,67 @@ public:
GRID_ASSERT(mgrid->_ndimension == nd+1);
int nr = mgrid->_fdimensions[0];
// Every slice carries the same vector; slice 0 is the answer. The
// slices are not a check here -- see CheckMrhsSlices, which does that
// once on a field chosen for the purpose.
Field min(mgrid), mout(mgrid);
for(int r=0;r<nr;r++) InsertSliceFast(in,min,r,0);
Op.M(min,mout);
ExtractSliceFast(out,mout,0,0);
}
////////////////////////////////////////////////////////////////////
// Do the multiRHS slices stay independent? Slice r carries (r+1) times
// one RANDOM field, so linearity says the results scale likewise, and a
// slot reading another slot's data is an O(1) error.
//
// The field is random on purpose. Run on the output of the inverse, as
// this check once was, it measures the cancellation in A applied to
// A^-1 x instead: at the production point that is 1.3e-4 in an fp32
// sector with an fp32 inversion, so the check then fires on any change
// of summation order while saying nothing about slice independence.
// Once per import, not per apply.
////////////////////////////////////////////////////////////////////
template<class CoarseOp>
void CheckMrhsSlices(CoarseOp &Op)
{
if ( Op.Grid() == grid ) return; // single rhs: nothing to mix
GridBase *mgrid = Op.Grid();
GRID_ASSERT(mgrid->_ndimension == nd+1);
int nr = mgrid->_fdimensions[0];
if ( nr < 2 ) return;
Field in(grid);
GridParallelRNG rng(grid); rng.SeedFixedIntegers(std::vector<int>({7,8,9,10}));
random(rng,in);
Field min(mgrid), mout(mgrid);
for(int r=0;r<nr;r++){
Field scaled(grid);
scaled = CoarseScalar(r+1.0,0.0)*in;
InsertSliceFast(scaled,min,r,0);
}
Op.M(min,mout);
ExtractSliceFast(out,mout,0,0);
Field s0(grid);
ExtractSliceFast(s0,mout,0,0);
RealD worst=0.0;
for(int r=1;r<nr;r++){
Field sr(grid),d(grid);
ExtractSliceFast(sr,mout,r,0);
d = sr - CoarseScalar(r+1.0,0.0)*out;
d = sr - CoarseScalar(r+1.0,0.0)*s0;
RealD rel = std::sqrt(norm2(d)/norm2(sr));
// Slices differ by rounding amplified by the operator's conditioning
// when the product cancels (A applied to A^-1 x): the tolerance follows
// the coarse precision (fp32 measured ~5e-6 here). A genuine rhs
// mix-up is O(1).
worst = std::max(worst,rel);
const RealD otol = (sizeof(CoarseScalar)==sizeof(ComplexF)) ? 1.0e-4 : 1.0e-6;
if ( rel >= otol ) {
std::cout << GridLogMessage << "DenseCoarseMatrix: oracle rhs "<<r
<<" inconsistent with rhs 0, rel "<<rel<<std::endl;
std::cout << GridLogMessage << "DenseCoarseMatrix: mrhs slice "<<r
<<" inconsistent with slice 0, rel "<<rel<<std::endl;
}
GRID_ASSERT( rel < otol );
}
std::cout << GridLogMessage << "DenseCoarseMatrix: mrhs slices independent to "
<< worst << " over " << nr << " right hand sides" << std::endl;
}
////////////////////////////////////////////////////////////////////
@@ -317,9 +353,8 @@ public:
for(int p=0; p<Op.Geometry().npoint; p++){
Coordinate shift = Op.Geometry().shifts[p];
// _A[p] is PADDED after ExchangeCoarseLinks (end of CoarsenOperator):
// extract the unpadded field before peeking with unpadded coordinates
// (exactly as MultiGeneralCoarsenedMatrix::CopyMatrix does).
// ExtractMatrix hands back the unpadded matrix field, which is what the
// unpadded coordinates below index.
CoarseMatrix Aun(grid); Op.ExtractMatrix(p,Aun);
autoView(Av, Aun, CpuRead);
thread_for(ss, lsites, {
@@ -34,13 +34,13 @@ NAMESPACE_BEGIN(Grid);
// Fine Object == (per site) type of fine field
// nbasis == number of deflation vectors
template<class Fobj,class CComplex,int nbasis>
class MultiGeneralCoarsenedOperatorV2 : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
class MultiGeneralCoarsenedOperator : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
public:
typedef typename CComplex::scalar_object SComplex;
// The BLAS scalar of this level follows the coefficient precision:
// ComplexF for an fp32 coarse space, ComplexD for fp64.
typedef typename GridTypeMapper<CComplex>::scalar_type CoarseBLASScalar;
typedef MultiGeneralCoarsenedOperatorV2<Fobj,CComplex,nbasis> MultiGeneralCoarseOp;
typedef MultiGeneralCoarsenedOperator<Fobj,CComplex,nbasis> MultiGeneralCoarseOp;
typedef iVector<CComplex,nbasis > siteVector;
typedef iMatrix<CComplex,nbasis > siteMatrix;
@@ -82,7 +82,13 @@ public:
deviceVector<calcVector> BLAS_B;
deviceVector<calcVector> BLAS_C;
std::vector<deviceVector<calcMatrix> > BLAS_A;
// The matrix, ONE buffer, site major: a site's npoint blocks are contiguous.
// Both organisations read it through their pointer tables, the folded one a
// site at a time and the unfolded one a point at a time with a stride of the
// whole stencil, so nothing is duplicated. Writing or reading one point
// goes through a single-point scratch.
deviceVector<calcMatrix> BLAS_Amat; // [site][point]
deviceVector<calcMatrix> BLAS_Ascr; // one point, for I/O
std::vector<deviceVector<CoarseBLASScalar *> > BLAS_AP;
std::vector<deviceVector<CoarseBLASScalar *> > BLAS_BP;
@@ -143,20 +149,106 @@ public:
for(int l=0;l<legs;l++){
int p = g*LegGroup+l;
for(int32_t ss=0;ss<sites;ss++){
CoarseBLASScalar *ptr = (CoarseBLASScalar *)&BLAS_A[p][ss];
CoarseBLASScalar *ptr = (CoarseBLASScalar *)&BLAS_Amat[(uint64_t)ss*geom.npoint+p];
acceleratorPut(BLAS_APg[g][l*sites+ss],ptr);
}
}
}
}
void SetLegGroup(int G)
///////////////////////////////////////////////////////////////////////////
// Stencil legs folded along the CONTRACTION instead of the batch.
//
// One GEMM per site of [nbasis x (KFold*nbasis)] . [(KFold*nbasis) x nrhs],
// accumulating straight into C, so there are no partial results to combine.
// It needs the site's matrices contiguous, which is a repack done once, and
// the neighbour vectors gathered into one block per site, which replaces the
// copy the ungrouped path already does.
//
// Measured on MI250X at nbasis 60, batch 1024, single precision: the short
// form (K=60) reaches 487 GB/s at one right-hand side and 625 at twelve,
// the folded form (K=1980) 912 and 823. Beyond about K=2000 it collapses:
// K=4860, the full 81 point box, gives 304 GB/s. So the fold is chunked to
// keep K near 2000, and the chunks accumulate.
//
// The gather costs one extra write of nbasis*nrhs per leg per site, which is
// 3% of the traffic at one right-hand side and 20% at twelve; the fold wins
// at both, and by most at one, which is the HMC case.
///////////////////////////////////////////////////////////////////////////
static const int TargetK = 2048;
int KFold = 0; // legs per GEMM; 0 = off
deviceVector<CoarseBLASScalar> BLAS_BK; // [site][rhs][leg][nbasis]
std::vector<deviceVector<CoarseBLASScalar *> > BLAS_AKP; // per chunk, per site
std::vector<deviceVector<CoarseBLASScalar *> > BLAS_BKP; // per chunk, per site
deviceVector<CoarseBLASScalar *> BLAS_BPflat; // [point][site] neighbours
int KChunks(void) const { return (geom.npoint + KFold - 1)/KFold; }
int KLegs(int g) const { return std::min(KFold, geom.npoint - g*KFold); }
///////////////////////////////////////////////////////////////////////////
// Repack the matrix so a site's stencil points are contiguous: the folded
// GEMM reads each site's block as one column-major nbasis x (legs*nbasis)
// matrix, which is also a single sequential stream rather than npoint
// interleaved ones. Costs one pass over the matrix, once per solve setup.
///////////////////////////////////////////////////////////////////////////
void PackFoldedMatrix(void)
{
GRID_ASSERT( G>=1 );
GRID_ASSERT( G<=geom.npoint ); // every partial must get an initialising call
LegGroup = G;
BuildGroupedA();
if ( _CoarseGridMulti ) SetGrid(_CoarseGridMulti); // rebuild the rest
// Nothing to pack: the matrix is already site major. Only the pointer
// table is built, one entry per site per chunk, at the chunk's first leg.
int32_t sites = _CoarseGrid->lSites();
const int np = geom.npoint;
BLAS_AKP.resize(KChunks());
for(int g=0;g<KChunks();g++){
BLAS_AKP[g].resize(sites);
for(int32_t ss=0;ss<sites;ss++){
CoarseBLASScalar *ptr = (CoarseBLASScalar *)&BLAS_Amat[(uint64_t)ss*np + (uint64_t)g*KFold];
acceleratorPut(BLAS_AKP[g][ss],ptr);
}
}
}
// Legs per GEMM that keeps K near where the kernel runs best.
int AutoKFold(void) const
{
int K = TargetK/nbasis;
if ( K < 1 ) K = 1;
if ( K > geom.npoint ) K = geom.npoint;
return K;
}
// Legs per call that brings the batch up to where the kernel saturates.
int AutoLegGroup(void) const
{
int32_t sites = _CoarseGrid->lSites();
int G = (TargetBatch + sites - 1)/sites;
if ( G < 1 ) G = 1;
if ( G > geom.npoint ) G = geom.npoint;
return G;
}
///////////////////////////////////////////////////////////////////////////
// The fold along K gathers the neighbours, which costs nbasis*nrhs per leg
// per site: 16 MB of traffic per apply at one right hand side and 195 at
// twelve, against a matrix of 973. Measured at the production point it is
// worth 7.6% at one right hand side and costs 7.6% at twelve, so the fold
// is used only for a small batch, and a large one carries its legs in the
// batch dimension instead.
///////////////////////////////////////////////////////////////////////////
static const int FoldNrhsMax = 4;
void ChooseOrganisation(int nrhs)
{
if ( nrhs <= FoldNrhsMax ) {
KFold = AutoKFold();
} else {
KFold = 0;
}
LegGroup = AutoLegGroup();
if ( KFold ) {
PackFoldedMatrix();
} else {
BuildGroupedA();
}
}
///////////////////////
@@ -171,7 +263,7 @@ public:
{
if ( _CoarseGridMulti == nullptr ) {
std::cout << GridLogError
<< "MultiGeneralCoarsenedOperatorV2: the multiRHS grid has not been set."
<< "MultiGeneralCoarsenedOperator: the multiRHS grid has not been set."
<< std::endl;
std::cout << GridLogError
<< " Call SetGrid(CoarseGridMulti) with the D+1 dimensional grid your"
@@ -184,33 +276,58 @@ public:
}
//////////////////////////////////////////////////////////////////////////
// Bilingual accessors, matching GeneralCoarsenedMatrix. Note Grid() is the
// D+1 multiRHS grid here, so a consumer wanting the space the elements live
// on must ask for CoarseGridD().
// Accessors. Note Grid() is the D+1 multiRHS grid, so a consumer wanting
// the space the matrix elements live on must ask for CoarseGridD().
//////////////////////////////////////////////////////////////////////////
// Resident device memory this operator holds (matrix elements and the
// Nrhs-dependent BLAS buffers). Not evictable.
uint64_t DeviceBytes(void)
{
uint64_t b = (uint64_t)(BLAS_B.capacity()+BLAS_C.capacity()+BLAS_Cg.capacity())*sizeof(calcVector);
for(int p=0;p<(int)BLAS_A.size();p++) b += (uint64_t)BLAS_A[p].capacity()*sizeof(calcMatrix);
uint64_t b = (uint64_t)(BLAS_B.capacity()+BLAS_C.capacity()+BLAS_Cg.capacity())*sizeof(calcVector)
+ (uint64_t)BLAS_BK.capacity()*sizeof(CoarseBLASScalar);
b += (uint64_t)(BLAS_Amat.capacity()+BLAS_Ascr.capacity())*sizeof(calcMatrix);
return b;
}
NonLocalStencilGeometry & Geometry(void) { return geom_srhs; };
void ExtractMatrix(int p,CoarseMatrix &A) { BLAStoGrid(A,BLAS_A[p]); };
// One stencil point in or out of the site-major buffer.
void MatrixPointIn(int p,const deviceVector<calcMatrix> &src)
{
int32_t sites = _CoarseGrid->lSites();
const int np = geom.npoint;
calcMatrix *dst = &BLAS_Amat[0];
const calcMatrix *sp = &src[0];
accelerator_for(ss,sites,1,{ dst[(uint64_t)ss*np+p] = sp[ss]; });
}
void MatrixPointOut(int p,deviceVector<calcMatrix> &dst)
{
int32_t sites = _CoarseGrid->lSites();
const int np = geom.npoint;
const calcMatrix *sp = &BLAS_Amat[0];
dst.resize(sites);
calcMatrix *dp = &dst[0];
accelerator_for(ss,sites,1,{ dp[ss] = sp[(uint64_t)ss*np+p]; });
}
void ExtractMatrix(int p,CoarseMatrix &A)
{
MatrixPointOut(p,BLAS_Ascr);
BLAStoGrid(A,BLAS_Ascr);
}
// I/O on the operator matrices, via the BLAS layout array. The parameter is
// a vector over the geometry points; the body indexes A[p].
void SetMatrix (int p,std::vector<CoarseMatrix> & A)
{
GRID_ASSERT(A.size()==geom_srhs.npoint);
GridtoBLAS(A[p],BLAS_A[p]);
GridtoBLAS(A[p],BLAS_Ascr);
MatrixPointIn(p,BLAS_Ascr);
}
void GetMatrix (int p,std::vector<CoarseMatrix> & A)
{
GRID_ASSERT(A.size()==geom_srhs.npoint);
BLAStoGrid(A[p],BLAS_A[p]);
MatrixPointOut(p,BLAS_Ascr);
BLAStoGrid(A[p],BLAS_Ascr);
}
///////////////////////////////////////////////////////////////////////////
@@ -218,7 +335,7 @@ public:
// is independent of Nrhs, in particular the matrix elements, which must
// survive a change of Nrhs untouched.
///////////////////////////////////////////////////////////////////////////
MultiGeneralCoarsenedOperatorV2(NonLocalStencilGeometry &_geom,GridCartesian *CoarseGrid) :
MultiGeneralCoarsenedOperator(NonLocalStencilGeometry &_geom,GridCartesian *CoarseGrid) :
_CoarseGrid(CoarseGrid),
geom_srhs(_geom),
geom(CoarseGrid,_geom.hops,_geom.skip),
@@ -233,28 +350,26 @@ public:
/////////////////////////////////////////////////
// Matrix elements and their pointer table
/////////////////////////////////////////////////
BLAS_A.resize(geom.npoint);
BLAS_Amat.resize((uint64_t)unpadded_sites*geom.npoint);
BLAS_Ascr.resize(unpadded_sites);
BLAS_AP.resize(geom.npoint);
for(int p=0;p<geom.npoint;p++){
BLAS_A[p].resize (unpadded_sites); // no ghost zone, npoint elements
BLAS_AP[p].resize(unpadded_sites);
}
// Site identity mapping for A
// Point p at site ss: the unfolded path strides by the whole stencil.
for(int p=0;p<geom.npoint;p++){
for(int ss=0;ss<unpadded_sites;ss++){
CoarseBLASScalar *ptr = (CoarseBLASScalar *)&BLAS_A[p][ss];
CoarseBLASScalar *ptr = (CoarseBLASScalar *)&BLAS_Amat[(uint64_t)ss*geom.npoint+p];
acceleratorPut(BLAS_AP[p][ss],ptr);
}
}
// Enough legs to bring the batch up to where the kernel saturates, and
// never more legs than the stencil has.
LegGroup = (TargetBatch + unpadded_sites - 1)/unpadded_sites;
if ( LegGroup < 1 ) LegGroup = 1;
if ( LegGroup > geom.npoint ) LegGroup = geom.npoint;
// The organisation is chosen in SetGrid, where Nrhs is known.
KFold = 0;
LegGroup = AutoLegGroup();
BuildGroupedA();
std::cout << GridLogMessage << "MultiGeneralCoarsenedOperatorV2: stencil "
std::cout << GridLogMessage << "MultiGeneralCoarsenedOperator: stencil "
<< geom.npoint << " points, local coarse volume " << unpadded_sites
<< ", legs per GEMM " << LegGroup
<< " (batch " << LegGroup*unpadded_sites << ", "
@@ -263,7 +378,7 @@ public:
<< ")" << std::endl;
}
virtual ~MultiGeneralCoarsenedOperatorV2()
virtual ~MultiGeneralCoarsenedOperator()
{
ReleaseGrid();
}
@@ -285,6 +400,10 @@ public:
BLAS_Cg.resize(0);
BLAS_CPg.resize(0);
BLAS_CPr.resize(0);
BLAS_BK.resize(0);
BLAS_BPflat.resize(0);
for(int g=0;g<(int)BLAS_BKP.size();g++){ BLAS_BKP[g].resize(0); }
BLAS_BKP.resize(0);
for(int g=0;g<(int)BLAS_BPg.size();g++){ BLAS_BPg[g].resize(0); }
BLAS_BPg.resize(0);
for(int p=0;p<BLAS_BP.size();p++){
@@ -342,6 +461,19 @@ public:
GRID_ASSERT(CellMulti->grids.back()->lSites() == nrhs*padded_sites);
GRID_ASSERT(_CoarseGridMulti->lSites() == nrhs*unpadded_sites);
/////////////////////////////////////////////////
// Choose the organisation FIRST: every table below is sized from
// LegGroup or KFold, and is indexed with them in the apply, so the two
// must be decided before anything is allocated.
/////////////////////////////////////////////////
ChooseOrganisation(nrhs);
std::cout << GridLogMessage << "MultiGeneralCoarsenedOperator: Nrhs " << nrhs
<< ", coarse apply "
<< ( KFold ? "folds "+std::to_string(KFold)+" legs along K (K = "
+std::to_string(KFold*nbasis)+", "+std::to_string(KChunks())+" chunk(s))"
: "carries "+std::to_string(LegGroup)+" legs in the batch" )
<< std::endl;
/////////////////////////////////////////////////
// Device data vector storage
/////////////////////////////////////////////////
@@ -382,6 +514,24 @@ public:
BLAS_BPg[g].resize(legs*unpadded_sites);
}
// Folded-along-K working set: one chunk of gathered neighbour vectors,
// laid out per site as a column-major (legs*nbasis) x nrhs block, plus a
// pointer table per chunk because a short last chunk has a shorter site
// stride. BLAS_BPflat is the neighbour table the gather kernel reads.
if ( KFold ) {
BLAS_BK.resize((uint64_t)unpadded_sites*nrhs*KFold*nbasis);
BLAS_BKP.resize(KChunks());
for(int g=0;g<KChunks();g++){
int legs = KLegs(g);
BLAS_BKP[g].resize(unpadded_sites);
for(int32_t ss=0;ss<unpadded_sites;ss++){
CoarseBLASScalar *ptr = &BLAS_BK[(uint64_t)ss*nrhs*legs*nbasis];
acceleratorPut(BLAS_BKP[g][ss],ptr);
}
}
BLAS_BPflat.resize((uint64_t)geom.npoint*unpadded_sites);
}
// Neighbour table is more complicated
int32_t j=0; // Interior point counter (unpadded)
for(int32_t s=0;s<padded_sites;s++){ // D volume, padded
@@ -402,6 +552,7 @@ public:
CoarseBLASScalar * ptr = (CoarseBLASScalar *)&BLAS_B[nbr];
acceleratorPut(BLAS_BP[point][j],ptr); // neighbour indexing in ghost zone volume
acceleratorPut(BLAS_BPg[point/LegGroup][(point%LegGroup)*unpadded_sites+j],ptr);
if ( KFold ) acceleratorPut(BLAS_BPflat[(uint64_t)point*unpadded_sites+j],ptr);
}
j++;
}
@@ -707,7 +858,7 @@ public:
// orthonormalised basis into it, and the caller keeps it for the solve.
// The overload below makes its own and is kept for the pre-2026 drivers;
// it costs a second full basis store, which is why nothing new should use
// it (documentation/MultiGridMemoryAudit.md).
// it.
void CoarsenOperator(LinearOperatorBase<Lattice<Fobj> > &linop,
GridCartesian *FineGridMulti,
std::vector<FineField> &Subspace,
@@ -815,8 +966,10 @@ public:
}
for(int p=0;p<npoint;p++){
GridtoBLAS(_A[p],BLAS_A[p]);
GridtoBLAS(_A[p],BLAS_Ascr);
MatrixPointIn(p,BLAS_Ascr);
}
if ( KFold ) PackFoldedMatrix(); // the matrix is only valid from here
CoarsenReport(tphase,tphaseBZ,tslice,tmat,tproj,tinv);
}
@@ -920,8 +1073,10 @@ public:
}
for(int p=0;p<npoint;p++){
GridtoBLAS(_A[p],BLAS_A[p]);
GridtoBLAS(_A[p],BLAS_Ascr);
MatrixPointIn(p,BLAS_Ascr);
}
if ( KFold ) PackFoldedMatrix(); // the matrix is only valid from here
CoarsenReport(tphase,tphaseBZ,tslice,tmat,tproj,tinv);
}
void Mdag(const CoarseVector &in, CoarseVector &out)
@@ -944,7 +1099,7 @@ public:
CheckGridSet();
if ( in.Grid() != _CoarseGridMulti ) {
std::cout << GridLogError
<< "MultiGeneralCoarsenedOperatorV2::M called with a field on a"
<< "MultiGeneralCoarsenedOperator::M called with a field on a"
<< std::endl;
std::cout << GridLogError
<< " different grid object from the one given to SetGrid(). Two"
@@ -955,14 +1110,14 @@ public:
GRID_ASSERT(in.Grid() == _CoarseGridMulti);
}
GRID_TRACE("CoarseV2Mult");
GRID_TRACE("CoarseMult");
t_tot=-usecond();
t_exch=-usecond();
// The exchange takes its input by const reference, so it reads the
// caller's field directly. lambda scope so the roctx range covers
// exactly the exchange; the PaddedCellFwd/BwdMPI markers inside it then
// nest properly.
CoarseVector pin = [&](){ GRID_TRACE("CoarseV2Exchange");
CoarseVector pin = [&](){ GRID_TRACE("CoarseExchange");
return CellMulti->ExchangePeriodic(in); }(); //padded input
t_exch+=usecond();
@@ -985,7 +1140,7 @@ public:
t_GtoB=-usecond();
{ GRID_TRACE("CoarseV2GridToBLAS");
{ GRID_TRACE("CoarseGridToBLAS");
GridtoBLAS(pin,BLAS_B);
}
t_GtoB+=usecond();
@@ -993,10 +1148,46 @@ public:
GridBLAS BLAS;
t_mult=-usecond();
{ GRID_TRACE("CoarseV2StencilGEMM");
{ GRID_TRACE("CoarseStencilGEMM");
// The scalar type selects the GEMM: Cgemm for an fp32 coarse space,
// Zgemm for fp64.
if ( LegGroup == 1 ) {
if ( KFold ) {
// Gather one chunk of neighbour vectors per site, then one GEMM whose
// contraction runs over the chunk's legs, accumulating into C. No
// partial results, so nothing to combine afterwards.
const int nb = nbasis;
const int64_t ns = (int64_t)_CoarseGrid->lSites(); // D dimensional: CoarseGrid() is D+1
GRID_ASSERT( (int64_t)BLAS_BPflat.size() == (int64_t)geom.npoint*ns );
CoarseBLASScalar * const bk = &BLAS_BK[0];
CoarseBLASScalar * const * const bp = &BLAS_BPflat[0];
for(int g=0;g<KChunks();g++){
const int64_t legs = KLegs(g);
const int64_t leg0 = (int64_t)g*KFold;
{ GRID_TRACE("CoarseLegGather");
accelerator_for(idx, ns*legs*nrhs, 1, {
int64_t j = idx % nrhs; // right hand side (column)
int64_t t = idx / nrhs;
int64_t l = t % legs; // leg within the chunk
int64_t ss = t / legs; // site
const CoarseBLASScalar *src = bp[(leg0+l)*ns + ss] + j*nb;
CoarseBLASScalar *dst = bk + ((ss*nrhs + j)*legs + l)*nb;
for(int i=0;i<nb;i++) dst[i] = src[i];
});
}
RealD c = (g==0) ? 0.0 : 1.0;
BLAS.gemmBatched(nbasis,nrhs,legs*nbasis,
CoarseBLASScalar(1.0),
BLAS_AKP[g],
BLAS_BKP[g],
CoarseBLASScalar(c),
BLAS_CP);
// One gather buffer serves every chunk and the multiply is issued
// asynchronously, so the next chunk's gather must not start until this
// multiply has read it. One synchronise per chunk, none when the whole
// stencil folds into a single call.
if ( g+1 < KChunks() ) BLAS.synchronise();
}
} else if ( LegGroup == 1 ) {
for(int p=0;p<geom.npoint;p++){
RealD c = (p==0) ? 0.0 : 1.0;
BLAS.gemmBatched(nbasis,nrhs,nbasis,
@@ -1037,7 +1228,7 @@ public:
}
t_mult+=usecond();
if ( LegGroup > 1 ) { GRID_TRACE("CoarseV2LegSum");
if ( !KFold && LegGroup > 1 ) { GRID_TRACE("CoarseLegSum");
// Sum the LegGroup partial results into the single output buffer.
// Flat over scalars: a calcVector is nbasis of them and the partials
// are contiguous, one whole C volume after another.
@@ -1053,7 +1244,7 @@ public:
}
t_BtoG=-usecond();
{ GRID_TRACE("CoarseV2BLASToGrid");
{ GRID_TRACE("CoarseBLASToGrid");
BLAStoGrid(out,BLAS_C);
}
t_BtoG+=usecond();
+2 -2
View File
@@ -34,7 +34,7 @@ NAMESPACE_BEGIN(Grid);
// HDCGCoarsening everything that is a function of the gauge
// field: the raw near-null basis (and its
// refinement), the transfer operators, the
// Galerkin coarse operator (V2, Hermitian),
// Galerkin coarse operator (Hermitian),
// the coarse eigenvectors for deflation
// HDCGSolver the solve chain on a borrowed coarsening:
// smoother, deflated coarse solve, ADEF-2
@@ -51,7 +51,7 @@ template<class Fobj,class CComplex,int nbasis>
class HDCGCoarsening {
public:
typedef Lattice<Fobj> FineField;
typedef MultiGeneralCoarsenedOperatorV2<Fobj,CComplex,nbasis> CoarseOperator;
typedef MultiGeneralCoarsenedOperator<Fobj,CComplex,nbasis> CoarseOperator;
typedef typename CoarseOperator::CoarseVector CoarseVector;
typedef typename GridTypeMapper<Fobj>::SinglePrecision FobjF;
typedef Lattice<FobjF> FineFieldF;
+181 -49
View File
@@ -50,23 +50,74 @@ public:
LinearOperatorBase<Field> &Linop;
MrhsLinearFunction<Field> &Preconditioner;
std::function<void(int)> OnStep; // called with the outer step count after every step
void Level(int lv){ name = "Level " + std::to_string(lv); level=lv; }
void Name(std::string n){ name = n; trace_op = name+" MrhsPGCR::vOp"; trace_orthog = name+" MrhsPGCR orthog"; }
void SetZeroGuess(int z){ ZeroGuess=z; }
void Level(int lv){
name = "Level " + std::to_string(lv);
level=lv;
}
void Name(std::string n){
name = n;
trace_op = name+" MrhsPGCR::vOp";
trace_orthog = name+" MrhsPGCR orthog";
}
void SetZeroGuess(int z)
{
ZeroGuess=z;
}
MrhsPGCRNonHermitian(RealD tol,Integer maxit,LinearOperatorBase<Field> &_Linop,MrhsLinearFunction<Field> &Prec,int _mmax,int _nstep)
: Tolerance(tol),MaxIterations(maxit),Linop(_Linop),Preconditioner(Prec),mmax(_mmax),nstep(_nstep){ level=1; }
static RealD vnorm2(std::vector<Field> &x){ RealD s=0; for(auto &f:x) s+=norm2(f); return s; }
static ComplexD vinnerProduct(std::vector<Field> &x,std::vector<Field> &y){ ComplexD s(0); for(int r=0;r<(int)x.size();r++) s+=innerProduct(x[r],y[r]); return s; }
static void vaxpy(std::vector<Field> &z,ComplexD a,std::vector<Field> &x,std::vector<Field> &y){ for(int r=0;r<(int)z.size();r++) axpy(z[r],a,x[r],y[r]); }
void vOp(std::vector<Field> &in,std::vector<Field> &out){ GRID_TRACE(trace_op.c_str()); for(int r=0;r<(int)in.size();r++) Linop.Op(in[r],out[r]); }
: Tolerance(tol),MaxIterations(maxit),Linop(_Linop),Preconditioner(Prec),mmax(_mmax),nstep(_nstep)
{
level=1;
}
static RealD vnorm2(std::vector<Field> &x)
{
RealD s=0;
for(auto &f:x) s+=norm2(f);
return s;
}
static ComplexD vinnerProduct(std::vector<Field> &x,std::vector<Field> &y){
ComplexD s(0);
for(int r=0;r<(int)x.size();r++) {
s+=innerProduct(x[r],y[r]);
}
return s;
}
static void vaxpy(std::vector<Field> &z,ComplexD a,std::vector<Field> &x,std::vector<Field> &y)
{
for(int r=0;r<(int)z.size();r++) {
axpy(z[r],a,x[r],y[r]);
}
}
void vOp(std::vector<Field> &in,std::vector<Field> &out)
{
GRID_TRACE(trace_op.c_str());
for(int r=0;r<(int)in.size();r++) {
Linop.Op(in[r],out[r]);
}
}
void operator()(std::vector<Field> &src,std::vector<Field> &psi){
RealD cp,ssq,rsq; int nrhs=src.size(); GridBase *grid=src[0].Grid();
ssq=vnorm2(src); rsq=Tolerance*Tolerance*ssq;
RealD cp,ssq,rsq;
int nrhs=src.size();
GridBase *grid=src[0].Grid();
ssq=vnorm2(src);
rsq=Tolerance*Tolerance*ssq;
std::vector<Field> r(nrhs,grid);
GridStopWatch T; T.Start(); steps=0; FirstCycle=1;
GridStopWatch T;
T.Start();
steps=0;
FirstCycle=1;
for(int k=0;k<MaxIterations;k++){
cp=GCRnStep(src,psi,rsq);
std::cout<<GridLogMessage<<std::string(level,'\t')<<" "<<name<<" MrhsPGCR("<<mmax<<","<<nstep<<") "<<steps<<" steps cp = "<<cp<<" target "<<rsq<<std::endl;
std::cout<<GridLogMessage<<std::string(level,'\t')<<" "<<name<<" MrhsPGCR("<<mmax<<","<<nstep<<") "<<steps
<<" steps cp = "<<cp<<" target "<<rsq<<std::endl;
if(cp<rsq){
T.Stop(); vOp(psi,r); for(int rr=0;rr<nrhs;rr++) axpy(r[rr],-1.0,src[rr],r[rr]);
RealD tr=vnorm2(r);
@@ -78,28 +129,61 @@ public:
}
std::cout<<GridLogMessage<<"MrhsPGCR: did not converge"<<std::endl;
}
RealD GCRnStep(std::vector<Field> &src,std::vector<Field> &psi,RealD rsq){
RealD cp; ComplexD a,rq; int nrhs=src.size(); GridBase *grid=src[0].Grid();
RealD GCRnStep(std::vector<Field> &src,std::vector<Field> &psi,RealD rsq)
{
RealD cp;
ComplexD a,rq;
int nrhs=src.size();
GridBase *grid=src[0].Grid();
std::vector<Field> r(nrhs,grid),Az(nrhs,grid); // Az: restart residual scratch only
std::vector< std::vector<Field> > q(mmax,std::vector<Field>(nrhs,grid));
std::vector< std::vector<Field> > p(mmax,std::vector<Field>(nrhs,grid));
std::vector<RealD> qq(mmax);
if (ZeroGuess && FirstCycle) { for(int rr=0;rr<nrhs;rr++){ psi[rr]=Zero(); r[rr]=src[rr]; } }
else { vOp(psi,Az); for(int rr=0;rr<nrhs;rr++) r[rr]=src[rr]-Az[rr]; }
if (ZeroGuess && FirstCycle) {
for(int rr=0;rr<nrhs;rr++){
psi[rr]=Zero();
r[rr]=src[rr];
}
} else {
vOp(psi,Az);
for(int rr=0;rr<nrhs;rr++) {
r[rr]=src[rr]-Az[rr];
}
}
FirstCycle=0;
// p[0]=Prec(r), q[0]=A p[0], produced directly in the history slots (no copies)
Preconditioner(r,p[0]); vOp(p[0],q[0]); qq[0]=vnorm2(q[0]); cp=vnorm2(r);
Preconditioner(r,p[0]);
vOp(p[0],q[0]);
qq[0]=vnorm2(q[0]);
cp=vnorm2(r);
for(int k=0;k<nstep;k++){
steps++; int kp=k+1, peri_k=k%mmax, peri_kp=kp%mmax;
steps++;
int kp=k+1, peri_k=k%mmax, peri_kp=kp%mmax;
if ( OnStep ) OnStep(steps);
rq=vinnerProduct(q[peri_k],r); a=rq/qq[peri_k];
vaxpy(psi,a,p[peri_k],psi); vaxpy(r,-a,q[peri_k],r); cp=vnorm2(r);
rq=vinnerProduct(q[peri_k],r);
a=rq/qq[peri_k];
vaxpy(psi,a,p[peri_k],psi);
vaxpy(r,-a,q[peri_k],r);
cp=vnorm2(r);
std::cout<<GridLogMessage<<std::string(level,'\t')<<" "<<name<<" MrhsPGCR step["<<steps<<"] resid "<<cp<<" target "<<rsq<<std::endl;
if((k==nstep-1)||(cp<rsq)) return cp;
// New direction straight into its history slot: p=Prec(r), q=A p.
Preconditioner(r,p[peri_kp]);
vOp(p[peri_kp],q[peri_kp]);
int northog=((kp)>(mmax-1))?(mmax-1):(kp);
{
GRID_TRACE(trace_orthog.c_str());
// Classical Gram-Schmidt: all coefficients against the UN-updated new q
@@ -110,15 +194,30 @@ public:
std::vector<ComplexD> bcoef(northog,ComplexD(0.0)), part;
for(int rr=0;rr<nrhs;rr++){
std::vector<const Field*> qwin(northog);
for(int back=0;back<northog;back++){ int peri_back=(k-back)%mmax; GRID_ASSERT((k-back)>=0); qwin[back]=&q[peri_back][rr]; }
for(int back=0;back<northog;back++){
int peri_back=(k-back)%mmax;
GRID_ASSERT((k-back)>=0);
qwin[back]=&q[peri_back][rr];
}
rankInnerProductMulti(part,qwin,q[peri_kp][rr]);
for(int back=0;back<northog;back++) bcoef[back]+=part[back];
}
if(northog) grid->GlobalSumVector(&bcoef[0],northog);
for(int back=0;back<northog;back++){ int peri_back=(k-back)%mmax; bcoef[back]=-bcoef[back]/qq[peri_back]; }
for(int back=0;back<northog;back++){
int peri_back=(k-back)%mmax;
bcoef[back]=-bcoef[back]/qq[peri_back];
}
for(int rr=0;rr<nrhs;rr++){
std::vector<const Field*> qwin(northog), pwin(northog);
for(int back=0;back<northog;back++){ int peri_back=(k-back)%mmax; qwin[back]=&q[peri_back][rr]; pwin[back]=&p[peri_back][rr]; }
for(int back=0;back<northog;back++){
int peri_back=(k-back)%mmax;
qwin[back]=&q[peri_back][rr];
pwin[back]=&p[peri_back][rr];
}
axpyMulti(p[peri_kp][rr],bcoef,pwin);
axpyMulti(q[peri_kp][rr],bcoef,qwin);
}
@@ -168,24 +267,42 @@ public:
_Coarse5d(Coarse5d), _CoarseCoarse5d(CoarseCoarse5d), _CoarseCoarseMrhs(CoarseCoarseMrhs), _nrhs(nrhs) {}
using LinearFunction<CoarseField>::operator();
virtual void operator()(const CoarseField &in, CoarseField &out) {
GRID_TRACE("MGCoarseVcycle");
CoarseField vec1(in.Grid());
CoarseField vec2(in.Grid());
out = in;
_CoarseOp.Op(out,vec1); sub(vec1,in,vec1);
// restrict, through the mixed blockProject: D+1 coarse in, D+1 cc out
CoarseCoarseField CCsrc(_CoarseCoarseMrhs);
CoarseCoarseField CCsol(_CoarseCoarseMrhs);
_Projector.blockProject(vec1,CCsrc);
_CoarseCoarseSolve(CCsrc,CCsol);
_Projector.blockPromote(vec1,CCsol);
add(out,out,vec1);
_CoarseOp.Op(out,vec1); sub(vec1,in,vec1);
_CoarseSmoother(vec1,vec2);
add(out,out,vec2);
// The cycle starts from x0 = in, one unit-step Richardson iteration ahead of
// a zero start; measured, it pays for its apply by cutting Couter steps 22%.
// x0 is never materialised in out: the operator is applied to in directly and
// the copy folds into the add that lands the coarse-coarse correction.
{
GRID_TRACE("MGCoarseResidual");
_CoarseOp.Op(in,vec1); sub(vec1,in,vec1);
}
// restrict, through the mixed blockProject: D+1 coarse in, D+1 cc out
{
GRID_TRACE("MGCCProject");
_Projector.blockProject(vec1,CCsrc);
}
{
GRID_TRACE("MGCCSolve");
_CoarseCoarseSolve(CCsrc,CCsol);
}
{
GRID_TRACE("MGCCPromote");
_Projector.blockPromote(vec1,CCsol);
add(out,in,vec1);
}
{
GRID_TRACE("MGCoarseResidual2");
_CoarseOp.Op(out,vec1); sub(vec1,in,vec1);
}
{
GRID_TRACE("MGCoarseSmooth");
_CoarseSmoother(vec1,vec2);
add(out,out,vec2);
}
}
};
@@ -211,37 +328,47 @@ public:
GridBase *_CoarseGrid, *_CoarseGridMrhs;
std::function<void(int)> SetSloppy = [](int){};
int SloppyComms = 0; // value passed to SetSloppy on entry
MrhsTwoLevelMG(LinearOperatorBase<FineField> &FineOp, FineSmoother &Post,
Projector_t &Projector, LinearFunction<CoarseVector> &CoarseSolve,
GridBase *CoarseGrid, GridBase *CoarseGridMrhs)
: _FineOperator(FineOp),_PostSmoother(Post),_Projector(Projector),_CoarseSolve(CoarseSolve),
_CoarseGrid(CoarseGrid),_CoarseGridMrhs(CoarseGridMrhs){}
virtual void operator()(std::vector<FineField> &in, std::vector<FineField> &out){
virtual void operator()(std::vector<FineField> &in, std::vector<FineField> &out)
{
GRID_TRACE("MGVcycle");
SetSloppy(SloppyComms);
int nrhs=in.size(); GridBase *fgrid=in[0].Grid();
std::vector<FineField> vec1(nrhs,fgrid),vec2(nrhs,fgrid);
for(int r=0;r<nrhs;r++) out[r]=in[r];
{ GRID_TRACE("MGFineResidual");
for(int r=0;r<nrhs;r++){ _FineOperator.Op(out[r],vec1[r]); sub(vec1[r],in[r],vec1[r]); }
// x0 = in, never materialised in out: the operator is applied to in directly
// and the copy folds into the add that lands the coarse correction.
{
GRID_TRACE("MGFineResidual");
for(int r=0;r<nrhs;r++){ _FineOperator.Op(in[r],vec1[r]); sub(vec1[r],in[r],vec1[r]); }
}
// fine vector -> D+1 coarse, via the mixed blockProject
CoarseVector CsrcMrhs(_CoarseGridMrhs), CsolMrhs(_CoarseGridMrhs);
{ GRID_TRACE("MGProject");
{
GRID_TRACE("MGProject");
_Projector.blockProject(vec1,CsrcMrhs);
}
CsolMrhs=Zero();
{ GRID_TRACE("MGCoarseSolve");
{
GRID_TRACE("MGCoarseSolve");
_CoarseSolve(CsrcMrhs,CsolMrhs);
}
{ GRID_TRACE("MGPromote");
{
GRID_TRACE("MGPromote");
_Projector.blockPromote(vec1,CsolMrhs);
for(int r=0;r<nrhs;r++) add(out[r],out[r],vec1[r]);
for(int r=0;r<nrhs;r++) add(out[r],in[r],vec1[r]);
}
{ GRID_TRACE("MGFineResidual2");
{
GRID_TRACE("MGFineResidual2");
for(int r=0;r<nrhs;r++){ _FineOperator.Op(out[r],vec1[r]); sub(vec1[r],in[r],vec1[r]); }
}
{ GRID_TRACE("MGPostSmooth");
{
GRID_TRACE("MGPostSmooth");
for(int r=0;r<nrhs;r++){
_PostSmoother(vec1[r],vec2[r]); add(out[r],out[r],vec2[r]);
}
@@ -271,15 +398,19 @@ public:
precisionChangeWorkspace _ws_f2d; // out fp64, in fp32
GridBase *_gridF;
std::vector<FieldF> _in_f, _out_f;
MrhsMixedPrecPreconditioner(MrhsPreconditioner<FieldF> &Inner, GridBase *gridD, GridBase *gridF, int nrhs)
: _Inner(Inner), _ws_d2f(gridF,gridD), _ws_f2d(gridD,gridF), _gridF(gridF) {}
void Scratch(int nrhs){
if ( (int)_in_f.size() == nrhs ) return;
_in_f.clear(); _in_f.reserve(nrhs);
_out_f.clear(); _out_f.reserve(nrhs);
for(int r=0;r<nrhs;r++){ _in_f.emplace_back(_gridF); _out_f.emplace_back(_gridF); }
}
virtual void operator()(std::vector<FieldD> &in, std::vector<FieldD> &out){
virtual void operator()(std::vector<FieldD> &in, std::vector<FieldD> &out)
{
GRID_TRACE("MGPrecisionSeam");
int nrhs=in.size(); Scratch(nrhs);
for(int r=0;r<nrhs;r++) precisionChange(_in_f[r],in[r],_ws_d2f);
@@ -287,7 +418,8 @@ public:
for(int r=0;r<nrhs;r++) precisionChange(out[r],_out_f[r],_ws_f2d);
}
// The start and the timers are the inner preconditioner's
virtual void Vstart(std::vector<FieldD> &x, std::vector<FieldD> &src){
virtual void Vstart(std::vector<FieldD> &x, std::vector<FieldD> &src)
{
GRID_TRACE("MGPrecisionSeamVstart");
int nrhs=src.size(); Scratch(nrhs);
for(int r=0;r<nrhs;r++) precisionChange(_in_f[r],src[r],_ws_d2f);
+6 -6
View File
@@ -30,12 +30,12 @@ Author: Peter Boyle <pboyle@bnl.gov>
#include <Grid/algorithms/multigrid/Aggregates.h>
#include <Grid/algorithms/multigrid/Geometry.h>
#include <Grid/algorithms/multigrid/CoarsenedMatrix.h>
#include <Grid/algorithms/multigrid/GeneralCoarsenedMatrixMultiRHSV2.h>
// DEPRECATED: the V1 coarse operators and the Aggregation-based single-RHS
// ADEF-2. Nothing in the library uses them; the PVdagM and HDCG chains are
// on V2 (PVdagMMultiGrid.h, HDCGMultiGrid.h). Kept so the pre-2026 drivers
// in tests/debug and examples still build; removing this block is the
// deletion gate for them.
#include <Grid/algorithms/multigrid/GeneralCoarsenedMatrixMultiRHS.h>
// DEPRECATED: the superseded coarse operators and the Aggregation-based
// single-RHS ADEF-2. Nothing in the library uses them; the PVdagM and HDCG
// chains are on MultiGeneralCoarsenedOperator (PVdagMMultiGrid.h,
// HDCGMultiGrid.h). Kept so the pre-2026 drivers in tests/debug and
// examples still build; removing this block is the deletion gate for them.
#include <Grid/algorithms/multigrid/deprecated/GeneralCoarsenedMatrix.h>
#include <Grid/algorithms/multigrid/deprecated/GeneralCoarsenedMatrixMultiRHS.h>
#include <Grid/algorithms/multigrid/deprecated/TwoLevelADEF2.h>
+3 -3
View File
@@ -212,14 +212,14 @@ template<class Fobj,class CComplex,int nbasis>
class PVdagMMultiGridCoarsening {
public:
typedef Lattice<Fobj> FineField;
typedef MultiGeneralCoarsenedOperatorV2<Fobj,CComplex,nbasis> CoarseOperator;
typedef MultiGeneralCoarsenedOperator<Fobj,CComplex,nbasis> CoarseOperator;
typedef typename CoarseOperator::CoarseVector CoarseVector;
typedef typename CoarseVector::vector_object CoarseSiteObj;
// Every coarse level carries the same site type iVector<CComplex,nbasis>:
// the coefficient scalar does not deepen with the level (the operator keeps
// its own deeper scratch type for the block inner products). Levels are
// told apart by their grids, not their C++ types.
typedef MultiGeneralCoarsenedOperatorV2<CoarseSiteObj,CComplex,nbasis> CoarseCoarseOperator;
typedef MultiGeneralCoarsenedOperator<CoarseSiteObj,CComplex,nbasis> CoarseCoarseOperator;
typedef typename CoarseCoarseOperator::CoarseVector CoarseCoarseVector;
typedef DenseCoarseMatrix<CComplex,nbasis> DenseBottom;
// The fp32 fine field, derived from the fp64 one: the fp32 fine level of
@@ -359,7 +359,7 @@ public:
rawPsi.reserve(nbasis);
for(int k=0;k<nbasis;k++) rawPsi.push_back(psi[k]);
// L2: the V2 L1 operator is natively multiRHS at the batch grid
// L2: the L1 operator is natively multiRHS at the batch grid
CoarseOpL2.SetGrid(Grids.CoarseCoarseBatch);
NonHermitianLinearOperator<CoarseOperator,CoarseVector> LinOpCoarse(CoarseOpPV);
std::cout << GridLogMessage << "PVdagMMultiGridCoarsening: L2 CoarsenOperator, batch "
+117 -26
View File
@@ -36,17 +36,54 @@ NAMESPACE_BEGIN(Grid);
//////////////////////////////////////////////////////////////////////
template<class Matrix,class Field>
class PVdagMLinearOperator : public LinearOperatorBase<Field> {
Matrix &_Mat; Matrix &_PV;
Matrix &_Mat;
Matrix &_PV;
public:
PVdagMLinearOperator(Matrix &Mat,Matrix &PV): _Mat(Mat),_PV(PV) {};
void SloppyComms(int sloppy) { _Mat.SloppyComms(sloppy); _PV.SloppyComms(sloppy); }
void OpDiag (const Field &in, Field &out) { GRID_ASSERT(0); }
void OpDir (const Field &in, Field &out,int dir,int disp) { GRID_ASSERT(0); }
void OpDirAll (const Field &in, std::vector<Field> &out){ GRID_ASSERT(0); };
void Op (const Field &in, Field &out){ Field tmp(in.Grid()); _Mat.M(in,tmp); _PV.Mdag(tmp,out); }
void AdjOp (const Field &in, Field &out){ Field tmp(in.Grid()); _PV.M(in,tmp); _Mat.Mdag(tmp,out); }
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2){ HermOp(in,out); ComplexD d=innerProduct(in,out); n1=real(d); n2=norm2(out); }
void HermOp(const Field &in, Field &out){ Field tmp(in.Grid()); Op(in,tmp); AdjOp(tmp,out); }
PVdagMLinearOperator(Matrix &Mat,Matrix &PV):
_Mat(Mat),_PV(PV)
{};
void SloppyComms(int sloppy)
{
_Mat.SloppyComms(sloppy);
_PV.SloppyComms(sloppy);
}
void OpDiag (const Field &in, Field &out)
{
GRID_ASSERT(0);
}
void OpDir (const Field &in, Field &out,int dir,int disp)
{
GRID_ASSERT(0);
}
void OpDirAll (const Field &in, std::vector<Field> &out)
{
GRID_ASSERT(0);
};
void Op (const Field &in, Field &out)
{
Field tmp(in.Grid());
_Mat.M(in,tmp);
_PV.Mdag(tmp,out);
}
void AdjOp (const Field &in, Field &out)
{
Field tmp(in.Grid());
_PV.M(in,tmp);
_Mat.Mdag(tmp,out);
}
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2)
{
HermOp(in,out);
ComplexD d=innerProduct(in,out);
n1=real(d);
n2=norm2(out);
}
void HermOp(const Field &in, Field &out)
{
Field tmp(in.Grid());
Op(in,tmp);
AdjOp(tmp,out);
}
};
//////////////////////////////////////////////////////////////////////
@@ -55,17 +92,46 @@ public:
//////////////////////////////////////////////////////////////////////
template<class Matrix,class Field>
class ShiftedPVdagMLinearOperator : public LinearOperatorBase<Field> {
Matrix &_Mat; Matrix &_PV;
Matrix &_Mat;
Matrix &_PV;
public:
RealD shift;
ShiftedPVdagMLinearOperator(RealD _shift,Matrix &Mat,Matrix &PV): shift(_shift),_Mat(Mat),_PV(PV){};
void OpDiag (const Field &in, Field &out) { GRID_ASSERT(0); }
void OpDir (const Field &in, Field &out,int dir,int disp) { GRID_ASSERT(0); }
void OpDirAll (const Field &in, std::vector<Field> &out){ GRID_ASSERT(0); };
void Op (const Field &in, Field &out){ Field tmp(in.Grid()); _Mat.M(in,tmp); _PV.Mdag(tmp,out); out = out + shift*in; }
void AdjOp (const Field &in, Field &out){ Field tmp(in.Grid()); _PV.M(in,tmp); _Mat.Mdag(tmp,out); out = out + shift*in; }
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2){ GRID_ASSERT(0); }
void HermOp(const Field &in, Field &out){ Field tmp(in.Grid()); Op(in,tmp); AdjOp(tmp,out); }
void OpDiag (const Field &in, Field &out)
{
GRID_ASSERT(0);
}
void OpDir (const Field &in, Field &out,int dir,int disp)
{
GRID_ASSERT(0);
}
void OpDirAll (const Field &in, std::vector<Field> &out)
{
GRID_ASSERT(0);
};
void Op (const Field &in, Field &out)
{
Field tmp(in.Grid());
_Mat.M(in,tmp);
_PV.Mdag(tmp,out);
out = out + shift*in;
}
void AdjOp (const Field &in, Field &out)
{
Field tmp(in.Grid());
_PV.M(in,tmp);
_Mat.Mdag(tmp,out);
out = out + shift*in;
}
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2)
{
GRID_ASSERT(0);
}
void HermOp(const Field &in, Field &out){
Field tmp(in.Grid());
Op(in,tmp);
AdjOp(tmp,out);
}
};
//////////////////////////////////////////////////////////////////////
@@ -73,16 +139,41 @@ public:
//////////////////////////////////////////////////////////////////////
template<class Field>
class ShiftedLinearOperator : public LinearOperatorBase<Field> {
LinearOperatorBase<Field> &_Op; RealD shift;
LinearOperatorBase<Field> &_Op;
RealD shift;
public:
ShiftedLinearOperator(RealD _shift, LinearOperatorBase<Field> &Op) : _Op(Op), shift(_shift) {}
void OpDiag (const Field &in, Field &out) { GRID_ASSERT(0); }
void OpDir (const Field &in, Field &out,int dir,int disp) { GRID_ASSERT(0); }
void OpDirAll (const Field &in, std::vector<Field> &out) { GRID_ASSERT(0); }
void Op (const Field &in, Field &out) { _Op.Op(in,out); out = out + shift*in; }
void AdjOp (const Field &in, Field &out) { _Op.AdjOp(in,out); out = out + shift*in; }
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2){ GRID_ASSERT(0); }
void HermOp (const Field &in, Field &out) { Field tmp(in.Grid()); Op(in,tmp); AdjOp(tmp,out); }
void OpDiag (const Field &in, Field &out)
{
GRID_ASSERT(0);
}
void OpDir (const Field &in, Field &out,int dir,int disp)
{
GRID_ASSERT(0);
}
void OpDirAll (const Field &in, std::vector<Field> &out)
{
GRID_ASSERT(0);
}
void Op (const Field &in, Field &out)
{
_Op.Op(in,out);
out = out + shift*in;
}
void AdjOp (const Field &in, Field &out) {
_Op.AdjOp(in,out);
out = out + shift*in;
}
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2)
{
GRID_ASSERT(0);
}
void HermOp (const Field &in, Field &out)
{
Field tmp(in.Grid());
Op(in,tmp);
AdjOp(tmp,out);
}
};
NAMESPACE_END(Grid);
@@ -38,10 +38,10 @@ NAMESPACE_BEGIN(Grid);
// Fine Object == (per site) type of fine field
// nbasis == number of deflation vectors
template<class Fobj,class CComplex,int nbasis>
class GeneralCoarsenedMatrix : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
class DeprecatedGeneralCoarsenedMatrix : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
public:
typedef GeneralCoarsenedMatrix<Fobj,CComplex,nbasis> GeneralCoarseOp;
typedef DeprecatedGeneralCoarsenedMatrix<Fobj,CComplex,nbasis> GeneralCoarseOp;
typedef iVector<CComplex,nbasis > siteVector;
typedef iMatrix<CComplex,nbasis > siteMatrix;
typedef Lattice<iScalar<CComplex> > CoarseComplexField;
@@ -103,7 +103,7 @@ public:
void ProjectNearestNeighbour(RealD shift, GeneralCoarseOp &CopyMe)
{
int nfound=0;
std::cout << GridLogMessage <<"GeneralCoarsenedMatrix::ProjectNearestNeighbour "<< CopyMe._A[0].Grid()<<std::endl;
std::cout << GridLogMessage <<"DeprecatedGeneralCoarsenedMatrix::ProjectNearestNeighbour "<< CopyMe._A[0].Grid()<<std::endl;
for(int p=0;p<geom.npoint;p++){
for(int pp=0;pp<CopyMe.geom.npoint;pp++){
// Search for the same relative shift
@@ -120,7 +120,7 @@ public:
}
*/
GeneralCoarsenedMatrix(NonLocalStencilGeometry &_geom,GridBase *FineGrid, GridCartesian * CoarseGrid,int _herm=1)
DeprecatedGeneralCoarsenedMatrix(NonLocalStencilGeometry &_geom,GridBase *FineGrid, GridCartesian * CoarseGrid,int _herm=1)
: geom(_geom),
_FineGrid(FineGrid),
_CoarseGrid(CoarseGrid),
@@ -34,11 +34,11 @@ NAMESPACE_BEGIN(Grid);
// Fine Object == (per site) type of fine field
// nbasis == number of deflation vectors
template<class Fobj,class CComplex,int nbasis>
class MultiGeneralCoarsenedMatrix : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
class DeprecatedMultiGeneralCoarsenedMatrix : public SparseMatrixBase<Lattice<iVector<CComplex,nbasis > > > {
public:
typedef typename CComplex::scalar_object SComplex;
typedef GeneralCoarsenedMatrix<Fobj,CComplex,nbasis> GeneralCoarseOp;
typedef MultiGeneralCoarsenedMatrix<Fobj,CComplex,nbasis> MultiGeneralCoarseOp;
typedef DeprecatedGeneralCoarsenedMatrix<Fobj,CComplex,nbasis> GeneralCoarseOp;
typedef DeprecatedMultiGeneralCoarsenedMatrix<Fobj,CComplex,nbasis> MultiGeneralCoarseOp;
typedef iVector<CComplex,nbasis > siteVector;
typedef iMatrix<CComplex,nbasis > siteMatrix;
@@ -78,13 +78,33 @@ public:
GridCartesian * CoarseGrid(void) { return _CoarseGridMulti; }; // this is all the linalg routines need to know
//////////////////////////////////////////////////////////////////////////
// Bilingual accessors, matching GeneralCoarsenedMatrix. Grid() here is the
// Accessors shared with DeprecatedGeneralCoarsenedMatrix. Grid() here is the
// D+1 multiRHS grid and this class never holds the D dimensional one, so
// ExtractMatrix writes into whatever grid the caller's lattice is on.
//////////////////////////////////////////////////////////////////////////
NonLocalStencilGeometry & Geometry(void) { return geom_srhs; };
void ExtractMatrix(int p,CoarseMatrix &A) { BLAStoGrid(A,BLAS_A[p]); };
// One stencil point in or out, matching MultiGeneralCoarsenedOperator, so a
// caller can read or write a point without knowing which of the two layouts
// (per-point here, site-major there) holds it.
void MatrixPointIn(int p,const deviceVector<calcMatrix> &src)
{
int64_t sites = BLAS_A[p].size();
GRID_ASSERT((int64_t)src.size() == sites);
calcMatrix *dst = &BLAS_A[p][0];
const calcMatrix *sp = &src[0];
accelerator_for(ss,sites,1,{ dst[ss] = sp[ss]; });
}
void MatrixPointOut(int p,deviceVector<calcMatrix> &dst)
{
int64_t sites = BLAS_A[p].size();
dst.resize(sites);
const calcMatrix *sp = &BLAS_A[p][0];
calcMatrix *dp = &dst[0];
accelerator_for(ss,sites,1,{ dp[ss] = sp[ss]; });
}
// I/O on the operator matrices, via the BLAS layout array. The parameter is
// a vector over the geometry points; the body indexes A[p].
void SetMatrix (int p,std::vector<CoarseMatrix> & A)
@@ -121,7 +141,7 @@ public:
}
*/
MultiGeneralCoarsenedMatrix(NonLocalStencilGeometry &_geom,GridCartesian *CoarseGridMulti) :
DeprecatedMultiGeneralCoarsenedMatrix(NonLocalStencilGeometry &_geom,GridCartesian *CoarseGridMulti) :
_CoarseGridMulti(CoarseGridMulti),
geom_srhs(_geom),
geom(_CoarseGridMulti,_geom.hops,_geom.skip+1),
@@ -27,15 +27,15 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
/* END LEGAL */
#pragma once
// DEPRECATED with the V1 coarse operators: the single-RHS ADEF-2 on an
// DEPRECATED with the other classes in deprecated/: the single-RHS ADEF-2 on an
// Aggregation (ProjectToSubspace/PromoteFromSubspace) and a D-dimensional
// coarse operator. The mrhs solver in AdefMrhs.h covers one right-hand
// side through its LinearFunction interface, on the V2 coarse operator.
// side through its LinearFunction interface, on MultiGeneralCoarsenedOperator.
NAMESPACE_BEGIN(Grid);
template<class Field, class CoarseField, class Aggregation>
class TwoLevelADEF2 : public TwoLevelCG<Field>
class DeprecatedTwoLevelADEF2 : public TwoLevelCG<Field>
{
public:
///////////////////////////////////////////////////////////////////////////////////
@@ -50,7 +50,7 @@ class TwoLevelADEF2 : public TwoLevelCG<Field>
///////////////////////////////////////////////////////////////////////////////////
// more most opertor functions
TwoLevelADEF2(RealD tol,
DeprecatedTwoLevelADEF2(RealD tol,
Integer maxit,
LinearOperatorBase<Field> &FineLinop,
LinearFunction<Field> &Smoother,
+1
View File
@@ -133,6 +133,7 @@ public:
template<class obj> void GlobalSumP2P(obj &o)
{
GRID_TRACE("GlobalSumP2P");
std::vector<obj> column;
obj accum = o;
int source,dest;
+43 -13
View File
@@ -898,7 +898,8 @@ template<int B,class vobj>
void axpyMultiChunk(Lattice<vobj> &z,const ComplexD *b,
const std::vector<const Lattice<vobj>*> &x,int m,
int do_norm,
decltype(innerProduct(vobj(),vobj())) *inner_tmp_v)
decltype(innerProduct(vobj(),vobj())) *inner_tmp_v,
const Lattice<vobj> *from=nullptr)
{
typedef decltype(z.View(AcceleratorRead)) View;
GRID_ASSERT(m>=1 && m<=B);
@@ -917,13 +918,29 @@ void axpyMultiChunk(Lattice<vobj> &z,const ComplexD *b,
}
for(int j=m;j<B;j++){ pack.set(j,h_v[0]); pack.b[j] = ComplexD(0.0); }
autoView(z_v,z,AcceleratorWrite);
accelerator_for(ss,sites,nsimd,{
auto acc = coalescedRead(z_v[ss]);
for(int j=0;j<B;j++) if ( j<m ) acc = acc + pack.b[j]*coalescedRead(pack.v(j)[ss]);
coalescedWrite(z_v[ss],acc);
if ( do_norm ) coalescedWrite(inner_tmp_v[ss],innerProduct(acc,acc));
});
if ( from == nullptr ) {
autoView(z_v,z,AcceleratorWrite);
accelerator_for(ss,sites,nsimd,{
auto acc = coalescedRead(z_v[ss]);
for(int j=0;j<B;j++) if ( j<m ) acc = acc + pack.b[j]*coalescedRead(pack.v(j)[ss]);
coalescedWrite(z_v[ss],acc);
if ( do_norm ) coalescedWrite(inner_tmp_v[ss],innerProduct(acc,acc));
});
} else {
// z = from + sum_j b[j] x[j]. The accumulator is seeded from another field,
// so z is written without being read and the caller needs no copy into it.
conformable(*from,z);
GRID_ASSERT(from!=&z);
const Lattice<vobj> &y = *from; // named: autoView is a macro, *p.View() misparses
autoView(y_v,y,AcceleratorRead);
autoView(z_v,z,AcceleratorWriteDiscard);
accelerator_for(ss,sites,nsimd,{
auto acc = coalescedRead(y_v[ss]);
for(int j=0;j<B;j++) if ( j<m ) acc = acc + pack.b[j]*coalescedRead(pack.v(j)[ss]);
coalescedWrite(z_v[ss],acc);
if ( do_norm ) coalescedWrite(inner_tmp_v[ss],innerProduct(acc,acc));
});
}
for(int j=0;j<m;j++) h_v[j].ViewClose();
}
@@ -931,7 +948,8 @@ void axpyMultiChunk(Lattice<vobj> &z,const ComplexD *b,
// (last) pass. Chunks of 16 for windows longer than 16.
template<class vobj>
RealD axpyMultiNormImpl(Lattice<vobj> &z,const std::vector<ComplexD> &b,
const std::vector<const Lattice<vobj>*> &x,int do_norm)
const std::vector<const Lattice<vobj>*> &x,int do_norm,
const Lattice<vobj> *from=nullptr)
{
typedef decltype(innerProduct(vobj(),vobj())) inner_t;
int m = x.size();
@@ -943,6 +961,7 @@ RealD axpyMultiNormImpl(Lattice<vobj> &z,const std::vector<ComplexD> &b,
inner_t *inner_tmp_v = &inner_tmp[0];
if ( m==0 ) {
if ( from ) z = *from;
if ( do_norm ) return norm2(z);
return 0.0;
}
@@ -951,10 +970,12 @@ RealD axpyMultiNormImpl(Lattice<vobj> &z,const std::vector<ComplexD> &b,
int last = (j0+mm>=m);
std::vector<const Lattice<vobj>*> sub(x.begin()+j0,x.begin()+j0+mm);
int dn = do_norm && last;
if ( mm<=2 ) axpyMultiChunk<2> (z,&b[j0],sub,mm,dn,inner_tmp_v);
else if ( mm<=4 ) axpyMultiChunk<4> (z,&b[j0],sub,mm,dn,inner_tmp_v);
else if ( mm<=8 ) axpyMultiChunk<8> (z,&b[j0],sub,mm,dn,inner_tmp_v);
else axpyMultiChunk<16>(z,&b[j0],sub,mm,dn,inner_tmp_v);
// Only the first chunk is seeded from `from`; later chunks accumulate into z.
const Lattice<vobj> *f = (j0==0) ? from : nullptr;
if ( mm<=2 ) axpyMultiChunk<2> (z,&b[j0],sub,mm,dn,inner_tmp_v,f);
else if ( mm<=4 ) axpyMultiChunk<4> (z,&b[j0],sub,mm,dn,inner_tmp_v,f);
else if ( mm<=8 ) axpyMultiChunk<8> (z,&b[j0],sub,mm,dn,inner_tmp_v,f);
else axpyMultiChunk<16>(z,&b[j0],sub,mm,dn,inner_tmp_v,f);
}
RealD nrm = 0.0;
@@ -969,6 +990,15 @@ void axpyMulti(Lattice<vobj> &z,const std::vector<ComplexD> &b,const std::vector
{
axpyMultiNormImpl(z,b,x,0);
}
// z = y + sum_j b[j] x[j]. Same single pass as axpyMulti, but the accumulator
// starts from y, so a caller whose z would otherwise have to be primed with a
// copy of y does not make that copy.
template<class vobj>
void axpyMultiFrom(Lattice<vobj> &z,const Lattice<vobj> &y,
const std::vector<ComplexD> &b,const std::vector<const Lattice<vobj>*> &x)
{
axpyMultiNormImpl(z,b,x,0,&y);
}
template<class vobj>
RealD axpyMultiNorm(Lattice<vobj> &z,const std::vector<ComplexD> &b,const std::vector<const Lattice<vobj>*> &x)
{
+36 -1
View File
@@ -800,6 +800,39 @@ void localCopyRegion(const Lattice<vobj> &From,Lattice<vobj> & To,Coordinate Fro
autoView(from_v,From,AcceleratorRead);
autoView(to_v,To,AcceleratorWrite);
if constexpr ( vobj::Nsimd() == 1 ) {
// Unvectorised: there is one lane, so the lane indices are identically zero
// and the internal word index is free to carry the thread. Threading over
// (site,word) with the word fastest makes consecutive threads read and write
// consecutive elements; one thread per site instead leaves them `words`
// apart, which is what made this uncoalesced. The site count and the word
// count are folded into one range so the large product lands on the only
// block dimension that can hold it.
accelerator_for(ss,nsite*words,1,{
int w = ss % words;
uint64_t idx = ss / words;
Coordinate from_coor, to_coor, base;
Lexicographic::CoorFromIndex(base,idx,RegionSize);
for(int i=0;i<nd;i++){
from_coor[i] = base[i] + FromLowerLeft[i];
to_coor[i] = base[i] + ToLowerLeft[i];
}
// rdimensions == ldimensions when unvectorised and the region is local,
// so the oSite sums need no modulo and there is no lane term.
int from_oidx = 0; for(int d=0;d<nd;d++) from_oidx+=f_ostride[d]*from_coor[d];
int to_oidx = 0; for(int d=0;d<nd;d++) to_oidx +=t_ostride[d]*to_coor[d];
const vector_type* from = (const vector_type *)&from_v[from_oidx];
vector_type* to = (vector_type *)&to_v[to_oidx];
to[w] = from[w];
});
} else {
accelerator_for(idx,nsite,1,{
Coordinate from_coor, to_coor, base;
@@ -815,13 +848,15 @@ void localCopyRegion(const Lattice<vobj> &From,Lattice<vobj> & To,Coordinate Fro
const vector_type* from = (const vector_type *)&from_v[from_oidx];
vector_type* to = (vector_type *)&to_v[to_oidx];
scalar_type stmp;
for(int w=0;w<words;w++){
stmp = getlane(from[w], from_lane);
putlane(to[w], stmp, to_lane);
}
});
}
}
template<class vobj>
+37 -19
View File
@@ -94,7 +94,7 @@ template<class vobj> inline void ScatterSlice(const deviceVector<vobj> &buf,
// FIXME -- can put internal indices into thread loop
auto buf_p = & buf[0];
autoView(lat_v, lat, AcceleratorWrite);
accelerator_for(ss, face_ovol/simd[dim],Nsimd,{
accelerator_forNB(ss, face_ovol/simd[dim],Nsimd,{
// scalar layout won't coalesce
#ifdef GRID_SIMT
@@ -180,7 +180,7 @@ template<class vobj> inline void GatherSlice(deviceVector<vobj> &buf,
//for cross platform
//For CPU perhaps just run a loop over Nsimd
auto buf_p = & buf[0];
accelerator_for(ss, face_ovol/simd[dim],Nsimd,{
accelerator_forNB(ss, face_ovol/simd[dim],Nsimd,{
// scalar layout won't coalesce
#ifdef GRID_SIMT
@@ -303,10 +303,17 @@ public:
template<class vobj>
inline Lattice<vobj> Exchange(const Lattice<vobj> &in, const CshiftImplBase<vobj> &cshift = CshiftImplDefault<vobj>()) const
{
GridBase *old_grid = in.Grid();
int dims = old_grid->Nd();
Lattice<vobj> tmp = in;
for(int d=0;d<dims;d++){
Coordinate processors=unpadded_grid->_processors;
int dims = in.Grid()->Nd();
// An undecomposed dimension is not padded, so Expand on it would only copy
// its input onto the same grid: those are skipped. The first decomposed
// dimension expands from `in` itself, so the chain needs no seed copy.
int first=-1;
for(int d=0;d<dims;d++) if ( processors[d] > 1 ) { first=d; break; }
if ( first < 0 ) return in; // nothing decomposed: the padded grid IS the input grid
Lattice<vobj> tmp = Expand(first,in,cshift);
for(int d=first+1;d<dims;d++){
if ( processors[d] == 1 ) continue;
tmp = Expand(d,tmp,cshift); // rvalue && assignment
}
return tmp;
@@ -314,10 +321,18 @@ public:
template<class vobj>
inline Lattice<vobj> ExchangePeriodic(const Lattice<vobj> &in) const
{
GridBase *old_grid = in.Grid();
int dims = old_grid->Nd();
Lattice<vobj> tmp = in;
for(int d=0;d<dims;d++){
Coordinate processors=unpadded_grid->_processors;
int dims = in.Grid()->Nd();
// As Exchange: undecomposed dimensions are pure copies and are skipped, and
// the first decomposed dimension expands from `in` so there is no seed copy.
// On the D+1 mrhs coarse grid (nrhs,1,x,y,z,t) this removes three full
// coarse-field copies per operator application.
int first=-1;
for(int d=0;d<dims;d++) if ( processors[d] > 1 ) { first=d; break; }
if ( first < 0 ) return in; // nothing decomposed: the padded grid IS the input grid
Lattice<vobj> tmp = ExpandPeriodic(first,in);
for(int d=first+1;d<dims;d++){
if ( processors[d] == 1 ) continue;
tmp = ExpandPeriodic(d,tmp); // rvalue && assignment
}
return tmp;
@@ -513,12 +528,18 @@ public:
RealD t_tot=-usecond();
int plane=0;
for ( int d=0;d < depth ; d ++ ) {
int tag = d*1024 + dimension*2+0;
t=usecond();
GatherSlice(send_buf,from,d,dimension,plane*buffer_size); plane++;
t_gather+=usecond()-t;
}
for ( int d=0;d < depth ; d ++ ) {
t=usecond();
GatherSlice(send_buf,from,ld-depth+d,dimension,plane*buffer_size); plane++;
t_gather+= usecond() - t;
}
accelerator_barrier();
for ( int d=0;d < depth ; d ++ ) {
int tag = d*1024 + dimension*2+0;
t=usecond();
if(d==0) fwd_trace = traceStart("PaddedCellFwdMPI");
@@ -533,14 +554,9 @@ public:
(void *)&hrecv_buf[d*buffer_size], recv_from_rank, bytes, tag);
#endif
t_comms+=usecond()-t;
}
}
for ( int d=0;d < depth ; d ++ ) {
int tag = d*1024 + dimension*2+1;
t=usecond();
GatherSlice(send_buf,from,ld-depth+d,dimension,plane*buffer_size); plane++;
t_gather+= usecond() - t;
t=usecond();
if (d==0) bwd_trace = traceStart("PaddedCellBwdMPI");
#ifdef ACCELERATOR_AWARE_MPI
@@ -603,9 +619,11 @@ public:
for ( int d=0;d < depth ; d ++ ) {
ScatterSlice(recv_buf,to,d,dimension,plane*buffer_size); plane++;
}
accelerator_barrier();
t_scatter+= usecond() - t;
t_tot+=usecond();
std::cout << GridLogPerformance << "PaddedCell::Expand new timings: gather :" << t_gather/1000 << "ms"<<std::endl;
std::cout << GridLogPerformance << "PaddedCell::Expand new timings: scatter:" << t_scatter/1000 << "ms"<<std::endl;
std::cout << GridLogPerformance << "PaddedCell::Expand new timings: copy :" << t_copy/1000 << "ms"<<std::endl;
+2 -2
View File
@@ -52,7 +52,7 @@ protected:
public:
//Define the action used to evolve the plaquettes
//(Lüscher: https://arxiv.org/pdf/1006.4518 eq. 1.4)
//(Luscher: https://arxiv.org/pdf/1006.4518 eq. 1.4)
//V'(t) = -g^2 * ( d/dVt S[Vt](g) ) * Vt
// = -g^2 * ( d/dVt (1/g^2 * sum_p Re tr{ 1 - Vt(p) } ) ) * Vt
// = - d/dVt ( sum_p ( Nc - Re tr Vt(p) ) * Vt
@@ -157,7 +157,7 @@ public:
////////////////////////////////////////////////////////////////////////////////
//Compute t^2 <E(t)> for time from the plaquette form
//(Lüscher: https://arxiv.org/pdf/1006.4518 eq. 3.1)
//(Luscher: https://arxiv.org/pdf/1006.4518 eq. 3.1)
//E(t) = 2 * sum_p Retr{ 1 - Vt(p) } =
// = 2 * sum_p ( Nc - Retr Vt(p) ) =
// = 2 * Nc * sum_p ( 1 - Retr Vt(p)/Nc )
+3 -3
View File
@@ -1002,10 +1002,10 @@ accelerator_inline void precisionChange(vRealD *out,const vRealF *in,int n
Optimization::PrecisionChange::StoD(in[m].v,out[n].v,out[n+1].v);
// Bug in gcc 10.0.1 and gcc 10.1 using fixed-size SVE ACLE data types CAS-159553-Y1K4C6
// function call results in compile-time error:
// In function ‘void Grid::precisionChange(Grid::vRealD*, Grid::vRealF*, int)’:
// In function 'void Grid::precisionChange(Grid::vRealD*, Grid::vRealF*, int)':
// .../Grid_vector_types.h:961:56: error:
// cannot bind non-const lvalue reference of type ‘vecd&’ {aka ‘svfloat64_t&’}
// to an rvalue of type ‘vecd’ {aka ‘svfloat64_t’}
// cannot bind non-const lvalue reference of type 'vecd&' {aka 'svfloat64_t&'}
// to an rvalue of type 'vecd' {aka 'svfloat64_t'}
// 961 | Optimization::PrecisionChange::StoD(in[m].v,out[n].v,out[n+1].v);
// | ~~~~~~~^
}
+127
View File
@@ -104,6 +104,119 @@ inline double usecond(void) {
gettimeofday(&tv,NULL);
return 1.0e6*tv.tv_sec + 1.0*tv.tv_usec;
}
/**************************************************************
* Point to point cost, per Cartesian direction.
*
* PaddedCell::Face_exchange issues one MPI_Sendrecv per dimension and waits,
* so the quantity that bounds a halo round is the sendrecv time for that
* direction, not a one way ping-pong. Sweeping the size gives the latency
* intercept and the bandwidth slope separately.
*
* Each direction is labelled on- or off-node, which is what decides whether a
* halo round rides the intra-node fabric or the network, and therefore which
* dimension ordering costs least. Times are reduced across ranks: the max is
* the one that bounds a collective halo round.
**************************************************************
*/
void PingPong(std::vector<int> cart_geom,bool use_device,int ncall)
{
int Nd=cart_geom.size();
std::vector<int> periodic(Nd,1);
std::vector<int> coor(Nd);
int rank;
MPI_Comm communicator;
MPI_Cart_create(WorldComm,Nd,&cart_geom[0],&periodic[0],0,&communicator);
MPI_Comm_rank(communicator,&rank);
MPI_Cart_coords(communicator,rank,Nd,&coor[0]);
// An identifier every rank on a node agrees on, so a neighbour can be
// classified by comparing it.
int node_id = WorldRank;
MPI_Bcast(&node_id,1,MPI_INT,0,WorldShmComm);
size_t max_bytes = 2*1024*1024;
void *xmit, *recv;
if ( use_device ) {
xmit = acceleratorAllocDevice(max_bytes);
recv = acceleratorAllocDevice(max_bytes);
} else {
xmit = malloc(max_bytes);
recv = malloc(max_bytes);
}
if ( !WorldRank ) {
printf("= dim dir bytes us(max) us(min) MB/s off-node ranks\n");
fflush(stdout);
}
for(int d=0;d<Nd;d++){
// An undecomposed dimension is a local wrap in PaddedCell and carries no
// message at all; a shift here would only time a rank talking to itself.
if ( cart_geom[d] == 1 ) continue;
for(int sign=-1;sign<=1;sign+=2){
int from,to;
MPI_Cart_shift(communicator,d,sign,&from,&to);
// Fetch the identifier of the neighbour we SEND to, which means sending
// ours the other way round the shift.
int peer_node=node_id;
MPI_Sendrecv(&node_id, 1,MPI_INT,from,rank,
&peer_node,1,MPI_INT,to, to,
communicator,MPI_STATUS_IGNORE);
int offnode = (peer_node != node_id) ? 1 : 0;
int offnode_count=0;
MPI_Reduce(&offnode,&offnode_count,1,MPI_INT,MPI_SUM,0,communicator);
for(size_t bytes=8; bytes<=max_bytes; bytes*=8){
// Latency needs the repeats; bandwidth does not, and large messages
// would otherwise dominate the run time.
int nc = (bytes<=32768) ? ncall : (ncall/20 > 10 ? ncall/20 : 10);
// Untimed warm-up: the first exchange on a fresh buffer pays connection
// setup and registration, and would otherwise land in the first row.
MPI_Sendrecv(xmit,bytes,MPI_CHAR,to,rank,
recv,bytes,MPI_CHAR,from,from,
communicator,MPI_STATUS_IGNORE);
MPI_Barrier(communicator);
double t0=usecond();
for(int i=0;i<nc;i++){
MPI_Sendrecv(xmit,bytes,MPI_CHAR,to,rank,
recv,bytes,MPI_CHAR,from,from,
communicator,MPI_STATUS_IGNORE);
}
double t1=usecond();
double us = (t1-t0)/(double)nc;
double us_max, us_min;
MPI_Reduce(&us,&us_max,1,MPI_DOUBLE,MPI_MAX,0,communicator);
MPI_Reduce(&us,&us_min,1,MPI_DOUBLE,MPI_MIN,0,communicator);
if ( !WorldRank ) {
printf("= %3d %s %10zu %11.2f %11.2f %10.1f %d\n",
d, sign>0?"+":"-", bytes, us_max, us_min,
(double)bytes/us_max, offnode_count);
fflush(stdout);
}
}
}
}
if ( use_device ) {
acceleratorFreeDevice(xmit);
acceleratorFreeDevice(recv);
} else {
free(xmit);
free(recv);
}
MPI_Comm_free(&communicator);
}
/**************************************************************
* Main benchmark routine
**************************************************************
@@ -303,6 +416,20 @@ int main(int argc, char **argv)
}
if( !WorldRank ) {
printf("=========================================================\n");
printf("= Point to point cost per direction, HOST memory \n");
printf("=========================================================\n");fflush(stdout);
}
PingPong(mpi,false,1000);
if( !WorldRank ) {
printf("=========================================================\n");
printf("= Point to point cost per direction, DEVICE memory \n");
printf("=========================================================\n");fflush(stdout);
}
PingPong(mpi,true,1000);
if( !WorldRank ) {
printf("=========================================================\n");
printf("= Benchmarking HOST memory MPI performance \n");
+214
View File
@@ -0,0 +1,214 @@
// Minimal reproducer for the CXI uncached-registration defect.
//
// Hypothesis: with the libfabric MR cache disabled, a device pointer at a non-zero offset
// into a hipMalloc'd allocation is registered as the whole allocation and the offset is
// discarded, so a send transmits from the base and a receive lands at the base.
//
// Each test therefore prints, for both the data and its landing place, what a correct
// implementation must give, what the hypothesis predicts instead, and what was observed.
//
// hipcc -O2 -o simple_reproducer_noncache simple_reproducer_noncache.cc \
// -I$MPICH_DIR/include -L$MPICH_DIR/lib -lmpi -L$MPICH_DIR/gtl/lib -lmpi_gtl_hsa
//
// MPICH_GPU_SUPPORT_ENABLED=1 FI_MR_CACHE_MAX_COUNT=0 \
// srun -N2 -n2 --ntasks-per-node=1 ./simple_reproducer_noncache
#include <mpi.h>
#include <hip/hip_runtime.h>
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <vector>
// The message must be large enough to go rendezvous: an eager message is staged through a host
// bounce buffer with hipMemcpy, which honours the offset, and never registers the user pointer.
static const int NWORDS = 262144; // 2 MB buffers
static int NMSG = 16384; // 128 KB messages (argv[1], in words)
static int OFF = 16384; // offending offset in words (argv[2]); clearest at >= NMSG,
// which keeps the intended and predicted windows disjoint
static int rank, size, peer;
static uint64_t *A, *B;
static std::vector<uint64_t> h(NWORDS);
#define HIP(cmd) do { hipError_t e=(cmd); if ( e != hipSuccess ) { \
printf("hip error %s at line %d\n",hipGetErrorString(e),__LINE__); \
MPI_Abort(MPI_COMM_WORLD,1); } } while(0)
static uint64_t word(uint64_t tag,int r,int n)
{
return (tag<<60) | ((uint64_t)r<<32) | (uint64_t)n;
}
static const char *verdict(int observed,int correct,int predicted)
{
if ( observed == correct ) return "CORRECT";
if ( observed == predicted ) return "WRONG, as predicted";
return "WRONG, and not as predicted";
}
static void refresh(void)
{
for(int n=0;n<NWORDS;n++) h[n] = word(0xa,rank,n);
HIP(hipMemcpy(A,h.data(),NWORDS*sizeof(uint64_t),hipMemcpyHostToDevice));
for(int n=0;n<NWORDS;n++) h[n] = word(0xb,rank,n);
HIP(hipMemcpy(B,h.data(),NWORDS*sizeof(uint64_t),hipMemcpyHostToDevice));
MPI_Barrier(MPI_COMM_WORLD);
}
static void exchange(int sendoff,int recvoff)
{
MPI_Request req[2];
MPI_Irecv(B+recvoff,NMSG*sizeof(uint64_t),MPI_BYTE,peer,0,MPI_COMM_WORLD,&req[0]);
MPI_Isend(A+sendoff,NMSG*sizeof(uint64_t),MPI_BYTE,peer,0,MPI_COMM_WORLD,&req[1]);
MPI_Waitall(2,req,MPI_STATUSES_IGNORE);
}
// Report the state of one NMSG-word window of B: how much of it was overwritten, and by what.
static void window(const char *label,int at)
{
int mod=0, m0=-1;
for(int i=0;i<NMSG && at+i<NWORDS;i++) {
if ( h[at+i] != word(0xb,rank,at+i) ) { if ( m0 < 0 ) m0 = at+i; mod++; }
}
if ( mod == 0 ) {
printf("rank %d %-9s B[%d..%d] : untouched\n",rank,label,at,at+NMSG-1);
return;
}
uint64_t w = h[m0];
printf("rank %d %-9s B[%d..%d] : %d of %d words overwritten, B[%d] holds A[%d] from rank %d\n",
rank,label,at,at+NMSG-1,mod,NMSG,m0,(int)(w & 0xffffffff),(int)((w>>32) & 0xfffffff));
}
// Name a word: still the receive pattern, or a word of the neighbour's send buffer.
static const char *describe(uint64_t w,int idx)
{
static char s[64];
if ( w == word(0xb,rank,idx) ) snprintf(s,sizeof(s),"untouched");
else if ( (w>>60) == 0xa ) snprintf(s,sizeof(s),"A[%d] of rank %d",
(int)(w & 0xffffffff),(int)((w>>32) & 0xfffffff));
else snprintf(s,sizeof(s),"unrecognised");
return s;
}
static void report(const char *name,int sendoff,int recvoff)
{
HIP(hipMemcpy(h.data(),B,NWORDS*sizeof(uint64_t),hipMemcpyDeviceToHost));
int first=-1, last=-1;
for(int n=0;n<NWORDS;n++) {
if ( h[n] != word(0xb,rank,n) ) { if ( first < 0 ) first = n; last = n; }
}
printf("rank %d %s : send &A[%d] -> recv &B[%d]\n",rank,name,sendoff,recvoff);
if ( first < 0 ) { printf("rank %d nothing was written\n",rank); return; }
uint64_t w = h[first];
int src = (int)(w & 0xffffffff);
int sr = (int)((w>>32) & 0xfffffff);
// The hypothesis predicts both offsets are discarded: the bytes at A[0], landing at B[0].
printf("rank %d data : want A[%d] predict A[0] got A[%d] from rank %d %s\n",
rank,sendoff,src,sr,verdict(src,sendoff,0));
printf("rank %d place : want B[%d] predict B[0] got B[%d] %s\n",
rank,recvoff,first,verdict(first,recvoff,0));
if ( src == sendoff && first == recvoff ) return;
// Both candidate landing sites in our receive buffer, and both candidate source words in the
// neighbour's send buffer, which the fill pattern fixes by construction.
printf("rank %d Boff %d\n",rank,recvoff);
printf("rank %d B[0] = 0x%016llx %s\n",
rank,(unsigned long long)h[0],describe(h[0],0));
if ( recvoff != 0 )
printf("rank %d B[%d]\t= 0x%016llx %s\n",
rank,recvoff,(unsigned long long)h[recvoff],describe(h[recvoff],recvoff));
printf("rank %d Neighbour rank %d send buffer holds\n",rank,peer);
printf("rank %d A[0]\t= 0x%016llx\n",
rank,(unsigned long long)word(0xa,peer,0));
if ( sendoff != 0 )
printf("rank %d A[%d]\t= 0x%016llx\n",
rank,sendoff,(unsigned long long)word(0xa,peer,sendoff));
window("intended",recvoff);
if ( recvoff != 0 ) window("predicted",0);
if ( first < recvoff || last >= recvoff+NMSG ) {
if ( first != 0 || last != NMSG-1 ) {
printf("rank %d stray : modified words span B[%d..%d], outside both windows\n",
rank,first,last);
}
}
int bad=0;
for(int i=0;i<NMSG && first+i<NWORDS;i++) {
if ( h[first+i] != word(0xa,peer,src+i) ) bad++;
}
if ( bad ) printf("rank %d body : %d of %d words are not a contiguous run from A[%d]\n",
rank,bad,NMSG,src);
}
// Let each rank print in turn: flush, then a barrier so rank r is on the page before r+1 starts.
static void ordered(void (*print)(void))
{
for(int r=0;r<size;r++) {
if ( r == rank ) { print(); fflush(stdout); }
MPI_Barrier(MPI_COMM_WORLD);
}
}
static const char *test_name;
static int test_sendoff, test_recvoff;
static void print_report(void) { report(test_name,test_sendoff,test_recvoff); }
static void print_setup (void)
{
printf("rank %d A %p B %p message %d words (%d bytes) offset %d words (%d bytes)\n",
rank,(void *)A,(void *)B,NMSG,(int)(NMSG*sizeof(uint64_t)),OFF,(int)(OFF*sizeof(uint64_t)));
}
int main(int argc,char **argv)
{
MPI_Init(&argc,&argv);
MPI_Comm_rank(MPI_COMM_WORLD,&rank);
MPI_Comm_size(MPI_COMM_WORLD,&size);
if ( size != 2 ) {
if ( rank == 0 ) printf("run with two ranks, one per node\n");
MPI_Finalize();
return 1;
}
peer = 1-rank;
if ( argc > 1 ) NMSG = atoi(argv[1]);
if ( argc > 2 ) OFF = atoi(argv[2]);
if ( NMSG < 1 || OFF < 0 || OFF+NMSG > NWORDS ) {
if ( rank == 0 ) printf("message %d and offset %d words do not fit in %d\n",NMSG,OFF,NWORDS);
MPI_Finalize();
return 1;
}
int ndev;
HIP(hipGetDeviceCount(&ndev));
if ( ndev < 1 ) {
printf("rank %d sees no GPU\n",rank);
MPI_Abort(MPI_COMM_WORLD,1);
}
HIP(hipSetDevice(rank%ndev));
HIP(hipMalloc(&A,NWORDS*sizeof(uint64_t)));
HIP(hipMalloc(&B,NWORDS*sizeof(uint64_t)));
ordered(print_setup);
const int offsets[4][2] = { {0,0}, {OFF,0}, {0,OFF}, {OFF,OFF} };
const char *names[4] = { "test 1","test 2","test 3","test 4" };
for(int t=0;t<4;t++) {
refresh();
exchange(offsets[t][0],offsets[t][1]);
test_name = names[t]; test_sendoff = offsets[t][0]; test_recvoff = offsets[t][1];
ordered(print_report);
}
HIP(hipFree(A));
HIP(hipFree(B));
MPI_Finalize();
return 0;
}
+2 -2
View File
@@ -87,7 +87,7 @@ int main (int argc, char ** argv)
HermOpAdaptor<LatticeFermionD> HermFineOp(MdagMOp);
// ── Coarse geometry ────────────────────────────────────────────────────
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
@@ -171,7 +171,7 @@ int main (int argc, char ** argv)
LatticeFermionD src(FGrid); random(RNG5, src);
LatticeFermionD result(FGrid); result = Zero();
TwoLevelADEF2<LatticeFermionD, CoarseVector, Subspace>
DeprecatedTwoLevelADEF2<LatticeFermionD, CoarseVector, Subspace>
HDCG(1.0e-8, 1000,
HermFineOp,
Smoother,
+2 -2
View File
@@ -31,7 +31,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
//
// Operator hierarchy:
// Fine: M†M, acted on via HermOpAdaptor so Op = HermOp = M†M
// Coarse: 33-point GeneralCoarsenedMatrix (NextToNearest, 2-hop M†M)
// Coarse: 33-point DeprecatedGeneralCoarsenedMatrix (NextToNearest, 2-hop M†M)
//
// Setup:
// 1. Chebyshev filter (T_600 x T_2500) on M†M to build near-null subspace {ψᵢ}
@@ -198,7 +198,7 @@ int main(int argc, char **argv)
///////////////////////////////////////////////////////////
// Coarse operator: 33-point stencil for 2-hop M†M
///////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector, vTComplex, nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector, vTComplex, nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNearestStencilGeometry5D geom(Coarse5d);
+2 -2
View File
@@ -364,7 +364,7 @@ void runMG(
) {
// typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
// typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
// typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
// typedef LittleDiracOperator::CoarseVector CoarseVector;
ParseEnvironment();
@@ -620,7 +620,7 @@ int main (int argc, char ** argv)
// assert(nbasis <= Nevecs); // need to have enough evecs
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNearestStencilGeometry5D geom(Coarse5d);
+23 -23
View File
@@ -190,8 +190,8 @@ public:
}
};
// Lüscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi† src.
// Luscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi^dag src.
template<class Field>
class LuscherGuesser : public LinearFunction<Field> {
const std::vector<Field> &psi;
@@ -342,7 +342,7 @@ void runMG(
TrivialPrecon<LatticeFermionD> simple_fine;
//////////////////////////////////////////////////////////////////////
// Level 0→1: coarsen PVdagM, build LinOpCoarse
// Level 0->1: coarsen PVdagM, build LinOpCoarse
//////////////////////////////////////////////////////////////////////
LittleDiracOperator LittleDiracOpPV(geom, FGrid, Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM, AggregatesPD);
@@ -366,7 +366,7 @@ void runMG(
//////////////////////////////////////////////////////////////////////
// psi_coarse: coarse projections of pre-GS fine null vectors.
// These are the Level 1 near-null vectors, promoted from Level 0.
// Used as the aggregation basis for Level 1→2 coarsening.
// Used as the aggregation basis for Level 1->2 coarsening.
//////////////////////////////////////////////////////////////////////
std::vector<CoarseVector> psi_coarse(nbasis, Coarse5d);
for (int k = 0; k < nbasis; k++)
@@ -396,22 +396,22 @@ void runMG(
RealD normC = C.norm();
RealD normCmCdag = (C - C.adjoint()).norm();
std::cout << GridLogMessage << "Coarse null matrix ||C|| = " << normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C†||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C^dag||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Galerkin check ||C||/||W|| = " << normC/normW << std::endl;
}
//////////////////////////////////////////////////////////////////////
// Level 1→2: set up aggregation using psi_coarse as subspace.
// Level 1->2: set up aggregation using psi_coarse as subspace.
// Block factor 2,2,3,2 (removes odd local sublattice in z given MPI
// geometry 3×6×4×4 where z-local at Level 1 is 6).
// geometry 3x6x4x4 where z-local at Level 1 is 6).
// psi_coarse are assigned directly; CoarsenOperator performs
// block-GS orthogonalisation before building LinOpCoarseCoarse.
//////////////////////////////////////////////////////////////////////
// innerProduct(CoarseSiteObj, CoarseSiteObj) returns iScalar<vTComplex>, so CComplex
// for the L1→L2 level must be iScalar<vTComplex>, not vTComplex.
// for the L1->L2 level must be iScalar<vTComplex>, not vTComplex.
typedef typename CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef typename LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,NB> SubspaceL2;
typedef MGPreconditioner<CoarseSiteObj,vTTComplex,NB> L1to2MG;
@@ -429,12 +429,12 @@ void runMG(
TrivialPrecon<CoarseCoarseVector> simpleCC;
//////////////////////////////////////////////////////////////////////
// Lüscher deflation guesser for L3PGCR.
// Luscher deflation guesser for L3PGCR.
// Step 1: project psi_coarse[k] (promoted fine null vectors) to
// CoarseCoarseVector space — these cover the zero-momentum
// CoarseCoarseVector space -- these cover the zero-momentum
// component of the near-null space of LinOpCC.
// Step 2: breed Nextra additional null vectors directly on LinOpCC
// using GCR with random sources — these pick up near-null
// using GCR with random sources -- these pick up near-null
// modes at all spatial frequencies not spanned by step 1.
// Step 3: build C_{st} = <psi_cc[s]|LinOpCC|psi_cc[t]> over the
// full augmented basis and invert directly via Eigen LU.
@@ -476,7 +476,7 @@ void runMG(
RealD normCcc = Ccc.norm();
RealD normCccmCdag = (Ccc - Ccc.adjoint()).norm();
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc|| = " << normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc†||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc^dag||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
}
Eigen::MatrixXcd Ccc_inv = Ccc.inverse();
LuscherGuesser<CoarseCoarseVector> CCDeflGuesser(psi_cc, Ccc_inv);
@@ -489,7 +489,7 @@ void runMG(
L3PGCR.Name("CCouter");
//////////////////////////////////////////////////////////////////////
// Coarse-level GCR smoother for Level 1→2 V-cycle.
// Coarse-level GCR smoother for Level 1->2 V-cycle.
// Mirrors fine-grid SmootherGCR: shifted operator + fixed step count.
// coarse_smoother_shift and coarse_smoother_nstep are the tuning knobs.
//////////////////////////////////////////////////////////////////////
@@ -505,7 +505,7 @@ void runMG(
CoarseSmootherGCR.Name("Csmoother");
//////////////////////////////////////////////////////////////////////
// Level 1→2 V-cycle preconditioner.
// Level 1->2 V-cycle preconditioner.
//////////////////////////////////////////////////////////////////////
L1to2MG L1to2Precon(AggregatesL2,
LinOpCoarse,
@@ -513,7 +513,7 @@ void runMG(
CoarseSmootherGCR, // post-smoother: 12 GCR steps
LinOpCC,
L3PGCR,
CCDeflGuesser); // Lüscher guesser: psi_cc C^{-1} psi_cc†
CCDeflGuesser); // Luscher guesser: psi_cc C^{-1} psi_cc^dag
//////////////////////////////////////////////////////////////////////
// Standalone Level 1 two-level solve test.
@@ -548,7 +548,7 @@ void runMG(
f_src = one;
// Pre-smoother: none (TrivialPrecon); post-smoother: shifted PGCR.
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1→2 V-cycle).
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1->2 V-cycle).
TwoLevelMG ThreeLevelPrecon(AggregatesPD,
PVdagM,
simple_fine,
@@ -592,7 +592,7 @@ int main (int argc, char ** argv)
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
// Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
// Level 1 coarse grid: block 2^4 from fine (48x48x48x96 -> 24x24x24x48, Ls=1)
Coordinate clatt = lat_size;
for (int d = 0; d < 4; d++) clatt[d] /= 2;
std::cout << GridLogMessage << "Level 1 coarse lattice: " << clatt << std::endl;
@@ -600,13 +600,13 @@ int main (int argc, char ** argv)
GridCartesian *Coarse4d = SpaceTimeGrid::makeFourDimGrid(clatt, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
GridCartesian *Coarse5d = SpaceTimeGrid::makeFiveDimGrid(1,Coarse4d);
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24×24×24×48 → 12×12×8×16, Ls=1).
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24x24x24x48 -> 12x12x8x16, Ls=1).
// MPI geometry 3.6.4.4 (288 ranks): fine local {16,8,12,24}.
// Level 1 local {8,4,6,12}; Level 2 local {4,2,2,4}.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) → must use 3.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) -> must use 3.
// t blocked by 3: t-Level1-local=12; 12/3=4 divisible by Nsimd=4 (gen-simd-width=64).
// t-block=2 gives t2-local=6, 6 mod 4 ≠ 0, fails Grid SIMD assertion. ✓
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4). ✓
// t-block=2 gives t2-local=6, 6 mod 4 != 0, fails Grid SIMD assertion.
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4).
Coordinate clatt2 = clatt;
clatt2[0] /= 2;
clatt2[1] /= 2;
@@ -635,7 +635,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis> TwoLevelMG;
@@ -2,7 +2,7 @@
Grid physics library, www.github.com/paboyle/Grid
Source file: ./examples/Example_pvdagm_v2_3level_DenseCoarseMatrix.cc
Source file: ./examples/Example_pvdagm_3level_DenseCoarseMatrix.cc
Copyright (C) 2026
@@ -27,7 +27,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
/* END LEGAL */
//
// PVdagM three level multigrid on the V2 coarse operator.
// PVdagM three level multigrid on MultiGeneralCoarsenedOperator.
//
// STAGES ONE AND TWO: grids, types, subspace, and the L1 and L2 coarsenings.
// The dense bottom and the solves are not here yet.
@@ -37,12 +37,13 @@ Author: Peter Boyle <pboyle@bnl.gov>
// * The coarse space is UNVECTORISED (sComplexD). The fine space stays
// vectorised. MultiRHSBlockProject carries the mixed layout.
//
// * One operator, not two. V1 needed GeneralCoarsenedMatrix to coarsen and
// MultiGeneralCoarsenedMatrix to apply, bridged by CopyMatrix. V2 does
// both, and single versus multiRHS is SetGrid on the same object with the
// matrix elements built once.
// * One operator, not two. The deprecated path needed
// DeprecatedGeneralCoarsenedMatrix to coarsen and
// DeprecatedMultiGeneralCoarsenedMatrix to apply, bridged by CopyMatrix.
// This operator does both, and single versus multiRHS is SetGrid on the
// same object with the matrix elements built once.
//
// * Nrhs is unconstrained. V1 required nrhs % vComplex::Nsimd() == 0 because
// * Nrhs is unconstrained. The deprecated path required nrhs % vComplex::Nsimd() == 0 because
// its multiRHS grid carried the SIMD in the rhs direction.
//
// * CoarsenOperator takes the subspace vectors, not an Aggregation. It block
@@ -54,7 +55,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
// ||<psi|psi> - I||_F guard below is what catches that.
//
// Env: LATT LS MASS NBASIS(compile time) NRHS BLOCK BLOCK2 COARSEN_BATCH
// HOT_START CONFIG SUBSPACE_FILE V1_CHECK MRHS_COARSEN
// HOT_START CONFIG SUBSPACE_FILE DEPRECATED_CHECK MRHS_COARSEN
//
#include <Grid/Grid.h>
@@ -205,7 +206,7 @@ void ParseEnvironment(void)
P("PolyRecordSelect",PolyRecordSelect); P("PolyRefresh",PolyRefresh); P("PolyVerbose",PolyVerbose);
P("SmootherCoeffLog",SmootherCoeffLog);
// library-side knobs, read straight from the environment as the library will
const char *envs[] = {"BLOCK","BLOCK2","SUBSPACE_FILE","CONFIG","HOT_START","MRHS_COARSEN","V1_CHECK",
const char *envs[] = {"BLOCK","BLOCK2","SUBSPACE_FILE","CONFIG","HOT_START","MRHS_COARSEN","DEPRECATED_CHECK",
"DENSE_CC","DENSE_SCHUR","DENSE_SCHUR2D","DENSE_DEVICE_SUM","DENSE_SPLITK","DENSE_NB",
"SCHUR2D_LEAF_SPAN","SCHUR2D_LEAF_LU","SCHUR2D_PROBE","SUMMA_HANDSHAKE","DENSE_APPLY_PROFILE","SLAB_FILE",
"SOLVE_SRHS"};
@@ -433,7 +434,7 @@ int main (int argc, char ** argv)
GridCartesian *Coarse5d = new GridCartesian(c5latt,c5simd,c5mpi);
// 6D coarse multiRHS grid: rhs is dim 0, undistributed and unvectorised.
// No divisibility constraint on nrhs, unlike V1.
// No divisibility constraint on nrhs, unlike the deprecated operators.
Coordinate cmlatt({nrhs,1,clatt[0],clatt[1],clatt[2],clatt[3]});
Coordinate cmsimd({1,1,1,1,1,1});
Coordinate cmmpi ({1,1,mpi[0],mpi[1],mpi[2],mpi[3]});
@@ -495,7 +496,7 @@ int main (int argc, char ** argv)
// Level 1 types: unvectorised coarse scalar
//////////////////////////////////////////////////////////////////////
typedef sTComplexD CComplexS;
typedef MultiGeneralCoarsenedOperatorV2<vSpinColourVector,CComplexS,nbasis> CoarseOperator;
typedef MultiGeneralCoarsenedOperator<vSpinColourVector,CComplexS,nbasis> CoarseOperator;
typedef CoarseOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,CComplexS,nbasis> Subspace;
@@ -556,23 +557,23 @@ int main (int argc, char ** argv)
// the projector imports it below -- the 60 fine vectors are 20 GB of host memory per
// rank and 60 LRU-eligible device fields competing with the solver's working set)
//////////////////////////////////////////////////////////////////////
// Optional cross check of the coarse matrix elements against the V1
// path, which needs a vectorised coarse space. Block Gram-Schmidt is
// idempotent, so V1 may re-orthonormalise the same vectors in place
// Optional cross check of the coarse matrix elements against the
// deprecated path, which needs a vectorised coarse space. Block Gram-Schmidt
// is idempotent, so the deprecated path may re-orthonormalise the same vectors in place
// without a second copy of the subspace.
//////////////////////////////////////////////////////////////////////
if ( getenv("V1_CHECK") ) {
if ( getenv("DEPRECATED_CHECK") ) {
typedef GeneralCoarsenedMatrix <vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix <vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef Aggregation <vSpinColourVector,vTComplex,nbasis> SubspaceV;
Coordinate v5latt({1,clatt[0],clatt[1],clatt[2],clatt[3]});
Coordinate v5simd({1,fsimd[0],fsimd[1],fsimd[2],fsimd[3]});
GridCartesian *Coarse5dV = new GridCartesian(v5latt,v5simd,c5mpi);
int nrhs_v1 = vComplex::Nsimd();
Coordinate vmlatt({nrhs_v1,1,clatt[0],clatt[1],clatt[2],clatt[3]});
int nrhs_dep = vComplex::Nsimd();
Coordinate vmlatt({nrhs_dep,1,clatt[0],clatt[1],clatt[2],clatt[3]});
Coordinate vmsimd({vComplex::Nsimd(),1,1,1,1,1});
GridCartesian *CoarseMrhsV = new GridCartesian(vmlatt,vmsimd,cmmpi);
@@ -582,25 +583,30 @@ int main (int argc, char ** argv)
for(int k=0;k<nbasis;k++) AggV.subspace[k]=AggregatesGCR.subspace[k];
LittleDiracOperator LittleDiracOpPV(geomV,FGrid,Coarse5dV);
std::cout << GridLogMessage << "*** V1 CoarsenOperator (cross check) ***" << std::endl;
std::cout << GridLogMessage << "*** deprecated CoarsenOperator (cross check) ***" << std::endl;
SetFineSloppy(FineSloppyComms);
LittleDiracOpPV.CoarsenOperator(PVdagM,AggV);
SetFineSloppy(0);
MrhsLittleDiracOperator mrhsV1(geomV,CoarseMrhsV);
mrhsV1.CopyMatrix(LittleDiracOpPV);
MrhsLittleDiracOperator mrhsDep(geomV,CoarseMrhsV);
mrhsDep.CopyMatrix(LittleDiracOpPV);
// BLAS_A is written by GridtoBLAS in lSite order and both sides carry
// the same scalar_object, so the two are directly comparable.
typedef MrhsLittleDiracOperator::calcMatrix calcMatrix;
int npoint = geom.npoint;
RealD num=0.0, den=0.0;
// The operator holds ONE site-major matrix buffer; MatrixPointOut hands back
// stencil point p in the same lSite order as the deprecated per-point
// array, and sizes a2.
deviceVector<calcMatrix> a2;
for(int p=0;p<npoint;p++){
int64_t sites = mrhsV1.BLAS_A[p].size();
GRID_ASSERT(sites == (int64_t)CoarseOpPV.BLAS_A[p].size());
int64_t sites = mrhsDep.BLAS_A[p].size();
CoarseOpPV.MatrixPointOut(p,a2);
GRID_ASSERT(sites == (int64_t)a2.size());
std::vector<calcMatrix> h1(sites),h2(sites);
acceleratorCopyFromDevice(&mrhsV1.BLAS_A[p][0], &h1[0],sites*sizeof(calcMatrix));
acceleratorCopyFromDevice(&CoarseOpPV.BLAS_A[p][0],&h2[0],sites*sizeof(calcMatrix));
acceleratorCopyFromDevice(&mrhsDep.BLAS_A[p][0],&h1[0],sites*sizeof(calcMatrix));
acceleratorCopyFromDevice(&a2[0], &h2[0],sites*sizeof(calcMatrix));
ComplexD *w1=(ComplexD *)&h1[0];
ComplexD *w2=(ComplexD *)&h2[0];
int64_t words = sites*sizeof(calcMatrix)/sizeof(ComplexD);
@@ -610,8 +616,8 @@ int main (int argc, char ** argv)
den += real(w1[i])*real(w1[i])+imag(w1[i])*imag(w1[i]);
}
}
std::cout << GridLogMessage << "V1_CHECK: |A_V1|^2 = " << den << std::endl;
std::cout << GridLogMessage << "V1_CHECK: |A_V1 - A_V2|^2 / |A_V1|^2 = " << num/den << std::endl;
std::cout << GridLogMessage << "DEPRECATED_CHECK: |A_dep|^2 = " << den << std::endl;
std::cout << GridLogMessage << "DEPRECATED_CHECK: |A_dep - A|^2 / |A_dep|^2 = " << num/den << std::endl;
GRID_ASSERT( den > 0.0 );
GRID_ASSERT( num/den < 1.0e-18 );
}
@@ -627,7 +633,7 @@ int main (int argc, char ** argv)
MrhsProjector.Allocate(nbasis,FGrid,Coarse5d);
MrhsProjector.ImportBasis(AggregatesGCR.subspace); // block orthonormal basis
// The basis now lives in the projector's BLAS_V (device); nothing after this
// point reads AggregatesGCR.subspace (V1_CHECK moved above). Free it: 60 fine
// point reads AggregatesGCR.subspace (DEPRECATED_CHECK moved above). Free it: 60 fine
// fields = 20 GB host per rank, and 60 LRU-eligible device copies.
AggregatesGCR.subspace.clear(); AggregatesGCR.subspace.shrink_to_fit();
@@ -640,7 +646,7 @@ int main (int argc, char ** argv)
//////////////////////////////////////////////////////////////////////
// STAGE TWO: L2 -> L3.
//
// The fine operator here is V2 at L1, which is natively multiRHS, so the
// The fine operator here is the L1 operator, which is natively multiRHS, so the
// multiRHS driver applies with no promotion adapter: its D+1 grid IS the
// batch grid the L1 operator is currently set to.
//////////////////////////////////////////////////////////////////////
@@ -662,7 +668,7 @@ int main (int argc, char ** argv)
// Coarsening deepens the tensor nest by one iScalar
typedef CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<CComplexS> CComplexS2;
typedef MultiGeneralCoarsenedOperatorV2<CoarseSiteObj,CComplexS2,nbasis> CoarseCoarseOperator;
typedef MultiGeneralCoarsenedOperator<CoarseSiteObj,CComplexS2,nbasis> CoarseCoarseOperator;
typedef CoarseCoarseOperator::CoarseVector CoarseCoarseVector;
NextToNearestStencilGeometry5D geom2(CoarseCoarse5d);
@@ -730,7 +736,7 @@ int main (int argc, char ** argv)
// STAGE THREE (part one): the dense bottom on L2.
//
// DenseCoarseMatrix is bilingual: it takes the elements through
// Geometry()/ExtractMatrix(), so the V2 operator serves directly. It does
// Geometry()/ExtractMatrix(), so the L2 operator serves directly. It does
// detect that a multiRHS op cannot apply on the D dimensional grid and
// skips its own certificate and VERIFY, so the equivalent check is done
// here instead, driving the L2 operator at Nrhs 1 through a slice.
@@ -740,7 +746,7 @@ int main (int argc, char ** argv)
if ( getenv("DENSE_CC")==nullptr || atoi(getenv("DENSE_CC")) ) {
std::cout << GridLogMessage << "*** L3 dense bottom: import from the V2 L2 operator ***" << std::endl;
std::cout << GridLogMessage << "*** L3 dense bottom: import from the L2 operator ***" << std::endl;
DenseCC.reset(new DenseCC_t(CoarseCoarse5d));
DenseCC->Import(CoarseOpL2);
@@ -790,7 +796,7 @@ int main (int argc, char ** argv)
auto RunSolve = [&](int nr)
{
std::cout << GridLogMessage << "**********************************************" << std::endl;
std::cout << GridLogMessage << " V2 THREE-level solve, Nrhs = " << nr << std::endl;
std::cout << GridLogMessage << " THREE-level solve, Nrhs = " << nr << std::endl;
// Device-memory budget BEFORE the solve (2026-08-28: NRHS=6 died in hipMalloc at the
// first fine-smoother history allocation -- reported as an asynchronous "memory
// access fault" unless AMD_SERIALIZE_KERNEL/COPY made it a clean OOM). The outer
@@ -952,7 +958,7 @@ int main (int argc, char ** argv)
GridStopWatch w; w.Start();
L1PGCR(src,sol);
w.Stop();
std::cout << GridLogMessage << "V2 3-level solve Nrhs "<<nr<<" total " << w.Elapsed()
std::cout << GridLogMessage << "3-level solve Nrhs "<<nr<<" total " << w.Elapsed()
<< " (per RHS: " << w.useconds()/1.0e6/nr << " s)" << std::endl;
// The outer operator is exact by policy; assert the state rather than
File diff suppressed because it is too large. Load diff
+12 -12
View File
@@ -212,8 +212,8 @@ public:
}
};
// Lüscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi† src.
// Luscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi^dag src.
template<class Field>
class LuscherGuesser : public LinearFunction<Field> {
const std::vector<Field> &psi;
@@ -704,7 +704,7 @@ void runMG(
TrivialPrecon<LatticeFermionD> simple_fine;
//////////////////////////////////////////////////////////////////////
// Level 0→1: coarsen PVdagM, build LinOpCoarse
// Level 0->1: coarsen PVdagM, build LinOpCoarse
//////////////////////////////////////////////////////////////////////
LittleDiracOperator LittleDiracOpPV(geom, FGrid, Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM, AggregatesPD);
@@ -736,7 +736,7 @@ void runMG(
// Optional sigma-ordering of psi_coarse (SVD_REORDER set): replace the crude
// first-NB_CC slice with the NB_CC most-null directions of span(psi_coarse)
// under LinOpCoarse. Nullness measure = singular values (eig of the
// Gram-whitened Psi†A†APsi), NOT the numerical range Q†AQ which
// Gram-whitened Psi^dag A^dag APsi), NOT the numerical range Q^dag AQ which
// non-normality contaminates. The printed sigma spectrum shows where the
// truncation cliff sits. Unset => raw first-30 (crude GS-ordered slice).
//////////////////////////////////////////////////////////////////////
@@ -787,7 +787,7 @@ void runMG(
}
//////////////////////////////////////////////////////////////////////
// Level 1→2: SUPERCOARSE aggregation using psi_coarse as subspace.
// Level 1->2: SUPERCOARSE aggregation using psi_coarse as subspace.
// Maximal block {8,4,3,6}: CC = [3,6,8,8] = the dense-invertible floor.
//
// UNBLOCKING TRUNCATION NB_CC = 30 (first-30 slice of psi_coarse):
@@ -811,7 +811,7 @@ void runMG(
typedef typename CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB_CC> LittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB_CC> LittleDiracOperatorL2;
typedef typename LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,NB_CC> SubspaceL2;
typedef MGPreconditioner<CoarseSiteObj,vTTComplex,NB_CC> L1to2MG;
@@ -830,7 +830,7 @@ void runMG(
//////////////////////////////////////////////////////////////////////
// CC solve: DENSE (exact, non-iterative) by default; DENSE_CC=0 gives
// the previous iterative L3PGCR + Lüscher-guesser path for A/B.
// the previous iterative L3PGCR + Luscher-guesser path for A/B.
//////////////////////////////////////////////////////////////////////
int use_dense = 1;
if (getenv("DENSE_CC")) use_dense = atoi(getenv("DENSE_CC"));
@@ -859,7 +859,7 @@ void runMG(
ccGuess = &simpleCC; // exact solve ignores/overwrites any guess
} else {
////////////////////////////////////////////////////////////////////
// Lüscher deflation guesser (arXiv:0706.2298 A.3) for the iterative CC
// Luscher deflation guesser (arXiv:0706.2298 A.3) for the iterative CC
// solve, as in the earlier supercoarse configuration.
////////////////////////////////////////////////////////////////////
psi_cc.resize(nbasis, CoarseCoarse5d);
@@ -900,7 +900,7 @@ void runMG(
}
//////////////////////////////////////////////////////////////////////
// Coarse-level GCR smoother for Level 1→2 V-cycle.
// Coarse-level GCR smoother for Level 1->2 V-cycle.
//////////////////////////////////////////////////////////////////////
RealD coarse_smoother_shift = 0.1;
int coarse_smoother_nstep = 2;
@@ -914,7 +914,7 @@ void runMG(
CoarseSmootherGCR.SetZeroGuess(1); // post-smoother slot: caller zeroes vec2 (NOT L2MGsolver: it takes the Luscher guess)
//////////////////////////////////////////////////////////////////////
// Level 1→2 V-cycle preconditioner.
// Level 1->2 V-cycle preconditioner.
//////////////////////////////////////////////////////////////////////
L1to2MG L1to2Precon(AggregatesL2,
LinOpCoarse,
@@ -998,7 +998,7 @@ int main (int argc, char ** argv)
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
// Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
// Level 1 coarse grid: block 2^4 from fine (48x48x48x96 -> 24x24x24x48, Ls=1)
Coordinate clatt = lat_size;
// Coordinate Block1({2,2,2,2});
// Coordinate Block2({8,4,3,6});
@@ -1047,7 +1047,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis> TwoLevelMG;
+17 -17
View File
@@ -323,7 +323,7 @@ void runMG(
TrivialPrecon<LatticeFermionD> simple_fine;
//////////////////////////////////////////////////////////////////////
// Level 0→1: coarsen PVdagM, build LinOpCoarse
// Level 0->1: coarsen PVdagM, build LinOpCoarse
//////////////////////////////////////////////////////////////////////
LittleDiracOperator LittleDiracOpPV(geom, FGrid, Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM, AggregatesPD);
@@ -348,7 +348,7 @@ void runMG(
//////////////////////////////////////////////////////////////////////
// psi_coarse: coarse projections of pre-GS fine null vectors.
// These are the Level 1 near-null vectors, promoted from Level 0.
// Used as the aggregation basis for Level 1→2 coarsening.
// Used as the aggregation basis for Level 1->2 coarsening.
//////////////////////////////////////////////////////////////////////
std::vector<CoarseVector> psi_coarse(nbasis, Coarse5d);
for (int k = 0; k < nbasis; k++)
@@ -378,22 +378,22 @@ void runMG(
RealD normC = C.norm();
RealD normCmCdag = (C - C.adjoint()).norm();
std::cout << GridLogMessage << "Coarse null matrix ||C|| = " << normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C†||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C^dag||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Galerkin check ||C||/||W|| = " << normC/normW << std::endl;
}
//////////////////////////////////////////////////////////////////////
// Level 1→2: set up aggregation using psi_coarse as subspace.
// Level 1->2: set up aggregation using psi_coarse as subspace.
// Block factor 2,2,3,2 (removes odd local sublattice in z given MPI
// geometry 3×6×4×4 where z-local at Level 1 is 6).
// geometry 3x6x4x4 where z-local at Level 1 is 6).
// psi_coarse are assigned directly; CoarsenOperator performs
// block-GS orthogonalisation before building LinOpCoarseCoarse.
//////////////////////////////////////////////////////////////////////
// innerProduct(CoarseSiteObj, CoarseSiteObj) returns iScalar<vTComplex>, so CComplex
// for the L1→L2 level must be iScalar<vTComplex>, not vTComplex.
// for the L1->L2 level must be iScalar<vTComplex>, not vTComplex.
typedef typename CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef typename LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,NB> SubspaceL2;
typedef MGPreconditioner<CoarseSiteObj,vTTComplex,NB> L1to2MG;
@@ -412,14 +412,14 @@ void runMG(
// Level 2 solver: plain GCR, no further coarsening
//////////////////////////////////////////////////////////////////////
TrivialPrecon<CoarseCoarseVector> simpleCC;
// L3PGCR is an inner solver inside the L1→2 V-cycle; does not need to converge
// L3PGCR is an inner solver inside the L1->2 V-cycle; does not need to converge
// to fine-grid precision. Loose tolerance (3e-2) and large restart (64) to allow
// the Krylov space to span enough of the near-null spectrum of LinOpCC per cycle.
PrecGeneralisedConjugateResidualNonHermitian<CoarseCoarseVector> L3PGCR(1.0e-4,5,LinOpCC,simpleCC,64,64);
L3PGCR.Level(3);
//////////////////////////////////////////////////////////////////////
// Coarse-level GCR smoother for Level 1→2 V-cycle.
// Coarse-level GCR smoother for Level 1->2 V-cycle.
// Mirrors fine-grid SmootherGCR: shifted operator + fixed step count.
// coarse_smoother_shift and coarse_smoother_nstep are the tuning knobs.
//////////////////////////////////////////////////////////////////////
@@ -433,7 +433,7 @@ void runMG(
CoarseSmootherGCR.Level(2);
//////////////////////////////////////////////////////////////////////
// Level 1→2 V-cycle preconditioner.
// Level 1->2 V-cycle preconditioner.
//////////////////////////////////////////////////////////////////////
L1to2MG L1to2Precon(AggregatesL2,
LinOpCoarse,
@@ -474,7 +474,7 @@ void runMG(
f_src = one;
// Pre-smoother: none (TrivialPrecon); post-smoother: shifted PGCR.
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1→2 V-cycle).
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1->2 V-cycle).
TwoLevelMG ThreeLevelPrecon(AggregatesPD,
PVdagM,
simple_fine,
@@ -518,7 +518,7 @@ int main (int argc, char ** argv)
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
// Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
// Level 1 coarse grid: block 2^4 from fine (48x48x48x96 -> 24x24x24x48, Ls=1)
Coordinate clatt = lat_size;
for (int d = 0; d < 4; d++) clatt[d] /= 2;
std::cout << GridLogMessage << "Level 1 coarse lattice: " << clatt << std::endl;
@@ -526,13 +526,13 @@ int main (int argc, char ** argv)
GridCartesian *Coarse4d = SpaceTimeGrid::makeFourDimGrid(clatt, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
GridCartesian *Coarse5d = SpaceTimeGrid::makeFiveDimGrid(1,Coarse4d);
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24×24×24×48 → 12×12×8×16, Ls=1).
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24x24x24x48 -> 12x12x8x16, Ls=1).
// MPI geometry 3.6.4.4 (288 ranks): fine local {16,8,12,24}.
// Level 1 local {8,4,6,12}; Level 2 local {4,2,2,4}.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) → must use 3.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) -> must use 3.
// t blocked by 3: t-Level1-local=12; 12/3=4 divisible by Nsimd=4 (gen-simd-width=64).
// t-block=2 gives t2-local=6, 6 mod 4 ≠ 0, fails Grid SIMD assertion. ✓
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4). ✓
// t-block=2 gives t2-local=6, 6 mod 4 != 0, fails Grid SIMD assertion.
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4).
Coordinate clatt2 = clatt;
clatt2[0] /= 2;
clatt2[1] /= 2;
@@ -566,7 +566,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis> TwoLevelMG;
+29 -29
View File
@@ -190,8 +190,8 @@ public:
}
};
// Lüscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi† src.
// Luscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi^dag src.
template<class Field>
class LuscherGuesser : public LinearFunction<Field> {
const std::vector<Field> &psi;
@@ -343,7 +343,7 @@ void runMG(
TrivialPrecon<LatticeFermionD> simple_fine;
//////////////////////////////////////////////////////////////////////
// Level 0→1: coarsen PVdagM, build LinOpCoarse
// Level 0->1: coarsen PVdagM, build LinOpCoarse
//////////////////////////////////////////////////////////////////////
LittleDiracOperator LittleDiracOpPV(geom, FGrid, Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM, AggregatesPD);
@@ -367,7 +367,7 @@ void runMG(
//////////////////////////////////////////////////////////////////////
// psi_coarse: coarse projections of pre-GS fine null vectors.
// These are the Level 1 near-null vectors, promoted from Level 0.
// Used as the aggregation basis for Level 1→2 coarsening.
// Used as the aggregation basis for Level 1->2 coarsening.
//////////////////////////////////////////////////////////////////////
std::vector<CoarseVector> psi_coarse(nbasis, Coarse5d);
for (int k = 0; k < nbasis; k++)
@@ -397,22 +397,22 @@ void runMG(
RealD normC = C.norm();
RealD normCmCdag = (C - C.adjoint()).norm();
std::cout << GridLogMessage << "Coarse null matrix ||C|| = " << normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C†||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C^dag||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Galerkin check ||C||/||W|| = " << normC/normW << std::endl;
}
//////////////////////////////////////////////////////////////////////
// Level 1→2: set up aggregation using psi_coarse as subspace.
// Level 1->2: set up aggregation using psi_coarse as subspace.
// Block factor 2,2,3,2 (removes odd local sublattice in z given MPI
// geometry 3×6×4×4 where z-local at Level 1 is 6).
// geometry 3x6x4x4 where z-local at Level 1 is 6).
// psi_coarse are assigned directly; CoarsenOperator performs
// block-GS orthogonalisation before building LinOpCoarseCoarse.
//////////////////////////////////////////////////////////////////////
// innerProduct(CoarseSiteObj, CoarseSiteObj) returns iScalar<vTComplex>, so CComplex
// for the L1→L2 level must be iScalar<vTComplex>, not vTComplex.
// for the L1->L2 level must be iScalar<vTComplex>, not vTComplex.
typedef typename CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef typename LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,NB> SubspaceL2;
typedef MGPreconditioner<CoarseSiteObj,vTTComplex,NB> L1to2MG;
@@ -430,12 +430,12 @@ void runMG(
TrivialPrecon<CoarseCoarseVector> simpleCC;
//////////////////////////////////////////////////////////////////////
// Lüscher deflation guesser for L3PGCR.
// Luscher deflation guesser for L3PGCR.
// Step 1: project psi_coarse[k] (promoted fine null vectors) to
// CoarseCoarseVector space — these cover the zero-momentum
// CoarseCoarseVector space -- these cover the zero-momentum
// component of the near-null space of LinOpCC.
// Step 2: breed Nextra additional null vectors directly on LinOpCC
// using GCR with random sources — these pick up near-null
// using GCR with random sources -- these pick up near-null
// modes at all spatial frequencies not spanned by step 1.
// Step 3: build C_{st} = <psi_cc[s]|LinOpCC|psi_cc[t]> over the
// full augmented basis and invert directly via Eigen LU.
@@ -477,13 +477,13 @@ void runMG(
RealD normCcc = Ccc.norm();
RealD normCccmCdag = (Ccc - Ccc.adjoint()).norm();
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc|| = " << normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc†||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc^dag||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
}
Eigen::MatrixXcd Ccc_inv = Ccc.inverse();
LuscherGuesser<CoarseCoarseVector> CCDeflGuesser(psi_cc, Ccc_inv);
//////////////////////////////////////////////////////////////////////
// Level 2→3: coarsen LinOpCC using the RAW promoted psi_cc as aggregation
// Level 2->3: coarsen LinOpCC using the RAW promoted psi_cc as aggregation
// to build the Level 4 (coarse-coarse-coarse) operator.
// psi_cc[0..nbasis-1] are the coarse-coarse near-null vectors, projected
// from the RAW psi_coarse (themselves projected from the RAW fine null
@@ -492,11 +492,11 @@ void runMG(
// so assign COPIES of psi_cc and keep psi_cc itself raw.
//
// Tensor depth deepens once more: innerProduct(CoarseCoarseSiteObj,...) returns
// iScalar<vTTComplex>, so CComplex for the L2→L3 level is iScalar<iScalar<vTComplex>>.
// iScalar<vTTComplex>, so CComplex for the L2->L3 level is iScalar<iScalar<vTComplex>>.
//////////////////////////////////////////////////////////////////////
typedef typename CoarseCoarseVector::vector_object CoarseCoarseSiteObj;
typedef iScalar<vTTComplex> vTTTComplex;
typedef GeneralCoarsenedMatrix<CoarseCoarseSiteObj,vTTTComplex,NB> LittleDiracOperatorL3;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseCoarseSiteObj,vTTTComplex,NB> LittleDiracOperatorL3;
typedef typename LittleDiracOperatorL3::CoarseVector CoarseCoarseCoarseVector;
typedef Aggregation<CoarseCoarseSiteObj,vTTTComplex,NB> SubspaceL3;
typedef MGPreconditioner<CoarseCoarseSiteObj,vTTTComplex,NB> L2to3MG;
@@ -531,7 +531,7 @@ void runMG(
L4PGCR.Name("CCCouter");
//////////////////////////////////////////////////////////////////////
// Level 2→3 V-cycle: depth-2 SHIFTED smoother on LinOpCC + L4 bottom.
// Level 2->3 V-cycle: depth-2 SHIFTED smoother on LinOpCC + L4 bottom.
// The shift slides the coarse-coarse field of values off the origin so a
// 2-step smoother has something to bite on a non-normal operator (IRS idea).
//////////////////////////////////////////////////////////////////////
@@ -556,7 +556,7 @@ void runMG(
simpleCCC); // trivial guesser at the bottom
//////////////////////////////////////////////////////////////////////
// Level 3 (coarse-coarse) solve: GCR preconditioned by the L2→L3 V-cycle.
// Level 3 (coarse-coarse) solve: GCR preconditioned by the L2->L3 V-cycle.
// Replaces the plain L3PGCR of the 3-level build -- the coarse-coarse level
// is now smoothed shallowly and recursed rather than solved deeply.
//////////////////////////////////////////////////////////////////////
@@ -565,7 +565,7 @@ void runMG(
L3MGsolver.Name("CCouter");
//////////////////////////////////////////////////////////////////////
// Coarse-level GCR smoother for Level 1→2 V-cycle.
// Coarse-level GCR smoother for Level 1->2 V-cycle.
// Mirrors fine-grid SmootherGCR: shifted operator + fixed step count.
// coarse_smoother_shift and coarse_smoother_nstep are the tuning knobs.
//////////////////////////////////////////////////////////////////////
@@ -581,15 +581,15 @@ void runMG(
CoarseSmootherGCR.Name("Csmoother");
//////////////////////////////////////////////////////////////////////
// Level 1→2 V-cycle preconditioner.
// Level 1->2 V-cycle preconditioner.
//////////////////////////////////////////////////////////////////////
L1to2MG L1to2Precon(AggregatesL2,
LinOpCoarse,
simpleC, // no pre-smoother (matches fine-grid setup)
CoarseSmootherGCR, // post-smoother: depth-2 shifted GCR
LinOpCC,
L3MGsolver, // coarse-coarse solve is now the L2→L3 V-cycle
CCDeflGuesser); // Lüscher guesser: psi_cc C^{-1} psi_cc†
L3MGsolver, // coarse-coarse solve is now the L2->L3 V-cycle
CCDeflGuesser); // Luscher guesser: psi_cc C^{-1} psi_cc^dag
//////////////////////////////////////////////////////////////////////
// Standalone Level 1 two-level solve test.
@@ -624,7 +624,7 @@ void runMG(
f_src = one;
// Pre-smoother: none (TrivialPrecon); post-smoother: shifted PGCR.
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1→2 V-cycle).
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1->2 V-cycle).
TwoLevelMG ThreeLevelPrecon(AggregatesPD,
PVdagM,
simple_fine,
@@ -668,7 +668,7 @@ int main (int argc, char ** argv)
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
// Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
// Level 1 coarse grid: block 2^4 from fine (48x48x48x96 -> 24x24x24x48, Ls=1)
Coordinate clatt = lat_size;
for (int d = 0; d < 4; d++) clatt[d] /= 2;
std::cout << GridLogMessage << "Level 1 coarse lattice: " << clatt << std::endl;
@@ -676,13 +676,13 @@ int main (int argc, char ** argv)
GridCartesian *Coarse4d = SpaceTimeGrid::makeFourDimGrid(clatt, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
GridCartesian *Coarse5d = SpaceTimeGrid::makeFiveDimGrid(1,Coarse4d);
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24×24×24×48 → 12×12×8×16, Ls=1).
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24x24x24x48 -> 12x12x8x16, Ls=1).
// MPI geometry 3.6.4.4 (288 ranks): fine local {16,8,12,24}.
// Level 1 local {8,4,6,12}; Level 2 local {4,2,2,4}.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) → must use 3.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) -> must use 3.
// t blocked by 3: t-Level1-local=12; 12/3=4 divisible by Nsimd=4 (gen-simd-width=64).
// t-block=2 gives t2-local=6, 6 mod 4 ≠ 0, fails Grid SIMD assertion. ✓
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4). ✓
// t-block=2 gives t2-local=6, 6 mod 4 != 0, fails Grid SIMD assertion.
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4).
Coordinate clatt2 = clatt;
clatt2[0] /= 2;
clatt2[1] /= 2;
@@ -728,7 +728,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis> TwoLevelMG;
+38 -38
View File
@@ -190,8 +190,8 @@ public:
}
};
// Lüscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi† src.
// Luscher deflated guesser (arXiv:0706.2298 Sec A.3) for a non-Hermitian solve.
// C_{st} = <psi[s] | LinOp | psi[t]>; guess = sum_s c_s psi[s] where c = C^{-1} psi^dag src.
template<class Field>
class LuscherGuesser : public LinearFunction<Field> {
const std::vector<Field> &psi;
@@ -344,7 +344,7 @@ void runMG(
TrivialPrecon<LatticeFermionD> simple_fine;
//////////////////////////////////////////////////////////////////////
// Level 0→1: coarsen PVdagM, build LinOpCoarse
// Level 0->1: coarsen PVdagM, build LinOpCoarse
//////////////////////////////////////////////////////////////////////
LittleDiracOperator LittleDiracOpPV(geom, FGrid, Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM, AggregatesPD);
@@ -368,7 +368,7 @@ void runMG(
//////////////////////////////////////////////////////////////////////
// psi_coarse: coarse projections of pre-GS fine null vectors.
// These are the Level 1 near-null vectors, promoted from Level 0.
// Used as the aggregation basis for Level 1→2 coarsening.
// Used as the aggregation basis for Level 1->2 coarsening.
//////////////////////////////////////////////////////////////////////
std::vector<CoarseVector> psi_coarse(nbasis, Coarse5d);
for (int k = 0; k < nbasis; k++)
@@ -398,22 +398,22 @@ void runMG(
RealD normC = C.norm();
RealD normCmCdag = (C - C.adjoint()).norm();
std::cout << GridLogMessage << "Coarse null matrix ||C|| = " << normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C†||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Coarse null matrix ||C - C^dag||/||C|| = " << normCmCdag/normC << std::endl;
std::cout << GridLogMessage << "Galerkin check ||C||/||W|| = " << normC/normW << std::endl;
}
//////////////////////////////////////////////////////////////////////
// Level 1→2: set up aggregation using psi_coarse as subspace.
// Level 1->2: set up aggregation using psi_coarse as subspace.
// Block factor 2,2,3,2 (removes odd local sublattice in z given MPI
// geometry 3×6×4×4 where z-local at Level 1 is 6).
// geometry 3x6x4x4 where z-local at Level 1 is 6).
// psi_coarse are assigned directly; CoarsenOperator performs
// block-GS orthogonalisation before building LinOpCoarseCoarse.
//////////////////////////////////////////////////////////////////////
// innerProduct(CoarseSiteObj, CoarseSiteObj) returns iScalar<vTComplex>, so CComplex
// for the L1→L2 level must be iScalar<vTComplex>, not vTComplex.
// for the L1->L2 level must be iScalar<vTComplex>, not vTComplex.
typedef typename CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,NB> LittleDiracOperatorL2;
typedef typename LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,NB> SubspaceL2;
typedef MGPreconditioner<CoarseSiteObj,vTTComplex,NB> L1to2MG;
@@ -431,12 +431,12 @@ void runMG(
TrivialPrecon<CoarseCoarseVector> simpleCC;
//////////////////////////////////////////////////////////////////////
// Lüscher deflation guesser for L3PGCR.
// Luscher deflation guesser for L3PGCR.
// Step 1: project psi_coarse[k] (promoted fine null vectors) to
// CoarseCoarseVector space — these cover the zero-momentum
// CoarseCoarseVector space -- these cover the zero-momentum
// component of the near-null space of LinOpCC.
// Step 2: breed Nextra additional null vectors directly on LinOpCC
// using GCR with random sources — these pick up near-null
// using GCR with random sources -- these pick up near-null
// modes at all spatial frequencies not spanned by step 1.
// Step 3: build C_{st} = <psi_cc[s]|LinOpCC|psi_cc[t]> over the
// full augmented basis and invert directly via Eigen LU.
@@ -478,13 +478,13 @@ void runMG(
RealD normCcc = Ccc.norm();
RealD normCccmCdag = (Ccc - Ccc.adjoint()).norm();
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc|| = " << normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc†||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
std::cout << GridLogMessage << "Coarse-coarse deflation matrix ||Ccc-Ccc^dag||/||Ccc|| = " << normCccmCdag/normCcc << std::endl;
}
Eigen::MatrixXcd Ccc_inv = Ccc.inverse();
LuscherGuesser<CoarseCoarseVector> CCDeflGuesser(psi_cc, Ccc_inv);
//////////////////////////////////////////////////////////////////////
// Level 2→3: coarsen LinOpCC using the RAW promoted psi_cc as aggregation
// Level 2->3: coarsen LinOpCC using the RAW promoted psi_cc as aggregation
// to build the Level 4 (coarse-coarse-coarse) operator.
// psi_cc[0..nbasis-1] are the coarse-coarse near-null vectors, projected
// from the RAW psi_coarse (themselves projected from the RAW fine null
@@ -493,11 +493,11 @@ void runMG(
// so assign COPIES of psi_cc and keep psi_cc itself raw.
//
// Tensor depth deepens once more: innerProduct(CoarseCoarseSiteObj,...) returns
// iScalar<vTTComplex>, so CComplex for the L2→L3 level is iScalar<iScalar<vTComplex>>.
// iScalar<vTTComplex>, so CComplex for the L2->L3 level is iScalar<iScalar<vTComplex>>.
//////////////////////////////////////////////////////////////////////
typedef typename CoarseCoarseVector::vector_object CoarseCoarseSiteObj;
typedef iScalar<vTTComplex> vTTTComplex;
typedef GeneralCoarsenedMatrix<CoarseCoarseSiteObj,vTTTComplex,NB> LittleDiracOperatorL3;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseCoarseSiteObj,vTTTComplex,NB> LittleDiracOperatorL3;
typedef typename LittleDiracOperatorL3::CoarseVector CoarseCoarseCoarseVector;
typedef Aggregation<CoarseCoarseSiteObj,vTTTComplex,NB> SubspaceL3;
typedef MGPreconditioner<CoarseCoarseSiteObj,vTTTComplex,NB> L2to3MG;
@@ -514,7 +514,7 @@ void runMG(
TrivialPrecon<CoarseCoarseCoarseVector> simpleCCC;
//////////////////////////////////////////////////////////////////////
// Level 3→4: coarsen LinOpCCC to build the Level 5 operator, using a
// Level 3->4: coarsen LinOpCCC to build the Level 5 operator, using a
// TRUNCATED basis of only the first NB5 (< nbasis) raw promoted null vectors.
// psi_ccc[k] = raw psi_cc projected through the (block-GS'd) L3 aggregation
// -- the pre-block-GS chain continued one level deeper. We keep only the
@@ -526,7 +526,7 @@ void runMG(
// NB: a positive result is conservative (sigma-ordering can only help); a
// negative one is inconclusive until the sigma-ordered NB5 is tried.
//
// Tensor depth deepens once more: CComplex for the L3→L4 level is
// Tensor depth deepens once more: CComplex for the L3->L4 level is
// iScalar<vTTTComplex>. NB5 (the coarse dimension) is independent of the
// depth -- it just makes the coarsest site vector NB5-dimensional.
//////////////////////////////////////////////////////////////////////
@@ -542,10 +542,10 @@ void runMG(
// Optional sigma-ordering of psi_ccc (SVD_REORDER set): replace the crude
// first-NB5 slice with the NB5 genuinely-most-null directions of span(psi_ccc)
// under LinOpCCC. For a NON-NORMAL operator the nullness measure is the
// singular value of A restricted to the span -- eig of Q†A†AQ -- NOT the
// numerical range Q†AQ (which non-normality contaminates). Robust route:
// singular value of A restricted to the span -- eig of Q^dag A^dag AQ -- NOT the
// numerical range Q^dag AQ (which non-normality contaminates). Robust route:
// whiten by the Gram (drop near-dependent directions), Hermitian-eig the
// whitened A†A, rotate. The printed singular spectrum IS the SVD study: where
// whitened A^dag A, rotate. The printed singular spectrum IS the SVD study: where
// it falls off tells you the natural NB5, and the same numbers illuminate why
// the earlier singular-subspace deflation re-entered. Safe here because we
// ORDER vectors that then feed a Galerkin projection, not REMOVE a subspace.
@@ -603,7 +603,7 @@ void runMG(
typedef typename CoarseCoarseCoarseVector::vector_object CoarseCoarseCoarseSiteObj;
typedef iScalar<vTTTComplex> vTTTTComplex;
typedef GeneralCoarsenedMatrix<CoarseCoarseCoarseSiteObj,vTTTTComplex,NB5> LittleDiracOperatorL4;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseCoarseCoarseSiteObj,vTTTTComplex,NB5> LittleDiracOperatorL4;
typedef typename LittleDiracOperatorL4::CoarseVector CoarseCoarseCoarseCoarseVector;
typedef Aggregation<CoarseCoarseCoarseSiteObj,vTTTTComplex,NB5> SubspaceL4;
typedef MGPreconditioner<CoarseCoarseCoarseSiteObj,vTTTTComplex,NB5> L3to4MG;
@@ -635,7 +635,7 @@ void runMG(
L5PGCR.Name("CCCCouter");
//////////////////////////////////////////////////////////////////////
// Level 3→4 V-cycle: depth-2 SHIFTED smoother on LinOpCCC + Level 5 bottom.
// Level 3->4 V-cycle: depth-2 SHIFTED smoother on LinOpCCC + Level 5 bottom.
// Level 4 is no longer the bottom -- it is smoothed shallowly and recursed to
// Level 5, mirroring how Level 3 recurses to Level 4.
//////////////////////////////////////////////////////////////////////
@@ -660,14 +660,14 @@ void runMG(
simpleCCCC); // trivial guesser at the bottom
//////////////////////////////////////////////////////////////////////
// Level 4 (coarse-coarse-coarse) solve: GCR preconditioned by the L3→L4 V-cycle.
// Level 4 (coarse-coarse-coarse) solve: GCR preconditioned by the L3->L4 V-cycle.
//////////////////////////////////////////////////////////////////////
PrecGeneralisedConjugateResidualNonHermitian<CoarseCoarseCoarseVector> L4MGsolver(1.0e-1,200,LinOpCCC,L3to4Precon,16,16);
L4MGsolver.Level(4);
L4MGsolver.Name("CCCouter");
//////////////////////////////////////////////////////////////////////
// Level 2→3 V-cycle: depth-2 SHIFTED smoother on LinOpCC + Level 4 solve.
// Level 2->3 V-cycle: depth-2 SHIFTED smoother on LinOpCC + Level 4 solve.
// The shift slides the coarse-coarse field of values off the origin so a
// 2-step smoother has something to bite on a non-normal operator (IRS idea).
//////////////////////////////////////////////////////////////////////
@@ -688,11 +688,11 @@ void runMG(
simpleCC, // no pre-smoother
CoarseCoarseSmootherGCR, // post-smoother: depth-2 shifted GCR
LinOpCCC,
L4MGsolver, // coarse solve is now the L3→L4 V-cycle
L4MGsolver, // coarse solve is now the L3->L4 V-cycle
simpleCCC); // trivial guesser
//////////////////////////////////////////////////////////////////////
// Level 3 (coarse-coarse) solve: GCR preconditioned by the L2→L3 V-cycle.
// Level 3 (coarse-coarse) solve: GCR preconditioned by the L2->L3 V-cycle.
// Replaces the plain L3PGCR of the 3-level build -- the coarse-coarse level
// is now smoothed shallowly and recursed rather than solved deeply.
//////////////////////////////////////////////////////////////////////
@@ -701,7 +701,7 @@ void runMG(
L3MGsolver.Name("CCouter");
//////////////////////////////////////////////////////////////////////
// Coarse-level GCR smoother for Level 1→2 V-cycle.
// Coarse-level GCR smoother for Level 1->2 V-cycle.
// Mirrors fine-grid SmootherGCR: shifted operator + fixed step count.
// coarse_smoother_shift and coarse_smoother_nstep are the tuning knobs.
//////////////////////////////////////////////////////////////////////
@@ -717,15 +717,15 @@ void runMG(
CoarseSmootherGCR.Name("Csmoother");
//////////////////////////////////////////////////////////////////////
// Level 1→2 V-cycle preconditioner.
// Level 1->2 V-cycle preconditioner.
//////////////////////////////////////////////////////////////////////
L1to2MG L1to2Precon(AggregatesL2,
LinOpCoarse,
simpleC, // no pre-smoother (matches fine-grid setup)
CoarseSmootherGCR, // post-smoother: depth-2 shifted GCR
LinOpCC,
L3MGsolver, // coarse-coarse solve is now the L2→L3 V-cycle
CCDeflGuesser); // Lüscher guesser: psi_cc C^{-1} psi_cc†
L3MGsolver, // coarse-coarse solve is now the L2->L3 V-cycle
CCDeflGuesser); // Luscher guesser: psi_cc C^{-1} psi_cc^dag
//////////////////////////////////////////////////////////////////////
// Standalone Level 1 two-level solve test.
@@ -760,7 +760,7 @@ void runMG(
f_src = one;
// Pre-smoother: none (TrivialPrecon); post-smoother: shifted PGCR.
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1→2 V-cycle).
// Coarse solver: L2MGsolver (PGCR preconditioned by Level 1->2 V-cycle).
TwoLevelMG ThreeLevelPrecon(AggregatesPD,
PVdagM,
simple_fine,
@@ -804,7 +804,7 @@ int main (int argc, char ** argv)
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
// Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
// Level 1 coarse grid: block 2^4 from fine (48x48x48x96 -> 24x24x24x48, Ls=1)
Coordinate clatt = lat_size;
for (int d = 0; d < 4; d++) clatt[d] /= 2;
std::cout << GridLogMessage << "Level 1 coarse lattice: " << clatt << std::endl;
@@ -812,13 +812,13 @@ int main (int argc, char ** argv)
GridCartesian *Coarse4d = SpaceTimeGrid::makeFourDimGrid(clatt, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
GridCartesian *Coarse5d = SpaceTimeGrid::makeFiveDimGrid(1,Coarse4d);
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24×24×24×48 → 12×12×8×16, Ls=1).
// Level 2 coarse-coarse grid: block 2,2,3,3 from Level 1 (24x24x24x48 -> 12x12x8x16, Ls=1).
// MPI geometry 3.6.4.4 (288 ranks): fine local {16,8,12,24}.
// Level 1 local {8,4,6,12}; Level 2 local {4,2,2,4}.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) → must use 3.
// z blocked by 3: z-Level1-local=6; 6/3=2 (even), 6/2=3 (odd) -> must use 3.
// t blocked by 3: t-Level1-local=12; 12/3=4 divisible by Nsimd=4 (gen-simd-width=64).
// t-block=2 gives t2-local=6, 6 mod 4 ≠ 0, fails Grid SIMD assertion. ✓
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4). ✓
// t-block=2 gives t2-local=6, 6 mod 4 != 0, fails Grid SIMD assertion.
// With {4,2,2,4}: Nsimd=4 goes into x or t (both =4).
Coordinate clatt2 = clatt;
clatt2[0] /= 2;
clatt2[1] /= 2;
@@ -880,7 +880,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis> TwoLevelMG;
+2 -2
View File
@@ -49,7 +49,7 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
// sigma_min << min|lambda| : non-normal near origin
// lambda_min(H) < 0 : half-plane condition violated
//
// Requires the dagger code path in GeneralCoarsenedMatrix:
// Requires the dagger code path in DeprecatedGeneralCoarsenedMatrix:
// _Adag allocated, PopulateAdag active, _Adag exchanged, hermitian=0.
//
// Env vars:
@@ -308,7 +308,7 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
+1 -1
View File
@@ -275,7 +275,7 @@ int main (int argc, char ** argv)
MobiusFermionD Dpv (Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,1.0, M5,b,c);
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
// Grid index contraction is positional (colour/spin/lorentz order is meaningful),
// so each MG projection adds one index to the tensor nest rather than reusing a slot:
+3 -3
View File
@@ -38,7 +38,7 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
// precon: V-cycle -- per-RHS fine post-smoother (16-step shifted GCR),
// batched restriction (MultiRHSBlockProject / GEMM),
// ONE coarse PGCR on the 6D mrhs coarse operator
// (MultiGeneralCoarsenedMatrix, GEMM mults -- the ~10x win),
// (DeprecatedMultiGeneralCoarsenedMatrix, GEMM mults -- the ~10x win),
// batched prolongation.
//
// The coarse operator is coarsened once with the standard single-RHS
@@ -260,8 +260,8 @@ int main (int argc, char ** argv)
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
+5 -5
View File
@@ -23,7 +23,7 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
// with L3_DEFL=0 (NO deflation), applied to the enlarged block-diagonal mRHS system:
// the coarse and coarse-coarse levels run a SINGLE Krylov (one GCR polynomial, inner
// products summed over rhs) on the packed 6D mrhs fields, so both coarse levels batch
// through GEMM (MultiGeneralCoarsenedMatrix) -- the valence throughput win at BOTH levels.
// through GEMM (DeprecatedMultiGeneralCoarsenedMatrix) -- the valence throughput win at BOTH levels.
//
// Level structure (each coarse level is a single-field PGCR on a packed 6D mrhs field):
// L1 (fine) : std::vector<LatticeFermionD>, MrhsPGCRNonHermitian on PVdagM,
@@ -227,16 +227,16 @@ int main (int argc, char ** argv)
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
// Level 1 tensor types
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
// Level 2 tensor types (coarsening deepens the nest by one iScalar -- see CLAUDE.md)
typedef CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef MultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef DeprecatedMultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,nbasis> SubspaceL2;
@@ -233,16 +233,16 @@ int main (int argc, char ** argv)
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
// Level 1 tensor types
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
// Level 2 tensor types (coarsening deepens the nest by one iScalar)
typedef CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef MultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef DeprecatedMultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,nbasis> SubspaceL2;
+5 -5
View File
@@ -237,7 +237,7 @@ public:
#endif
// NB: the apply GEMM (Y = slab^dag X) is a tiny-output/huge-K shape that
// under-fills the GPU (~13ms). The fix is a software split-K via
// GridBLAS.gemmBatched (see MultiRHSBlockCGLinalg.h / 2409.03904 Fig 11) —
// GridBLAS.gemmBatched (see MultiRHSBlockCGLinalg.h / 2409.03904 Fig 11) --
// NOT a raw rocblas strided-batched batch, which hung on Frontier and was
// removed. TODO: reimplement through GridBLAS when the ~1.4 s/RHS is wanted.
@@ -891,16 +891,16 @@ int main (int argc, char ** argv)
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
// Level 1 tensor types
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MrhsLittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
typedef Aggregation<vSpinColourVector,vTComplex,nbasis> Subspace;
// Level 2 tensor types (coarsening deepens the nest by one iScalar)
typedef CoarseVector::vector_object CoarseSiteObj;
typedef iScalar<vTComplex> vTTComplex;
typedef GeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef MultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef DeprecatedGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> LittleDiracOperatorL2;
typedef DeprecatedMultiGeneralCoarsenedMatrix<CoarseSiteObj,vTTComplex,nbasis> MrhsLittleDiracOperatorL2;
typedef LittleDiracOperatorL2::CoarseVector CoarseCoarseVector;
typedef Aggregation<CoarseSiteObj,vTTComplex,nbasis> SubspaceL2;
+9 -3
View File
@@ -2,13 +2,15 @@ SUBDIRS = .
include Make.inc
# Precision variants of the two multigrid drivers. The coarse-sector and
# Compile-time variants of the two multigrid drivers. The coarse-sector and
# dense-inversion precisions are compile-time instantiations, so each is its
# own binary built from the same source; the fine-level precision is a
# run-time parameter and needs no variant. Per-target flags give each
# variant its own object file.
# run-time parameter and needs no variant. The recur variant additionally
# takes the GCR residual norm from the recurrence rather than a reduction.
# Per-target flags give each variant its own object file.
bin_PROGRAMS += Example_pvdagm_multigrid_fp32coarse \
Example_pvdagm_multigrid_fp32coarse_fp32dense \
Example_pvdagm_multigrid_fp32coarse_fp32dense_recur \
Example_hdcg_multigrid_fp32coarse
Example_pvdagm_multigrid_fp32coarse_SOURCES = Example_pvdagm_multigrid.cc
@@ -19,6 +21,10 @@ Example_pvdagm_multigrid_fp32coarse_fp32dense_SOURCES = Example_pvdagm_multigri
Example_pvdagm_multigrid_fp32coarse_fp32dense_CPPFLAGS = -DCOARSE_SINGLE -DGRID_DENSE_INVERSE_SINGLE
Example_pvdagm_multigrid_fp32coarse_fp32dense_LDADD = $(top_builddir)/Grid/libGrid.a
Example_pvdagm_multigrid_fp32coarse_fp32dense_recur_SOURCES = Example_pvdagm_multigrid.cc
Example_pvdagm_multigrid_fp32coarse_fp32dense_recur_CPPFLAGS = -DCOARSE_SINGLE -DGRID_DENSE_INVERSE_SINGLE -DGRID_GCR_RESIDUAL_RECURRENCE
Example_pvdagm_multigrid_fp32coarse_fp32dense_recur_LDADD = $(top_builddir)/Grid/libGrid.a
Example_hdcg_multigrid_fp32coarse_SOURCES = Example_hdcg_multigrid.cc
Example_hdcg_multigrid_fp32coarse_CPPFLAGS = -DCOARSE_SINGLE
Example_hdcg_multigrid_fp32coarse_LDADD = $(top_builddir)/Grid/libGrid.a
+142
View File
@@ -0,0 +1,142 @@
#!/usr/bin/env python3
"""
Reconstruct the GCR smoother polynomial from PGCR coefficient logs
(SmootherCoeffLog=1 in Example_pvdagm_3level_DenseCoarseMatrix, or
LogCoefficients(1) on any PrecGeneralisedConjugateResidualNonHermitian).
GCR with z = r (trivial preconditioner) builds p_k = P_k(A) r0,
r_k = R_k(A) r0:
R_0 = P_0 = 1
R_{k+1}(l) = R_k(l) - a_k l P_k(l)
P_{k+1}(l) = R_{k+1}(l) + sum_j b_{k,j} P_{k-j}(l)
so after m steps x_m = S(A) r0 with S = sum_k a_k P_k and residual
polynomial R_m = 1 - l S(l), R_m(0) = 1.
Usage
gcr_polynomial.py LOG [--name Fsmoother] [--lmin 0] [--lmax 3] [--n 600]
writes <name>_poly.dat : l |R_mean(l)| Re R_mean Im R_mean |R|min |R|max |S_mean(l)|
where mean = polynomial built from the per-step MEAN coefficients over
all calls, and min/max are the envelope over the individual calls.
Prints the mean coefficients and their relative spread per step.
gcr_polynomial.py --selftest
runs GCR in numpy on a random non-normal matrix, logs a,b the same way,
and checks R_m(A) r0 == the actual residual to rounding.
gnuplot:
set logscale y; plot 'Fsmoother_poly.dat' u 1:2 w l t '|R_m| mean', \
'' u 1:5 w l t 'min', '' u 1:6 w l t 'max'
"""
import re, sys, argparse
import numpy as np
from numpy.polynomial import polynomial as Poly # coefficient arrays, low->high
A_RE = re.compile(r'(\S+)\s+coeff\[(\d+)\]\s+a=\(([^,]+),([^)]+)\)')
B_RE = re.compile(r'(\S+)\s+coeff\[(\d+)\]((?:\s+b\[\d+\]=\([^)]*\))+)')
BJ_RE = re.compile(r'b\[(\d+)\]=\(([^,]+),([^)]+)\)')
def parse(path, name):
"""-> list of calls; each call = list of (a_k, [b_k0, b_k1, ...]) in step order."""
calls, cur = [], None
for line in open(path, errors='replace'):
m = A_RE.search(line)
if m and m.group(1) == name:
k = int(m.group(2)); a = complex(float(m.group(3)), float(m.group(4)))
if k == 0:
if cur: calls.append(cur)
cur = []
if cur is not None:
cur.append([a, []])
continue
m = B_RE.search(line)
if m and m.group(1) == name and cur is not None:
k = int(m.group(2))
bs = [complex(float(x[1]), float(x[2])) for x in BJ_RE.findall(m.group(3))]
if k < len(cur): cur[k][1] = bs
if cur: calls.append(cur)
return calls
def polynomials(coeffs):
"""coeffs: list of (a_k, [b_kj]) -> (S, R) as low->high complex coefficient arrays."""
R = [np.array([1.0+0j])]; P = [np.array([1.0+0j])]
S = np.array([0.0+0j])
for k, (a, bs) in enumerate(coeffs):
S = Poly.polyadd(S, a*P[k])
Rn = Poly.polysub(R[k], a*Poly.polymulx(P[k])) # R_{k+1} = R_k - a l P_k
Pn = Rn.copy()
for j, b in enumerate(bs): # P_{k+1} = R_{k+1} + sum b_kj P_{k-j}
Pn = Poly.polyadd(Pn, b*P[k-j])
R.append(Rn); P.append(Pn)
return S, R[-1]
def mean_coeffs(calls):
m = min(len(c) for c in calls)
out, spread = [], []
for k in range(m):
a = np.array([c[k][0] for c in calls])
nb = min(len(c[k][1]) for c in calls)
bs = [np.array([c[k][1][j] for c in calls]) for j in range(nb)]
out.append((a.mean(), [b.mean() for b in bs]))
sa = a.std()/max(abs(a.mean()),1e-300)
sb = [b.std()/max(abs(b.mean()),1e-300) for b in bs]
spread.append((sa, sb))
return out, spread
def selftest(n=40, m=6, mmax=6, seed=1):
rng = np.random.default_rng(seed)
A = np.eye(n)*1.5 + 0.4*rng.standard_normal((n,n)) + 0.2j*rng.standard_normal((n,n)) # non-normal
r0 = rng.standard_normal(n) + 1j*rng.standard_normal(n)
r = r0.copy(); x = np.zeros(n, complex)
p = [r.copy()]; q = [A@r]; qq = [np.vdot(q[0],q[0]).real]
log = []
for k in range(m):
a = np.vdot(q[k], r)/qq[k] # <q,r>/<q,q>
x = x + a*p[k]; r = r - a*q[k]
z = r; Az = A@z; pn = z.copy(); qn = Az.copy(); bs = []
for back in range(min(k+1, mmax-1)):
b = -(np.vdot(q[k-back], Az)/qq[k-back]).real # real part, as Grid does
pn = pn + b*p[k-back]; qn = qn + b*q[k-back]; bs.append(complex(b))
p.append(pn); q.append(qn); qq.append(np.vdot(qn,qn).real)
log.append((a, bs))
S, R = polynomials(log)
def apply(poly, v):
out = np.zeros_like(v); Ak = v.copy()
for c in poly: out = out + c*Ak; Ak = A@Ak
return out
err_r = np.linalg.norm(apply(R, r0) - r)/np.linalg.norm(r)
err_x = np.linalg.norm(apply(S, r0) - x)/np.linalg.norm(x)
print(f"selftest: |R_m(A)r0 - r_m|/|r_m| = {err_r:.2e} |S(A)r0 - x_m|/|x_m| = {err_x:.2e}")
return err_r < 1e-10 and err_x < 1e-10
if __name__ == "__main__":
ap = argparse.ArgumentParser()
ap.add_argument("log", nargs="?")
ap.add_argument("--name", default="Fsmoother")
ap.add_argument("--lmin", type=float, default=0.0)
ap.add_argument("--lmax", type=float, default=3.0)
ap.add_argument("--n", type=int, default=600)
ap.add_argument("--selftest", action="store_true")
args = ap.parse_args()
if args.selftest:
sys.exit(0 if selftest() else 1)
calls = parse(args.log, args.name)
if not calls:
sys.exit(f"no '{args.name} coeff[...]' lines found in {args.log}")
mean, spread = mean_coeffs(calls)
print(f"{args.name}: {len(calls)} calls, {len(mean)} steps")
for k,(a,bs) in enumerate(mean):
print(f" step {k}: a = ({a.real:+.5f},{a.imag:+.5f}) rel spread {spread[k][0]:.3f} "
+ " ".join(f"b[{j}]={b.real:+.5f} (spread {spread[k][1][j]:.3f})" for j,b in enumerate(bs)))
S, R = polynomials(mean)
print(" S coefficients (low->high):", np.array2string(S, precision=5))
lam = np.linspace(args.lmin, args.lmax, args.n)
Rm = Poly.polyval(lam, R); Sm = Poly.polyval(lam, S)
per = np.array([np.abs(Poly.polyval(lam, polynomials(c)[1])) for c in calls])
out = f"{args.name}_poly.dat"
with open(out, "w") as f:
f.write("# lambda |R_mean| ReR ImR |R|min |R|max |S_mean|\n")
for i,l in enumerate(lam):
f.write(f"{l:.6f} {abs(Rm[i]):.6e} {Rm[i].real:.6e} {Rm[i].imag:.6e} "
f"{per[:,i].min():.6e} {per[:,i].max():.6e} {abs(Sm[i]):.6e}\n")
print(f" wrote {out}")
+3 -3
View File
@@ -221,11 +221,11 @@ judgement.
the exact inverse of whatever matrix it was handed. Both certify the import
and inversion, not the accuracy of `A_c`.
The instruments that DO see it: coarsen twice and compare `BLAS_A` element by
element against an exact run (the comparison `Test_coarse_v2_coarsen` already
performs between V1 and V2), or the outer iteration count and time to
element against an exact run (the comparison `Test_coarse_coarsen` already
performs against the deprecated operator), or the outer iteration count and time to
solution.
- **Does not reach the coarse levels.** `SloppyComms` sets a flag on
`Stencil`/`StencilEven`/`StencilOdd`; V2's `M` uses `PaddedCell::Face_exchange`,
`Stencil`/`StencilEven`/`StencilOdd`; the coarse operator.s `M` uses `PaddedCell::Face_exchange`,
which sends raw `vobj` bytes with no precision option. Coarse levels are less
bandwidth sensitive, so this matters less than it sounds.
- **Object sharing caveat:** the flag is per fermion operator. A single
+3 -3
View File
@@ -66,7 +66,7 @@ export MPICH_SMP_SINGLE_COPY_MODE=CMA
export MPICH_OFI_NIC_POLICY=GPU
module load libfabric
BIN=$root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix
BIN=$root/examples/Example_pvdagm_3level_DenseCoarseMatrix
OPTS1="--accelerator-threads 8 --shm 4096 --shm-mpi 1"
# LOCAL VOLUME MATCHED to production: 48.48.48.96 on 3.6.4.4 is local 16.8.12.24; on 8 GCDs
# (1.1.2.4) that is global 16.8.24.96 with the production blockings 2.2.3.3 / 4.4.2.4:
@@ -112,7 +112,7 @@ run_cell () {
logopt=""; [ -n "$GRID_LOG_EXTRA" ] && logopt="--log Error,Warning,Message,$GRID_LOG_EXTRA"
env "$@" srun -N1 -n8 --kill-on-bad-exit=1 ./select_gpu $BIN --mpi ${MPI_GEOM} --grid $vol $OPTS1 --device-mem $devmem --comms-overlap $logopt > $fname 2>&1
echo " exit $?"; sleep 10
grep -h "V2 3-level solve\|Memory access fault\|GRID_ASSERT\|Converged on iteration" $fname | sed 's/^Grid : Message : [0-9.]* s : //' | tail -4 | cut -c1-140
grep -h "3-level solve\|Memory access fault\|GRID_ASSERT\|Converged on iteration" $fname | sed 's/^Grid : Message : [0-9.]* s : //' | tail -4 | cut -c1-140
grep -n -B3 "Memory access fault" $fname | tail -5 | cut -c1-160
}
@@ -124,5 +124,5 @@ run_cell F3_nrhs6_nodense 16000 NRHS=6 DENSE_CC=0
run_cell F4_nrhs6_bigmem 48000 NRHS=6 # no MemoryManager eviction
echo "========================================================="
for f in log.fault.F*; do echo "$f: $(grep -c 'Memory access fault' $f) faults; $(grep -h 'V2 3-level solve' $f | tail -1 | sed 's/.*V2/V2/' | cut -c1-60)"; done
for f in log.fault.F*; do echo "$f: $(grep -c 'Memory access fault' $f) faults; $(grep -h '3-level solve' $f | tail -1 | sed 's/.*3-level/3-level/' | cut -c1-60)"; done
echo "========================================================="
+2 -2
View File
@@ -67,7 +67,7 @@ export MPICH_SMP_SINGLE_COPY_MODE=CMA
export MPICH_OFI_NIC_POLICY=GPU
module load libfabric
BIN=$root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix
BIN=$root/examples/Example_pvdagm_3level_DenseCoarseMatrix
OPTS1="--accelerator-threads 8 --shm 4096 --shm-mpi 1 --device-mem 32000"
OVERLAP="--comms-overlap" # set empty for a no-overlap job
vol=48.48.48.96
@@ -112,7 +112,7 @@ run_cell () {
tail -15 $(dirname $f)/Grid.stdout.$rank | cut -c1-160
else
echo " no 'Memory access fault' in any Grid.stderr.*"
grep -h "V2 3-level solve\|Fouter MrhsPGCR: Converged" $GRID_STDOUT_ROOT/0/Grid.stdout.0 | tail -3 | cut -c1-120
grep -h "3-level solve\|Fouter MrhsPGCR: Converged" $GRID_STDOUT_ROOT/0/Grid.stdout.0 | tail -3 | cut -c1-120
fi
}
+206
View File
@@ -0,0 +1,206 @@
#!/bin/bash -l
#SBATCH --job-name=pvdagm-mixed-precision
#SBATCH --nodes=36
#SBATCH --ntasks-per-node=8
#SBATCH --cpus-per-task=7
#SBATCH --gpus-per-node=8
#SBATCH --time=2:00:00
#SBATCH --account=phy157_dwf
#SBATCH --gpu-bind=none
#SBATCH --exclusive
#SBATCH --mem=0
#SBATCH -S 0
##############################################################################
# The three precision axes of the PVdagM multigrid, on the banked point of
# pvdagm_multigrid.job (48^3x96, Ls=24, nbasis=60, 288 GCDs). The outer
# Krylov is fp64 and exact in every cell, so fp32 below it costs convergence
# rate only, never correctness -- the FINAL exact-halo true residual is the
# proof and is printed by every cell.
#
# 1. fine level inside the preconditioner run time <FinePrecision>
# 2. coarse + coarse-coarse sector compile -DCOARSE_SINGLE
# 3. distributed dense coarse-coarse invert compile -DGRID_DENSE_INVERSE_SINGLE
#
# The compile-time axes are separate binaries from the same source, all
# plain `make` targets (examples/Makefile.am):
# Example_pvdagm_multigrid fp64 coarse, fp64 dense
# Example_pvdagm_multigrid_fp32coarse fp32 coarse, fp64 dense
# Example_pvdagm_multigrid_fp32coarse_fp32dense fp32 coarse, fp32 dense
# The default basis is 60. The subspace load reads the first 60 vectors
# of the file.
##############################################################################
LUSTRE=/lustre/orion/phy157/proj-shared/phy157_dwf/paboyle
RUNDIR=$LUSTRE/runs/${SLURM_JOB_NAME}_${SLURM_JOB_ID}
mkdir -p $RUNDIR && cd $RUNDIR && echo "RUNDIR $RUNDIR"
cat << EOF > select_gpu
#!/bin/bash
export GPU_MAP=(0 1 2 3 7 6 5 4)
export NUMA_MAP=(3 3 1 1 2 2 0 0)
export GPU=\${GPU_MAP[\$SLURM_LOCALID]}
export NUMA=\${NUMA_MAP[\$SLURM_LOCALID]}
export HIP_VISIBLE_DEVICES=\$GPU
unset ROCR_VISIBLE_DEVICES
if [ \$SLURM_PROCID = "0" ]; then echo \$*; fi
exec numactl -m \$NUMA -N \$NUMA \$*
EOF
chmod +x ./select_gpu
root=/lustre/orion/phy157/proj-shared/phy157_dwf/paboyle/MGrewrite/Grid/systems/Frontier
source $root/sourceme-rocm7.2.sh
# sourceme sets FI_HMEM_ROCR_USE_DMABUF=0 -- the standing avoidance of the CXI
# NO_TRANSLATION fault at NRHS>=12 (libfabric #12775).
export OMP_NUM_THREADS=7
ulimit -c 0 # no 22 GB GPU core dumps
export FI_MR_CACHE_MONITOR=kdreg2 # site default; device-buffer MPI on Slingshot
export MPICH_GPU_SUPPORT_ENABLED=1
export MPICH_SMP_SINGLE_COPY_MODE=CMA
export MPICH_OFI_NIC_POLICY=GPU
export GRID_ALLOC_NCACHE_LARGE=64 # allocator large-ring depth, not a solver knob
module load libfabric
C64D64=$root/examples/Example_pvdagm_multigrid
C32D64=$root/examples/Example_pvdagm_multigrid_fp32coarse
C32D32=$root/examples/Example_pvdagm_multigrid_fp32coarse_fp32dense
# 40000, not 32000: at Nrhs 12 the working set is about 37 GB (outer restart
# history 24.5, fine smoother history 8.2, sources and solutions 4.1), so a
# 32 GB cache evicts. Measured: 808 ms of device-to-host copies in a 20 s
# window, and 9.706 s per rhs against 8.748 with the larger cache. The
# non-evictable total is 6.3 GiB in the fp32 sector, so 40 GB of cache plus
# comms and shm still fits a 64 GB device.
OPTS="--accelerator-threads 8 --shm 4096 --shm-mpi 1 --device-mem 40000 --comms-overlap"
vol=48.48.48.96
MPI_GEOM=3.6.4.4
SUBSPACE=$LUSTRE/subspace_nb64.scidac
CONFIG=/ccs/home/poare/ckpoint_lat.1000
##############################################################################
# write_params <file> <FinePrecision> <FineSloppyComms>
#
# Everything else is the banked adaptive operating point of
# pvdagm_multigrid.job, unchanged, so the cells differ only in precision.
##############################################################################
write_params () {
local f=$1 fprec=$2 sloppy=$3
cat << EOF > $f
<?xml version="1.0"?>
<grid>
<PVdagMDriver>
<Ls>24</Ls>
<Mass>0.00078</Mass>
<M5>1.8</M5>
<MobiusB>1.5</MobiusB>
<MobiusC>0.5</MobiusC>
<Config>$CONFIG</Config>
<Nrhs>12</Nrhs>
<SolveSingleRHS>1</SolveSingleRHS>
<MultiGrid>
<Setup>
<Block1><elem>2</elem><elem>2</elem><elem>3</elem><elem>3</elem></Block1>
<Block2><elem>4</elem><elem>4</elem><elem>2</elem><elem>4</elem></Block2>
<CoarsenBatch>9</CoarsenBatch>
<SubspaceFile>$SUBSPACE</SubspaceFile>
<FineSloppyComms>$sloppy</FineSloppyComms>
<FinePrecision>$fprec</FinePrecision>
<RetainSubspace>0</RetainSubspace>
</Setup>
<FineSmoother>
<Shift>0.1</Shift><Nstep>6</Nstep><Mmax>4</Mmax>
</FineSmoother>
<CoarseSmoother>
<Shift>2.0</Shift><Nstep>2</Nstep><Mmax>2</Mmax>
</CoarseSmoother>
<CoarseSolver>
<Tol>0.05</Tol><Order>200</Order><Mmax>8</Mmax>
</CoarseSolver>
<Outer>
<Tol>1e-8</Tol><MaxIterations>1000</MaxIterations><Mmax>6</Mmax><Nstep>12</Nstep>
</Outer>
<Dense>
<LeafSpan>9</LeafSpan>
</Dense>
</MultiGrid>
</PVdagMDriver>
</grid>
EOF
}
##############################################################################
# run_cell <name> <binary> <FinePrecision> <FineSloppyComms>
##############################################################################
run_cell () {
name=$1; bin=$2; fprec=$3; sloppy=$4
if [ ! -x "$bin" ]; then
echo "----- $name : SKIPPED, no binary $bin"
echo " build it with a plain make in examples/"
return
fi
write_params params.$name.xml $fprec $sloppy
echo "----- $name : fine=$fprec sloppy=$sloppy bin=$(basename $bin) -----"
export GRID_STDOUT_ROOT=$RUNDIR/mg_$name
srun -N36 -n288 --kill-on-bad-exit=1 ./select_gpu $bin --mpi ${MPI_GEOM} --grid $vol $OPTS \
--pvdagm-params params.$name.xml --debug-stdout --log Error,Warning,Message,Performance \
> log.mg.$name 2>&1
echo " exit $?"; sleep 30
# Device OOM is reported on stderr as "hipMalloc failed for <bytes> out of
# memory" and nowhere else, so it must be in this pattern or a cell that
# died of it looks like a cell that printed nothing.
f=$(grep -l "Memory access fault\|NO_TRANSLATION\|hipMalloc failed\|out of memory\|illegal memory access" $GRID_STDOUT_ROOT/*/Grid.stderr.* 2>/dev/null | head -1)
if [ -n "$f" ]; then
echo " FAULT in $f (rank $(basename $f | sed 's/Grid.stderr.//'))"; tail -20 $f | cut -c1-160
else
echo " certificates, convergence, timing:"
grep -h "Coarse sector precision\|GALERKIN CERTIFICATE\|IMPORT CERTIFICATE\|inversion-source import\|SCHUR .* distributed invert\|VERIFY\|PVdagMMultiGridSolver: Nrhs\|MrhsPGCR: Converged\|FINAL Nrhs .*worst" \
$GRID_STDOUT_ROOT/0/Grid.stdout.0 2>/dev/null | cut -c1-150
fi
}
##############################################################################
# Cells. P0 is the regression gate: with fp64 everywhere it must reproduce
# the banked outer count and residual of pvdagm_multigrid.job. If it does
# not, nothing downstream means anything, and the Galerkin certificate
# (2.3e-15 expected in fp64) separates a coarsening change from a solve one.
##############################################################################
# name binary fineprec sloppy
run_cell P0_fp64 $C64D64 fp64 1
run_cell P1_fp32fine $C64D64 fp32 1
run_cell P2_fp32coarse $C32D64 fp64 1
run_cell P3_fp32both $C32D64 fp32 1
run_cell P4_denseF32 $C32D32 fp32 1
# Halo wire format is orthogonal to arithmetic: on an fp32 operator
# FineSloppyComms is bf16 compression, on an fp64 operator it is fp32. One
# exact-halo cell to price it inside the preconditioner.
run_cell P5_exacthalo $C32D64 fp32 0
##############################################################################
# What to read.
#
# P0: the gate. Outer count and s/RHS against the banked numbers.
#
# P1 vs P0: fp32 fine level. Laptop 8^4 saw the same outer count and ~12%
# less wall; here the fine level is the dominant cost and the halo is real,
# so this is where the fp32 case is actually made or lost.
#
# P2 vs P0: fp32 coarse sector. Watch the Galerkin certificate move from
# ~2e-15 to ~3e-7 -- that is fp32 rounding in the coarsening, not an error --
# and check the outer count does not move with it. Halves the coarse
# operator's storage and its comms.
#
# P4 vs P3: fp32 dense coarse-coarse inversion. The reading is the VERIFY
# certificate ||A Ainv x - x||/||x||, which scales as kappa(A_cc) * eps32 and
# so measures the conditioning of the coarse-coarse operator at production
# size. The laptop's N=128 was too small to say anything (5.3e-6 -> 6.7e-6);
# at N=69120 it is a real test. If VERIFY degrades but the outer count holds,
# the fp32 factorisation is fine and halves that footprint; if the outer count
# moves, the fallback is fp32 factorise plus one fp64 refinement step, which
# is two extra dense applies and is not yet written.
#
# P5 vs P3: what the reduced-precision halo is worth once the arithmetic is
# already fp32 (bf16 wire on an fp32 operator).
##############################################################################
+1 -1
View File
@@ -70,7 +70,7 @@ export GRID_ALLOC_NCACHE_LARGE=0 # allocator large-ring depth (Grid alloca
module load libfabric
BIN=$root/examples/Example_pvdagm_multigrid
OPTS="--accelerator-threads 8 --shm 4096 --shm-mpi 0 --device-mem 32000 --comms-overlap"
OPTS="--accelerator-threads 8 --shm 4096 --shm-mpi 0 --device-mem 40000 --comms-overlap"
vol=48.48.48.96
MPI_GEOM=3.6.4.4
+2 -2
View File
@@ -131,14 +131,14 @@ export NRHS=4
echo "----- F4a: 1D baseline (DENSE_SCHUR2D unset) -----"
unset DENSE_SCHUR2D
srun -N36 -n288 ./select_gpu $root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix \
srun -N36 -n288 ./select_gpu $root/examples/Example_pvdagm_3level_DenseCoarseMatrix \
--mpi ${MPI_GEOM} --grid $vol $OPTS1 --comms-overlap
if [ $F3RC -eq 0 ]
then
echo "----- F4b: 2D block-cyclic (DENSE_SCHUR2D=1) -----"
export DENSE_SCHUR2D=1
srun -N36 -n288 ./select_gpu $root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix \
srun -N36 -n288 ./select_gpu $root/examples/Example_pvdagm_3level_DenseCoarseMatrix \
--mpi ${MPI_GEOM} --grid $vol $OPTS1 --comms-overlap
else
echo "----- F4b SKIPPED: F3 failed (rc=$F3RC), not spending the setup on it -----"
+2 -2
View File
@@ -121,10 +121,10 @@ run_cell () {
export FineSmootherMode=$5; export CoarseSmootherMode=$6
echo "----- $name : Fso=$FineSmootherOrder Fss=$FineSmootherShift Csn=$CoarseSmootherNstep fine=$FineSmootherMode coarse=$CoarseSmootherMode -----"
fname=log.ladder.$name
srun -N36 -n288 --kill-on-bad-exit=1 ./select_gpu $root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix \
srun -N36 -n288 --kill-on-bad-exit=1 ./select_gpu $root/examples/Example_pvdagm_3level_DenseCoarseMatrix \
--mpi ${MPI_GEOM} --grid $vol $OPTS1 --comms-overlap > $fname 2>&1
echo " exit $?"; sleep 60
echo " $(grep -h 'V2 3-level solve Nrhs' $fname | sed 's/.*V2/V2/' | tr '\n' ' ')"
echo " $(grep -h '3-level solve Nrhs' $fname | sed 's/.*3-level/3-level/' | tr '\n' ' ')"
echo " $(grep -h 'Fouter MrhsPGCR: Converged' $fname | sed 's/.*Converged/Converged/' | cut -c1-60 | tr '\n' ' ')"
echo " replay per-call |r|/|r0| (Nrhs=1 solve): $(awk '/THREE-level solve, Nrhs = 1/{s=1} s && /Fsmoother replay \|r\|/{v=$NF; n++; t+=v; if(v>mx)mx=v} END{if(n) printf "mean %.4f max %.4f over %d calls", t/n, mx, n}' $fname)"
grep -h "SCHUR fp64 distributed invert took\|GB/s/rank\|BIG LEAVES\|ring histogram\|^Grid : Message : [0-9.]* s : >=" $fname | sed 's/^Grid : Message : [0-9.]* s : //' | cut -c1-150 | head -12
+168 -48
View File
@@ -9,7 +9,7 @@ Contents:
3. Profiling
************************
* 1. INTERCONNECT + MPI
* 0. MPI IO
************************
--------------------------------------------------------------------
@@ -27,6 +27,89 @@ https://github.com/paboyle/Grid/issues/381
https://github.com/pmodels/mpich/commit/3a479ab0
--------------------------------------------------------------------
MPI2-IO performance: introduced All to all based aggregation no changes required.
--------------------------------------------------------------------
https://github.com/pmodels/mpich/issues/7935
MPI_File_write_all / read_all collective buffering does not scale for 4D block-decomposed subarray views (Lustre): 6x slower than a user-level MPI_Alltoallv + POSIX transposition
/* phase 1: redistribute within the row subcommunicator.
sendcounts/senddispls derived analytically from the plan; iodata untouched. */
MPI_Alltoallv(iodata, sendcounts, senddispls, MPI_BYTE,
aggregated, recvcounts, recvdispls, MPI_BYTE, rowcomm);
/* phase 2: a handful of large contiguous writes per rank, no MPI-IO involved */
for (e = 0; e < nextent; e++) {
lseek(fd, offset + extentGsite[e]*sizeof(fobj), SEEK_SET);
write(fd, aggregated + extentLocal[e], extentSites[e]*sizeof(fobj));
}
************************
* 1. INTERCONNECT + MPI
************************
============================================================================
Piz Daint
============================================================================
https://github.com/ofiwg/libfabric/issues/11451
SEGFAULT in first MPI_Comm_dup on more than one node, after MPI_Init
Cause: symptom (ALPS/CSCS, aarch64 Grace/H200, cray-mpich 8.1.32, libfabric 1.22): the
memhooks monitor intercepts munmap by PATCHING THE PLT at MPI_Init
(ofi_memhooks_start -> ofi_write_patch with a garbage data_size); the write overruns
munmap@plt into the NEXT PLT entry (MPI_Comm_dup in Grid's binary)
leaving
`br x15` where an adrp belongs ->
Diagnosed with a hardware watchpoint on the PLT entry during MPI_Init; see the issue for the
gdb transcript. Reproduce on one node with MPICH_SINGLE_HOST_ENABLED=0.
Fix (either):
export FI_MR_CACHE_MONITOR=kdreg2
OR
export FI_MR_CACHE_MONITOR=disabled
export FI_MR_CACHE_MAX_COUNT=0
On ALPS this is a CORRECTNESS requirement for MPI on Slingshot, not a tuning.
============================================================================
Frontier
============================================================================
a) FI_MR_CACHE_MONITOR=kdreg2 (site default) leads to Runtime MPI errors with "NO_TRANSLATION"
userfaultfd and memhooks also fail.
https://github.com/ofiwg/libfabric/issues/12775
Symptom (Frontier, x86, Cray MPICH 8.1.x, ROCm 7.2): device-buffer MPI fails with
MPI_Waitall ... MPIDI_OFI_handle_cq_error: OFI poll failed
(ofi_events.c:MPIDI_OFI_handle_cq_error: Input/output error - NO_TRANSLATION)
on a LIVE, never-freed hipMalloc buffer
b) export FI_MR_CACHE_MONITOR=disabled
export FI_MR_CACHE_MAX_COUNT=0
Leads to INCORRECT RESULTS, see
https://github.com/ofiwg/libfabric/issues/12773
Fix:
export FI_HMEM_ROCR_USE_DMABUF=0
and use site default
FI_MR_CACHE_MONITOR=kdreg2
Every systems/Frontier job and sourceme now sets kdreg2 on OLCF's recommendation;
whether it resolves the Frontier NO_TRANSLATION is the pending test above.
--------------------------------------------------------------------
Slingshot: Frontier and Perlmutter libfabric slow down
and physical memory fragmentation
@@ -187,7 +270,6 @@ unset MPIR_CVAR_COLL_SELECTION_TUNING_JSON_FILE
unset MPIR_CVAR_CH4_POSIX_COLL_SELECTION_TUNING_JSON_FILE
************************
* 3. Visual profile tools
************************
@@ -196,59 +278,97 @@ unset MPIR_CVAR_CH4_POSIX_COLL_SELECTION_TUNING_JSON_FILE
Frontier/rocprof
--------------------------------------------------------------------
cat << EOF > select_gpu
#!/bin/bash
export GPU_MAP=(0 1 2 3 7 6 5 4)
export NUMA_MAP=(3 3 1 1 2 2 0 0)
export GPU=\${GPU_MAP[\$SLURM_LOCALID]}
export NUMA=\${NUMA_MAP[\$SLURM_LOCALID]}
export HIP_VISIBLE_DEVICES=\$GPU
unset ROCR_VISIBLE_DEVICES
if [ \$SLURM_PROCID = "0" ]; then echo \$*; fi
if [ "1" = "1" ]
then
if [ \$SLURM_PROCID = "0" ]
then
exec numactl -m \$NUMA -N \$NUMA rocprofv3 -d $TRACE -o mgtrace.\$SLURM_JOBID \
-r -f csv pftrace --stats -D \
-P 140:20:1 \
-- \$*
else
exec numactl -m \$NUMA -N \$NUMA \$*
fi
else
exec numactl -m \$NUMA -N \$NUMA \$*
fi
EOF
chmod +x ./select_gpu
--------------------------------------------------------------------
Aurora/unitrace
--------------------------------------------------------------------
select_gpu:
#!/bin/bash
export NUMA_PMAP=(0 0 0 1 1 1 0 0 0 1 1 1 );
export NUMA_HMAP=(2 2 2 3 3 3 3 2 2 2 2 3 3 3 );
export GPU_MAP=(0.0 1.0 2.0 3.0 4.0 5.0 0.1 1.1 2.1 3.1 4.1 5.1 )
export NUMAP=${NUMA_PMAP[$PALS_LOCAL_RANKID]}
export NUMAH=${NUMA_HMAP[$PALS_LOCAL_RANKID]}
export gpu_id=${GPU_MAP[$PALS_LOCAL_RANKID]}
unset EnableWalkerPartition
export EnableImplicitScaling=0
export ZE_AFFINITY_MASK=$gpu_id
export ONEAPI_DEVICE_FILTER=gpu,level_zero
export SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=0
export SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1
export SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE=0:4
export SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE_FOR_D2D_COPY=1
echo "rank $PALS_RANKID ; local rank $PALS_LOCAL_RANKID ; ZE_AFFINITY_MASK=$ZE_AFFINITY_MASK ; NUMA $NUMA "
if [ $PALS_RANKID = "0" ]
then
numactl -p $NUMAP -N $NUMAP unitrace --chrome-kernel-logging --chrome-mpi-logging --chrome-sycl-logging --demangle "$@"
else
numactl -p $NUMAP -N $NUMAP "$@"
fi
--------------------------------------------------------------------
Tursa/nsight-sys
--------------------------------------------------------------------
============================================================================
2026-08-28 libfabric memory-registration-cache "memhooks" monitor -- DEFECTIVE
(libfabric issue #11451, filed by PB, HPE JIRA opened)
============================================================================
STATUS: the ALPS/aarch64 case below is DEMONSTRATED (issue + both fixes verified
there). The Frontier case is a HYPOTHESIS as of 2026-08-28: the NO_TRANSLATION
failure is observed, its attribution to the memhooks MR cache is by mechanism and
by analogy, and the fix is UNTESTED here. Test ladder, one knob per run, NRHS=12:
(1) FineSloppyComms=0 (2) FI_MR_CACHE_MONITOR=kdreg2 (3) FI_MR_CACHE_MAX_COUNT=0
(4) --disable-accelerator-aware-mpi build.
REVISIT this entry with the result; if (2) does not cure it, remove kdreg2 from the
job scripts' justification (it stays as OLCF's recommendation regardless).
2026-08-28 22:18, job 5372414: step (2) DID NOT cure it -- kdreg2 set (and the
environment already carried FI_MR_CACHE_MAX_COUNT=786432), sloppy comms OFF, still
NO_TRANSLATION at outer step 60 (Couter, Waitall count=2).
CORRECTION 2026-08-28 23:xx (fi_mr(3) man page): FI_MR_CACHE_MONITOR governs SYSTEM
memory only; device (HMEM_ROCR) registrations are monitored by
FI_MR_ROCR_CACHE_MONITOR_ENABLED=0|1. Grid's comms window is hipMalloc'd device
memory, so the kdreg2 step tested nothing relevant -- the ALPS PLT bug (#11451) and
the Frontier NO_TRANSLATION are the same *library* but different monitors.
Attributed by gdb on a hung rank (5371826): main thread in MPI_Waitall <- Grid
0x5d539c3 (StencilSendToRecvFromComplete/CommsComplete class), i.e. the HALO
EXCHANGE; the two aborts' request counts (16 = fine Dhop, 2 = coarse PaddedCell
direction) agree. Next: FI_MR_CACHE_MAX_COUNT=0 (all memory) and
FI_MR_ROCR_CACHE_MONITOR_ENABLED=0, one per cell.
cat << EOF > mpiwrapper.sh
#!/bin/bash
lrank=\$OMPI_COMM_WORLD_LOCAL_RANK
numa1=\$(( 2 * \$lrank ))
numa2=\$(( 2 * \$lrank + 1 ))
export CUDA_VISIBLE_DEVICES=\$lrank
export UCX_NET_DEVICES=mlx5_\${lrank}:1
if [ \$OMPI_COMM_WORLD_RANK = "0" ]; then echo \$*; fi
Symptom (Frontier, x86, Cray MPICH 8.1.x, ROCm 7.2): device-buffer MPI fails
with
MPI_Waitall ... MPIDI_OFI_handle_cq_error: OFI poll failed
(ofi_events.c:MPIDI_OFI_handle_cq_error: Input/output error - NO_TRANSLATION)
on a LIVE, never-freed hipMalloc buffer (the sloppy-comms compressed halo buffer,
a static deviceVector) once the process does sustained hipMalloc/hipFree churn
(MemoryManager eviction at NRHS>=12, EvictAll/DropCache). Translations cached by
the provider go stale.
Symptom (ALPS/CSCS, aarch64 Grace/H200, cray-mpich 8.1.32, libfabric 1.22): the
memhooks monitor intercepts munmap by PATCHING THE PLT at MPI_Init
(ofi_memhooks_start -> ofi_write_patch with a garbage data_size); the write overruns
munmap@plt into the NEXT PLT entry (MPI_Comm_dup in Grid's binary), leaving
`br x15` where an adrp belongs -> segfault on first MPI_Comm_dup. Diagnosed with
a hardware watchpoint on the PLT entry during MPI_Init; see the issue for the
gdb transcript. Reproduce on one node with MPICH_SINGLE_HOST_ENABLED=0.
Fix (either): export FI_MR_CACHE_MONITOR=kdreg2 (kernel-driven invalidation; OLCF's
own recommendation, NOT the default)
export FI_MR_CACHE_MAX_COUNT=0 (no registration cache at all)
On ALPS this is a CORRECTNESS requirement for MPI on Slingshot, not a tuning.
Every systems/Frontier job and sourceme now sets kdreg2 on OLCF's recommendation;
whether it resolves the Frontier NO_TRANSLATION is the pending test above.
if [ "1" = "1" ]
then
if [ \$OMPI_COMM_WORLD_RANK = "0" ]
then
exec numactl --interleave=\$numa1,\$numa2 nsys profile \
-o $TRACE/nsys.\$SLURM_JOB_ID --force-overwrite=true \
--trace=cuda,nvtx,mpi,osrt --mpi-impl=op
--sample=none --cpuctxsw=none \
--delay=140 --duration=20 --stats=true \
\$*
else
exec numactl --interleave=\$numa1,\$numa2 \$*
fi
else
exec numactl --interleave=\$numa1,\$numa2 \
fi
EOF
chmod +x ./mpiwrapper.sh
mpirun -np $SLURM_NTASKS -x LD_LIBRARY_PATH --bind-to none ./mpiwrapper.sh <binary> <args>
+7 -7
View File
@@ -8,7 +8,7 @@
then repeatedly applies FFT_all_dim to the same propagator 400 times.
If PlannedFFT is working correctly the RSS should remain flat after the first
iteration — no new plans, no new deviceVector allocations beyond the per-call
iteration -- no new plans, no new deviceVector allocations beyond the per-call
pencil buffer which is freed at the end of each FFT_dim_execute call.
Build exactly like any other Grid test, e.g.:
@@ -63,7 +63,7 @@ static long getGPUUsedMb()
}
// ============================================================
// Convenience struct — one snapshot of both sides
// Convenience struct -- one snapshot of both sides
// ============================================================
struct MemSnapshot {
long cpu_rss_kb; // host RSS in kB (-1 if unavailable)
@@ -113,7 +113,7 @@ int main(int argc, char **argv)
<< "Grid is setup to use " << threads << " threads" << std::endl;
// ------------------------------------------------------------------
// Grid setup — use whatever lattice/mpi/simd was passed on the CLI,
// Grid setup -- use whatever lattice/mpi/simd was passed on the CLI,
// e.g. --grid 8.8.8.8 --mpi 1.1.1.1
// ------------------------------------------------------------------
Coordinate latt_size = GridDefaultLatt();
@@ -174,13 +174,13 @@ int main(int argc, char **argv)
<< std::endl;
// ------------------------------------------------------------------
// Create the PlannedFFT — plans are allocated here ONCE for all
// Create the PlannedFFT -- plans are allocated here ONCE for all
// dimensions and stored inside the object.
// ------------------------------------------------------------------
PlannedFFT<iSpinColourMatrix<vComplexD>> plannedFFT(&GRID);
// ------------------------------------------------------------------
// Snapshot AFTER plan construction — this is the true baseline
// Snapshot AFTER plan construction -- this is the true baseline
// for the loop, because cufftPlanMany itself grabs device memory.
// ------------------------------------------------------------------
MemSnapshot snap_after_plan = takeSnapshot();
@@ -197,7 +197,7 @@ int main(int argc, char **argv)
// ------------------------------------------------------------------
// 400-iteration loop.
// Each iteration computes the full 4d forward FFT of `prop`.
// We deliberately do NOT cache the result — we always start from
// We deliberately do NOT cache the result -- we always start from
// the same `prop` so the FFT is recomputed identically each time.
// The point is to watch memory, not correctness.
// ------------------------------------------------------------------
@@ -229,7 +229,7 @@ int main(int argc, char **argv)
}
// cudaMemGetInfo reflects the state *after* any pooled frees have
// been committed, so this is accurate without an explicit sync —
// been committed, so this is accurate without an explicit sync --
// FFT_dim_execute already calls accelerator_barrier() internally.
MemSnapshot snap_now = takeSnapshot();
printRow(i, snap_now, snap_prev);
+1 -1
View File
@@ -156,7 +156,7 @@ int main (int argc, char ** argv)
}
////////////////////////////////////////////////////
// Dwf matrix — verify Fourier representation using PlannedFFT<LatticeFermionD>
// Dwf matrix -- verify Fourier representation using PlannedFFT<LatticeFermionD>
////////////////////////////////////////////////////
{
std::cout<<"****************************************"<<std::endl;
+1 -1
View File
@@ -20,7 +20,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
//////////////////////////////////////////////////////////////////////////////
// Regression gate for BlockCyclicLayout -- stage 1 of the 2D distributed
// dense inverse (documentation/DistributedDenseInverse2D.tex).
// dense inverse.
//
// The layout is pure index arithmetic, so this test is EXHAUSTIVE rather
// than statistical: every stage sweeps a battery of (N, nb, Pr, Pc)
@@ -2,7 +2,7 @@
Grid physics library, www.github.com/paboyle/Grid
Source file: ./tests/debug/Test_coarse_v2.cc
Source file: ./tests/debug/Test_coarse.cc
Copyright (C) 2026
@@ -27,17 +27,18 @@ Author: Peter Boyle <pboyle@bnl.gov>
/* END LEGAL */
//
// MultiGeneralCoarsenedOperatorV2 against the existing mrhs coarse operator.
// MultiGeneralCoarsenedOperator against the existing mrhs coarse operator.
//
// V1 is constructed on the D+1 grid as now; V2 on the D dimensional grid,
// with SetGrid() adopting the caller owned D+1 grid and building its padded
// cell and neighbour table from the D dimensional stencil, with the Nrhs
// factor multiplied in.
// The reference (the deprecated operator) is constructed on the D+1 grid; the
// operator under test on the D dimensional grid, with SetGrid() adopting the
// caller owned D+1 grid and building its padded cell and neighbour table from
// the D dimensional stencil, with the Nrhs factor multiplied in.
//
// Both are given identical matrix elements, so any difference in the apply is
// the restructured neighbour table. The same geometry object is passed to
// both: V1 adds one to skip for the rhs direction, V2 uses it as is on the D
// dimensional grid, so both describe the same stencil over the D dimensions.
// both: the reference adds one to skip for the rhs direction, the operator
// under test uses it as is on the D dimensional grid, so both describe the
// same stencil over the D dimensions.
//
#include <Grid/Grid.h>
@@ -48,8 +49,8 @@ const int nbasis = 8;
typedef vSpinColourVector FineObj;
typedef sTComplexD CComplexT; // unvectorised coarse space
typedef MultiGeneralCoarsenedMatrix <FineObj,CComplexT,nbasis> MrhsV1;
typedef MultiGeneralCoarsenedOperatorV2<FineObj,CComplexT,nbasis> MrhsV2;
typedef DeprecatedMultiGeneralCoarsenedMatrix <FineObj,CComplexT,nbasis> RefOperator;
typedef MultiGeneralCoarsenedOperator<FineObj,CComplexT,nbasis> TestOperator;
////////////////////////////////////////////////////////////////////////
// Identical random matrix elements into both operators
@@ -59,11 +60,16 @@ void SeedMatrixElements(OpA &A,OpB &B,int npoint,GridSerialRNG &sRNG)
{
typedef typename OpA::calcMatrix calcMatrix;
deviceVector<calcMatrix> bufA,bufB;
for(int p=0;p<npoint;p++){
GRID_ASSERT(A.BLAS_A[p].size() == B.BLAS_A[p].size());
// MatrixPointOut/In carry one stencil point in lSite order whichever
// internal layout the operator uses.
A.MatrixPointOut(p,bufA);
B.MatrixPointOut(p,bufB);
GRID_ASSERT(bufA.size() == bufB.size());
int64_t sites = A.BLAS_A[p].size();
int64_t sites = bufA.size();
std::vector<calcMatrix> host(sites);
ComplexD *w = (ComplexD *)&host[0];
@@ -75,8 +81,9 @@ void SeedMatrixElements(OpA &A,OpB &B,int npoint,GridSerialRNG &sRNG)
w[i] = ComplexD(re-0.5,im-0.5);
}
acceleratorCopyToDevice(&host[0],&A.BLAS_A[p][0],sites*sizeof(calcMatrix));
acceleratorCopyToDevice(&host[0],&B.BLAS_A[p][0],sites*sizeof(calcMatrix));
acceleratorCopyToDevice(&host[0],&bufA[0],sites*sizeof(calcMatrix));
A.MatrixPointIn(p,bufA);
B.MatrixPointIn(p,bufA);
}
}
@@ -85,10 +92,12 @@ RealD MatrixChecksum(Op &O,int npoint)
{
typedef typename Op::calcMatrix calcMatrix;
RealD sum=0.0;
deviceVector<calcMatrix> buf;
for(int p=0;p<npoint;p++){
int64_t sites = O.BLAS_A[p].size();
O.MatrixPointOut(p,buf);
int64_t sites = buf.size();
std::vector<calcMatrix> host(sites);
acceleratorCopyFromDevice(&O.BLAS_A[p][0],&host[0],sites*sizeof(calcMatrix));
acceleratorCopyFromDevice(&buf[0],&host[0],sites*sizeof(calcMatrix));
ComplexD *w = (ComplexD *)&host[0];
int64_t words = sites*sizeof(calcMatrix)/sizeof(ComplexD);
for(int64_t i=0;i<words;i++) sum += real(w[i])*real(w[i]) + imag(w[i])*imag(w[i]);
@@ -107,7 +116,7 @@ int main (int argc, char ** argv)
Coordinate cmpi = GridDefaultMpi();
////////////////////////////////////////////////
// D dimensional coarse grid, and D+1 for V1
// D dimensional coarse grid, and D+1 for the reference
////////////////////////////////////////////////
GridCartesian *CoarseD = new GridCartesian(clatt,csimd,cmpi);
@@ -129,35 +138,35 @@ int main (int argc, char ** argv)
}
GridCartesian *CoarseMulti = new GridCartesian(mlatt,msimd,mmpi);
MrhsV2 OpV2(geom,CoarseD);
MrhsV1 OpV1(geom,CoarseMulti);
OpV2.SetGrid(CoarseMulti);
TestOperator OpTest(geom,CoarseD);
RefOperator OpRef(geom,CoarseMulti);
OpTest.SetGrid(CoarseMulti);
std::cout << GridLogMessage << "coarse D+1 grid nrhs " << nrhs
<< " Nsimd " << CoarseMulti->Nsimd() << std::endl;
std::cout << GridLogMessage << "npoint V1 " << OpV1.geom.npoint
<< " npoint V2 " << OpV2.geom.npoint << std::endl;
GRID_ASSERT(OpV1.geom.npoint == OpV2.geom.npoint);
std::cout << GridLogMessage << "npoint ref " << OpRef.geom.npoint
<< " npoint test " << OpTest.geom.npoint << std::endl;
GRID_ASSERT(OpRef.geom.npoint == OpTest.geom.npoint);
int npoint = OpV1.geom.npoint;
int npoint = OpRef.geom.npoint;
////////////////////////////////////////////////
// Identical matrix elements
////////////////////////////////////////////////
GridSerialRNG sRNG; sRNG.SeedFixedIntegers(std::vector<int>({7,8,9,10}));
SeedMatrixElements(OpV1,OpV2,npoint,sRNG);
SeedMatrixElements(OpRef,OpTest,npoint,sRNG);
RealD ckV1 = MatrixChecksum(OpV1,npoint);
RealD ckV2 = MatrixChecksum(OpV2,npoint);
std::cout << GridLogMessage << "matrix element checksum V1 " << ckV1
<< " V2 " << ckV2 << std::endl;
GRID_ASSERT( ckV1 == ckV2 );
RealD ckRef = MatrixChecksum(OpRef,npoint);
RealD ckTest = MatrixChecksum(OpTest,npoint);
std::cout << GridLogMessage << "matrix element checksum ref " << ckRef
<< " test " << ckTest << std::endl;
GRID_ASSERT( ckRef == ckTest );
////////////////////////////////////////////////
// Same input, compare the applies
////////////////////////////////////////////////
typedef MrhsV1::CoarseVector CoarseVector;
typedef RefOperator::CoarseVector CoarseVector;
// RNG on the D dimensional grid fills any D+1 field: the rhs direction is
// undistributed and divides cleanly. One RNG serves every Nrhs.
GridParallelRNG pRNG(CoarseD); pRNG.SeedFixedIntegers(std::vector<int>({1,2,3,4}));
@@ -167,22 +176,22 @@ int main (int argc, char ** argv)
CoarseVector out2(CoarseMulti);
CoarseVector err (CoarseMulti);
OpV1.M(in,out1);
OpV2.M(in,out2);
OpRef.M(in,out1);
OpTest.M(in,out2);
err = out1 - out2;
std::cout << GridLogMessage << "|V1 out|^2 = " << norm2(out1)
<< " |V2 out|^2 = " << norm2(out2) << std::endl;
std::cout << GridLogMessage << "|V1 - V2|^2 = " << norm2(err) << std::endl;
std::cout << GridLogMessage << "|ref out|^2 = " << norm2(out1)
<< " |test out|^2 = " << norm2(out2) << std::endl;
std::cout << GridLogMessage << "|ref - test|^2 = " << norm2(err) << std::endl;
GRID_ASSERT( norm2(out1) > 0.0 );
GRID_ASSERT( norm2(err) == 0.0 );
////////////////////////////////////////////////
// SetGrid is idempotent on pointer identity
////////////////////////////////////////////////
OpV2.SetGrid(CoarseMulti);
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 );
OpV2.M(in,out2);
OpTest.SetGrid(CoarseMulti);
GRID_ASSERT( MatrixChecksum(OpTest,npoint) == ckTest );
OpTest.M(in,out2);
err = out1 - out2;
GRID_ASSERT( norm2(err) == 0.0 );
std::cout << GridLogMessage << "SetGrid idempotent on identity" << std::endl;
@@ -194,7 +203,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////
// Nrhs 1 is the single RHS case through the multiRHS path, and each slice
// of the Nrhs 4 apply must come back unchanged.
OpV2.M(in,out2);
OpTest.M(in,out2);
for(int nr=2;nr>=1;nr--){
Coordinate latt2(1,nr), simd2(1,1), mpi2(1,1);
for(int d=0;d<Nd;d++){
@@ -204,9 +213,9 @@ int main (int argc, char ** argv)
}
GridCartesian *CoarseMulti2 = new GridCartesian(latt2,simd2,mpi2);
OpV2.SetGrid(CoarseMulti2);
GRID_ASSERT( OpV2.Nrhs() == nr );
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 );
OpTest.SetGrid(CoarseMulti2);
GRID_ASSERT( OpTest.Nrhs() == nr );
GRID_ASSERT( MatrixChecksum(OpTest,npoint) == ckTest );
CoarseVector in2 (CoarseMulti2);
CoarseVector out(CoarseMulti2);
@@ -215,7 +224,7 @@ int main (int argc, char ** argv)
ExtractSliceFast(slice,in,r,0);
InsertSliceFast(slice,in2,r,0);
}
OpV2.M(in2,out);
OpTest.M(in2,out);
RealD sdiff=0.0;
for(int r=0;r<nr;r++){
@@ -231,22 +240,22 @@ int main (int argc, char ** argv)
GRID_ASSERT( norm2(out) > 0.0 );
GRID_ASSERT( sdiff/norm2(out) < 1.0e-20 );
OpV2.ReleaseGrid();
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 ); // survives release
OpTest.ReleaseGrid();
GRID_ASSERT( MatrixChecksum(OpTest,npoint) == ckTest ); // survives release
delete CoarseMulti2;
}
OpV2.SetGrid(CoarseMulti);
GRID_ASSERT( OpV2.Nrhs() == nrhs );
OpTest.SetGrid(CoarseMulti);
GRID_ASSERT( OpTest.Nrhs() == nrhs );
OpV2.M(in,out2);
OpTest.M(in,out2);
err = out1 - out2;
std::cout << GridLogMessage << "after Nrhs 4 -> 2 -> 1 -> release -> 4, |V1 - V2|^2 = "
std::cout << GridLogMessage << "after Nrhs 4 -> 2 -> 1 -> release -> 4, |ref - test|^2 = "
<< norm2(err) << std::endl;
GRID_ASSERT( norm2(err) == 0.0 );
std::cout << GridLogMessage << "Test_coarse_v2: ALL PASS" << std::endl;
std::cout << GridLogMessage << "Test_coarse: ALL PASS" << std::endl;
Grid_finalize();
}
@@ -2,7 +2,7 @@
Grid physics library, www.github.com/paboyle/Grid
Source file: ./tests/debug/Test_coarse_v2_coarsen.cc
Source file: ./tests/debug/Test_coarse_coarsen.cc
Copyright (C) 2026
@@ -29,11 +29,13 @@ Author: Peter Boyle <pboyle@bnl.gov>
//
// Coarsen the same fine operator two ways and compare the matrix elements.
//
// V1 : existing mrhs CoarsenOperator, vectorised coarse space matched to
// the fine SIMD layout, single RHS fine applications
// V2 : D+1 CoarsenOperator, unvectorised sComplexD coarse space, the batch
// of phased basis vectors carried in the rhs direction and applied
// through MrhsPromotedOperator
// ref : DeprecatedMultiGeneralCoarsenedMatrix::CoarsenOperator, vectorised
// coarse space matched to the fine SIMD layout, single RHS fine
// applications
// test : MultiGeneralCoarsenedOperator::CoarsenOperator on the D+1 grid,
// unvectorised sComplexD coarse space, the batch of phased basis
// vectors carried in the rhs direction and applied through
// MrhsPromotedOperator
//
// BLAS_A is written by GridtoBLAS in lSite order, which does not depend on
// the SIMD layout, so the two are directly comparable.
@@ -46,21 +48,25 @@ const int nbasis = 8;
const int batch = 9;
typedef vSpinColourVector FineObj;
typedef vTComplex CComplexV; // vectorised coarse space, V1 reference
typedef sTComplexD CComplexS; // unvectorised coarse space, V2
typedef vTComplex CComplexV; // vectorised coarse space, reference
typedef sTComplexD CComplexS; // unvectorised coarse space, under test
typedef MultiGeneralCoarsenedMatrix <FineObj,CComplexV,nbasis> MrhsV1;
typedef MultiGeneralCoarsenedOperatorV2<FineObj,CComplexS,nbasis> MrhsV2;
typedef DeprecatedMultiGeneralCoarsenedMatrix <FineObj,CComplexV,nbasis> RefOperator;
typedef MultiGeneralCoarsenedOperator<FineObj,CComplexS,nbasis> TestOperator;
template<class Op>
void ReadMatrix(Op &O,int npoint,std::vector<std::vector<typename Op::calcMatrix> > &host)
{
typedef typename Op::calcMatrix calcMatrix;
host.resize(npoint);
// MatrixPointOut carries one stencil point in lSite order whichever internal
// layout the operator uses.
deviceVector<calcMatrix> buf;
for(int p=0;p<npoint;p++){
int64_t sites = O.BLAS_A[p].size();
O.MatrixPointOut(p,buf);
int64_t sites = buf.size();
host[p].resize(sites);
acceleratorCopyFromDevice(&O.BLAS_A[p][0],&host[p][0],sites*sizeof(calcMatrix));
acceleratorCopyFromDevice(&buf[0],&host[p][0],sites*sizeof(calcMatrix));
}
}
@@ -83,24 +89,24 @@ int main (int argc, char ** argv)
GridRedBlackCartesian *FrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(FineGrid);
////////////////////////////////////////////////
// V1 coarse space: SIMD layout matched to the fine
// Reference coarse space: SIMD layout matched to the fine
////////////////////////////////////////////////
Coordinate cvsimd = GridDefaultSimd(Nd,CComplexV::Nsimd());
GridCartesian *CoarseV = new GridCartesian(clatt,cvsimd,fmpi);
// V1 puts all of the SIMD in the rhs direction. CoarsenOperator does not
// use the multiRHS grid; it only sizes BLAS_A, so one lane of rhs suffices.
int nrhs_v1 = CComplexV::Nsimd();
Coordinate v1latt(1,nrhs_v1),v1simd(1,CComplexV::Nsimd()),v1mpi(1,1);
// The reference puts all of the SIMD in the rhs direction. CoarsenOperator
// does not use the multiRHS grid; it only sizes BLAS_A, so one lane of rhs suffices.
int nrhs_ref = CComplexV::Nsimd();
Coordinate reflatt(1,nrhs_ref),refsimd(1,CComplexV::Nsimd()),refmpi(1,1);
for(int d=0;d<Nd;d++){
v1latt.push_back(clatt[d]);
v1simd.push_back(1);
v1mpi .push_back(fmpi[d]);
reflatt.push_back(clatt[d]);
refsimd.push_back(1);
refmpi .push_back(fmpi[d]);
}
GridCartesian *CoarseVMulti = new GridCartesian(v1latt,v1simd,v1mpi);
GridCartesian *CoarseVMulti = new GridCartesian(reflatt,refsimd,refmpi);
////////////////////////////////////////////////
// V2 coarse space: unvectorised
// Coarse space under test: unvectorised
////////////////////////////////////////////////
Coordinate cssimd(Nd,1);
GridCartesian *CoarseS = new GridCartesian(clatt,cssimd,fmpi);
@@ -124,8 +130,8 @@ int main (int argc, char ** argv)
std::cout << GridLogMessage << "fine "<<flatt<<" coarse "<<clatt<<" batch "<<batch<<std::endl;
std::cout << GridLogMessage << "Nsimd fine "<<FineGrid->Nsimd()
<< " coarse V1 "<<CoarseV->Nsimd()
<< " coarse V2 "<<CoarseS->Nsimd()<<std::endl;
<< " coarse ref "<<CoarseV->Nsimd()
<< " coarse test "<<CoarseS->Nsimd()<<std::endl;
////////////////////////////////////////////////
// Fine operator
@@ -151,32 +157,32 @@ int main (int argc, char ** argv)
NextToNearestStencilGeometry4D geomS(CoarseS);
////////////////////////////////////////////////
// V1 coarsening, matched layouts
// Reference coarsening, matched layouts
////////////////////////////////////////////////
MrhsV1 OpV1(geomV,CoarseVMulti);
std::cout << GridLogMessage << "V1 CoarsenOperator" << std::endl;
OpV1.CoarsenOperator(HermOp,Subspace,CoarseV);
RefOperator OpRef(geomV,CoarseVMulti);
std::cout << GridLogMessage << "reference CoarsenOperator" << std::endl;
OpRef.CoarsenOperator(HermOp,Subspace,CoarseV);
////////////////////////////////////////////////
// V2 coarsening, D+1 fine applications, unvectorised coarse
// Coarsening under test, D+1 fine applications, unvectorised coarse
////////////////////////////////////////////////
MrhsV2 OpV2(geomS,CoarseS);
OpV2.SetGrid(CoarseSMulti);
TestOperator OpTest(geomS,CoarseS);
OpTest.SetGrid(CoarseSMulti);
MrhsPromotedOperator<LatticeFermionD> MrhsHermOp(HermOp,FineGrid,batch);
std::cout << GridLogMessage << "V2 CoarsenOperator (D+1)" << std::endl;
OpV2.CoarsenOperator(MrhsHermOp,FineGridMulti,subspace,CoarseS);
std::cout << GridLogMessage << "test CoarsenOperator (D+1)" << std::endl;
OpTest.CoarsenOperator(MrhsHermOp,FineGridMulti,subspace,CoarseS);
////////////////////////////////////////////////
// Compare matrix elements
////////////////////////////////////////////////
int npoint = OpV1.geom.npoint;
GRID_ASSERT(npoint == OpV2.geom.npoint);
typedef MrhsV1::calcMatrix calcMatrix;
int npoint = OpRef.geom.npoint;
GRID_ASSERT(npoint == OpTest.geom.npoint);
typedef RefOperator::calcMatrix calcMatrix;
std::vector<std::vector<calcMatrix> > A1,A2;
ReadMatrix(OpV1,npoint,A1);
ReadMatrix(OpV2,npoint,A2);
ReadMatrix(OpRef,npoint,A1);
ReadMatrix(OpTest,npoint,A2);
RealD num=0.0, den=0.0;
for(int p=0;p<npoint;p++){
@@ -190,8 +196,8 @@ int main (int argc, char ** argv)
den += real(w1[i])*real(w1[i])+imag(w1[i])*imag(w1[i]);
}
}
std::cout << GridLogMessage << "|A_V1|^2 = " << den << std::endl;
std::cout << GridLogMessage << "|A_V1 - A_V2|^2 / |A_V1|^2 = " << num/den << std::endl;
std::cout << GridLogMessage << "|A_ref|^2 = " << den << std::endl;
std::cout << GridLogMessage << "|A_ref - A_test|^2 / |A_ref|^2 = " << num/den << std::endl;
GRID_ASSERT( den > 0.0 );
GRID_ASSERT( num/den < 1.0e-20 );
@@ -200,14 +206,14 @@ int main (int argc, char ** argv)
// batch assembled on the coarse side by the mixed blockProject. Block
// Gram-Schmidt is idempotent so the subspace may be reused in place.
////////////////////////////////////////////////
MrhsV2 OpV2s(geomS,CoarseS);
OpV2s.SetGrid(CoarseSMulti);
TestOperator OpTestSrhs(geomS,CoarseS);
OpTestSrhs.SetGrid(CoarseSMulti);
std::cout << GridLogMessage << "V2 CoarsenOperator (single RHS fine op)" << std::endl;
OpV2s.CoarsenOperator(HermOp,subspace,CoarseS,batch);
std::cout << GridLogMessage << "test CoarsenOperator (single RHS fine op)" << std::endl;
OpTestSrhs.CoarsenOperator(HermOp,subspace,CoarseS,batch);
std::vector<std::vector<calcMatrix> > A3;
ReadMatrix(OpV2s,npoint,A3);
ReadMatrix(OpTestSrhs,npoint,A3);
RealD nums=0.0;
for(int p=0;p<npoint;p++){
@@ -220,7 +226,7 @@ int main (int argc, char ** argv)
nums += real(d)*real(d)+imag(d)*imag(d);
}
}
std::cout << GridLogMessage << "|A_V1 - A_V2srhs|^2 / |A_V1|^2 = " << nums/den << std::endl;
std::cout << GridLogMessage << "|A_ref - A_test_srhs|^2 / |A_ref|^2 = " << nums/den << std::endl;
GRID_ASSERT( nums/den < 1.0e-20 );
////////////////////////////////////////////////
@@ -229,15 +235,15 @@ int main (int argc, char ** argv)
// needs no SetGrid since the matrix elements are Nrhs independent.
////////////////////////////////////////////////
{
typedef MrhsV2::CoarseMatrix CoarseMatrixS;
typedef TestOperator::CoarseMatrix CoarseMatrixS;
std::vector<CoarseMatrixS> Aget(npoint,CoarseS);
for(int p=0;p<npoint;p++) OpV2.GetMatrix(p,Aget);
for(int p=0;p<npoint;p++) OpTest.GetMatrix(p,Aget);
MrhsV2 OpV2rt(geomS,CoarseS);
for(int p=0;p<npoint;p++) OpV2rt.SetMatrix(p,Aget);
TestOperator OpTestRoundTrip(geomS,CoarseS);
for(int p=0;p<npoint;p++) OpTestRoundTrip.SetMatrix(p,Aget);
std::vector<std::vector<calcMatrix> > A4;
ReadMatrix(OpV2rt,npoint,A4);
ReadMatrix(OpTestRoundTrip,npoint,A4);
RealD numrt=0.0;
for(int p=0;p<npoint;p++){
@@ -255,19 +261,19 @@ int main (int argc, char ** argv)
}
////////////////////////////////////////////////
// and V2 applies the matrix it just built
// and the operator under test applies the matrix it just built
////////////////////////////////////////////////
typedef MrhsV2::CoarseVector CoarseVectorS;
typedef TestOperator::CoarseVector CoarseVectorS;
GridParallelRNG cRNG(CoarseS); cRNG.SeedFixedIntegers(std::vector<int>({5,6,7,8}));
CoarseVectorS in(CoarseSMulti); random(cRNG,in);
CoarseVectorS out(CoarseSMulti);
OpV2.M(in,out);
OpTest.M(in,out);
std::cout << GridLogMessage << "|in|^2 = " << norm2(in)
<< " |M_V2 in|^2 = " << norm2(out) << std::endl;
<< " |M_test in|^2 = " << norm2(out) << std::endl;
GRID_ASSERT( norm2(out) > 0.0 );
std::cout << GridLogMessage << "Test_coarse_v2_coarsen: ALL PASS" << std::endl;
std::cout << GridLogMessage << "Test_coarse_coarsen: ALL PASS" << std::endl;
Grid_finalize();
}
+2 -2
View File
@@ -114,7 +114,7 @@ int main (int argc, char ** argv)
///////////////////////////////////////////////////
std::cout<<GridLogMessage << "Building little Dirac operator"<< std::endl;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -223,7 +223,7 @@ int main (int argc, char ** argv)
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
#if 0
MultiGeneralCoarsenedMatrix mrhs(LittleDiracOp,CoarseMrhs);
DeprecatedMultiGeneralCoarsenedMatrix mrhs(LittleDiracOp,CoarseMrhs);
typedef decltype(mrhs) MultiGeneralCoarsenedMatrix_t;
//////////////////////////////////////////
+2 -2
View File
@@ -96,7 +96,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -140,7 +140,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
mrhs.CoarsenOperator(FineHermOp,Aggregates,Coarse5d);
+2 -2
View File
@@ -189,7 +189,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -371,7 +371,7 @@ slurm-1482367.out:Grid : Message : 6169.469330 s : HDCG: Pcg converged in 487 it
//////////////////////////////////////////
// Build a HDCG solver
//////////////////////////////////////////
TwoLevelADEF2<LatticeFermion,CoarseVector,Subspace>
DeprecatedTwoLevelADEF2<LatticeFermion,CoarseVector,Subspace>
HDCG(1.0e-8, 700,
FineHermOp,
CGsmooth,
@@ -218,7 +218,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -269,7 +269,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
std::cout << "**************************************"<<std::endl;
@@ -407,7 +407,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -456,7 +456,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
std::cout << "**************************************"<<std::endl;
@@ -406,7 +406,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -455,7 +455,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
std::cout << "**************************************"<<std::endl;
@@ -215,7 +215,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -161,7 +161,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -210,7 +210,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
std::cout << "**************************************"<<std::endl;
@@ -162,7 +162,7 @@ int main (int argc, char ** argv)
////////////////////////////////////////////////////////////
///////////// Coarse basis and Little Dirac Operator ///////
////////////////////////////////////////////////////////////
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
@@ -217,7 +217,7 @@ int main (int argc, char ** argv)
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
///////////////////////
+1 -1
View File
@@ -254,7 +254,7 @@ int main (int argc, char ** argv)
const int cb = 0 ;
LatticeFermion prom(FGrid);
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNearestStencilGeometry5D geom(Coarse5d);
@@ -376,10 +376,10 @@ int main (int argc, char ** argv)
std::cout <<bl<<" "<<br<<"\t"<<innerProduct(U.subspace[bl],src)<<std::endl;
}}
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
typedef LittleDiracOperatorV::CoarseVector CoarseVectorV;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
V.Orthogonalise();
@@ -375,10 +375,10 @@ int main (int argc, char ** argv)
std::cout <<bl<<" "<<br<<"\t"<<innerProduct(U.subspace[bl],src)<<std::endl;
}}
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
typedef LittleDiracOperatorV::CoarseVector CoarseVectorV;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
V.Orthogonalise();
@@ -369,8 +369,8 @@ int main (int argc, char ** argv)
// }
// typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
// typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
LittleDiracOperator LittleDiracOpPV(geom,FGrid,Coarse5d);
LittleDiracOpPV.CoarsenOperator(PVdagM,V,V);
+1 -1
View File
@@ -185,7 +185,7 @@ int main (int argc, char ** argv)
const int cb = 0 ;
LatticeFermion prom(FGrid);
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NearestStencilGeometry4D geom(Coarse4d);
@@ -185,7 +185,7 @@ int main (int argc, char ** argv)
const int cb = 0 ;
LatticeFermion prom(FGrid);
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NearestStencilGeometry4D geom(Coarse4d);
@@ -182,7 +182,7 @@ int main (int argc, char ** argv)
const int cb = 0 ;
LatticeFermion prom(FGrid);
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NearestStencilGeometry4D geom(Coarse4d);
@@ -182,7 +182,7 @@ int main (int argc, char ** argv)
const int cb = 0 ;
LatticeFermion prom(FGrid);
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NearestStencilGeometry4D geom(Coarse4d);
+5 -5
View File
@@ -3,9 +3,9 @@
*
* Tests three orderings with an empty rocFFT cache to find which GPU
* operation before plan creation triggers the failure:
* A) hipMalloc only — hypothesis: passes (no async GPU work)
* B) hipMalloc + hipMemset — hypothesis: fails (async GPU work in flight)
* C) hipMalloc + hipMemset — hypothesis: passes (work completed before plan)
* A) hipMalloc only -- hypothesis: passes (no async GPU work)
* B) hipMalloc + hipMemset -- hypothesis: fails (async GPU work in flight)
* C) hipMalloc + hipMemset -- hypothesis: passes (work completed before plan)
* + hipDeviceSynchronize
*
* Compile:
@@ -59,13 +59,13 @@ int main(void) {
hipfftResult rvB = makePlan(G, howmany);
printf("G=%-4d B) hipMalloc + hipMemset : %s\n", G, res(rvB));
// C: hipMalloc + hipMemset + sync — does syncing before plan creation fix it?
// C: hipMalloc + hipMemset + sync -- does syncing before plan creation fix it?
hipMemset(buf, 0, nelems * sizeof(hipfftDoubleComplex));
hipDeviceSynchronize();
hipfftResult rvC = makePlan(G, howmany);
printf("G=%-4d C) hipMalloc + hipMemset + sync: %s\n", G, res(rvC));
// A last: hipMalloc only, no async GPU work — should always pass
// A last: hipMalloc only, no async GPU work -- should always pass
hipfftResult rvA = makePlan(G, howmany);
printf("G=%-4d A) hipMalloc only : %s\n\n", G, res(rvA));
+1 -1
View File
@@ -33,7 +33,7 @@ int main(void) {
int n[] = {G};
long nelems = (long)G * howmany;
// Plan created BEFORE hipMalloc — succeeds for all G
// Plan created BEFORE hipMalloc -- succeeds for all G
hipfftHandle p;
size_t workSize = 0;
hipfftCreate(&p);
+2 -2
View File
@@ -38,8 +38,8 @@ static const char *hipfftResultString(hipfftResult r) {
// Plan creation + execution for (G, howmany).
// Tests two orderings to isolate whether a prior hipMalloc poisons hipfft
// plan creation for small G on ROCm 7:
// A) plan BEFORE hipMalloc — hypothesis: succeeds
// B) hipMalloc BEFORE plan — hypothesis: fails for G < 32
// A) plan BEFORE hipMalloc -- hypothesis: succeeds
// B) hipMalloc BEFORE plan -- hypothesis: fails for G < 32
static void tryPlanAndExec(int G, long howmany) {
int n[] = {G};
long nelems = (long)G * howmany;
+3 -3
View File
@@ -58,12 +58,12 @@ static void tryPlanAndExec(int G, long howmany) {
hipfftDoubleComplex *dbuf = nullptr;
hipError_t herr = hipMalloc(&dbuf, nelems * sizeof(hipfftDoubleComplex));
if (herr != hipSuccess) {
printf(" hipMalloc failed (%d) for %ld elems — skipping\n\n", (int)herr, nelems);
printf(" hipMalloc failed (%d) for %ld elems -- skipping\n\n", (int)herr, nelems);
return;
}
hipMemset(dbuf, 0, nelems * sizeof(hipfftDoubleComplex));
// 1. hipfftPlanMany (one-step, nullptr embed) — current Grid path
// 1. hipfftPlanMany (one-step, nullptr embed) -- current Grid path
{
hipfftHandle p;
hipfftResult rv = hipfftPlanMany(&p, 1, n,
@@ -79,7 +79,7 @@ static void tryPlanAndExec(int G, long howmany) {
}
}
// 2. hipfftCreate + hipfftMakePlanMany (two-step) — also current Grid path
// 2. hipfftCreate + hipfftMakePlanMany (two-step) -- also current Grid path
{
hipfftHandle p;
size_t workSize = 0;
+89
View File
@@ -0,0 +1,89 @@
/*************************************************************************************
Grid physics library, www.github.com/paboyle/Grid
Source file: ./tests/debug/Test_mrhs_deflation.cc
Copyright (C) 2026
Author: Peter Boyle <pboyle@bnl.gov>
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.
See the full license in the file "LICENSE" in the top level distribution
directory
*************************************************************************************/
/* END LEGAL */
#include <Grid/Grid.h>
// MultiRHSDeflation: the D+1 multiRHS interface (one permutation pass each
// way) must reproduce the vector-of-D-fields interface bit for bit up to
// the summation order of the GEMMs. Random "eigenvectors" and values: the
// deflation formula G = E (E^dag R)/lambda does not care that they are not
// eigenpairs of anything.
using namespace Grid;
int main (int argc, char ** argv)
{
Grid_init(&argc,&argv);
const int nbasis = 8;
const int nev = 12;
const int nrhs = 5;
typedef iVector<sTComplexD,nbasis> siteVector;
typedef Lattice<siteVector> CoarseVector;
// Unvectorised D and D+1 coarse grids, as MGCoarseGrids builds them
Coordinate latt({1,4,4,4,4});
Coordinate simd({1,1,1,1,1});
Coordinate mpi = GridDefaultMpi();
Coordinate mpi5({1,mpi[0],mpi[1],mpi[2],mpi[3]});
GridCartesian Coarse5d(latt,simd,mpi5);
Coordinate latt6({nrhs,1,4,4,4,4});
Coordinate simd6({1,1,1,1,1,1});
Coordinate mpi6({1,1,mpi[0],mpi[1],mpi[2],mpi[3]});
GridCartesian CoarseMrhs(latt6,simd6,mpi6);
GridParallelRNG RNG(&Coarse5d); RNG.SeedFixedIntegers(std::vector<int>({1,2,3,4}));
std::vector<CoarseVector> evec(nev,&Coarse5d);
std::vector<RealD> eval(nev);
for(int e=0;e<nev;e++){ random(RNG,evec[e]); eval[e] = 0.5 + 0.1*e; }
std::vector<CoarseVector> src(nrhs,&Coarse5d), guess(nrhs,&Coarse5d);
for(int r=0;r<nrhs;r++) random(RNG,src[r]);
MultiRHSDeflation<CoarseVector> Deflator;
Deflator.ImportEigenBasis(evec,eval);
// Reference: the vector interface
Deflator.DeflateSources(src,guess);
// The D+1 interface on the same sources
CoarseVector src_mrhs(&CoarseMrhs), guess_mrhs(&CoarseMrhs);
for(int r=0;r<nrhs;r++) InsertSliceFast(src[r],src_mrhs,r,0);
guess_mrhs = Zero();
Deflator.DeflateSources(src_mrhs,guess_mrhs);
RealD worst = 0.0;
CoarseVector g(&Coarse5d), d(&Coarse5d);
for(int r=0;r<nrhs;r++){
ExtractSliceFast(g,guess_mrhs,r,0);
d = g - guess[r];
RealD rel = std::sqrt(norm2(d)/norm2(guess[r]));
std::cout << GridLogMessage << "rhs " << r << " ||guess_mrhs - guess||/||guess|| = " << rel << std::endl;
worst = std::max(worst,rel);
}
std::cout << GridLogMessage << "MultiRHSDeflation D+1 vs vector interface: worst " << worst << std::endl;
GRID_ASSERT( worst < 1.0e-12 );
std::cout << GridLogMessage << "PASS" << std::endl;
Grid_finalize();
return 0;
}
+2 -2
View File
@@ -30,7 +30,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
// The DenseCoarseMatrix GLUE test, on a real (tiny) lattice coarse
// operator, CPU laptop build.
//
// Builds a genuine GeneralCoarsenedMatrix (DWF MdagM + 0.5 shift for a
// Builds a genuine DeprecatedGeneralCoarsenedMatrix (DWF MdagM + 0.5 shift for a
// guaranteed-invertible Galerkin coarse op, random aggregation basis)
// and runs the whole Import certificate chain through the glue:
// - fresh ImportDense + IMPORT CERTIFICATE vs Op.M
@@ -143,7 +143,7 @@ int main (int argc, char ** argv)
Aggregates.CreateSubspaceRandom(RNG5);
std::cout << GridLogMessage << "Coarsening shifted MdagM" << std::endl;
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
typedef LittleDiracOperator::CoarseVector CoarseVector;
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
LittleDiracOperator LittleDiracOp(geom,FGrid,Coarse5d);
+207
View File
@@ -0,0 +1,207 @@
/*************************************************************************************
Grid physics library, www.github.com/paboyle/Grid
Source file: ./tests/debug/Test_sloppy_dagger.cc
Copyright (C) 2026
Author: Peter Boyle <pboyle@bnl.gov>
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.
See the full license in the file "LICENSE" in the top level distribution
directory
*************************************************************************************/
/* END LEGAL */
//
// Isolate a sloppy-comms dagger defect: on a CPU/GEN multi-rank build the
// sloppy dagger halo can be wrong by ~5% of norm2, deterministically, which
// Benchmark_dwf's sloppy pass catches as a failed dagger Cshift check.
//
// The sloppy compressed-buffer pool (StencilBuffer::DeviceCommBuf) is a
// SHARED STATIC across all stencil objects, so scenarios contaminate each
// other inside one process: this program runs exactly ONE scenario per
// invocation, selected with --seq <name>:
//
// fresh-dag sloppy dagger is the operator's first call
// nodag-dag sloppy nodag x3 (one source), then sloppy dagger (same source)
// nodag-dag-fresh nodag x3 (fresh source each), then dagger (fresh source)
// nodag-dag-newsrc nodag x3 (one source), then dagger (different source)
// dag-nodag-dag dagger, nodag x3, dagger (dagger-first history)
//
// Every call is checked against a separate never-sloppy reference
// operator (exact Dhop == the Cshift construction at 1e-31, certified by
// Benchmark_dwf's exact pass). The last dagger's error is fingerprinted
// by slice along each MPI-decomposed direction.
//
// OMP_NUM_THREADS=1 mpirun -n 2 ./Test_sloppy_dagger --grid 8.8.8.16 \
// --mpi 1.1.1.2 --seq nodag-dag
//
#include <Grid/Grid.h>
using namespace std;
using namespace Grid;
int main (int argc, char ** argv)
{
Grid_init(&argc,&argv);
std::string seq("nodag-dag");
if( GridCmdOptionExists(argv,argv+argc,"--seq") )
seq = GridCmdOptionPayload(argv,argv+argc,"--seq");
const int Ls=8;
GridCartesian * UGrid = SpaceTimeGrid::makeFourDimGrid(GridDefaultLatt(),
GridDefaultSimd(Nd,vComplex::Nsimd()),
GridDefaultMpi());
GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
GridParallelRNG RNG4(UGrid); RNG4.SeedFixedIntegers({1,2,3,4});
GridParallelRNG RNG5(FGrid); RNG5.SeedFixedIntegers({5,6,7,8});
LatticeGaugeField Umu(UGrid);
SU<Nc>::HotConfiguration(RNG4,Umu);
RealD mass=0.1, M5=1.8;
DomainWallFermionD Dref(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5); // never sloppy
DomainWallFermionD Dw (Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5); // under test
Dw.SloppyComms(1);
Coordinate mpi = GridDefaultMpi();
LatticeFermionD src(FGrid), exact(FGrid), sloppy(FGrid), err(FGrid);
// One call: fresh or reused source, dag or not; always checked.
LatticeFermionD srcA(FGrid); gaussian(RNG5,srcA);
auto call = [&](int dag, int fresh, const char *tag){
if ( fresh ) gaussian(RNG5,src); else src = srcA;
Dref.Dhop(src,exact,dag);
Dw.Dhop(src,sloppy,dag);
err = sloppy - exact;
std::cout << GridLogMessage << "SEQ[" << seq << "] " << tag
<< (dag?" DAG ":" NODAG ") << " norm2(err) " << norm2(err) << std::endl;
};
// First bad site: coordinate + exact-vs-sloppy hex words (PB: look for
// expected/actual mismatching in the high-order half word).
auto firstbad = [&](LatticeFermionD &ex, LatticeFermionD &sl){
typedef LatticeFermionD::scalar_object sobj;
Coordinate gdims = FGrid->GlobalDimensions();
int64_t gsites = FGrid->gSites();
for(int64_t g=0; g<gsites; g++){
Coordinate gcoor(5);
Lexicographic::CoorFromIndex(gcoor,g,gdims);
sobj se, ss_;
peekSite(se,ex,gcoor);
peekSite(ss_,sl,gcoor);
uint64_t *we=(uint64_t *)&se;
uint64_t *ws=(uint64_t *)&ss_;
int nw = sizeof(sobj)/8;
for(int w=0;w<nw;w++){
double de=((double *)we)[w], ds=((double *)ws)[w];
if ( fabs(de-ds) > 1.0e-4 ) {
std::cout << GridLogMessage << "FIRSTBAD site " << gcoor
<< " word " << w << " (spin "<<(w/6)%4<<" col "<<(w/2)%3<<" reim "<<(w%2)<<")"
<< " exact " << std::hex << we[w]
<< " sloppy " << ws[w] << std::dec << std::endl;
for(int w2=0;w2<8&&w2<nw;w2++)
std::cout << GridLogMessage << " word["<<w2<<"] exact "<<std::hex<<we[w2]
<<" sloppy "<<ws[w2]<<std::dec<<std::endl;
return;
}
}
}
std::cout << GridLogMessage << "FIRSTBAD: none above 1e-4" << std::endl;
};
auto fingerprint = [&](void){
for(int mu=0;mu<Nd;mu++){
if ( mpi[mu] == 1 ) continue;
std::vector<RealD> sn;
sliceNorm(sn,err,mu+1);
std::cout << GridLogMessage << "SEQ[" << seq << "] err by slice of dim " << mu << ":";
for(int t=0;t<(int)sn.size();t++) std::cout << " " << sn[t];
std::cout << std::endl;
}
};
if ( seq == "fresh-dag" ) {
call(1,0,"call0");
} else if ( seq == "nodag-dag" ) {
call(0,0,"call0"); call(0,0,"call1"); call(0,0,"call2");
call(1,0,"call3");
} else if ( seq == "nodag-dag-fresh" ) {
call(0,1,"call0"); call(0,1,"call1"); call(0,1,"call2");
call(1,1,"call3");
} else if ( seq == "nodag-dag-newsrc" ) {
call(0,0,"call0"); call(0,0,"call1"); call(0,0,"call2");
call(1,1,"call3");
} else if ( seq == "dag-nodag-dag" ) {
call(1,0,"call0");
call(0,0,"call1"); call(0,0,"call2"); call(0,0,"call3");
call(1,0,"call4");
} else if ( seq == "noleave" ) {
// NO exact-operator call between sloppy calls: references precomputed.
LatticeFermionD refnodag(FGrid), refdag(FGrid);
Dref.Dhop(srcA,refnodag,DaggerNo);
Dref.Dhop(srcA,refdag,DaggerYes);
for(int i=0;i<2;i++){
Dw.Dhop(srcA,sloppy,DaggerNo);
err = sloppy - refnodag;
std::cout << GridLogMessage << "SEQ[noleave] call" << i << " NODAG norm2(err) " << norm2(err) << std::endl;
if ( i==0 ) firstbad(refnodag,sloppy);
}
Dw.Dhop(srcA,sloppy,DaggerYes);
err = sloppy - refdag;
std::cout << GridLogMessage << "SEQ[noleave] call3 DAG norm2(err) " << norm2(err) << std::endl;
} else if ( seq == "twoop" ) {
// A SECOND sloppy operator shares the static pool; sequence runs on it
// with interleaved exact references (as the honest runs had).
DomainWallFermionD Dw2(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5);
Dw2.SloppyComms(1);
Dw.Dhop(srcA,sloppy,DaggerNo); // prime the FIRST sloppy op's state
for(int i=0;i<3;i++){
Dref.Dhop(srcA,exact,DaggerNo);
Dw2.Dhop(srcA,sloppy,DaggerNo);
err = sloppy - exact;
std::cout << GridLogMessage << "SEQ[twoop] call" << i << " NODAG norm2(err) " << norm2(err) << std::endl;
}
Dref.Dhop(srcA,exact,DaggerYes);
Dw2.Dhop(srcA,sloppy,DaggerYes);
err = sloppy - exact;
std::cout << GridLogMessage << "SEQ[twoop] call3 DAG norm2(err) " << norm2(err) << std::endl;
} else if ( seq == "self-exact" ) {
// Exact call on the SAME object (Dw), opposite sense, then sloppy:
// distinguishes per-object from shared state.
LatticeFermionD refnodag(FGrid);
Dref.Dhop(srcA,refnodag,DaggerNo); // reference only, early
Dw.SloppyComms(0);
Dw.Dhop(srcA,sloppy,DaggerYes); // exact DAG on Dw itself
Dw.SloppyComms(1);
Dw.Dhop(srcA,sloppy,DaggerNo); // sloppy NODAG: sense mismatch
err = sloppy - refnodag;
std::cout << GridLogMessage << "SEQ[self-exact] sloppy NODAG after own exact DAG: " << norm2(err) << std::endl;
} else if ( seq == "cold" ) {
// NO exact Dhop anywhere before the sloppy calls.
LatticeFermionD outn(FGrid), outd(FGrid);
Dw.Dhop(srcA,outn,DaggerNo);
Dw.Dhop(srcA,outd,DaggerYes);
LatticeFermionD refnodag(FGrid), refdag(FGrid);
Dref.Dhop(srcA,refnodag,DaggerNo);
Dref.Dhop(srcA,refdag,DaggerYes);
err = outn - refnodag;
std::cout << GridLogMessage << "SEQ[cold] first-ever sloppy NODAG: " << norm2(err) << std::endl;
err = outd - refdag;
std::cout << GridLogMessage << "SEQ[cold] then sloppy DAG: " << norm2(err) << std::endl;
} else {
std::cout << GridLogMessage << "unknown --seq " << seq << std::endl;
}
fingerprint();
Grid_finalize();
}
+6 -6
View File
@@ -397,7 +397,7 @@ public:
MdagMLinearOperator<CoarseDiracMatrix, CoarseVector> coarseMdagMOp(_CoarseMatrix);
std::cout << GridLogMG << " Level " << _CurrentLevel << ": **************************************************" << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": MG correctness check: 0 == (M - (Mdiag + Σ_μ Mdir_μ)) * v" << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": MG correctness check: 0 == (M - (Mdiag + Sum_mu Mdir_mu)) * v" << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": **************************************************" << std::endl;
random(_LevelInfo.PRNGs[_CurrentLevel], fineTmps[0]);
@@ -406,22 +406,22 @@ public:
fineMdagMOp.OpDiag(fineTmps[0], fineTmps[2]); // Mdiag * v
fineTmps[4] = zero;
for(int dir = 0; dir < 4; dir++) { // Σ_μ Mdir_μ * v
for(int dir = 0; dir < 4; dir++) { // Sum_mu Mdir_mu * v
for(auto disp : {+1, -1}) {
fineMdagMOp.OpDir(fineTmps[0], fineTmps[3], dir, disp);
fineTmps[4] = fineTmps[4] + fineTmps[3];
}
}
fineTmps[5] = fineTmps[2] + fineTmps[4]; // (Mdiag + Σ_μ Mdir_μ) * v
fineTmps[5] = fineTmps[2] + fineTmps[4]; // (Mdiag + Sum_mu Mdir_mu) * v
fineTmps[6] = fineTmps[1] - fineTmps[5];
auto deviation = std::sqrt(norm2(fineTmps[6]) / norm2(fineTmps[1]));
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(M * v) = " << norm2(fineTmps[1]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Mdiag * v) = " << norm2(fineTmps[2]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Σ_μ Mdir_μ * v) = " << norm2(fineTmps[4]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2((Mdiag + Σ_μ Mdir_μ) * v) = " << norm2(fineTmps[5]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Sum_mu Mdir_mu * v) = " << norm2(fineTmps[4]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2((Mdiag + Sum_mu Mdir_mu) * v) = " << norm2(fineTmps[5]) << std::endl;
std::cout << GridLogMG << " Level " << _CurrentLevel << ": relative deviation = " << deviation;
if(deviation > tolerance) {
@@ -664,7 +664,7 @@ createMGInstance(WilsonMGParams &mgParams, LevelInfo &levelInfo, Matrix &FineMat
CASE_FOR_N_LEVELS(3);
CASE_FOR_N_LEVELS(4);
default:
std::cout << GridLogError << "We currently only support nLevels ∈ {2, 3, 4}" << std::endl;
std::cout << GridLogError << "We currently only support nLevels in {2, 3, 4}" << std::endl;
exit(EXIT_FAILURE);
break;
}