From 63c2cdb7123a21425cddb64a311cf0151fe2b812 Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Fri, 28 Aug 2026 11:20:43 -0400 Subject: [PATCH] Debug ulimit as core files driving me crazy --- benchmarks/Benchmark_allreduce.cc | 16 ++++++++++++---- skills/gpu-runtime-correctness.md | 30 ++++++++++++++++++++++++++++++ systems/Frontier/nrhs_fault.job | 13 ++++++++++++- systems/Frontier/nrhs_fault_36.job | 13 ++++++++++++- 4 files changed, 66 insertions(+), 6 deletions(-) diff --git a/benchmarks/Benchmark_allreduce.cc b/benchmarks/Benchmark_allreduce.cc index 034bad28b..cdda67fb9 100644 --- a/benchmarks/Benchmark_allreduce.cc +++ b/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 sP2P, sBare, sVec, sCart, sFlat, sP2Pdirect; + Stats sP2P, sBare, sVec, sCart, sFlat, sP2Pdirect, sP2Phead; RealD x; deviceVector d(1); std::vector h(1); // gap: a bandwidth-bound device kernel sized to take ~gap_us (calibrated once) uint64_t gapN = 0; deviceVector 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;rBarrier(); { 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 &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); diff --git a/skills/gpu-runtime-correctness.md b/skills/gpu-runtime-correctness.md index 24e696d07..4f9b4bd70 100644 --- a/skills/gpu-runtime-correctness.md +++ b/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.` (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. diff --git a/systems/Frontier/nrhs_fault.job b/systems/Frontier/nrhs_fault.job index f2949a82c..d6b6ba756 100644 --- a/systems/Frontier/nrhs_fault.job +++ b/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., ~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 diff --git a/systems/Frontier/nrhs_fault_36.job b/systems/Frontier/nrhs_fault_36.job index 257568b98..ef827a25e 100644 --- a/systems/Frontier/nrhs_fault_36.job +++ b/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., ~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