Debug ulimit as core files driving me crazy

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) [head: 1st after flat ring]", sP2Phead);
-  line("GlobalSum(RealD) = GlobalSumP2Pdirect",   sP2Pdirect);
+  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);

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

@@ -32,6 +32,13 @@
 #         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)
@@ -49,6 +56,10 @@ 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
@@ -70,7 +81,7 @@ export LATT=$vol
 export LS=24
 
 # The subspace is generated on first use and saved to SUBSPACE_FILE (a few minutes).
-export SUBSPACE_FILE=$PWD/subspace_16x8x24x96_nb64.scidac
+export SUBSPACE_FILE=$LUSTRE/subspace_16x8x24x96_nb64.scidac   # persistent across jobs (generated once)
 unset  SLAB_FILE
 export HOT_START=1
 export MASS=0.00078

systems/Frontier/nrhs_fault_36.job

@@ -33,6 +33,13 @@
 #   (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)
@@ -50,6 +57,10 @@ 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
@@ -86,7 +97,7 @@ export NRHS=6
 run_cell () {
   name=$1; shift
   echo "----- $name : $* -----"
-  export GRID_STDOUT_ROOT=$PWD/fault36_$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