Solving memory pressure in NRHS=12 solver

Grid/allocator/MemoryManagerShared.cc

@@ -24,6 +24,7 @@ void  MemoryManager::PrintState(void* CpuPtr)
 std::cout << GridLogMessage << "Host<->Device memory movement not currently managed by Grid." << std::endl;
 };
 void  MemoryManager::Print(void){};
+void  MemoryManager::EvictAll(void){};   // nothing to evict in unified memory
 void  MemoryManager::PrintAll(void){};
 void  MemoryManager::NotifyDeletion(void *ptr){};
 

benchmarks/Benchmark_allreduce.cc

@@ -160,7 +160,7 @@ void Run(GridCartesian *grid, const char *tname, uint64_t nbytes_lo, uint64_t nb
 void Scalar(GridCartesian *grid, int reps, double gap_us)
 {
   int me=grid->ThisRank();
-  Stats<RealD> sP2P, sBare, sVec, sCart, sFlat, sP2Pdirect;
+  Stats<RealD> sP2P, sBare, sVec, sCart, sFlat, sP2Pdirect, sP2Phead;
   RealD x; deviceVector<RealD> d(1); std::vector<RealD> h(1);
   // gap: a bandwidth-bound device kernel sized to take ~gap_us (calibrated once)
   uint64_t gapN = 0; deviceVector<RealD> gapbuf(1);
@@ -179,7 +179,14 @@ void Scalar(GridCartesian *grid, int reps, double gap_us)
   // A Barrier before each timed call so that node skew (from the gap kernel or
   // anything else) is not charged to the reduction: the timer measures the
   // collective from a synchronised start.
+  // The same GlobalSum is timed twice per rep: first thing after the previous
+  // rep's flat ring (574 Sendrecvs per rank at P=288, staggered exit), and again
+  // after the MPI calls.  2026-08-28 at P=288: first-position mean 198 us vs 51 us
+  // for the identical primitive one line later -- position, not the wrapper
+  // (StepLog is a pointer assignment).  If "head" stays slow and "P2P" matches
+  // the direct call, MPI is draining the ring's backlog into the next P2P call.
   for(int r=0;r<reps;r++){
+    gap(); x=1.0+me;   grid->Barrier(); { double t0=usecond(); grid->GlobalSum(x);                 sP2Phead.add(usecond()-t0); }
     gap(); x=1.0+me;   grid->Barrier(); { double t0=usecond(); grid->GlobalSum(x);                 sP2P.add(usecond()-t0); }
     gap(); x=1.0+me;   grid->Barrier(); { double t0=usecond(); grid->GlobalSumP2P(x);              sP2Pdirect.add(usecond()-t0); }
     gap(); h[0]=1.0+me;grid->Barrier(); { double t0=usecond(); BareAllreduce(grid,&h[0],1);        sBare.add(usecond()-t0); }
@@ -190,11 +197,12 @@ void Scalar(GridCartesian *grid, int reps, double gap_us)
   auto mx=[&](double v){ RealD y=v; grid->GlobalMax(y); return (double)y; };
   auto line=[&](const char *nm, Stats<RealD> &st){
     double tmin=mx(st.tmin), tmean=mx(st.tsum/st.n);
-    if ( me==0 ) std::cout << GridLogMessage << "  " << std::setw(34) << std::left << nm << std::right
+    if ( me==0 ) std::cout << GridLogMessage << "  " << std::setw(42) << std::left << nm << std::right
                            << " min " << std::setw(8) << tmin << " us   mean " << std::setw(8) << tmean << " us" << std::endl; };
   if ( me==0 ) std::cout << GridLogMessage << "==== SCALAR latency (RealD, " << reps << " reps, slowest rank, Barrier before each timed call, gap " << gap_us << " us)  P=" << grid->ProcessorCount() << std::endl;
-  line("GlobalSum(RealD) = GlobalSumP2P",   sP2P);
-  line("GlobalSum(RealD) = GlobalSumP2Pdirect",   sP2Pdirect);
+  line("GlobalSum(RealD) [head: 1st after flat ring]", sP2Phead);
+  line("GlobalSum(RealD) = GlobalSumP2P [2nd]", sP2P);
+  line("GlobalSumP2P direct [3rd]",          sP2Pdirect);
   line("bare MPI_Allreduce(1 double)",       sBare);
   line("GlobalSumVector(double*,1) [MPI]",   sVec);
   line("CartesianRingAllReduce n=1 (device)",sCart);

examples/Example_pvdagm_v2_3level_DenseCoarseMatrix.cc

@@ -708,23 +708,9 @@ int main (int argc, char ** argv)
   }
   SetFineSloppy(0);
 
-  // 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
-
-  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);
-
+  // (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 V1
   // path, which needs a vectorised coarse space. Block Gram-Schmidt is
@@ -786,6 +772,27 @@ int main (int argc, char ** argv)
     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 (V1_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.
   //
@@ -940,6 +947,29 @@ int main (int argc, char ** argv)
   {
     std::cout << GridLogMessage << "**********************************************" << std::endl;
     std::cout << GridLogMessage << " V2 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();
+      MemoryManager::PrintBytes();
+      acceleratorMem();
+    }
     std::cout << GridLogMessage << "**********************************************" << std::endl;
 
     Coordinate cml({nr,1,clatt[0],clatt[1],clatt[2],clatt[3]});

skills/gpu-runtime-correctness.md

@@ -99,3 +99,33 @@ Grid uses `computeStream` (CUDA/HIP) or `theGridAccelerator` (SYCL) consistently
 3. If results are non-deterministic across runs, insert a second barrier and observe whether reproducibility improves.
 4. For correctness-critical operations (reductions that will be compared against reference values), add the double-run checksum test from `correctness-verification.md`.
 5. If the process hangs at 100% CPU in a runtime library function, this is a driver/runtime bug — there is no application-level fix beyond scheduling a node reboot.
+
+## Making an asynchronous GPU fault attributable
+
+A "Memory access fault by GPU node-N ... Reason: Unknown" is raised by the
+runtime, asynchronously, possibly long after the offending kernel or copy
+was *queued*.  `--debug-signals` sees nothing useful: the host is elsewhere
+by then.  Before exhaustive logging or bisection, make the fault land on the
+call that caused it:
+
+| runtime | serialise launches/copies | name kernels as they launch |
+|---|---|---|
+| HIP/ROCm | `AMD_SERIALIZE_KERNEL=3 AMD_SERIALIZE_COPY=3` (every launch and copy synchronous) | `AMD_LOG_LEVEL=3` (4 is very verbose) |
+| CUDA | `CUDA_LAUNCH_BLOCKING=1` | `compute-sanitizer --tool memcheck ./bin ...` (names the kernel and the bad address) |
+| SYCL/L0 | `SYCL_PI_TRACE=2` / `ZE_DEBUG=1` | (Level Zero validation layer: `ZE_ENABLE_VALIDATION_LAYER=1`) |
+
+5-10x slower, so run on the smallest reproducer (one node, small lattice).
+The last kernel name printed before the fault is the culprit; combined with
+Grid's `--log ...,Memory` (every MemoryManager transfer with size, direction,
+device and host pointers, in program order) it separates "stale pointer after
+eviction" from "out-of-range index" in one run.  First used 2026-08-28 for the
+NRHS>=6 fault in the 3-level example (`systems/Frontier/nrhs_fault.job`).
+
+**GPU core dumps (ROCm):** on a fault the runtime writes `gpucore.<pid>` (a
+GCD's memory, ~22 GB on MI250X) into the process's working directory; it obeys
+`ulimit -c` exactly as the kernel core does, so `ulimit -c 0` in the batch
+script suppresses it (Slurm propagates rlimits to `srun` tasks by default) and
+`ulimit -c unlimited` + `HSA_ENABLE_DEBUG=1` produces one ROCgdb can load.
+Deleted-while-open dumps on NFS persist as hidden `.nfs*` files against quota
+until the holder exits.  Run fault hunts from a Lustre directory.
+Source: ROCgdb documentation, "AMD GPU" chapter, core-dump section.

systems/Frontier/nrhs_fault.job

@@ -0,0 +1,127 @@
+#!/bin/bash -l
+#SBATCH --job-name=nrhs-fault
+#SBATCH --nodes=1
+#SBATCH --ntasks-per-node=8
+#SBATCH --cpus-per-task=7
+#SBATCH --gpus-per-node=8
+#SBATCH --time=0:45:00
+#SBATCH --account=phy157_dwf
+#SBATCH --gpu-bind=none
+#SBATCH --exclusive
+#SBATCH --mem=0
+#SBATCH -S 0
+#SBATCH -q debug
+
+##############################################################################
+# GPU "Memory access fault ... Reason: Unknown" in the 3-level example,
+# 2026-08-27/28: NRHS=12 and NRHS=6 (adaptive gcr/gcr) deep in the solve; once
+# at Nrhs=1 after an Nrhs=4 solve (replay cell S6).  NRHS=4 has never faulted.
+# --debug-signals catches nothing: the fault is raised by the ROCm runtime,
+# asynchronously, possibly long after the offending kernel/copy was queued.
+#
+# Plan: a 1-node reproducer (debug queue, minutes), then make the fault
+# ATTRIBUTABLE.  AMD_SERIALIZE_KERNEL=3 / AMD_SERIALIZE_COPY=3 make every
+# kernel launch and copy synchronous, so the fault surfaces at the call that
+# caused it; AMD_LOG_LEVEL=3 prints each kernel's name as it launches, so the
+# last name before the fault is the culprit.  ~5-10x slower: fine at this size.
+#
+# Cells:  F0 NRHS=4 control  |  F1 NRHS=6 (reproduce?)  |  F2 NRHS=6 serialised
+#         F3 NRHS=6 DENSE_CC=0 (no dense bottom)  |  F4 NRHS=6 --device-mem 48000 (no LRU eviction)
+# Read:   grep -n -B8 "Memory access fault" log.fault.F*  ;  for F2 the last
+#         "ShaderName"/kernel lines before the fault; F3/F4 say whether the
+#         dense bottom or MemoryManager eviction is involved.
+##############################################################################
+
+# Run in a per-job subdirectory on Lustre: logs, per-rank --debug-stdout trees and GPU
+# core files are large and must not land in $HOME (quota).  Everything below is
+# relative to RUNDIR; the job's own slurm-*.out stays where sbatch was run.
+LUSTRE=/lustre/orion/phy157/proj-shared/phy157_dwf/paboyle
+RUNDIR=$LUSTRE/runs/${SLURM_JOB_NAME}_${SLURM_JOB_ID}
+mkdir -p $RUNDIR && cd $RUNDIR && echo "RUNDIR $RUNDIR"
+
+cat << EOF > select_gpu
+#!/bin/bash
+export GPU_MAP=(0 1 2 3 7 6 5 4)
+export NUMA_MAP=(3 3 1 1 2 2 0 0)
+export GPU=\${GPU_MAP[\$SLURM_LOCALID]}
+export NUMA=\${NUMA_MAP[\$SLURM_LOCALID]}
+export HIP_VISIBLE_DEVICES=\$GPU
+unset ROCR_VISIBLE_DEVICES
+if [ \$SLURM_PROCID = "0" ]; then echo \$*; fi
+exec numactl -m \$NUMA -N \$NUMA \$*
+EOF
+chmod +x ./select_gpu
+
+root=$HOME/ParallelIO/systems/Frontier
+source $root/sourceme-rocm7.2.sh
+
+export OMP_NUM_THREADS=7
+# No core dumps: the ROCm runtime's GPU core dump (gpucore.<pid>, ~22 GB per faulting GCD, the
+# .nfs quota disaster of 2026-08-25) obeys ulimit -c like the kernel's (ROCgdb docs, "AMD GPU
+# core dumps"); Slurm propagates the batch shell's limits to srun tasks by default.
+ulimit -c 0
+export MPICH_GPU_SUPPORT_ENABLED=1
+export MPICH_SMP_SINGLE_COPY_MODE=CMA
+export MPICH_OFI_NIC_POLICY=GPU
+module load libfabric
+
+BIN=$root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix
+OPTS1="--accelerator-threads 8 --shm 4096 --shm-mpi 1"
+# LOCAL VOLUME MATCHED to production: 48.48.48.96 on 3.6.4.4 is local 16.8.12.24; on 8 GCDs
+# (1.1.2.4) that is global 16.8.24.96 with the production blockings 2.2.3.3 / 4.4.2.4:
+# coarse local 8.4.4.8 and coarse-coarse local 2.1.2.2 exactly as at 288 ranks.  The global
+# lattice is oddly shaped and the dense bottom is tiny (N = 2.1.4.8 x 60 = 3840, one block),
+# but every solver-side buffer has the production shape.  (First submission, 2026-08-28,
+# used 16.16.16.32 with 2^4 blocks: local 16.16.8.8 -- NOT matched.)
+vol=16.8.24.96
+MPI_GEOM=1.1.2.4
+# The example sizes its lattice from LATT/LS (env), NOT from --grid (first submission
+# asserted heap_bytes<heap_size trying to build 48^3x96 on one node for exactly this reason).
+export LATT=$vol
+export LS=24
+
+# The subspace is generated on first use and saved to SUBSPACE_FILE (a few minutes).
+export SUBSPACE_FILE=$LUSTRE/subspace_16x8x24x96_nb64.scidac   # persistent across jobs (generated once)
+unset  SLAB_FILE
+export HOT_START=1
+export MASS=0.00078
+export BLOCK=2.2.3.3
+export BLOCK2=4.4.2.4
+export DENSE_SCHUR=1
+export DENSE_SCHUR2D=1
+export DENSE_CC=1
+export DENSE_SPLITK=1
+export DENSE_DEVICE_SUM=4
+export SCHUR2D_LEAF_SPAN=1
+export SCHUR2D_LEAF_LU=1
+export SCHUR2D_PROBE=0
+export GRID_ALLOC_NCACHE_LARGE=64
+# banked adaptive solver settings
+export FineSmootherShift=0.1  FineSmootherOrder=6 FineSmootherMmax=4
+export CoarseSmootherShift=2.0 CoarseSmootherNstep=2 CoarseSmootherMmax=2
+export CoarseSolverTol=0.05 CoarseSolverOrder=200 CoarseSolverMmax=8
+export OuterTol=1e-8 OuterMmax=6 OuterNstep=12
+export PowerIterations=0 SmootherCoeffLog=0 PolyVerbose=0
+
+run_cell () {
+  name=$1; devmem=$2; shift 2
+  GRID_LOG_EXTRA=""; for kv in "$@"; do case $kv in GRID_LOG_EXTRA=*) GRID_LOG_EXTRA=${kv#GRID_LOG_EXTRA=};; esac; done
+  echo "----- $name : --device-mem $devmem  $* -----"
+  fname=log.fault.$name
+  logopt=""; [ -n "$GRID_LOG_EXTRA" ] && logopt="--log Error,Warning,Message,$GRID_LOG_EXTRA"
+  env "$@" srun -N1 -n8 --kill-on-bad-exit=1 ./select_gpu $BIN --mpi ${MPI_GEOM} --grid $vol $OPTS1 --device-mem $devmem --comms-overlap $logopt > $fname 2>&1
+  echo "  exit $?"; sleep 10
+  grep -h "V2 3-level solve\|Memory access fault\|GRID_ASSERT\|Converged on iteration" $fname | sed 's/^Grid : Message : [0-9.]* s : //' | tail -4 | cut -c1-140
+  grep -n -B3 "Memory access fault" $fname | tail -5 | cut -c1-160
+}
+
+#          name         devmem  environment
+run_cell   F0_nrhs4     16000   NRHS=4                                                     # control: never faulted at 288
+run_cell   F1_nrhs6     16000   NRHS=6                                                     # reproduce on 1 node?
+run_cell   F2_nrhs6_ser 16000   NRHS=6 AMD_SERIALIZE_KERNEL=3 AMD_SERIALIZE_COPY=3 AMD_LOG_LEVEL=3 GRID_LOG_EXTRA=Memory   # attributable fault
+run_cell   F3_nrhs6_nodense 16000 NRHS=6 DENSE_CC=0                                        # dense bottom out of the loop
+run_cell   F4_nrhs6_bigmem 48000 NRHS=6                                                    # no MemoryManager eviction
+
+echo "========================================================="
+for f in log.fault.F*; do echo "$f: $(grep -c 'Memory access fault' $f) faults; $(grep -h 'V2 3-level solve' $f | tail -1 | sed 's/.*V2/V2/' | cut -c1-60)"; done
+echo "========================================================="

systems/Frontier/nrhs_fault_36.job

@@ -0,0 +1,118 @@
+#!/bin/bash -l
+#SBATCH --job-name=nrhs-fault-36
+#SBATCH --nodes=36
+#SBATCH --ntasks-per-node=8
+#SBATCH --cpus-per-task=7
+#SBATCH --gpus-per-node=8
+#SBATCH --time=1:30:00
+#SBATCH --account=phy157_dwf
+#SBATCH --gpu-bind=none
+#SBATCH --exclusive
+#SBATCH --mem=0
+#SBATCH -S 0
+
+##############################################################################
+# The NRHS>=6 GPU memory access fault, on the ACTUAL target (48^3x96, 288 GCDs),
+# made attributable in ONE queue traversal:
+#
+#   AMD_SERIALIZE_KERNEL=3 AMD_SERIALIZE_COPY=3  every launch/copy synchronous:
+#                                the fault is raised at the call that caused it
+#   AMD_LOG_LEVEL=3               the runtime names each kernel as it launches
+#   --debug-stdout                every rank's stdout/stderr to its own file:
+#                                GRID_STDOUT_ROOT/<32*(rank/32)>/Grid.std{out,err}.<rank>
+#                                (the runtime's log goes to stderr, so the faulting
+#                                rank's Grid.stderr.<rank> ends with the culprit)
+#   --log ...,Memory              MemoryManager transfers in program order
+#
+# Cell A: serialised (5-10x slower: setup ~10 min + solve until the fault).
+# Cell B: NOT serialised, same logging -- if serialisation makes the fault
+#         vanish (a race), B still captures the faulting rank's Memory trace.
+#
+# Readout:
+#   f=$(grep -l "Memory access fault" fault36_A/*/Grid.stderr.* | head -1); echo $f; tail -40 $f
+#   (the last kernel name / MemoryManager line before the fault); same for fault36_B.
+##############################################################################
+
+# Run in a per-job subdirectory on Lustre: logs, per-rank --debug-stdout trees and GPU
+# core files are large and must not land in $HOME (quota).  Everything below is
+# relative to RUNDIR; the job's own slurm-*.out stays where sbatch was run.
+LUSTRE=/lustre/orion/phy157/proj-shared/phy157_dwf/paboyle
+RUNDIR=$LUSTRE/runs/${SLURM_JOB_NAME}_${SLURM_JOB_ID}
+mkdir -p $RUNDIR && cd $RUNDIR && echo "RUNDIR $RUNDIR"
+
+cat << EOF > select_gpu
+#!/bin/bash
+export GPU_MAP=(0 1 2 3 7 6 5 4)
+export NUMA_MAP=(3 3 1 1 2 2 0 0)
+export GPU=\${GPU_MAP[\$SLURM_LOCALID]}
+export NUMA=\${NUMA_MAP[\$SLURM_LOCALID]}
+export HIP_VISIBLE_DEVICES=\$GPU
+unset ROCR_VISIBLE_DEVICES
+if [ \$SLURM_PROCID = "0" ]; then echo \$*; fi
+exec numactl -m \$NUMA -N \$NUMA \$*
+EOF
+chmod +x ./select_gpu
+
+root=$HOME/ParallelIO/systems/Frontier
+source $root/sourceme-rocm7.2.sh
+
+export OMP_NUM_THREADS=7
+# No core dumps: the ROCm runtime's GPU core dump (gpucore.<pid>, ~22 GB per faulting GCD, the
+# .nfs quota disaster of 2026-08-25) obeys ulimit -c like the kernel's (ROCgdb docs, "AMD GPU
+# core dumps"); Slurm propagates the batch shell's limits to srun tasks by default.
+ulimit -c 0
+export MPICH_GPU_SUPPORT_ENABLED=1
+export MPICH_SMP_SINGLE_COPY_MODE=CMA
+export MPICH_OFI_NIC_POLICY=GPU
+module load libfabric
+
+BIN=$root/examples/Example_pvdagm_v2_3level_DenseCoarseMatrix
+OPTS1="--accelerator-threads 8 --shm 4096 --shm-mpi 1 --device-mem 32000"
+vol=48.48.48.96
+MPI_GEOM=3.6.4.4
+
+# production setup (as smoother_modes.job, banked adaptive point)
+export SUBSPACE_FILE=/lustre/orion/phy157/proj-shared/phy157_dwf/paboyle/subspace_nb64.scidac
+unset  SLAB_FILE
+export MASS=0.00078
+export BLOCK=2.2.3.3
+export BLOCK2=4.4.2.4
+export DENSE_SCHUR=1
+export DENSE_SCHUR2D=1
+export DENSE_CC=1
+export DENSE_APPLY_PROFILE=0
+export DENSE_SPLITK=128
+export DENSE_DEVICE_SUM=4
+export SCHUR2D_LEAF_SPAN=9
+export SCHUR2D_LEAF_LU=1
+export SCHUR2D_PROBE=0
+export GRID_ALLOC_NCACHE_LARGE=64
+export FineSmootherShift=0.1  FineSmootherOrder=6 FineSmootherMmax=4
+export CoarseSmootherShift=2.0 CoarseSmootherNstep=2 CoarseSmootherMmax=2
+export CoarseSolverTol=0.05 CoarseSolverOrder=200 CoarseSolverMmax=8
+export OuterTol=1e-8 OuterMmax=6 OuterNstep=12
+export PowerIterations=0 SmootherCoeffLog=0 PolyVerbose=0
+export NRHS=6
+
+run_cell () {
+  name=$1; shift
+  echo "----- $name : $* -----"
+  export GRID_STDOUT_ROOT=$RUNDIR/fault36_$name
+  env "$@" srun -N36 -n288 --kill-on-bad-exit=1 ./select_gpu $BIN --mpi ${MPI_GEOM} --grid $vol $OPTS1 --comms-overlap \
+       --debug-stdout --log Error,Warning,Message,Memory > log.fault36.$name 2>&1
+  echo "  exit $?"; sleep 30
+  f=$(grep -l "Memory access fault" $GRID_STDOUT_ROOT/*/Grid.stderr.* 2>/dev/null | head -1)
+  if [ -n "$f" ]; then
+    echo "  FAULT in $f  (rank $(basename $f | sed 's/Grid.stderr.//'))"
+    rank=$(basename $f | sed 's/Grid.stderr.//')
+    echo "  --- last 30 lines of stderr (runtime log) ---";  tail -30 $f | cut -c1-160
+    echo "  --- last 15 lines of that rank's stdout (Grid Message/Memory) ---"
+    tail -15 $(dirname $f)/Grid.stdout.$rank | cut -c1-160
+  else
+    echo "  no 'Memory access fault' in any Grid.stderr.*"
+    grep -h "V2 3-level solve\|Fouter MrhsPGCR: Converged" $GRID_STDOUT_ROOT/0/Grid.stdout.0 | tail -3 | cut -c1-120
+  fi
+}
+
+run_cell   A_serialised   AMD_SERIALIZE_KERNEL=3 AMD_SERIALIZE_COPY=3 AMD_LOG_LEVEL=3
+run_cell   B_plain        AMD_LOG_LEVEL=1