mirror of
https://github.com/paboyle/Grid.git
synced 2026-10-10 01:38:06 +01:00
Updates
This commit is contained in:
Peter Boyle committed
1 parent
3bd5e8883f
commit
3a3a20b8c9
11 files changed
+4506
-41
No files matched your search
@@ -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
|
||||
|
||||
File diff suppressed because it is too large.
Load diff
@@ -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;
|
||||
}
|
||||
@@ -0,0 +1,990 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./examples/Example_pvdagm_3level_DenseCoarseMatrix.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.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
|
||||
//
|
||||
// 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.
|
||||
//
|
||||
// Differences from Example_pvdagm_mrhs_3level_DenseCoarseMatrix.cc:
|
||||
//
|
||||
// * The coarse space is UNVECTORISED (sComplexD). The fine space stays
|
||||
// vectorised. MultiRHSBlockProject carries the mixed layout.
|
||||
//
|
||||
// * 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. 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
|
||||
// orthonormalises them IN PLACE -- the vectors are far too large to copy
|
||||
// defensively -- so rawNull is taken first and the RAW vectors are what
|
||||
// define the L2 null space. Do not insert an Orthogonalise() anywhere:
|
||||
// projecting a block-orthonormal vector onto its own block-orthonormalised
|
||||
// aggregation gives e_k, and the near null content is silently gone. The
|
||||
// ||<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 DEPRECATED_CHECK MRHS_COARSEN
|
||||
//
|
||||
|
||||
#include <Grid/Grid.h>
|
||||
#include <algorithm> // std::sort, for the runtime-environment dump in ParseEnvironment
|
||||
#include <Grid/lattice/PaddedCell.h>
|
||||
#include <Grid/stencil/GeneralLocalStencil.h>
|
||||
#include <Grid/algorithms/iterative/PrecGeneralisedConjugateResidualNonHermitian.h>
|
||||
#include <Grid/algorithms/multigrid/MrhsMultiGrid.h>
|
||||
#include <Grid/algorithms/multigrid/DenseCoarseMatrix.h>
|
||||
|
||||
#include <memory>
|
||||
|
||||
using namespace std;
|
||||
using namespace Grid;
|
||||
|
||||
// Compile time so it can be cut down for laptop runs: -DNBASIS=8
|
||||
#ifndef NBASIS
|
||||
#define NBASIS 60
|
||||
#endif
|
||||
|
||||
RealD mass = 0.00078;
|
||||
int Nrhs = 12;
|
||||
int Ls = 24;
|
||||
int CoarsenBatch = 9;
|
||||
std::vector<int> lat_size({48,48,48,96});
|
||||
|
||||
// Solver tuning. PRINCIPLE (PB, 2026-08-24): the defaults ARE the current
|
||||
// optimum, so an unset environment reproduces the best banked result; they
|
||||
// are updated as and when a better point is found, and every change is
|
||||
// dated here. Environment variables of the same names override for sweeps.
|
||||
//
|
||||
// Current optimum: 2026-08-24, slurm-5335492 F4, 48^3x96 Ls=24 on 288 GCDs,
|
||||
// 17.2 s/RHS at Nrhs=4, 32.2 s at Nrhs=1 (exact-halo FINAL ~1e-8 pending
|
||||
// the exact-outer rerun). Smoother mmax == order (full GCR history);
|
||||
// PB's mmax=1 trial gave 72 vs ~60 outer iterations and was slower.
|
||||
RealD FineSmootherShift = 0.1;
|
||||
int FineSmootherOrder = 6;
|
||||
int FineSmootherMmax = 6;
|
||||
RealD CoarseSmootherShift = 0.1;
|
||||
int PowerIterations = 0; // >0: power-iterate the smoother operators before the solves (spectral edge)
|
||||
// Smoother implementation per level (Smoothers.h):
|
||||
// gcr : the adaptive PGCR (default, as always)
|
||||
// replay : run the PGCR with a coefficient recorder for the first
|
||||
// PolyRecordIters outer steps, then switch to GCRReplaySmoother
|
||||
// (same polynomial, no inner products) -- 1402.2585 p.13 revisited
|
||||
// cheb : ChebyshevNonHermitianSmoother, 1/x on [ChebLo,ChebHi], order
|
||||
// = the GCR step count of that level
|
||||
std::string FineSmootherMode = "gcr";
|
||||
std::string CoarseSmootherMode = "gcr";
|
||||
int PolyRecordIters = 4;
|
||||
int PolyRecordStart = 0; // outer step at which recording begins (0: from the first step)
|
||||
std::string PolyRecordSelect = "last"; // which recorded call to replay: last|first|mean (mean is the bad one)
|
||||
int PolyRefresh = 0; // >0: every PolyRefresh outer steps, one adaptive step re-records the polynomial (HDCG: every 10)
|
||||
int PolyVerbose = 0; // 1: fixed-polynomial smoothers print |r_m|/|r_0| per call
|
||||
RealD FineChebLo = 3.0, FineChebHi = 137.0; // from the harvested polynomial and PowerIterations edge
|
||||
RealD CoarseChebLo = 8.0, CoarseChebHi = 45.0;
|
||||
int CoarseSmootherNstep = 2;
|
||||
int CoarseSmootherMmax = 2;
|
||||
RealD CoarseSolverTol = 0.05;
|
||||
int CoarseSolverOrder = 200;
|
||||
int CoarseSolverMmax = 16;
|
||||
RealD OuterTol = 1.0e-8;
|
||||
int OuterMmax = 4;
|
||||
int OuterNstep = 8;
|
||||
|
||||
// "It's legal to get the same answer faster, not to get a less correct
|
||||
// answer." (PB, 2026-08-24)
|
||||
//
|
||||
// Halo-precision POLICY: reduced-precision (fp32 wire)
|
||||
// halos belong in the PRECONDITIONER -- the smoother, the V-cycle's own
|
||||
// residuals, and the coarsening -- and NEVER in the outer Krylov. The
|
||||
// outer operator's applications define what "converged" means; making them
|
||||
// sloppy turns the stopping criterion into a statement about the wrong
|
||||
// operator (measured: solver stops at computed 9.8e-9 while the true
|
||||
// residual is 3.4e-8). Exactness costs one exact fine matvec per outer
|
||||
// iteration, a few percent of the solve.
|
||||
//
|
||||
// Stencil::SloppyComms is a free runtime setter (Stencil.h:303), so the
|
||||
// policy is implemented by SCOPED toggling: SetFineSloppy(1) on entering
|
||||
// the preconditioner / coarsening, SetFineSloppy(0) on leaving. The
|
||||
// operators default to EXACT. FineSloppyComms therefore now means
|
||||
// "sloppy inside the preconditioner"; =0 makes everything exact.
|
||||
int FineSloppyComms = 1;
|
||||
// SmootherCoeffLog=1 : print the GCR step lengths a_k and orthogonalisation
|
||||
// coefficients b_kj of BOTH smoothers every call -- the harvest for a fixed
|
||||
// polynomial smoother (stable coefficients => stationary p(A), no reductions).
|
||||
int SmootherCoeffLog = 0;
|
||||
std::function<void(int)> SetFineSloppy = [](int){};
|
||||
|
||||
void ParseEnvironment(void)
|
||||
{
|
||||
if(getenv("MASS")) mass = atof(getenv("MASS"));
|
||||
if(getenv("NRHS")) Nrhs = atoi(getenv("NRHS"));
|
||||
if(getenv("LS")) Ls = atoi(getenv("LS"));
|
||||
if(getenv("COARSEN_BATCH")) CoarsenBatch= atoi(getenv("COARSEN_BATCH"));
|
||||
if(getenv("FineSmootherShift")) FineSmootherShift = atof(getenv("FineSmootherShift"));
|
||||
if(getenv("FineSmootherOrder")) FineSmootherOrder = atoi(getenv("FineSmootherOrder"));
|
||||
if(getenv("FineSmootherMmax")) FineSmootherMmax = atoi(getenv("FineSmootherMmax"));
|
||||
if(getenv("CoarseSmootherShift"))CoarseSmootherShift= atof(getenv("CoarseSmootherShift"));
|
||||
if(getenv("PowerIterations")) PowerIterations = atoi(getenv("PowerIterations"));
|
||||
if(getenv("FineSmootherMode")) FineSmootherMode = getenv("FineSmootherMode");
|
||||
if(getenv("CoarseSmootherMode")) CoarseSmootherMode = getenv("CoarseSmootherMode");
|
||||
if(getenv("PolyRecordIters")) PolyRecordIters = atoi(getenv("PolyRecordIters"));
|
||||
if(getenv("PolyRecordStart")) PolyRecordStart = atoi(getenv("PolyRecordStart"));
|
||||
if(getenv("PolyRecordSelect")) PolyRecordSelect = getenv("PolyRecordSelect");
|
||||
if(getenv("PolyRefresh")) PolyRefresh = atoi(getenv("PolyRefresh"));
|
||||
if(getenv("PolyVerbose")) PolyVerbose = atoi(getenv("PolyVerbose"));
|
||||
if(getenv("FineChebLo")) FineChebLo = atof(getenv("FineChebLo"));
|
||||
if(getenv("FineChebHi")) FineChebHi = atof(getenv("FineChebHi"));
|
||||
if(getenv("CoarseChebLo")) CoarseChebLo = atof(getenv("CoarseChebLo"));
|
||||
if(getenv("CoarseChebHi")) CoarseChebHi = atof(getenv("CoarseChebHi"));
|
||||
if(getenv("CoarseSmootherNstep"))CoarseSmootherNstep= atoi(getenv("CoarseSmootherNstep"));
|
||||
if(getenv("CoarseSmootherMmax")) CoarseSmootherMmax = atoi(getenv("CoarseSmootherMmax"));
|
||||
if(getenv("CoarseSolverTol")) CoarseSolverTol = atof(getenv("CoarseSolverTol"));
|
||||
if(getenv("CoarseSolverOrder")) CoarseSolverOrder = atoi(getenv("CoarseSolverOrder"));
|
||||
if(getenv("CoarseSolverMmax")) CoarseSolverMmax = atoi(getenv("CoarseSolverMmax"));
|
||||
if(getenv("OuterTol")) OuterTol = atof(getenv("OuterTol"));
|
||||
if(getenv("OuterMmax")) OuterMmax = atoi(getenv("OuterMmax"));
|
||||
if(getenv("FineSloppyComms")) FineSloppyComms = atoi(getenv("FineSloppyComms"));
|
||||
if(getenv("SmootherCoeffLog")) SmootherCoeffLog = atoi(getenv("SmootherCoeffLog"));
|
||||
if(getenv("OuterNstep")) OuterNstep = atoi(getenv("OuterNstep"));
|
||||
if(getenv("LATT")){
|
||||
Coordinate l;
|
||||
GridCmdOptionIntVector(std::string(getenv("LATT")),l);
|
||||
GRID_ASSERT(l.size()==4);
|
||||
for(int d=0;d<4;d++) lat_size[d]=l[d];
|
||||
}
|
||||
|
||||
std::cout << GridLogMessage << "PARAM: LATT "
|
||||
<< lat_size[0]<<"."<<lat_size[1]<<"."<<lat_size[2]<<"."<<lat_size[3] << std::endl;
|
||||
std::cout << GridLogMessage << "PARAM: LS " << Ls << std::endl;
|
||||
std::cout << GridLogMessage << "PARAM: MASS " << mass << std::endl;
|
||||
std::cout << GridLogMessage << "PARAM: NBASIS " << NBASIS << std::endl;
|
||||
std::cout << GridLogMessage << "PARAM: NRHS " << Nrhs << std::endl;
|
||||
std::cout << GridLogMessage << "PARAM: COARSEN_BATCH " << CoarsenBatch << std::endl;
|
||||
// EVERY knob this programme parsed, so a log identifies its own run (2026-08-28:
|
||||
// four jobs in the queue differing in FineSloppyComms/NRHS/mmax and none of it
|
||||
// printed). Interim until the Serializable parameter struct replaces all of this.
|
||||
auto P = [](const char *n, auto v){ std::cout << GridLogMessage << "PARAM: " << std::left << std::setw(20) << n << std::right << " " << v << std::endl; };
|
||||
P("FineSmootherShift",FineSmootherShift); P("FineSmootherOrder",FineSmootherOrder); P("FineSmootherMmax",FineSmootherMmax);
|
||||
P("FineSmootherMode",FineSmootherMode); P("FineChebLo",FineChebLo); P("FineChebHi",FineChebHi);
|
||||
P("FineSloppyComms",FineSloppyComms);
|
||||
P("CoarseSmootherShift",CoarseSmootherShift); P("CoarseSmootherNstep",CoarseSmootherNstep); P("CoarseSmootherMmax",CoarseSmootherMmax);
|
||||
P("CoarseSmootherMode",CoarseSmootherMode); P("CoarseChebLo",CoarseChebLo); P("CoarseChebHi",CoarseChebHi);
|
||||
P("CoarseSolverTol",CoarseSolverTol); P("CoarseSolverOrder",CoarseSolverOrder); P("CoarseSolverMmax",CoarseSolverMmax);
|
||||
P("OuterTol",OuterTol); P("OuterMmax",OuterMmax); P("OuterNstep",OuterNstep);
|
||||
P("PowerIterations",PowerIterations); P("PolyRecordIters",PolyRecordIters); P("PolyRecordStart",PolyRecordStart);
|
||||
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","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"};
|
||||
for(auto e : envs) P(e, getenv(e) ? std::string(getenv(e)) : std::string("(unset)"));
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////////////
|
||||
// The COMPLETE fabric/runtime environment, scanned from environ rather than a curated
|
||||
// list. 2026-08-30: we could not tell from a failing log whether FI_CXI_ATS was set,
|
||||
// because it was not in the list -- and a module or the submitting shell can set
|
||||
// anything. A log must describe its own run without reference to the job script.
|
||||
////////////////////////////////////////////////////////////////////////////////////////
|
||||
{
|
||||
extern char **environ;
|
||||
const char *prefixes[] = {"FI_","OFI_","MPICH_","MPIR_","PMI_","CXI_","HSA_","HIP_","ROCR_",
|
||||
"ROCM_","AMD_","GPU_","NUMA_","OMP_","GOMP_","GRID_","CRAY_","LIBFABRIC"};
|
||||
std::vector<std::string> hits;
|
||||
for(char **e = environ; e && *e; e++){
|
||||
std::string s(*e);
|
||||
for(auto p : prefixes){
|
||||
if( s.compare(0,strlen(p),p)==0 ){
|
||||
if(s.size()>200) s = s.substr(0,197)+"..."; // paths can be enormous
|
||||
hits.push_back(s);
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
std::sort(hits.begin(),hits.end());
|
||||
std::cout << GridLogMessage << "PARAM: ---- runtime environment: "<<hits.size()<<" variables set ----"<<std::endl;
|
||||
for(auto &s : hits) std::cout << GridLogMessage << "PARAM: ENV " << s << std::endl;
|
||||
// Variables whose ABSENCE is as informative as their value (defaults bite: an unset
|
||||
// FI_MR_CACHE_MONITOR means libfabric's defective memhooks, see systems/WorkArounds.txt).
|
||||
const char *notable[] = {"FI_MR_CACHE_MONITOR","FI_MR_CACHE_MAX_COUNT","FI_MR_ROCR_CACHE_MONITOR_ENABLED",
|
||||
"FI_CXI_ATS","FI_CXI_RDZV_THRESHOLD","FI_CXI_DISABLE_HMEM_DEV_REGISTER",
|
||||
"FI_HMEM_ROCR_USE_DMABUF","MPICH_GPU_SUPPORT_ENABLED","MPICH_OFI_NIC_POLICY",
|
||||
"MPICH_SMP_SINGLE_COPY_MODE","OMP_NUM_THREADS","GRID_ALLOC_NCACHE_LARGE"};
|
||||
for(auto n : notable) if(!getenv(n))
|
||||
std::cout << GridLogMessage << "PARAM: ENV " << n << " (UNSET - provider/runtime default applies)" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
template <class Field>
|
||||
void saveSubspace(std::vector<Field> &subspace, std::string const fname){
|
||||
#ifdef HAVE_LIME
|
||||
Grid::emptyUserRecord record;
|
||||
Grid::ScidacWriter SW(subspace[0].Grid()->IsBoss());
|
||||
SW.open(fname);
|
||||
for (int k = 0; k < (int)subspace.size(); k++) SW.writeScidacFieldRecord(subspace[k], record);
|
||||
SW.close();
|
||||
#endif
|
||||
}
|
||||
template <class Field>
|
||||
void loadSubspace(std::vector<Field> &subspace, std::string const fname){
|
||||
#ifdef HAVE_LIME
|
||||
Grid::emptyUserRecord record;
|
||||
Grid::ScidacReader SR;
|
||||
SR.open(fname);
|
||||
for (int k = 0; k < (int)subspace.size(); k++) SR.readScidacFieldRecord(subspace[k], record);
|
||||
SR.close();
|
||||
#endif
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// A = PV^dag M (non-Hermitian)
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
template<class Matrix,class Field>
|
||||
class PVdagMLinearOperator : public LinearOperatorBase<Field> {
|
||||
Matrix &_Mat; Matrix &_PV;
|
||||
public:
|
||||
PVdagMLinearOperator(Matrix &Mat,Matrix &PV): _Mat(Mat),_PV(PV) {};
|
||||
void OpDiag (const Field &in, Field &out) { assert(0); }
|
||||
void OpDir (const Field &in, Field &out,int dir,int disp) { assert(0); }
|
||||
void OpDirAll (const Field &in, std::vector<Field> &out){ 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); }
|
||||
};
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ||<v|v> - I||_F over a set of coarse vectors. Small means the raw near null
|
||||
// content survived the projection; see GramGuard for where a leak lands.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
template<class CoarseField>
|
||||
RealD GramDefect(std::vector<CoarseField> &v)
|
||||
{
|
||||
RealD s2=0.0;
|
||||
for(int i=0;i<(int)v.size();i++){
|
||||
for(int j=0;j<(int)v.size();j++){
|
||||
ComplexD sij=TensorRemove(innerProduct(v[i],v[j]));
|
||||
ComplexD d=sij-(i==j?ComplexD(1.0):ComplexD(0.0));
|
||||
s2+=real(d)*real(d)+imag(d)*imag(d);
|
||||
}
|
||||
}
|
||||
return std::sqrt(s2);
|
||||
}
|
||||
|
||||
// On a leak every image collapses to the block unit e_k, the Gram becomes
|
||||
// N*I, and the defect lands at (N-1)*sqrt(nbasis) -- orders above the ~0.2
|
||||
// of a content preserving projection. Trip well below that so a mis-set
|
||||
// threshold costs a log line rather than the run.
|
||||
template<class CoarseField>
|
||||
void GramGuard(const std::string &name,std::vector<CoarseField> &v,GridBase *grid)
|
||||
{
|
||||
RealD defect = GramDefect(v);
|
||||
RealD N = (RealD)grid->gSites();
|
||||
RealD leak = (N-1.0)*std::sqrt((RealD)v.size());
|
||||
RealD trip = std::sqrt(N);
|
||||
std::cout << GridLogMessage << "GUARD: ||<"<<name<<"|"<<name<<"> - I||_F = " << defect
|
||||
<< " (e_k leak would be " << leak << ", trip at " << trip << ")" << std::endl;
|
||||
GRID_ASSERT( defect < trip );
|
||||
}
|
||||
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Shifted variants for the smoothers
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
template<class Matrix,class Field>
|
||||
class ShiftedPVdagMLinearOperator : public LinearOperatorBase<Field> {
|
||||
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) { assert(0); }
|
||||
void OpDir (const Field &in, Field &out,int dir,int disp) { assert(0); }
|
||||
void OpDirAll (const Field &in, std::vector<Field> &out){ 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(tmp,out); _Mat.Mdag(in,tmp); out = out + shift*in; }
|
||||
void HermOpAndNorm(const Field &in, Field &out,RealD &n1,RealD &n2){ assert(0); }
|
||||
void HermOp(const Field &in, Field &out){ Field tmp(in.Grid()); Op(in,tmp); AdjOp(tmp,out); }
|
||||
};
|
||||
|
||||
template<class Field>
|
||||
class ShiftedLinearOperator : public LinearOperatorBase<Field> {
|
||||
LinearOperatorBase<Field> &_Op; RealD shift;
|
||||
public:
|
||||
ShiftedLinearOperator(RealD _shift, LinearOperatorBase<Field> &Op) : _Op(Op), shift(_shift) {}
|
||||
void OpDiag (const Field &in, Field &out) { assert(0); }
|
||||
void OpDir (const Field &in, Field &out,int dir,int disp) { assert(0); }
|
||||
void OpDirAll (const Field &in, std::vector<Field> &out) { 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){ assert(0); }
|
||||
void HermOp (const Field &in, Field &out) { Field tmp(in.Grid()); Op(in,tmp); AdjOp(tmp,out); }
|
||||
};
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Power iteration on a (non-Hermitian) operator: the spectral edge the
|
||||
// smoother polynomial must not exceed. Reports
|
||||
// step 0 : |A v|/|v| on a RANDOM unit v -- a one-sample lower bound on
|
||||
// sigma_max(A). If this and the converged value agree, the
|
||||
// operator is near-normal and the spectral picture (R_m(lambda)
|
||||
// on the spectrum) is trustworthy; if not, the field of values
|
||||
// sets the safe interval and the spectrum understates it.
|
||||
// step k : |A v_k|/|v_k| -> |lambda_max| as v_k -> the dominant
|
||||
// eigenvector; the complex Rayleigh quotient <v,Av> gives its
|
||||
// phase (real => on the axis). A non-converging oscillation
|
||||
// means a complex-conjugate pair of equal modulus at the top.
|
||||
// Uses Op(), not HermOp(): this is the operator the smoother sees.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
template<class Field>
|
||||
void PowerIteration(const std::string &name, LinearOperatorBase<Field> &Op, GridBase *grid, int iters)
|
||||
{
|
||||
GRID_TRACE("PowerIteration");
|
||||
GridParallelRNG RNG(grid); RNG.SeedFixedIntegers(std::vector<int>({7,11,13,17}));
|
||||
Field v(grid), Av(grid);
|
||||
gaussian(RNG,v);
|
||||
RealD nv = std::sqrt(norm2(v)); v = v*(1.0/nv);
|
||||
RealD ratio=0.0, ratio0=0.0; ComplexD rq(0.0);
|
||||
for(int i=0;i<iters;i++){
|
||||
Op.Op(v,Av);
|
||||
RealD nAv = std::sqrt(norm2(Av));
|
||||
rq = innerProduct(v,Av); // v is unit
|
||||
ratio = nAv;
|
||||
if ( i==0 ) {
|
||||
ratio0 = ratio;
|
||||
std::cout << GridLogMessage << "PowerIteration " << name << " step 0 (random v): |Av|/|v| = " << ratio
|
||||
<< " [lower bound on sigma_max]" << std::endl;
|
||||
}
|
||||
if ( (i%10==0) || (i==iters-1) )
|
||||
std::cout << GridLogMessage << "PowerIteration " << name << " step " << i << " |Av|/|v| = " << ratio
|
||||
<< " Rayleigh = (" << real(rq) << "," << imag(rq) << ")" << std::endl;
|
||||
v = Av*(1.0/nAv);
|
||||
}
|
||||
std::cout << GridLogMessage << "PowerIteration " << name << " SUMMARY: |lambda_max| ~ " << ratio
|
||||
<< " Rayleigh (" << real(rq) << "," << imag(rq) << ")"
|
||||
<< " phase " << std::atan2(imag(rq),real(rq)) << " rad"
|
||||
<< " step-0 ratio / converged = " << ratio0/ratio
|
||||
<< (ratio0/ratio > 1.2 ? " ** non-normal: sigma_max well above |lambda_max| **" : " (near-normal)")
|
||||
<< std::endl;
|
||||
}
|
||||
|
||||
|
||||
int main (int argc, char ** argv)
|
||||
{
|
||||
Grid_init(&argc,&argv);
|
||||
ParseEnvironment();
|
||||
|
||||
RealD M5=1.8, b=1.5, c=0.5;
|
||||
const int nbasis=NBASIS;
|
||||
const int nrhs=Nrhs;
|
||||
const int batch=CoarsenBatch;
|
||||
|
||||
Coordinate mpi = GridDefaultMpi();
|
||||
Coordinate fsimd= GridDefaultSimd(Nd,vComplex::Nsimd());
|
||||
|
||||
GridCartesian * UGrid = SpaceTimeGrid::makeFourDimGrid(lat_size,fsimd,mpi);
|
||||
GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
|
||||
GridCartesian * FGrid = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
|
||||
GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
|
||||
|
||||
// Level 1 blocking (default 2^4)
|
||||
Coordinate clatt = lat_size;
|
||||
Coordinate Block({2,2,3,3}); // L1->L2 blocking; banked optimum 2026-08-24 (env BLOCK overrides)
|
||||
if ( getenv("BLOCK") ){ GridCmdOptionIntVector(std::string(getenv("BLOCK")),Block); GRID_ASSERT(Block.size()==4); }
|
||||
for(int d=0;d<4;d++){ GRID_ASSERT(lat_size[d]%Block[d]==0); clatt[d]=lat_size[d]/Block[d]; }
|
||||
std::cout << GridLogMessage << "Block " << Block << " coarse lattice " << clatt << std::endl;
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// The coarse space is unvectorised. The 5D coarse grid is built here
|
||||
// rather than through SpaceTimeGrid so the SIMD layout is ours.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
Coordinate c5latt({1,clatt[0],clatt[1],clatt[2],clatt[3]});
|
||||
Coordinate c5simd({1,1,1,1,1});
|
||||
Coordinate c5mpi ({1,mpi[0],mpi[1],mpi[2],mpi[3]});
|
||||
GridCartesian *Coarse5d = new GridCartesian(c5latt,c5simd,c5mpi);
|
||||
|
||||
// 6D coarse multiRHS grid: rhs is dim 0, undistributed and unvectorised.
|
||||
// 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]});
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(cmlatt,cmsimd,cmmpi);
|
||||
|
||||
// 6D coarse grid at the coarsening batch, used only while CoarsenOperator
|
||||
// runs. The matrix elements survive the change back to nrhs.
|
||||
Coordinate cblatt({batch,1,clatt[0],clatt[1],clatt[2],clatt[3]});
|
||||
GridCartesian *CoarseBatch = new GridCartesian(cblatt,cmsimd,cmmpi);
|
||||
|
||||
// 6D fine grid carrying the coarsening batch: fine SIMD layout preserved
|
||||
Coordinate fmlatt({batch,Ls,lat_size[0],lat_size[1],lat_size[2],lat_size[3]});
|
||||
Coordinate fmsimd({1,1,fsimd[0],fsimd[1],fsimd[2],fsimd[3]});
|
||||
Coordinate fmmpi ({1,1,mpi[0],mpi[1],mpi[2],mpi[3]});
|
||||
GridCartesian *FineMrhs = new GridCartesian(fmlatt,fmsimd,fmmpi);
|
||||
|
||||
std::cout << GridLogMessage << "Nsimd fine " << FGrid->Nsimd()
|
||||
<< " coarse " << Coarse5d->Nsimd() << std::endl;
|
||||
|
||||
GridParallelRNG RNG4(UGrid); RNG4.SeedFixedIntegers({1,2,3,4});
|
||||
GridParallelRNG RNG5(FGrid); RNG5.SeedFixedIntegers({5,6,7,8});
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Gauge field
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
LatticeGaugeField Umu(UGrid);
|
||||
if ( getenv("HOT_START") ) {
|
||||
std::cout << GridLogMessage << "Hot start gauge field" << std::endl;
|
||||
SU<Nc>::HotConfiguration(RNG4,Umu);
|
||||
} else {
|
||||
std::string file("/ccs/home/poare/ckpoint_lat.1000");
|
||||
if ( getenv("CONFIG") ) file = std::string(getenv("CONFIG"));
|
||||
std::cout << GridLogMessage << "Reading gauge field " << file << std::endl;
|
||||
FieldMetaData header;
|
||||
NerscIO::readConfiguration(Umu,header,file);
|
||||
}
|
||||
|
||||
MobiusFermionD Ddwf(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5,b,c);
|
||||
MobiusFermionD Dpv (Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,1.0, M5,b,c);
|
||||
|
||||
// PVdagM and ShiftedPVdagM are thin wrappers over these same two objects,
|
||||
// so this one callback controls every fine halo in the program. Default
|
||||
// EXACT; the preconditioner and the coarsening turn sloppiness on for
|
||||
// their own scope only (policy note at FineSloppyComms).
|
||||
SetFineSloppy = [&Ddwf,&Dpv](int sloppy){
|
||||
Ddwf.SloppyComms(sloppy);
|
||||
Dpv .SloppyComms(sloppy);
|
||||
};
|
||||
SetFineSloppy(0);
|
||||
std::cout << GridLogMessage << "Fine halo policy: preconditioner+coarsening "
|
||||
<< (FineSloppyComms ? "SLOPPY (fp32 wire)" : "exact")
|
||||
<< ", outer Krylov EXACT" << std::endl;
|
||||
|
||||
typedef PVdagMLinearOperator<MobiusFermionD,LatticeFermionD> PVdagM_t;
|
||||
typedef ShiftedPVdagMLinearOperator<MobiusFermionD,LatticeFermionD> ShiftedPVdagM_t;
|
||||
PVdagM_t PVdagM(Ddwf,Dpv);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Level 1 types: unvectorised coarse scalar
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
typedef sTComplexD CComplexS;
|
||||
typedef MultiGeneralCoarsenedOperator<vSpinColourVector,CComplexS,nbasis> CoarseOperator;
|
||||
typedef CoarseOperator::CoarseVector CoarseVector;
|
||||
typedef Aggregation<vSpinColourVector,CComplexS,nbasis> Subspace;
|
||||
|
||||
NextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Subspace: load RAW (no Orthogonalise!), or generate.
|
||||
//
|
||||
// The Aggregation is scaffolding for CreateSubspaceGCR only. That runs
|
||||
// entirely on the fine grid and ends in GlobalOrthonormalise, which is a
|
||||
// whole-lattice Gram-Schmidt, so the coarse grid it holds is never
|
||||
// dereferenced and may be the unvectorised one.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
std::string subspace_file = "subspace_nb" + std::to_string(nbasis) + ".scidac";
|
||||
if ( getenv("SUBSPACE_FILE") ) subspace_file = std::string(getenv("SUBSPACE_FILE"));
|
||||
uint64_t file_exists=0;
|
||||
if ( UGrid->IsBoss() ){ std::ifstream f(subspace_file); file_exists=f.good()?1:0; }
|
||||
UGrid->GlobalSum(file_exists);
|
||||
|
||||
const int cb=0;
|
||||
Subspace AggregatesGCR(Coarse5d,FGrid,cb);
|
||||
if ( file_exists ){
|
||||
std::cout << GridLogMessage << "*** Loading subspace from disk (kept RAW) ***" << std::endl;
|
||||
loadSubspace(AggregatesGCR.subspace, subspace_file);
|
||||
} else {
|
||||
std::cout << GridLogMessage << "*** GCR subspace generation ***" << std::endl;
|
||||
AggregatesGCR.CreateSubspaceGCR(RNG5,PVdagM,nbasis);
|
||||
saveSubspace(AggregatesGCR.subspace, subspace_file);
|
||||
}
|
||||
|
||||
// RAW copy BEFORE CoarsenOperator block-orthonormalises in place.
|
||||
std::vector<LatticeFermionD> rawNull(nbasis,FGrid);
|
||||
for(int k=0;k<nbasis;k++) rawNull[k]=AggregatesGCR.subspace[k];
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// L1 coarsening. The fine operator is single RHS, so it is promoted to
|
||||
// the 6D batch grid; a natively multiRHS fine operator would substitute
|
||||
// here with no other change.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
CoarseOperator CoarseOpPV(geom,Coarse5d);
|
||||
CoarseOpPV.SetGrid(CoarseBatch);
|
||||
|
||||
std::cout << GridLogMessage << "*** L1 CoarsenOperator, batch "<<batch<<" ***" << std::endl;
|
||||
SetFineSloppy(FineSloppyComms); // coarsening builds the PRECONDITIONER
|
||||
if ( getenv("MRHS_COARSEN") ) {
|
||||
// Promote the single RHS operator and pack the batch: costs an
|
||||
// ExtractSlice/InsertSlice pair per rhs. Here for the A/B only; this is
|
||||
// the path a natively multiRHS fine operator would take.
|
||||
MrhsPromotedOperator<LatticeFermionD> MrhsPVdagM(PVdagM,FGrid,batch);
|
||||
CoarseOpPV.CoarsenOperator(MrhsPVdagM,FineMrhs,AggregatesGCR.subspace,Coarse5d);
|
||||
} else {
|
||||
// PVdagM is single RHS: apply it directly, batch on the coarse side.
|
||||
CoarseOpPV.CoarsenOperator(PVdagM,AggregatesGCR.subspace,Coarse5d,batch);
|
||||
}
|
||||
SetFineSloppy(0);
|
||||
|
||||
// (moved here 2026-08-28: it needs AggregatesGCR.subspace, which is freed right after
|
||||
// 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
|
||||
// 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("DEPRECATED_CHECK") ) {
|
||||
|
||||
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_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);
|
||||
|
||||
NextToNearestStencilGeometry5D geomV(Coarse5dV);
|
||||
|
||||
SubspaceV AggV(Coarse5dV,FGrid,cb);
|
||||
for(int k=0;k<nbasis;k++) AggV.subspace[k]=AggregatesGCR.subspace[k];
|
||||
|
||||
LittleDiracOperator LittleDiracOpPV(geomV,FGrid,Coarse5dV);
|
||||
std::cout << GridLogMessage << "*** deprecated CoarsenOperator (cross check) ***" << std::endl;
|
||||
SetFineSloppy(FineSloppyComms);
|
||||
LittleDiracOpPV.CoarsenOperator(PVdagM,AggV);
|
||||
SetFineSloppy(0);
|
||||
|
||||
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 = mrhsDep.BLAS_A[p].size();
|
||||
CoarseOpPV.MatrixPointOut(p,a2);
|
||||
GRID_ASSERT(sites == (int64_t)a2.size());
|
||||
std::vector<calcMatrix> h1(sites),h2(sites);
|
||||
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);
|
||||
for(int64_t i=0;i<words;i++){
|
||||
ComplexD d=w1[i]-w2[i];
|
||||
num += real(d)*real(d)+imag(d)*imag(d);
|
||||
den += real(w1[i])*real(w1[i])+imag(w1[i])*imag(w1[i]);
|
||||
}
|
||||
}
|
||||
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 );
|
||||
}
|
||||
|
||||
// Stay on the batch grid: the L2 coarsening drives this operator at the
|
||||
// batch. It is switched to the solve Nrhs once L2 is built.
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// psi_coarse = P^dag (RAW fine null) -> Galerkin images that carry the
|
||||
// near null content, and are free: A_c (P psi) = P A psi.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
MultiRHSBlockProject<LatticeFermionD> MrhsProjector;
|
||||
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 (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();
|
||||
|
||||
std::vector<CoarseVector> psi_coarse(nbasis,Coarse5d);
|
||||
MrhsProjector.blockProject(rawNull,psi_coarse); // RAW vectors in
|
||||
rawNull.clear(); rawNull.shrink_to_fit();
|
||||
|
||||
GramGuard("psi_coarse",psi_coarse,Coarse5d);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// STAGE TWO: L2 -> L3.
|
||||
//
|
||||
// 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.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
Coordinate cclatt = clatt;
|
||||
Coordinate Block2({4,4,2,4}); // L2->L3 blocking; banked optimum 2026-08-24 (env BLOCK2 overrides)
|
||||
if ( getenv("BLOCK2") ){ GridCmdOptionIntVector(std::string(getenv("BLOCK2")),Block2); GRID_ASSERT(Block2.size()==4); }
|
||||
for(int d=0;d<4;d++){ GRID_ASSERT(clatt[d]%Block2[d]==0); cclatt[d]=clatt[d]/Block2[d]; }
|
||||
std::cout << GridLogMessage << "Block2 " << Block2 << " coarse-coarse lattice " << cclatt << std::endl;
|
||||
|
||||
Coordinate cc5latt({1,cclatt[0],cclatt[1],cclatt[2],cclatt[3]});
|
||||
GridCartesian *CoarseCoarse5d = new GridCartesian(cc5latt,c5simd,c5mpi);
|
||||
|
||||
Coordinate ccmlatt({nrhs,1,cclatt[0],cclatt[1],cclatt[2],cclatt[3]});
|
||||
GridCartesian *CoarseCoarseMrhs = new GridCartesian(ccmlatt,cmsimd,cmmpi);
|
||||
|
||||
Coordinate ccblatt({batch,1,cclatt[0],cclatt[1],cclatt[2],cclatt[3]});
|
||||
GridCartesian *CoarseCoarseBatch = new GridCartesian(ccblatt,cmsimd,cmmpi);
|
||||
|
||||
// Coarsening deepens the tensor nest by one iScalar
|
||||
typedef CoarseVector::vector_object CoarseSiteObj;
|
||||
typedef iScalar<CComplexS> CComplexS2;
|
||||
typedef MultiGeneralCoarsenedOperator<CoarseSiteObj,CComplexS2,nbasis> CoarseCoarseOperator;
|
||||
typedef CoarseCoarseOperator::CoarseVector CoarseCoarseVector;
|
||||
|
||||
NextToNearestStencilGeometry5D geom2(CoarseCoarse5d);
|
||||
|
||||
// RAW copy of the coarse null vectors, for the same reason as rawNull:
|
||||
// the L2 CoarsenOperator block-orthonormalises its subspace in place, and
|
||||
// the L3 basis must be defined by the vectors that still carry content.
|
||||
std::vector<CoarseVector> rawPsi(nbasis,Coarse5d);
|
||||
for(int k=0;k<nbasis;k++) rawPsi[k]=psi_coarse[k];
|
||||
|
||||
CoarseCoarseOperator CoarseOpL2(geom2,CoarseCoarse5d);
|
||||
CoarseOpL2.SetGrid(CoarseCoarseBatch);
|
||||
|
||||
NonHermitianLinearOperator<CoarseOperator,CoarseVector> LinOpCoarse(CoarseOpPV);
|
||||
|
||||
std::cout << GridLogMessage << "*** L2 CoarsenOperator, batch "<<batch<<" ***" << std::endl;
|
||||
CoarseOpL2.CoarsenOperator(LinOpCoarse,CoarseBatch,psi_coarse,CoarseCoarse5d);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Both operators to the solve Nrhs. The matrix elements are Nrhs
|
||||
// independent and survive the change.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
CoarseOpPV.SetGrid(CoarseMrhs);
|
||||
CoarseOpL2.SetGrid(CoarseCoarseMrhs);
|
||||
std::cout << GridLogMessage << "L1 operator at Nrhs " << CoarseOpPV.Nrhs()
|
||||
<< ", L2 operator at Nrhs " << CoarseOpL2.Nrhs() << std::endl;
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// psi_cc from the RAW coarse null vectors, and the same guard
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
MultiRHSBlockProject<CoarseVector> MrhsProjectorL2;
|
||||
MrhsProjectorL2.Allocate(nbasis,Coarse5d,CoarseCoarse5d);
|
||||
MrhsProjectorL2.ImportBasis(psi_coarse); // block orthonormal basis
|
||||
|
||||
{
|
||||
std::vector<CoarseCoarseVector> psi_cc(nbasis,CoarseCoarse5d);
|
||||
MrhsProjectorL2.blockProject(rawPsi,psi_cc); // RAW vectors in
|
||||
|
||||
GramGuard("psi_cc",psi_cc,CoarseCoarse5d);
|
||||
}
|
||||
rawPsi.clear(); rawPsi.shrink_to_fit();
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Both coarse operators apply on their solve grids
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
{
|
||||
GridParallelRNG cRNG(Coarse5d); cRNG.SeedFixedIntegers({3,4,5,6});
|
||||
CoarseVector cin(CoarseMrhs), cout_(CoarseMrhs);
|
||||
random(cRNG,cin);
|
||||
CoarseOpPV.M(cin,cout_);
|
||||
std::cout << GridLogMessage << "L1 apply |in|^2 = " << norm2(cin)
|
||||
<< " |M in|^2 = " << norm2(cout_) << std::endl;
|
||||
GRID_ASSERT( norm2(cout_) > 0.0 );
|
||||
|
||||
GridParallelRNG ccRNG(CoarseCoarse5d); ccRNG.SeedFixedIntegers({7,8,9,10});
|
||||
CoarseCoarseVector ccin(CoarseCoarseMrhs), ccout(CoarseCoarseMrhs);
|
||||
random(ccRNG,ccin);
|
||||
CoarseOpL2.M(ccin,ccout);
|
||||
std::cout << GridLogMessage << "L2 apply |in|^2 = " << norm2(ccin)
|
||||
<< " |M in|^2 = " << norm2(ccout) << std::endl;
|
||||
GRID_ASSERT( norm2(ccout) > 0.0 );
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// STAGE THREE (part one): the dense bottom on L2.
|
||||
//
|
||||
// DenseCoarseMatrix is bilingual: it takes the elements through
|
||||
// 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.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
typedef DenseCoarseMatrix<CComplexS2,nbasis> DenseCC_t;
|
||||
std::unique_ptr<DenseCC_t> DenseCC;
|
||||
|
||||
if ( getenv("DENSE_CC")==nullptr || atoi(getenv("DENSE_CC")) ) {
|
||||
|
||||
std::cout << GridLogMessage << "*** L3 dense bottom: import from the L2 operator ***" << std::endl;
|
||||
DenseCC.reset(new DenseCC_t(CoarseCoarse5d));
|
||||
DenseCC->Import(CoarseOpL2);
|
||||
|
||||
////////////////////////////////////////////////////////////////////
|
||||
// ||A Ainv x - x|| / ||x||, the check Import could not run itself
|
||||
////////////////////////////////////////////////////////////////////
|
||||
Coordinate cc1latt({1,1,cclatt[0],cclatt[1],cclatt[2],cclatt[3]});
|
||||
GridCartesian *CoarseCoarseOne = new GridCartesian(cc1latt,cmsimd,cmmpi);
|
||||
|
||||
CoarseOpL2.SetGrid(CoarseCoarseOne);
|
||||
|
||||
CoarseCoarseVector x(CoarseCoarse5d),y(CoarseCoarse5d),z(CoarseCoarse5d);
|
||||
GridParallelRNG dRNG(CoarseCoarse5d); dRNG.SeedFixedIntegers({11,12,13,14});
|
||||
random(dRNG,x);
|
||||
|
||||
(*DenseCC)(x,y); // y = Ainv x
|
||||
|
||||
CoarseCoarseVector y1(CoarseCoarseOne),z1(CoarseCoarseOne);
|
||||
InsertSliceFast(y,y1,0,0);
|
||||
CoarseOpL2.M(y1,z1); // z = A y
|
||||
ExtractSliceFast(z,z1,0,0);
|
||||
|
||||
z = z - x;
|
||||
RealD rel = std::sqrt(norm2(z)/norm2(x));
|
||||
std::cout << GridLogMessage << "L3 dense: ||A Ainv x - x||/||x|| = " << rel << std::endl;
|
||||
GRID_ASSERT( rel < 1.0e-2 );
|
||||
|
||||
CoarseOpL2.SetGrid(CoarseCoarseMrhs);
|
||||
delete CoarseCoarseOne;
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// STAGE THREE (part two): the solves.
|
||||
//
|
||||
// Both operators are driven from the SAME objects at whatever Nrhs is
|
||||
// asked for -- the matrix elements were built once and survive SetGrid --
|
||||
// so single RHS and multiRHS are the same code path with a different grid.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
GRID_ASSERT(DenseCC != nullptr); // the PGCR bottom is not ported yet
|
||||
|
||||
typedef PrecGeneralisedConjugateResidualNonHermitian<LatticeFermionD> FineSmoother_t;
|
||||
|
||||
ShiftedPVdagM_t ShiftedPVdagM(FineSmootherShift,Ddwf,Dpv);
|
||||
TrivialPrecon<LatticeFermionD> simple_fine;
|
||||
TrivialPrecon<CoarseVector> simpleC;
|
||||
|
||||
auto RunSolve = [&](int nr)
|
||||
{
|
||||
std::cout << GridLogMessage << "**********************************************" << 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
|
||||
// mRHS GCR holds src, sol, r, Az and OuterMmax x (p,q) fine fields PER RHS; the fine
|
||||
// smoother adds FineSmootherMmax x (p,q) once. The MemoryManager LRU cap
|
||||
// (--device-mem) must be BELOW what is physically left after the non-LRU allocations
|
||||
// (comms buffers, dense slab, stencil buffers), or the device fills before anything
|
||||
// is evicted. Print the estimate, the LRU state, and the device's own free count.
|
||||
{
|
||||
uint64_t fieldBytes = (uint64_t)FGrid->lSites()*sizeof(typename LatticeFermionD::scalar_object);
|
||||
double outerGB = (double)nr*(4 + 2*OuterMmax)*fieldBytes/1.0e9;
|
||||
double smthGB = (double)(2*FineSmootherMmax + 4)*fieldBytes/1.0e9;
|
||||
std::cout << GridLogMessage << "Device budget: fine field " << fieldBytes/1.0e6 << " MB; outer GCR history "
|
||||
<< nr << " x (4 + 2 x " << OuterMmax << ") fields = " << outerGB << " GB; fine smoother history + temps ~ "
|
||||
<< smthGB << " GB; MemoryManager device LRU " << MemoryManager::DeviceCacheBytes()/1.0e9 << " GB now, cap "
|
||||
<< MemoryManager::DeviceMaxBytes/1.0e9 << " GB" << std::endl;
|
||||
// Empty the device LRU: every setup-era Lattice copy (coarse null vectors,
|
||||
// coarsening temporaries) goes back to host, so the solve's working set
|
||||
// starts from a clean device and the cap applies to it alone.
|
||||
MemoryManager::EvictAll();
|
||||
// ...and release the allocation caches' held blocks (setup-era deviceVector
|
||||
// scratch that is "free" to the caller but not to hipMalloc).
|
||||
MemoryManager::DropCache();
|
||||
MemoryManager::PrintBytes();
|
||||
acceleratorMem();
|
||||
}
|
||||
std::cout << GridLogMessage << "**********************************************" << std::endl;
|
||||
|
||||
Coordinate cml({nr,1,clatt[0],clatt[1],clatt[2],clatt[3]});
|
||||
Coordinate ccml({nr,1,cclatt[0],cclatt[1],cclatt[2],cclatt[3]});
|
||||
GridCartesian *CMrhs = new GridCartesian(cml, cmsimd,cmmpi);
|
||||
GridCartesian *CCMrhs = new GridCartesian(ccml,cmsimd,cmmpi);
|
||||
|
||||
CoarseOpPV.SetGrid(CMrhs);
|
||||
CoarseOpL2.SetGrid(CCMrhs);
|
||||
|
||||
NonHermitianLinearOperator<CoarseOperator,CoarseVector> LinOpC (CoarseOpPV);
|
||||
NonHermitianLinearOperator<CoarseCoarseOperator,CoarseCoarseVector> LinOpCC(CoarseOpL2);
|
||||
|
||||
MrhsDenseCCSolve<DenseCC_t,CoarseCoarseVector> ccSolve(*DenseCC,nr);
|
||||
|
||||
ShiftedLinearOperator<CoarseVector> ShiftedC(CoarseSmootherShift, LinOpC);
|
||||
if ( PowerIterations > 0 ) {
|
||||
// Spectral edges the smoother polynomials must respect (see
|
||||
// scripts/gcr_polynomial.py: |R_m|>1 beyond the edge = amplification).
|
||||
PowerIteration<CoarseVector> ("CoarseSmootherOp(shift="+std::to_string(CoarseSmootherShift)+")", ShiftedC, CMrhs, PowerIterations);
|
||||
PowerIteration<CoarseVector> ("CoarseOp(unshifted)", LinOpC, CMrhs, PowerIterations);
|
||||
PowerIteration<LatticeFermionD>("FineSmootherOp(shift="+std::to_string(FineSmootherShift)+")", ShiftedPVdagM, Ddwf.FermionGrid(), PowerIterations);
|
||||
}
|
||||
PrecGeneralisedConjugateResidualNonHermitian<CoarseVector>
|
||||
CoarseSmootherGCR(0.01,1,ShiftedC,simpleC,CoarseSmootherMmax,CoarseSmootherNstep);
|
||||
CoarseSmootherGCR.Level(2); CoarseSmootherGCR.Name("Csmoother"); CoarseSmootherGCR.SetZeroGuess(1);
|
||||
|
||||
SwitchableSmoother<CoarseVector> CoarseSmootherSlot(CoarseSmootherGCR,"Csmoother GCR");
|
||||
MrhsCoarseThreeLevelPrec<CoarseVector,CoarseCoarseVector>
|
||||
L2to3Precon(LinOpC, CoarseSmootherSlot, MrhsProjectorL2, ccSolve,
|
||||
Coarse5d, CoarseCoarse5d, CCMrhs, nr);
|
||||
|
||||
PrecGeneralisedConjugateResidualNonHermitian<CoarseVector>
|
||||
L2PGCR(CoarseSolverTol, CoarseSolverOrder/16, LinOpC, L2to3Precon, CoarseSolverMmax, 16);
|
||||
L2PGCR.Level(2); L2PGCR.Name("Couter"); L2PGCR.SetZeroGuess(1);
|
||||
|
||||
FineSmoother_t SmootherGCR(0.0,1,ShiftedPVdagM,simple_fine,FineSmootherMmax,FineSmootherOrder);
|
||||
SmootherGCR.Level(1); SmootherGCR.Name("Fsmoother"); SmootherGCR.SetZeroGuess(1);
|
||||
SmootherGCR.LogCoefficients(SmootherCoeffLog);
|
||||
CoarseSmootherGCR.LogCoefficients(SmootherCoeffLog);
|
||||
|
||||
SwitchableSmoother<LatticeFermionD> FineSmootherSlot(SmootherGCR,"Fsmoother GCR");
|
||||
MrhsTwoLevelMG<LatticeFermionD,CoarseVector,SwitchableSmoother<LatticeFermionD> >
|
||||
ThreeLevelPrecon(PVdagM, FineSmootherSlot, MrhsProjector, L2PGCR, Coarse5d, CMrhs);
|
||||
ThreeLevelPrecon.SetSloppy = SetFineSloppy;
|
||||
ThreeLevelPrecon.SloppyComms = FineSloppyComms;
|
||||
|
||||
MrhsPGCRNonHermitian<LatticeFermionD>
|
||||
L1PGCR(OuterTol,1000,PVdagM,ThreeLevelPrecon,OuterMmax,OuterNstep);
|
||||
L1PGCR.Level(1); L1PGCR.Name("Fouter"); L1PGCR.SetZeroGuess(1);
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Smoother modes. Objects live for the duration of this RunSolve.
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
std::unique_ptr<ChebyshevNonHermitianSmoother<LatticeFermionD> > FineCheb;
|
||||
std::unique_ptr<ChebyshevNonHermitianSmoother<CoarseVector> > CoarseCheb;
|
||||
std::unique_ptr<GCRReplaySmoother<LatticeFermionD> > FineReplay;
|
||||
std::unique_ptr<GCRReplaySmoother<CoarseVector> > CoarseReplay;
|
||||
GCRCoefficients recF, recC;
|
||||
{
|
||||
GCRCoefficients::Select sel = GCRCoefficients::Last;
|
||||
if ( PolyRecordSelect=="first" ) sel = GCRCoefficients::First;
|
||||
if ( PolyRecordSelect=="mean" ) sel = GCRCoefficients::Mean;
|
||||
recF.select = sel; recC.select = sel;
|
||||
}
|
||||
if ( FineSmootherMode == "cheb" ) {
|
||||
FineCheb.reset(new ChebyshevNonHermitianSmoother<LatticeFermionD>(FineChebLo,FineChebHi,FineSmootherOrder,ShiftedPVdagM));
|
||||
FineCheb->Verbose = PolyVerbose; FineCheb->name = "Fsmoother";
|
||||
FineSmootherSlot.Set(*FineCheb,"Fsmoother Chebyshev");
|
||||
}
|
||||
if ( CoarseSmootherMode == "cheb" ) {
|
||||
CoarseCheb.reset(new ChebyshevNonHermitianSmoother<CoarseVector>(CoarseChebLo,CoarseChebHi,CoarseSmootherNstep,ShiftedC));
|
||||
CoarseCheb->Verbose = PolyVerbose; CoarseCheb->name = "Csmoother";
|
||||
CoarseSmootherSlot.Set(*CoarseCheb,"Csmoother Chebyshev");
|
||||
}
|
||||
// Recording window [PolyRecordStart, PolyRecordStart+PolyRecordIters).
|
||||
// M3 (2026-08-26): the GCR polynomial changes fast over the first outer
|
||||
// steps (per-call |r|/|r0| 0.0034 -> 0.017 over steps 1-4) and the MEAN
|
||||
// of those is a poor smoother (replay 0.033-0.049 per call); record a
|
||||
// settled window instead.
|
||||
if ( PolyRecordStart == 0 ) {
|
||||
if ( FineSmootherMode == "replay" ) SmootherGCR.SetCoefficientRecorder(&recF);
|
||||
if ( CoarseSmootherMode == "replay" ) CoarseSmootherGCR.SetCoefficientRecorder(&recC);
|
||||
}
|
||||
// Record -> replay, with optional periodic re-recording ("re-record, not
|
||||
// fade away": HDCG refreshed its polynomial every 10 steps, tracking the
|
||||
// evolving spectral content of the residual). Schedule on outer steps:
|
||||
// [PolyRecordStart, +PolyRecordIters) adaptive GCR, recording
|
||||
// then replay of the selected recorded call;
|
||||
// if PolyRefresh>0: every PolyRefresh steps, ONE adaptive recording
|
||||
// step, then replay of that call.
|
||||
auto BuildReplays = [&](void){
|
||||
if ( FineSmootherMode == "replay" ) {
|
||||
SmootherGCR.SetCoefficientRecorder(nullptr);
|
||||
recF.Flush(); recF.Report("Fsmoother");
|
||||
FineReplay.reset(new GCRReplaySmoother<LatticeFermionD>(ShiftedPVdagM,recF));
|
||||
FineReplay->Verbose = PolyVerbose; FineReplay->name = "Fsmoother";
|
||||
FineSmootherSlot.Set(*FineReplay,"Fsmoother replay");
|
||||
SmootherGCR.ReleaseHistory();
|
||||
}
|
||||
if ( CoarseSmootherMode == "replay" ) {
|
||||
CoarseSmootherGCR.SetCoefficientRecorder(nullptr);
|
||||
recC.Flush(); recC.Report("Csmoother");
|
||||
CoarseReplay.reset(new GCRReplaySmoother<CoarseVector>(ShiftedC,recC));
|
||||
CoarseReplay->Verbose = PolyVerbose; CoarseReplay->name = "Csmoother";
|
||||
CoarseSmootherSlot.Set(*CoarseReplay,"Csmoother replay");
|
||||
CoarseSmootherGCR.ReleaseHistory();
|
||||
}
|
||||
};
|
||||
auto StartRecording = [&](int step){
|
||||
if ( FineSmootherMode == "replay" ) { { auto sel=recF.select; recF = GCRCoefficients(); recF.select=sel; } SmootherGCR.SetCoefficientRecorder(&recF); FineSmootherSlot.Set(SmootherGCR,"Fsmoother GCR (recording)"); }
|
||||
if ( CoarseSmootherMode == "replay" ) { { auto sel=recC.select; recC = GCRCoefficients(); recC.select=sel; } CoarseSmootherGCR.SetCoefficientRecorder(&recC); CoarseSmootherSlot.Set(CoarseSmootherGCR,"Csmoother GCR (recording)"); }
|
||||
std::cout << GridLogMessage << "Smoother coefficient recording starts at outer step " << step << std::endl;
|
||||
};
|
||||
int switchStep = PolyRecordStart + PolyRecordIters;
|
||||
L1PGCR.OnStep = [&](int step){
|
||||
if ( FineSmootherMode != "replay" && CoarseSmootherMode != "replay" ) return;
|
||||
if ( step == PolyRecordStart && PolyRecordStart > 0 ) StartRecording(step);
|
||||
if ( step == switchStep ) { BuildReplays(); return; }
|
||||
if ( PolyRefresh > 0 && step > switchStep ) {
|
||||
int since = step - switchStep;
|
||||
if ( since % PolyRefresh == 0 ) { StartRecording(step); } // one adaptive, recorded step
|
||||
if ( since % PolyRefresh == 1 ) { BuildReplays(); } // then replay it
|
||||
}
|
||||
};
|
||||
std::cout << GridLogMessage << "Smoother modes: fine " << FineSmootherMode << " coarse " << CoarseSmootherMode
|
||||
<< (FineSmootherMode=="replay"||CoarseSmootherMode=="replay" ? " (record outer steps "+std::to_string(PolyRecordStart)+".."+std::to_string(PolyRecordStart+PolyRecordIters)+")" : "")
|
||||
<< std::endl;
|
||||
|
||||
std::vector<LatticeFermionD> src(nr,FGrid), sol(nr,FGrid);
|
||||
for(int r=0;r<nr;r++){ gaussian(RNG5,src[r]); sol[r]=Zero(); }
|
||||
|
||||
GridStopWatch w; w.Start();
|
||||
L1PGCR(src,sol);
|
||||
w.Stop();
|
||||
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
|
||||
// trust it -- a preconditioner that failed to restore would surface here.
|
||||
SetFineSloppy(0);
|
||||
{ LatticeFermionD Ax(FGrid); RealD worst=0.0;
|
||||
for(int r=0;r<nr;r++){ PVdagM.Op(sol[r],Ax); Ax=Ax-src[r];
|
||||
RealD rn=std::sqrt(norm2(Ax)/norm2(src[r]));
|
||||
std::cout << GridLogMessage << "FINAL Nrhs "<<nr<<": rhs["<<r<<"] true residual = " << rn << std::endl;
|
||||
worst=std::max(worst,rn); }
|
||||
std::cout << GridLogMessage << "FINAL Nrhs "<<nr<<": worst-case residual = " << worst
|
||||
<< " (exact-halo verification)" << std::endl;
|
||||
}
|
||||
|
||||
|
||||
// The operators borrow these grids and build a PaddedCell on them, so
|
||||
// they must let go before the grids are destroyed.
|
||||
CoarseOpPV.ReleaseGrid();
|
||||
CoarseOpL2.ReleaseGrid();
|
||||
delete CMrhs; delete CCMrhs;
|
||||
};
|
||||
|
||||
RunSolve(nrhs);
|
||||
if ( getenv("SOLVE_SRHS")==nullptr || atoi(getenv("SOLVE_SRHS")) ) RunSolve(1);
|
||||
|
||||
std::cout << GridLogMessage << "*** stage three complete: solves done ***" << std::endl;
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
File diff suppressed because it is too large.
Load diff
+47
-38
@@ -206,49 +206,58 @@ Tursa/nsight-sys
|
||||
--------------------------------------------------------------------
|
||||
|
||||
============================================================================
|
||||
2026-08-28 libfabric memory-registration-cache "memhooks" monitor -- DEFECTIVE
|
||||
(libfabric issue #11451, filed by PB, HPE JIRA opened)
|
||||
Piz Daint
|
||||
============================================================================
|
||||
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.
|
||||
|
||||
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
|
||||
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 -> segfault on first MPI_Comm_dup. Diagnosed with
|
||||
a hardware watchpoint on the PLT entry during MPI_Init; see the issue for the
|
||||
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 (kernel-driven invalidation; OLCF's
|
||||
own recommendation, NOT the default)
|
||||
export FI_MR_CACHE_MAX_COUNT=0 (no registration cache at all)
|
||||
|
||||
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 leads to Runtime MPI errors with "NO_TRANSLATION"
|
||||
|
||||
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
|
||||
|
||||
https://github.com/ofiwg/libfabric/issues/12773
|
||||
|
||||
https://github.com/ofiwg/libfabric/issues/12775
|
||||
|
||||
Fix:
|
||||
|
||||
export FI_HMEM_ROCR_USE_DMABUF=0
|
||||
|
||||
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.
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -0,0 +1,261 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./tests/debug/Test_coarse.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.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
|
||||
//
|
||||
// MultiGeneralCoarsenedOperator against the existing mrhs coarse operator.
|
||||
//
|
||||
// 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: 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>
|
||||
|
||||
using namespace Grid;
|
||||
|
||||
const int nbasis = 8;
|
||||
|
||||
typedef vSpinColourVector FineObj;
|
||||
typedef sTComplexD CComplexT; // unvectorised coarse space
|
||||
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix <FineObj,CComplexT,nbasis> RefOperator;
|
||||
typedef MultiGeneralCoarsenedOperator<FineObj,CComplexT,nbasis> TestOperator;
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// Identical random matrix elements into both operators
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
template<class OpA,class OpB>
|
||||
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++){
|
||||
|
||||
// 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 = bufA.size();
|
||||
std::vector<calcMatrix> host(sites);
|
||||
|
||||
ComplexD *w = (ComplexD *)&host[0];
|
||||
int64_t words = sites*sizeof(calcMatrix)/sizeof(ComplexD);
|
||||
for(int64_t i=0;i<words;i++){
|
||||
RealD re,im;
|
||||
random(sRNG,re);
|
||||
random(sRNG,im);
|
||||
w[i] = ComplexD(re-0.5,im-0.5);
|
||||
}
|
||||
|
||||
acceleratorCopyToDevice(&host[0],&bufA[0],sites*sizeof(calcMatrix));
|
||||
A.MatrixPointIn(p,bufA);
|
||||
B.MatrixPointIn(p,bufA);
|
||||
}
|
||||
}
|
||||
|
||||
template<class Op>
|
||||
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++){
|
||||
O.MatrixPointOut(p,buf);
|
||||
int64_t sites = buf.size();
|
||||
std::vector<calcMatrix> host(sites);
|
||||
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]);
|
||||
}
|
||||
return sum;
|
||||
}
|
||||
|
||||
int main (int argc, char ** argv)
|
||||
{
|
||||
Grid_init(&argc,&argv);
|
||||
|
||||
const int nrhs = 4;
|
||||
|
||||
Coordinate clatt = GridDefaultLatt();
|
||||
Coordinate csimd(Nd,1); // the coarse space is unvectorised
|
||||
Coordinate cmpi = GridDefaultMpi();
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// D dimensional coarse grid, and D+1 for the reference
|
||||
////////////////////////////////////////////////
|
||||
GridCartesian *CoarseD = new GridCartesian(clatt,csimd,cmpi);
|
||||
|
||||
std::cout << GridLogMessage << "coarse D grid "; for(int d=0;d<Nd;d++) std::cout<<clatt[d]<<" ";
|
||||
std::cout << " Nsimd " << CoarseD->Nsimd() << std::endl;
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// One geometry object for both, and one D+1 grid
|
||||
// owned here and shared by both operators: fields
|
||||
// conform only across a shared grid object.
|
||||
////////////////////////////////////////////////
|
||||
NextToNearestStencilGeometry4D geom(CoarseD);
|
||||
|
||||
Coordinate mlatt(1,nrhs), msimd(1,1), mmpi(1,1);
|
||||
for(int d=0;d<Nd;d++){
|
||||
mlatt.push_back(clatt[d]);
|
||||
msimd.push_back(csimd[d]);
|
||||
mmpi .push_back(cmpi[d]);
|
||||
}
|
||||
GridCartesian *CoarseMulti = new GridCartesian(mlatt,msimd,mmpi);
|
||||
|
||||
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 ref " << OpRef.geom.npoint
|
||||
<< " npoint test " << OpTest.geom.npoint << std::endl;
|
||||
GRID_ASSERT(OpRef.geom.npoint == OpTest.geom.npoint);
|
||||
|
||||
int npoint = OpRef.geom.npoint;
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Identical matrix elements
|
||||
////////////////////////////////////////////////
|
||||
GridSerialRNG sRNG; sRNG.SeedFixedIntegers(std::vector<int>({7,8,9,10}));
|
||||
SeedMatrixElements(OpRef,OpTest,npoint,sRNG);
|
||||
|
||||
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 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}));
|
||||
|
||||
CoarseVector in (CoarseMulti); random(pRNG,in);
|
||||
CoarseVector out1(CoarseMulti);
|
||||
CoarseVector out2(CoarseMulti);
|
||||
CoarseVector err (CoarseMulti);
|
||||
|
||||
OpRef.M(in,out1);
|
||||
OpTest.M(in,out2);
|
||||
|
||||
err = out1 - out2;
|
||||
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
|
||||
////////////////////////////////////////////////
|
||||
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;
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Move to a different Nrhs and back. The matrix
|
||||
// elements are Nrhs independent and must survive
|
||||
// both the release and the rebuild.
|
||||
////////////////////////////////////////////////
|
||||
// Nrhs 1 is the single RHS case through the multiRHS path, and each slice
|
||||
// of the Nrhs 4 apply must come back unchanged.
|
||||
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++){
|
||||
latt2.push_back(clatt[d]);
|
||||
simd2.push_back(csimd[d]);
|
||||
mpi2 .push_back(cmpi[d]);
|
||||
}
|
||||
GridCartesian *CoarseMulti2 = new GridCartesian(latt2,simd2,mpi2);
|
||||
|
||||
OpTest.SetGrid(CoarseMulti2);
|
||||
GRID_ASSERT( OpTest.Nrhs() == nr );
|
||||
GRID_ASSERT( MatrixChecksum(OpTest,npoint) == ckTest );
|
||||
|
||||
CoarseVector in2 (CoarseMulti2);
|
||||
CoarseVector out(CoarseMulti2);
|
||||
for(int r=0;r<nr;r++){
|
||||
CoarseVector slice(CoarseD);
|
||||
ExtractSliceFast(slice,in,r,0);
|
||||
InsertSliceFast(slice,in2,r,0);
|
||||
}
|
||||
OpTest.M(in2,out);
|
||||
|
||||
RealD sdiff=0.0;
|
||||
for(int r=0;r<nr;r++){
|
||||
CoarseVector a(CoarseD),b(CoarseD),e(CoarseD);
|
||||
ExtractSliceFast(a,out ,r,0);
|
||||
ExtractSliceFast(b,out2,r,0);
|
||||
e = a-b;
|
||||
sdiff += norm2(e);
|
||||
}
|
||||
// Not bit exact: a different Nrhs is a different GEMM shape
|
||||
std::cout << GridLogMessage << "Nrhs " << nr << " slices agree with Nrhs "
|
||||
<< nrhs << " : |diff|^2/|out|^2 = " << sdiff/norm2(out) << std::endl;
|
||||
GRID_ASSERT( norm2(out) > 0.0 );
|
||||
GRID_ASSERT( sdiff/norm2(out) < 1.0e-20 );
|
||||
|
||||
OpTest.ReleaseGrid();
|
||||
GRID_ASSERT( MatrixChecksum(OpTest,npoint) == ckTest ); // survives release
|
||||
|
||||
delete CoarseMulti2;
|
||||
}
|
||||
|
||||
OpTest.SetGrid(CoarseMulti);
|
||||
GRID_ASSERT( OpTest.Nrhs() == nrhs );
|
||||
|
||||
OpTest.M(in,out2);
|
||||
err = out1 - out2;
|
||||
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: ALL PASS" << std::endl;
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
@@ -0,0 +1,279 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./tests/debug/Test_coarse_coarsen.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.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
|
||||
//
|
||||
// Coarsen the same fine operator two ways and compare the matrix elements.
|
||||
//
|
||||
// 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.
|
||||
//
|
||||
#include <Grid/Grid.h>
|
||||
|
||||
using namespace Grid;
|
||||
|
||||
const int nbasis = 8;
|
||||
const int batch = 9;
|
||||
|
||||
typedef vSpinColourVector FineObj;
|
||||
typedef vTComplex CComplexV; // vectorised coarse space, reference
|
||||
typedef sTComplexD CComplexS; // unvectorised coarse space, under test
|
||||
|
||||
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++){
|
||||
O.MatrixPointOut(p,buf);
|
||||
int64_t sites = buf.size();
|
||||
host[p].resize(sites);
|
||||
acceleratorCopyFromDevice(&buf[0],&host[p][0],sites*sizeof(calcMatrix));
|
||||
}
|
||||
}
|
||||
|
||||
int main (int argc, char ** argv)
|
||||
{
|
||||
Grid_init(&argc,&argv);
|
||||
|
||||
Coordinate flatt = GridDefaultLatt();
|
||||
Coordinate fsimd = GridDefaultSimd(Nd,vComplexD::Nsimd());
|
||||
Coordinate fmpi = GridDefaultMpi();
|
||||
|
||||
Coordinate block({2,2,2,2});
|
||||
Coordinate clatt(Nd);
|
||||
for(int d=0;d<Nd;d++){
|
||||
GRID_ASSERT(flatt[d]%block[d]==0);
|
||||
clatt[d] = flatt[d]/block[d];
|
||||
}
|
||||
|
||||
GridCartesian *FineGrid = new GridCartesian(flatt,fsimd,fmpi);
|
||||
GridRedBlackCartesian *FrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(FineGrid);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Reference coarse space: SIMD layout matched to the fine
|
||||
////////////////////////////////////////////////
|
||||
Coordinate cvsimd = GridDefaultSimd(Nd,CComplexV::Nsimd());
|
||||
GridCartesian *CoarseV = new GridCartesian(clatt,cvsimd,fmpi);
|
||||
|
||||
// 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++){
|
||||
reflatt.push_back(clatt[d]);
|
||||
refsimd.push_back(1);
|
||||
refmpi .push_back(fmpi[d]);
|
||||
}
|
||||
GridCartesian *CoarseVMulti = new GridCartesian(reflatt,refsimd,refmpi);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Coarse space under test: unvectorised
|
||||
////////////////////////////////////////////////
|
||||
Coordinate cssimd(Nd,1);
|
||||
GridCartesian *CoarseS = new GridCartesian(clatt,cssimd,fmpi);
|
||||
|
||||
Coordinate cmlatt(1,batch),cmsimd(1,1),cmmpi(1,1);
|
||||
for(int d=0;d<Nd;d++){
|
||||
cmlatt.push_back(clatt[d]);
|
||||
cmsimd.push_back(1);
|
||||
cmmpi .push_back(fmpi[d]);
|
||||
}
|
||||
GridCartesian *CoarseSMulti = new GridCartesian(cmlatt,cmsimd,cmmpi);
|
||||
|
||||
// D+1 fine grid carrying the batch
|
||||
Coordinate fmlatt(1,batch), fmsimd(1,1), fmmpi(1,1);
|
||||
for(int d=0;d<Nd;d++){
|
||||
fmlatt.push_back(flatt[d]);
|
||||
fmsimd.push_back(fsimd[d]);
|
||||
fmmpi .push_back(fmpi[d]);
|
||||
}
|
||||
GridCartesian *FineGridMulti = new GridCartesian(fmlatt,fmsimd,fmmpi);
|
||||
|
||||
std::cout << GridLogMessage << "fine "<<flatt<<" coarse "<<clatt<<" batch "<<batch<<std::endl;
|
||||
std::cout << GridLogMessage << "Nsimd fine "<<FineGrid->Nsimd()
|
||||
<< " coarse ref "<<CoarseV->Nsimd()
|
||||
<< " coarse test "<<CoarseS->Nsimd()<<std::endl;
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Fine operator
|
||||
////////////////////////////////////////////////
|
||||
GridParallelRNG pRNG(FineGrid); pRNG.SeedFixedIntegers(std::vector<int>({1,2,3,4}));
|
||||
LatticeGaugeFieldD Umu(FineGrid); SU<Nc>::HotConfiguration(pRNG,Umu);
|
||||
|
||||
RealD mass = 0.1;
|
||||
WilsonFermionD Dw(Umu,*FineGrid,*FrbGrid,mass);
|
||||
MdagMLinearOperator<WilsonFermionD,LatticeFermionD> HermOp(Dw);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// One random subspace, two copies: CoarsenOperator orthogonalises in place
|
||||
////////////////////////////////////////////////
|
||||
Aggregation<FineObj,CComplexV,nbasis> Subspace(CoarseV,FineGrid,0);
|
||||
std::vector<LatticeFermionD> subspace(nbasis,FineGrid);
|
||||
for(int i=0;i<nbasis;i++){
|
||||
random(pRNG,Subspace.subspace[i]);
|
||||
subspace[i] = Subspace.subspace[i];
|
||||
}
|
||||
|
||||
NextToNearestStencilGeometry4D geomV(CoarseV);
|
||||
NextToNearestStencilGeometry4D geomS(CoarseS);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Reference coarsening, matched layouts
|
||||
////////////////////////////////////////////////
|
||||
RefOperator OpRef(geomV,CoarseVMulti);
|
||||
std::cout << GridLogMessage << "reference CoarsenOperator" << std::endl;
|
||||
OpRef.CoarsenOperator(HermOp,Subspace,CoarseV);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Coarsening under test, D+1 fine applications, unvectorised coarse
|
||||
////////////////////////////////////////////////
|
||||
TestOperator OpTest(geomS,CoarseS);
|
||||
OpTest.SetGrid(CoarseSMulti);
|
||||
|
||||
MrhsPromotedOperator<LatticeFermionD> MrhsHermOp(HermOp,FineGrid,batch);
|
||||
|
||||
std::cout << GridLogMessage << "test CoarsenOperator (D+1)" << std::endl;
|
||||
OpTest.CoarsenOperator(MrhsHermOp,FineGridMulti,subspace,CoarseS);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Compare matrix elements
|
||||
////////////////////////////////////////////////
|
||||
int npoint = OpRef.geom.npoint;
|
||||
GRID_ASSERT(npoint == OpTest.geom.npoint);
|
||||
typedef RefOperator::calcMatrix calcMatrix;
|
||||
std::vector<std::vector<calcMatrix> > A1,A2;
|
||||
ReadMatrix(OpRef,npoint,A1);
|
||||
ReadMatrix(OpTest,npoint,A2);
|
||||
|
||||
RealD num=0.0, den=0.0;
|
||||
for(int p=0;p<npoint;p++){
|
||||
GRID_ASSERT(A1[p].size()==A2[p].size());
|
||||
ComplexD *w1 = (ComplexD *)&A1[p][0];
|
||||
ComplexD *w2 = (ComplexD *)&A2[p][0];
|
||||
int64_t words = A1[p].size()*sizeof(calcMatrix)/sizeof(ComplexD);
|
||||
for(int64_t i=0;i<words;i++){
|
||||
ComplexD d = w1[i]-w2[i];
|
||||
num += real(d)*real(d)+imag(d)*imag(d);
|
||||
den += real(w1[i])*real(w1[i])+imag(w1[i])*imag(w1[i]);
|
||||
}
|
||||
}
|
||||
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 );
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Same coarsening through the single RHS variant: no mrhs packing, the
|
||||
// batch assembled on the coarse side by the mixed blockProject. Block
|
||||
// Gram-Schmidt is idempotent so the subspace may be reused in place.
|
||||
////////////////////////////////////////////////
|
||||
TestOperator OpTestSrhs(geomS,CoarseS);
|
||||
OpTestSrhs.SetGrid(CoarseSMulti);
|
||||
|
||||
std::cout << GridLogMessage << "test CoarsenOperator (single RHS fine op)" << std::endl;
|
||||
OpTestSrhs.CoarsenOperator(HermOp,subspace,CoarseS,batch);
|
||||
|
||||
std::vector<std::vector<calcMatrix> > A3;
|
||||
ReadMatrix(OpTestSrhs,npoint,A3);
|
||||
|
||||
RealD nums=0.0;
|
||||
for(int p=0;p<npoint;p++){
|
||||
GRID_ASSERT(A1[p].size()==A3[p].size());
|
||||
ComplexD *w1 = (ComplexD *)&A1[p][0];
|
||||
ComplexD *w3 = (ComplexD *)&A3[p][0];
|
||||
int64_t words = A1[p].size()*sizeof(calcMatrix)/sizeof(ComplexD);
|
||||
for(int64_t i=0;i<words;i++){
|
||||
ComplexD d = w1[i]-w3[i];
|
||||
nums += real(d)*real(d)+imag(d)*imag(d);
|
||||
}
|
||||
}
|
||||
std::cout << GridLogMessage << "|A_ref - A_test_srhs|^2 / |A_ref|^2 = " << nums/den << std::endl;
|
||||
GRID_ASSERT( nums/den < 1.0e-20 );
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// GetMatrix/SetMatrix round trip through the BLAS layout array. This is
|
||||
// how a bilingual DenseCoarseMatrix will retrieve the elements, and it
|
||||
// needs no SetGrid since the matrix elements are Nrhs independent.
|
||||
////////////////////////////////////////////////
|
||||
{
|
||||
typedef TestOperator::CoarseMatrix CoarseMatrixS;
|
||||
std::vector<CoarseMatrixS> Aget(npoint,CoarseS);
|
||||
for(int p=0;p<npoint;p++) OpTest.GetMatrix(p,Aget);
|
||||
|
||||
TestOperator OpTestRoundTrip(geomS,CoarseS);
|
||||
for(int p=0;p<npoint;p++) OpTestRoundTrip.SetMatrix(p,Aget);
|
||||
|
||||
std::vector<std::vector<calcMatrix> > A4;
|
||||
ReadMatrix(OpTestRoundTrip,npoint,A4);
|
||||
|
||||
RealD numrt=0.0;
|
||||
for(int p=0;p<npoint;p++){
|
||||
GRID_ASSERT(A2[p].size()==A4[p].size());
|
||||
ComplexD *w2 = (ComplexD *)&A2[p][0];
|
||||
ComplexD *w4 = (ComplexD *)&A4[p][0];
|
||||
int64_t words = A2[p].size()*sizeof(calcMatrix)/sizeof(ComplexD);
|
||||
for(int64_t i=0;i<words;i++){
|
||||
ComplexD d = w2[i]-w4[i];
|
||||
numrt += real(d)*real(d)+imag(d)*imag(d);
|
||||
}
|
||||
}
|
||||
std::cout << GridLogMessage << "GetMatrix/SetMatrix round trip |diff|^2 = " << numrt << std::endl;
|
||||
GRID_ASSERT( numrt == 0.0 );
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// and the operator under test applies the matrix it just built
|
||||
////////////////////////////////////////////////
|
||||
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);
|
||||
|
||||
OpTest.M(in,out);
|
||||
std::cout << GridLogMessage << "|in|^2 = " << norm2(in)
|
||||
<< " |M_test in|^2 = " << norm2(out) << std::endl;
|
||||
GRID_ASSERT( norm2(out) > 0.0 );
|
||||
|
||||
std::cout << GridLogMessage << "Test_coarse_coarsen: ALL PASS" << std::endl;
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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();
|
||||
}
|
||||
Reference in new issue
Block a user