From e8a82fa683c750b41f242d5fac6ee2d9bb23cb35 Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Tue, 8 Sep 2026 15:52:15 -0400 Subject: [PATCH] Multigrid clean up phase, and fix shell variables for libfabric bug work around on Frontier --- Grid/algorithms/multigrid/MultiGrid.h | 5 +++++ Grid/stencil/Stencil.h | 30 +++++++++++++++++++++++++++ systems/Frontier/sourceme-rocm6.4.sh | 5 +++++ systems/Frontier/sourceme-rocm7.0.sh | 5 +++++ systems/Frontier/sourceme-rocm7.2.sh | 29 +++++++++++++++----------- 5 files changed, 62 insertions(+), 12 deletions(-) diff --git a/Grid/algorithms/multigrid/MultiGrid.h b/Grid/algorithms/multigrid/MultiGrid.h index 88a369a90..4e092040b 100644 --- a/Grid/algorithms/multigrid/MultiGrid.h +++ b/Grid/algorithms/multigrid/MultiGrid.h @@ -35,3 +35,8 @@ Author: Peter Boyle #include #include #include +#include +// PVdagMOperators.h / MrhsMultiGrid.h / PVdagMMultiGrid.h / +// DenseCoarseMatrix.h are NOT in this umbrella: consumers of the PVdagM +// chain include PVdagMMultiGrid.h explicitly (it pulls the dense stack +// and BLAS). diff --git a/Grid/stencil/Stencil.h b/Grid/stencil/Stencil.h index 8c1a3f2e6..97c7b04c8 100644 --- a/Grid/stencil/Stencil.h +++ b/Grid/stencil/Stencil.h @@ -453,9 +453,23 @@ public: double *dbuf =(double *) packet.recv_buf; float *fbuf =(float *) packet.compressed_recv_buf; + // BUG FIX 2026-09-07: the lane structure is a carefully designed + // GPU optimisation -- lane = threadIdx.y puts adjacent threads on + // adjacent words, fully coalesced. On CPU builds there are no + // SIMT lanes and acceleratorSIMTlane() is identically zero, so + // this loop converted only 1/nsimd of the words and the rest of + // the halo silently kept STALE buffer content (deterministic + // wrong answers whenever the previous occupant differed; caught + // on the 2-rank laptop build via Benchmark_dwf's sloppy Cshift + // check). GRID_SIMT keeps the optimisation; CPU loops the lanes. accelerator_forNB(ss,outer,nsimd,{ +#ifdef GRID_SIMT int lane = acceleratorSIMTlane(nsimd); dbuf[ss*nsimd+lane] = fbuf[ss*nsimd+lane]; //conversion +#else + for(int lane=0;lane>16; // convert as in Bagel/BFM ; bfloat16 ; s7e8 Intel patent +#else + for(int lane=0;lane>16; // convert as in Bagel/BFM ; bfloat16 ; s7e8 Intel patent +#endif }); } else { diff --git a/systems/Frontier/sourceme-rocm6.4.sh b/systems/Frontier/sourceme-rocm6.4.sh index 66fff35ae..227351be4 100644 --- a/systems/Frontier/sourceme-rocm6.4.sh +++ b/systems/Frontier/sourceme-rocm6.4.sh @@ -14,3 +14,8 @@ export LD_LIBRARY_PATH=$CRAY_LD_LIBRARY_PATH:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=/opt/rocm-6.4.0/lib/llvm/lib/:$LD_LIBRARY_PATH module load emacs export FI_MR_CACHE_MONITOR=kdreg2 # REQUIRED for device-buffer MPI on Slingshot: libfabric memhooks monitor (default) is defective, see systems/WorkArounds.txt (libfabric #11451) + +# Standing avoidance of the CXI NO_TRANSLATION fault on persistent device +# buffers under hipMalloc/hipFree churn (libfabric issue #12775; ruling +# 2026-09-07). +export FI_HMEM_ROCR_USE_DMABUF=0 diff --git a/systems/Frontier/sourceme-rocm7.0.sh b/systems/Frontier/sourceme-rocm7.0.sh index 872ef1be5..a73952afa 100644 --- a/systems/Frontier/sourceme-rocm7.0.sh +++ b/systems/Frontier/sourceme-rocm7.0.sh @@ -12,3 +12,8 @@ export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$CRAY_LD_LIBRARY_PATH module load emacs export FI_MR_CACHE_MONITOR=kdreg2 # REQUIRED for device-buffer MPI on Slingshot: libfabric memhooks monitor (default) is defective, see systems/WorkArounds.txt (libfabric #11451) + +# Standing avoidance of the CXI NO_TRANSLATION fault on persistent device +# buffers under hipMalloc/hipFree churn (libfabric issue #12775; ruling +# 2026-09-07). +export FI_HMEM_ROCR_USE_DMABUF=0 diff --git a/systems/Frontier/sourceme-rocm7.2.sh b/systems/Frontier/sourceme-rocm7.2.sh index dcfc2f7d1..712d27bd1 100644 --- a/systems/Frontier/sourceme-rocm7.2.sh +++ b/systems/Frontier/sourceme-rocm7.2.sh @@ -3,17 +3,17 @@ module unload slate 2>/dev/null echo spack . /autofs/nccs-svm1_home1/paboyle/spack-frontier/share/spack/setup-env.sh -export SLATE=`spack find --paths slate | grep ^slate | awk '{print $2}' ` -export BLASPP=`spack find --paths blaspp | grep ^blaspp | awk '{print $2}' ` -export LAPACKPP=`spack find --paths lapackpp | grep ^lapackpp | awk '{print $2}' ` +#export SLATE=`spack find --paths slate | grep ^slate | awk '{print $2}' ` +#export BLASPP=`spack find --paths blaspp | grep ^blaspp | awk '{print $2}' ` +#export LAPACKPP=`spack find --paths lapackpp | grep ^lapackpp | awk '{print $2}' ` -echo SLATE $SLATE -echo LAPACKPP $LAPACKPP -echo BLASPP $BLASPP +#echo SLATE $SLATE +#echo LAPACKPP $LAPACKPP +#echo BLASPP $BLASPP -ls $SLATE/lib64/libslate.so $LAPACKPP/lib64/liblapackpp.so 2>/dev/null || ls $SLATE/lib $LAPACKPP/lib -ldd $SLATE/lib64/libslate.so | grep -E "blaspp|lapackpp|rocblas|amdhip|mpi_cray|libsci" -ldd $LAPACKPP/lib64/liblapackpp.so | grep -E "blaspp|libsci" +#ls $SLATE/lib64/libslate.so $LAPACKPP/lib64/liblapackpp.so 2>/dev/null || ls $SLATE/lib $LAPACKPP/lib +#ldd $SLATE/lib64/libslate.so | grep -E "blaspp|lapackpp|rocblas|amdhip|mpi_cray|libsci" +#ldd $LAPACKPP/lib64/liblapackpp.so | grep -E "blaspp|libsci" export CLIME=`spack find --paths c-lime | grep ^c-lime | awk '{print $2}' ` export MPFR=`spack find --paths mpfr | grep ^mpfr | awk '{print $2}' ` @@ -24,9 +24,14 @@ module load cce/21.0.0 module load cpe/26.03 module load rocm/7.2.0 export LD_LIBRARY_PATH=/opt/rocm-7.2.0/lib/llvm/lib/:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=$LAPACKPP/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=$BLASPP/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=$SLATE/lib64:$LD_LIBRARY_PATH +#export LD_LIBRARY_PATH=$LAPACKPP/lib64:$LD_LIBRARY_PATH +#export LD_LIBRARY_PATH=$BLASPP/lib64:$LD_LIBRARY_PATH +#export LD_LIBRARY_PATH=$SLATE/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$CRAY_LD_LIBRARY_PATH module load emacs + +# Standing avoidance of the CXI NO_TRANSLATION fault on persistent device +# buffers under hipMalloc/hipFree churn (libfabric issue #12775; ruling +# 2026-09-07). export FI_MR_CACHE_MONITOR=kdreg2 # REQUIRED for device-buffer MPI on Slingshot: libfabric memhooks monitor (default) is defective, see systems/WorkArounds.txt (libfabric #11451) +export FI_HMEM_ROCR_USE_DMABUF=0