mirror of
https://github.com/paboyle/Grid.git
synced 2026-09-10 03:33:16 +01:00
Multigrid clean up phase, and fix shell variables for libfabric bug work around on Frontier
This commit is contained in:
Peter Boyle
@@ -35,3 +35,8 @@ Author: Peter Boyle <pboyle@bnl.gov>
|
||||
#include <Grid/algorithms/multigrid/GeneralCoarsenedMatrixMultiRHSV2.h>
|
||||
#include <Grid/algorithms/multigrid/MrhsPromotedOperator.h>
|
||||
#include <Grid/algorithms/multigrid/Smoothers.h>
|
||||
#include <Grid/algorithms/multigrid/PVdagMMultiGridParams.h>
|
||||
// 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).
|
||||
|
||||
@@ -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<nsimd;lane++)
|
||||
dbuf[ss*nsimd+lane] = fbuf[ss*nsimd+lane]; //conversion
|
||||
#endif
|
||||
});
|
||||
|
||||
} else if ( sizeof(word)==4){
|
||||
@@ -467,8 +481,13 @@ public:
|
||||
uint16_t *hbuf =(uint16_t *) packet.compressed_recv_buf;
|
||||
|
||||
accelerator_forNB(ss,outer,nsimd,{
|
||||
#ifdef GRID_SIMT
|
||||
int lane = acceleratorSIMTlane(nsimd);
|
||||
fbuf[ss*nsimd+lane] = ((uint32_t)hbuf[ss*nsimd+lane])<<16; //copy back and pad each word with zeroes
|
||||
#else
|
||||
for(int lane=0;lane<nsimd;lane++)
|
||||
fbuf[ss*nsimd+lane] = ((uint32_t)hbuf[ss*nsimd+lane])<<16; //copy back and pad each word with zeroes
|
||||
#endif
|
||||
});
|
||||
|
||||
} else {
|
||||
@@ -515,9 +534,15 @@ public:
|
||||
double *dbuf =(double *) packet.send_buf;
|
||||
float *fbuf =(float *) packet.compressed_send_buf;
|
||||
|
||||
// BUG FIX 2026-09-07: CPU lane coverage -- see DecompressPacket.
|
||||
accelerator_forNB(ss,outer,nsimd,{
|
||||
#ifdef GRID_SIMT
|
||||
int lane = acceleratorSIMTlane(nsimd);
|
||||
fbuf[ss*nsimd+lane] = dbuf[ss*nsimd+lane]; // convert fp64 to fp32
|
||||
#else
|
||||
for(int lane=0;lane<nsimd;lane++)
|
||||
fbuf[ss*nsimd+lane] = dbuf[ss*nsimd+lane]; // convert fp64 to fp32
|
||||
#endif
|
||||
});
|
||||
|
||||
} else if ( sizeof(word)==4){
|
||||
@@ -526,8 +551,13 @@ public:
|
||||
uint16_t *hbuf =(uint16_t *) packet.compressed_send_buf;
|
||||
|
||||
accelerator_forNB(ss,outer,nsimd,{
|
||||
#ifdef GRID_SIMT
|
||||
int lane = acceleratorSIMTlane(nsimd);
|
||||
hbuf[ss*nsimd+lane] = fbuf[ss*nsimd+lane]>>16; // convert as in Bagel/BFM ; bfloat16 ; s7e8 Intel patent
|
||||
#else
|
||||
for(int lane=0;lane<nsimd;lane++)
|
||||
hbuf[ss*nsimd+lane] = fbuf[ss*nsimd+lane]>>16; // convert as in Bagel/BFM ; bfloat16 ; s7e8 Intel patent
|
||||
#endif
|
||||
});
|
||||
|
||||
} else {
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user