mirror of
https://github.com/paboyle/Grid.git
synced 2026-10-12 10:48:06 +01:00
Synch multigrid rewrite to repository and
broad unicode elimination effort
This commit is contained in:
1 parent
865c338dd2
commit
3bd5e8883f
73 files changed
+1018
-3219
No files matched your search
@@ -8,7 +8,7 @@
|
||||
then repeatedly applies FFT_all_dim to the same propagator 400 times.
|
||||
|
||||
If PlannedFFT is working correctly the RSS should remain flat after the first
|
||||
iteration — no new plans, no new deviceVector allocations beyond the per-call
|
||||
iteration -- no new plans, no new deviceVector allocations beyond the per-call
|
||||
pencil buffer which is freed at the end of each FFT_dim_execute call.
|
||||
|
||||
Build exactly like any other Grid test, e.g.:
|
||||
@@ -63,7 +63,7 @@ static long getGPUUsedMb()
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Convenience struct — one snapshot of both sides
|
||||
// Convenience struct -- one snapshot of both sides
|
||||
// ============================================================
|
||||
struct MemSnapshot {
|
||||
long cpu_rss_kb; // host RSS in kB (-1 if unavailable)
|
||||
@@ -113,7 +113,7 @@ int main(int argc, char **argv)
|
||||
<< "Grid is setup to use " << threads << " threads" << std::endl;
|
||||
|
||||
// ------------------------------------------------------------------
|
||||
// Grid setup — use whatever lattice/mpi/simd was passed on the CLI,
|
||||
// Grid setup -- use whatever lattice/mpi/simd was passed on the CLI,
|
||||
// e.g. --grid 8.8.8.8 --mpi 1.1.1.1
|
||||
// ------------------------------------------------------------------
|
||||
Coordinate latt_size = GridDefaultLatt();
|
||||
@@ -174,13 +174,13 @@ int main(int argc, char **argv)
|
||||
<< std::endl;
|
||||
|
||||
// ------------------------------------------------------------------
|
||||
// Create the PlannedFFT — plans are allocated here ONCE for all
|
||||
// Create the PlannedFFT -- plans are allocated here ONCE for all
|
||||
// dimensions and stored inside the object.
|
||||
// ------------------------------------------------------------------
|
||||
PlannedFFT<iSpinColourMatrix<vComplexD>> plannedFFT(&GRID);
|
||||
|
||||
// ------------------------------------------------------------------
|
||||
// Snapshot AFTER plan construction — this is the true baseline
|
||||
// Snapshot AFTER plan construction -- this is the true baseline
|
||||
// for the loop, because cufftPlanMany itself grabs device memory.
|
||||
// ------------------------------------------------------------------
|
||||
MemSnapshot snap_after_plan = takeSnapshot();
|
||||
@@ -197,7 +197,7 @@ int main(int argc, char **argv)
|
||||
// ------------------------------------------------------------------
|
||||
// 400-iteration loop.
|
||||
// Each iteration computes the full 4d forward FFT of `prop`.
|
||||
// We deliberately do NOT cache the result — we always start from
|
||||
// We deliberately do NOT cache the result -- we always start from
|
||||
// the same `prop` so the FFT is recomputed identically each time.
|
||||
// The point is to watch memory, not correctness.
|
||||
// ------------------------------------------------------------------
|
||||
@@ -229,7 +229,7 @@ int main(int argc, char **argv)
|
||||
}
|
||||
|
||||
// cudaMemGetInfo reflects the state *after* any pooled frees have
|
||||
// been committed, so this is accurate without an explicit sync —
|
||||
// been committed, so this is accurate without an explicit sync --
|
||||
// FFT_dim_execute already calls accelerator_barrier() internally.
|
||||
MemSnapshot snap_now = takeSnapshot();
|
||||
printRow(i, snap_now, snap_prev);
|
||||
|
||||
@@ -156,7 +156,7 @@ int main (int argc, char ** argv)
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////
|
||||
// Dwf matrix — verify Fourier representation using PlannedFFT<LatticeFermionD>
|
||||
// Dwf matrix -- verify Fourier representation using PlannedFFT<LatticeFermionD>
|
||||
////////////////////////////////////////////////////
|
||||
{
|
||||
std::cout<<"****************************************"<<std::endl;
|
||||
|
||||
@@ -1,252 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./tests/debug/Test_coarse_v2.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 */
|
||||
|
||||
//
|
||||
// MultiGeneralCoarsenedOperatorV2 against the existing mrhs coarse operator.
|
||||
//
|
||||
// V1 is constructed on the D+1 grid as now; V2 on the D dimensional grid,
|
||||
// with SetGrid() adopting the caller owned D+1 grid and building its padded
|
||||
// cell and neighbour table from the D dimensional stencil, with the Nrhs
|
||||
// factor multiplied in.
|
||||
//
|
||||
// Both are given identical matrix elements, so any difference in the apply is
|
||||
// the restructured neighbour table. The same geometry object is passed to
|
||||
// both: V1 adds one to skip for the rhs direction, V2 uses it as is on the D
|
||||
// dimensional grid, so both describe the same stencil over the D dimensions.
|
||||
//
|
||||
#include <Grid/Grid.h>
|
||||
|
||||
using namespace Grid;
|
||||
|
||||
const int nbasis = 8;
|
||||
|
||||
typedef vSpinColourVector FineObj;
|
||||
typedef sTComplexD CComplexT; // unvectorised coarse space
|
||||
|
||||
typedef MultiGeneralCoarsenedMatrix <FineObj,CComplexT,nbasis> MrhsV1;
|
||||
typedef MultiGeneralCoarsenedOperatorV2<FineObj,CComplexT,nbasis> MrhsV2;
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// 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;
|
||||
|
||||
for(int p=0;p<npoint;p++){
|
||||
|
||||
GRID_ASSERT(A.BLAS_A[p].size() == B.BLAS_A[p].size());
|
||||
|
||||
int64_t sites = A.BLAS_A[p].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],&A.BLAS_A[p][0],sites*sizeof(calcMatrix));
|
||||
acceleratorCopyToDevice(&host[0],&B.BLAS_A[p][0],sites*sizeof(calcMatrix));
|
||||
}
|
||||
}
|
||||
|
||||
template<class Op>
|
||||
RealD MatrixChecksum(Op &O,int npoint)
|
||||
{
|
||||
typedef typename Op::calcMatrix calcMatrix;
|
||||
RealD sum=0.0;
|
||||
for(int p=0;p<npoint;p++){
|
||||
int64_t sites = O.BLAS_A[p].size();
|
||||
std::vector<calcMatrix> host(sites);
|
||||
acceleratorCopyFromDevice(&O.BLAS_A[p][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 V1
|
||||
////////////////////////////////////////////////
|
||||
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);
|
||||
|
||||
MrhsV2 OpV2(geom,CoarseD);
|
||||
MrhsV1 OpV1(geom,CoarseMulti);
|
||||
OpV2.SetGrid(CoarseMulti);
|
||||
|
||||
std::cout << GridLogMessage << "coarse D+1 grid nrhs " << nrhs
|
||||
<< " Nsimd " << CoarseMulti->Nsimd() << std::endl;
|
||||
|
||||
std::cout << GridLogMessage << "npoint V1 " << OpV1.geom.npoint
|
||||
<< " npoint V2 " << OpV2.geom.npoint << std::endl;
|
||||
GRID_ASSERT(OpV1.geom.npoint == OpV2.geom.npoint);
|
||||
|
||||
int npoint = OpV1.geom.npoint;
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Identical matrix elements
|
||||
////////////////////////////////////////////////
|
||||
GridSerialRNG sRNG; sRNG.SeedFixedIntegers(std::vector<int>({7,8,9,10}));
|
||||
SeedMatrixElements(OpV1,OpV2,npoint,sRNG);
|
||||
|
||||
RealD ckV1 = MatrixChecksum(OpV1,npoint);
|
||||
RealD ckV2 = MatrixChecksum(OpV2,npoint);
|
||||
std::cout << GridLogMessage << "matrix element checksum V1 " << ckV1
|
||||
<< " V2 " << ckV2 << std::endl;
|
||||
GRID_ASSERT( ckV1 == ckV2 );
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Same input, compare the applies
|
||||
////////////////////////////////////////////////
|
||||
typedef MrhsV1::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);
|
||||
|
||||
OpV1.M(in,out1);
|
||||
OpV2.M(in,out2);
|
||||
|
||||
err = out1 - out2;
|
||||
std::cout << GridLogMessage << "|V1 out|^2 = " << norm2(out1)
|
||||
<< " |V2 out|^2 = " << norm2(out2) << std::endl;
|
||||
std::cout << GridLogMessage << "|V1 - V2|^2 = " << norm2(err) << std::endl;
|
||||
GRID_ASSERT( norm2(out1) > 0.0 );
|
||||
GRID_ASSERT( norm2(err) == 0.0 );
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// SetGrid is idempotent on pointer identity
|
||||
////////////////////////////////////////////////
|
||||
OpV2.SetGrid(CoarseMulti);
|
||||
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 );
|
||||
OpV2.M(in,out2);
|
||||
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.
|
||||
OpV2.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);
|
||||
|
||||
OpV2.SetGrid(CoarseMulti2);
|
||||
GRID_ASSERT( OpV2.Nrhs() == nr );
|
||||
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 );
|
||||
|
||||
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);
|
||||
}
|
||||
OpV2.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 );
|
||||
|
||||
OpV2.ReleaseGrid();
|
||||
GRID_ASSERT( MatrixChecksum(OpV2,npoint) == ckV2 ); // survives release
|
||||
|
||||
delete CoarseMulti2;
|
||||
}
|
||||
|
||||
OpV2.SetGrid(CoarseMulti);
|
||||
GRID_ASSERT( OpV2.Nrhs() == nrhs );
|
||||
|
||||
OpV2.M(in,out2);
|
||||
err = out1 - out2;
|
||||
std::cout << GridLogMessage << "after Nrhs 4 -> 2 -> 1 -> release -> 4, |V1 - V2|^2 = "
|
||||
<< norm2(err) << std::endl;
|
||||
GRID_ASSERT( norm2(err) == 0.0 );
|
||||
|
||||
std::cout << GridLogMessage << "Test_coarse_v2: ALL PASS" << std::endl;
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
@@ -1,273 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./tests/debug/Test_coarse_v2_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.
|
||||
//
|
||||
// V1 : existing mrhs CoarsenOperator, vectorised coarse space matched to
|
||||
// the fine SIMD layout, single RHS fine applications
|
||||
// V2 : D+1 CoarsenOperator, unvectorised sComplexD coarse space, the batch
|
||||
// of phased basis vectors carried in the rhs direction and applied
|
||||
// through MrhsPromotedOperator
|
||||
//
|
||||
// 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, V1 reference
|
||||
typedef sTComplexD CComplexS; // unvectorised coarse space, V2
|
||||
|
||||
typedef MultiGeneralCoarsenedMatrix <FineObj,CComplexV,nbasis> MrhsV1;
|
||||
typedef MultiGeneralCoarsenedOperatorV2<FineObj,CComplexS,nbasis> MrhsV2;
|
||||
|
||||
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);
|
||||
for(int p=0;p<npoint;p++){
|
||||
int64_t sites = O.BLAS_A[p].size();
|
||||
host[p].resize(sites);
|
||||
acceleratorCopyFromDevice(&O.BLAS_A[p][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);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// V1 coarse space: SIMD layout matched to the fine
|
||||
////////////////////////////////////////////////
|
||||
Coordinate cvsimd = GridDefaultSimd(Nd,CComplexV::Nsimd());
|
||||
GridCartesian *CoarseV = new GridCartesian(clatt,cvsimd,fmpi);
|
||||
|
||||
// V1 puts all of the SIMD in the rhs direction. CoarsenOperator does not
|
||||
// use the multiRHS grid; it only sizes BLAS_A, so one lane of rhs suffices.
|
||||
int nrhs_v1 = CComplexV::Nsimd();
|
||||
Coordinate v1latt(1,nrhs_v1),v1simd(1,CComplexV::Nsimd()),v1mpi(1,1);
|
||||
for(int d=0;d<Nd;d++){
|
||||
v1latt.push_back(clatt[d]);
|
||||
v1simd.push_back(1);
|
||||
v1mpi .push_back(fmpi[d]);
|
||||
}
|
||||
GridCartesian *CoarseVMulti = new GridCartesian(v1latt,v1simd,v1mpi);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// V2 coarse space: 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 V1 "<<CoarseV->Nsimd()
|
||||
<< " coarse V2 "<<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);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// V1 coarsening, matched layouts
|
||||
////////////////////////////////////////////////
|
||||
MrhsV1 OpV1(geomV,CoarseVMulti);
|
||||
std::cout << GridLogMessage << "V1 CoarsenOperator" << std::endl;
|
||||
OpV1.CoarsenOperator(HermOp,Subspace,CoarseV);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// V2 coarsening, D+1 fine applications, unvectorised coarse
|
||||
////////////////////////////////////////////////
|
||||
MrhsV2 OpV2(geomS,CoarseS);
|
||||
OpV2.SetGrid(CoarseSMulti);
|
||||
|
||||
MrhsPromotedOperator<LatticeFermionD> MrhsHermOp(HermOp,FineGrid,batch);
|
||||
|
||||
std::cout << GridLogMessage << "V2 CoarsenOperator (D+1)" << std::endl;
|
||||
OpV2.CoarsenOperator(MrhsHermOp,FineGridMulti,subspace,CoarseS);
|
||||
|
||||
////////////////////////////////////////////////
|
||||
// Compare matrix elements
|
||||
////////////////////////////////////////////////
|
||||
int npoint = OpV1.geom.npoint;
|
||||
GRID_ASSERT(npoint == OpV2.geom.npoint);
|
||||
typedef MrhsV1::calcMatrix calcMatrix;
|
||||
std::vector<std::vector<calcMatrix> > A1,A2;
|
||||
ReadMatrix(OpV1,npoint,A1);
|
||||
ReadMatrix(OpV2,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_V1|^2 = " << den << std::endl;
|
||||
std::cout << GridLogMessage << "|A_V1 - A_V2|^2 / |A_V1|^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.
|
||||
////////////////////////////////////////////////
|
||||
MrhsV2 OpV2s(geomS,CoarseS);
|
||||
OpV2s.SetGrid(CoarseSMulti);
|
||||
|
||||
std::cout << GridLogMessage << "V2 CoarsenOperator (single RHS fine op)" << std::endl;
|
||||
OpV2s.CoarsenOperator(HermOp,subspace,CoarseS,batch);
|
||||
|
||||
std::vector<std::vector<calcMatrix> > A3;
|
||||
ReadMatrix(OpV2s,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_V1 - A_V2srhs|^2 / |A_V1|^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 MrhsV2::CoarseMatrix CoarseMatrixS;
|
||||
std::vector<CoarseMatrixS> Aget(npoint,CoarseS);
|
||||
for(int p=0;p<npoint;p++) OpV2.GetMatrix(p,Aget);
|
||||
|
||||
MrhsV2 OpV2rt(geomS,CoarseS);
|
||||
for(int p=0;p<npoint;p++) OpV2rt.SetMatrix(p,Aget);
|
||||
|
||||
std::vector<std::vector<calcMatrix> > A4;
|
||||
ReadMatrix(OpV2rt,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 V2 applies the matrix it just built
|
||||
////////////////////////////////////////////////
|
||||
typedef MrhsV2::CoarseVector CoarseVectorS;
|
||||
GridParallelRNG cRNG(CoarseS); cRNG.SeedFixedIntegers(std::vector<int>({5,6,7,8}));
|
||||
CoarseVectorS in(CoarseSMulti); random(cRNG,in);
|
||||
CoarseVectorS out(CoarseSMulti);
|
||||
|
||||
OpV2.M(in,out);
|
||||
std::cout << GridLogMessage << "|in|^2 = " << norm2(in)
|
||||
<< " |M_V2 in|^2 = " << norm2(out) << std::endl;
|
||||
GRID_ASSERT( norm2(out) > 0.0 );
|
||||
|
||||
std::cout << GridLogMessage << "Test_coarse_v2_coarsen: ALL PASS" << std::endl;
|
||||
|
||||
Grid_finalize();
|
||||
}
|
||||
@@ -114,7 +114,7 @@ int main (int argc, char ** argv)
|
||||
///////////////////////////////////////////////////
|
||||
std::cout<<GridLogMessage << "Building little Dirac operator"<< std::endl;
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -223,7 +223,7 @@ int main (int argc, char ** argv)
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
|
||||
#if 0
|
||||
MultiGeneralCoarsenedMatrix mrhs(LittleDiracOp,CoarseMrhs);
|
||||
DeprecatedMultiGeneralCoarsenedMatrix mrhs(LittleDiracOp,CoarseMrhs);
|
||||
typedef decltype(mrhs) MultiGeneralCoarsenedMatrix_t;
|
||||
|
||||
//////////////////////////////////////////
|
||||
|
||||
@@ -96,7 +96,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -140,7 +140,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
mrhs.CoarsenOperator(FineHermOp,Aggregates,Coarse5d);
|
||||
|
||||
@@ -189,7 +189,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -371,7 +371,7 @@ slurm-1482367.out:Grid : Message : 6169.469330 s : HDCG: Pcg converged in 487 it
|
||||
//////////////////////////////////////////
|
||||
// Build a HDCG solver
|
||||
//////////////////////////////////////////
|
||||
TwoLevelADEF2<LatticeFermion,CoarseVector,Subspace>
|
||||
DeprecatedTwoLevelADEF2<LatticeFermion,CoarseVector,Subspace>
|
||||
HDCG(1.0e-8, 700,
|
||||
FineHermOp,
|
||||
CGsmooth,
|
||||
|
||||
@@ -218,7 +218,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -269,7 +269,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
std::cout << "**************************************"<<std::endl;
|
||||
|
||||
@@ -407,7 +407,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -456,7 +456,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
std::cout << "**************************************"<<std::endl;
|
||||
|
||||
@@ -406,7 +406,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -455,7 +455,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
std::cout << "**************************************"<<std::endl;
|
||||
|
||||
@@ -215,7 +215,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
|
||||
@@ -161,7 +161,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -210,7 +210,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
std::cout << "**************************************"<<std::endl;
|
||||
|
||||
@@ -162,7 +162,7 @@ int main (int argc, char ** argv)
|
||||
////////////////////////////////////////////////////////////
|
||||
///////////// Coarse basis and Little Dirac Operator ///////
|
||||
////////////////////////////////////////////////////////////
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
@@ -217,7 +217,7 @@ int main (int argc, char ** argv)
|
||||
Coordinate rhSimd({vComplex::Nsimd(),1, 1,1,1,1});
|
||||
|
||||
GridCartesian *CoarseMrhs = new GridCartesian(rhLatt,rhSimd,rhMpi);
|
||||
typedef MultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
typedef DeprecatedMultiGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> MultiGeneralCoarsenedMatrix_t;
|
||||
MultiGeneralCoarsenedMatrix_t mrhs(geom,CoarseMrhs);
|
||||
|
||||
///////////////////////
|
||||
|
||||
@@ -254,7 +254,7 @@ int main (int argc, char ** argv)
|
||||
const int cb = 0 ;
|
||||
LatticeFermion prom(FGrid);
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
|
||||
@@ -376,10 +376,10 @@ int main (int argc, char ** argv)
|
||||
std::cout <<bl<<" "<<br<<"\t"<<innerProduct(U.subspace[bl],src)<<std::endl;
|
||||
}}
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
|
||||
typedef LittleDiracOperatorV::CoarseVector CoarseVectorV;
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
V.Orthogonalise();
|
||||
|
||||
@@ -375,10 +375,10 @@ int main (int argc, char ** argv)
|
||||
std::cout <<bl<<" "<<br<<"\t"<<innerProduct(U.subspace[bl],src)<<std::endl;
|
||||
}}
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperatorV;
|
||||
typedef LittleDiracOperatorV::CoarseVector CoarseVectorV;
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
V.Orthogonalise();
|
||||
|
||||
@@ -369,8 +369,8 @@ int main (int argc, char ** argv)
|
||||
// }
|
||||
|
||||
|
||||
// typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
// typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
LittleDiracOperator LittleDiracOpPV(geom,FGrid,Coarse5d);
|
||||
LittleDiracOpPV.CoarsenOperator(PVdagM,V,V);
|
||||
|
||||
@@ -185,7 +185,7 @@ int main (int argc, char ** argv)
|
||||
const int cb = 0 ;
|
||||
LatticeFermion prom(FGrid);
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NearestStencilGeometry4D geom(Coarse4d);
|
||||
|
||||
@@ -185,7 +185,7 @@ int main (int argc, char ** argv)
|
||||
const int cb = 0 ;
|
||||
LatticeFermion prom(FGrid);
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NearestStencilGeometry4D geom(Coarse4d);
|
||||
|
||||
@@ -182,7 +182,7 @@ int main (int argc, char ** argv)
|
||||
const int cb = 0 ;
|
||||
LatticeFermion prom(FGrid);
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,2*nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NearestStencilGeometry4D geom(Coarse4d);
|
||||
|
||||
@@ -182,7 +182,7 @@ int main (int argc, char ** argv)
|
||||
const int cb = 0 ;
|
||||
LatticeFermion prom(FGrid);
|
||||
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
|
||||
NearestStencilGeometry4D geom(Coarse4d);
|
||||
|
||||
@@ -3,9 +3,9 @@
|
||||
*
|
||||
* Tests three orderings with an empty rocFFT cache to find which GPU
|
||||
* operation before plan creation triggers the failure:
|
||||
* A) hipMalloc only — hypothesis: passes (no async GPU work)
|
||||
* B) hipMalloc + hipMemset — hypothesis: fails (async GPU work in flight)
|
||||
* C) hipMalloc + hipMemset — hypothesis: passes (work completed before plan)
|
||||
* A) hipMalloc only -- hypothesis: passes (no async GPU work)
|
||||
* B) hipMalloc + hipMemset -- hypothesis: fails (async GPU work in flight)
|
||||
* C) hipMalloc + hipMemset -- hypothesis: passes (work completed before plan)
|
||||
* + hipDeviceSynchronize
|
||||
*
|
||||
* Compile:
|
||||
@@ -59,13 +59,13 @@ int main(void) {
|
||||
hipfftResult rvB = makePlan(G, howmany);
|
||||
printf("G=%-4d B) hipMalloc + hipMemset : %s\n", G, res(rvB));
|
||||
|
||||
// C: hipMalloc + hipMemset + sync — does syncing before plan creation fix it?
|
||||
// C: hipMalloc + hipMemset + sync -- does syncing before plan creation fix it?
|
||||
hipMemset(buf, 0, nelems * sizeof(hipfftDoubleComplex));
|
||||
hipDeviceSynchronize();
|
||||
hipfftResult rvC = makePlan(G, howmany);
|
||||
printf("G=%-4d C) hipMalloc + hipMemset + sync: %s\n", G, res(rvC));
|
||||
|
||||
// A last: hipMalloc only, no async GPU work — should always pass
|
||||
// A last: hipMalloc only, no async GPU work -- should always pass
|
||||
hipfftResult rvA = makePlan(G, howmany);
|
||||
printf("G=%-4d A) hipMalloc only : %s\n\n", G, res(rvA));
|
||||
|
||||
|
||||
@@ -33,7 +33,7 @@ int main(void) {
|
||||
int n[] = {G};
|
||||
long nelems = (long)G * howmany;
|
||||
|
||||
// Plan created BEFORE hipMalloc — succeeds for all G
|
||||
// Plan created BEFORE hipMalloc -- succeeds for all G
|
||||
hipfftHandle p;
|
||||
size_t workSize = 0;
|
||||
hipfftCreate(&p);
|
||||
|
||||
@@ -38,8 +38,8 @@ static const char *hipfftResultString(hipfftResult r) {
|
||||
// Plan creation + execution for (G, howmany).
|
||||
// Tests two orderings to isolate whether a prior hipMalloc poisons hipfft
|
||||
// plan creation for small G on ROCm 7:
|
||||
// A) plan BEFORE hipMalloc — hypothesis: succeeds
|
||||
// B) hipMalloc BEFORE plan — hypothesis: fails for G < 32
|
||||
// A) plan BEFORE hipMalloc -- hypothesis: succeeds
|
||||
// B) hipMalloc BEFORE plan -- hypothesis: fails for G < 32
|
||||
static void tryPlanAndExec(int G, long howmany) {
|
||||
int n[] = {G};
|
||||
long nelems = (long)G * howmany;
|
||||
|
||||
@@ -58,12 +58,12 @@ static void tryPlanAndExec(int G, long howmany) {
|
||||
hipfftDoubleComplex *dbuf = nullptr;
|
||||
hipError_t herr = hipMalloc(&dbuf, nelems * sizeof(hipfftDoubleComplex));
|
||||
if (herr != hipSuccess) {
|
||||
printf(" hipMalloc failed (%d) for %ld elems — skipping\n\n", (int)herr, nelems);
|
||||
printf(" hipMalloc failed (%d) for %ld elems -- skipping\n\n", (int)herr, nelems);
|
||||
return;
|
||||
}
|
||||
hipMemset(dbuf, 0, nelems * sizeof(hipfftDoubleComplex));
|
||||
|
||||
// 1. hipfftPlanMany (one-step, nullptr embed) — current Grid path
|
||||
// 1. hipfftPlanMany (one-step, nullptr embed) -- current Grid path
|
||||
{
|
||||
hipfftHandle p;
|
||||
hipfftResult rv = hipfftPlanMany(&p, 1, n,
|
||||
@@ -79,7 +79,7 @@ static void tryPlanAndExec(int G, long howmany) {
|
||||
}
|
||||
}
|
||||
|
||||
// 2. hipfftCreate + hipfftMakePlanMany (two-step) — also current Grid path
|
||||
// 2. hipfftCreate + hipfftMakePlanMany (two-step) -- also current Grid path
|
||||
{
|
||||
hipfftHandle p;
|
||||
size_t workSize = 0;
|
||||
|
||||
@@ -30,7 +30,7 @@ Author: Peter Boyle <pboyle@bnl.gov>
|
||||
// The DenseCoarseMatrix GLUE test, on a real (tiny) lattice coarse
|
||||
// operator, CPU laptop build.
|
||||
//
|
||||
// Builds a genuine GeneralCoarsenedMatrix (DWF MdagM + 0.5 shift for a
|
||||
// Builds a genuine DeprecatedGeneralCoarsenedMatrix (DWF MdagM + 0.5 shift for a
|
||||
// guaranteed-invertible Galerkin coarse op, random aggregation basis)
|
||||
// and runs the whole Import certificate chain through the glue:
|
||||
// - fresh ImportDense + IMPORT CERTIFICATE vs Op.M
|
||||
@@ -143,7 +143,7 @@ int main (int argc, char ** argv)
|
||||
Aggregates.CreateSubspaceRandom(RNG5);
|
||||
|
||||
std::cout << GridLogMessage << "Coarsening shifted MdagM" << std::endl;
|
||||
typedef GeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef DeprecatedGeneralCoarsenedMatrix<vSpinColourVector,vTComplex,nbasis> LittleDiracOperator;
|
||||
typedef LittleDiracOperator::CoarseVector CoarseVector;
|
||||
NextToNextToNextToNearestStencilGeometry5D geom(Coarse5d);
|
||||
LittleDiracOperator LittleDiracOp(geom,FGrid,Coarse5d);
|
||||
|
||||
@@ -397,7 +397,7 @@ public:
|
||||
MdagMLinearOperator<CoarseDiracMatrix, CoarseVector> coarseMdagMOp(_CoarseMatrix);
|
||||
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": **************************************************" << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": MG correctness check: 0 == (M - (Mdiag + Σ_μ Mdir_μ)) * v" << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": MG correctness check: 0 == (M - (Mdiag + Sum_mu Mdir_mu)) * v" << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": **************************************************" << std::endl;
|
||||
|
||||
random(_LevelInfo.PRNGs[_CurrentLevel], fineTmps[0]);
|
||||
@@ -406,22 +406,22 @@ public:
|
||||
fineMdagMOp.OpDiag(fineTmps[0], fineTmps[2]); // Mdiag * v
|
||||
|
||||
fineTmps[4] = zero;
|
||||
for(int dir = 0; dir < 4; dir++) { // Σ_μ Mdir_μ * v
|
||||
for(int dir = 0; dir < 4; dir++) { // Sum_mu Mdir_mu * v
|
||||
for(auto disp : {+1, -1}) {
|
||||
fineMdagMOp.OpDir(fineTmps[0], fineTmps[3], dir, disp);
|
||||
fineTmps[4] = fineTmps[4] + fineTmps[3];
|
||||
}
|
||||
}
|
||||
|
||||
fineTmps[5] = fineTmps[2] + fineTmps[4]; // (Mdiag + Σ_μ Mdir_μ) * v
|
||||
fineTmps[5] = fineTmps[2] + fineTmps[4]; // (Mdiag + Sum_mu Mdir_mu) * v
|
||||
|
||||
fineTmps[6] = fineTmps[1] - fineTmps[5];
|
||||
auto deviation = std::sqrt(norm2(fineTmps[6]) / norm2(fineTmps[1]));
|
||||
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(M * v) = " << norm2(fineTmps[1]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Mdiag * v) = " << norm2(fineTmps[2]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Σ_μ Mdir_μ * v) = " << norm2(fineTmps[4]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2((Mdiag + Σ_μ Mdir_μ) * v) = " << norm2(fineTmps[5]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2(Sum_mu Mdir_mu * v) = " << norm2(fineTmps[4]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": norm2((Mdiag + Sum_mu Mdir_mu) * v) = " << norm2(fineTmps[5]) << std::endl;
|
||||
std::cout << GridLogMG << " Level " << _CurrentLevel << ": relative deviation = " << deviation;
|
||||
|
||||
if(deviation > tolerance) {
|
||||
@@ -664,7 +664,7 @@ createMGInstance(WilsonMGParams &mgParams, LevelInfo &levelInfo, Matrix &FineMat
|
||||
CASE_FOR_N_LEVELS(3);
|
||||
CASE_FOR_N_LEVELS(4);
|
||||
default:
|
||||
std::cout << GridLogError << "We currently only support nLevels ∈ {2, 3, 4}" << std::endl;
|
||||
std::cout << GridLogError << "We currently only support nLevels in {2, 3, 4}" << std::endl;
|
||||
exit(EXIT_FAILURE);
|
||||
break;
|
||||
}
|
||||
|
||||
Reference in new issue
Block a user