From 2264afda845cfb2c33f2ac7c1f1d0525835223cc Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Thu, 27 Aug 2026 22:11:06 -0400 Subject: [PATCH] Allreduce benchmarking --- benchmarks/Benchmark_allreduce.cc | 191 +++++++++++++++++++++++++++ systems/Frontier/allreduce_bench.job | 75 +++++++++++ 2 files changed, 266 insertions(+) create mode 100644 benchmarks/Benchmark_allreduce.cc create mode 100644 systems/Frontier/allreduce_bench.job diff --git a/benchmarks/Benchmark_allreduce.cc b/benchmarks/Benchmark_allreduce.cc new file mode 100644 index 000000000..0a57dac6d --- /dev/null +++ b/benchmarks/Benchmark_allreduce.cc @@ -0,0 +1,191 @@ +/************************************************************************************* + + Grid physics library, www.github.com/paboyle/Grid + + Source file: ./benchmarks/Benchmark_allreduce.cc + + Copyright (C) 2026 + +Author: Peter Boyle + + This program is free software; you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 2 of the License, or + (at your option) any later version. + + See the full license in the file "LICENSE" in the top level distribution + directory +*************************************************************************************/ +/* END LEGAL */ + +////////////////////////////////////////////////////////////////////////////// +// Vector all-reduce shoot-out: MPI_Allreduce vs the P2P rings, across sizes, +// so GlobalSumVector's dispatch (if any) can be decided from a table rather +// than a guess. For each size, each type, the SAME data are reduced by: +// +// MPI-host GlobalSumVector on a host buffer (today's path) +// MPI-bare MPI_Allreduce(MPI_IN_PLACE,..) directly on grid->communicator, +// host buffer: GlobalSumVector minus its FlightRecorder step log +// MPI-dev GlobalSumVector on a device buffer (GPU-aware MPI; +// Cray MPICH aborts +// above ~8 MB: capped +// by BENCH_MPI_DEV_MAX_MB) +// cart-dev CartesianRingAllReduce, device buffer (~sum_d 2N bytes) +// flat-dev RingAllReduce, device buffer (2N(P-1)/P bytes) +// cart-host H2D copy + cart-dev + D2H (the host path when +// flat-host H2D copy + flat-dev + D2H the caller holds a +// std::vector) +// +// Every method is checked against MPI-host; the deterministic rings are also +// checked bitwise against themselves on a repeat. Reported per method: time, +// payload rate (N bytes / t), and "wire" rate (bytes actually moved by that +// method / t), min over reps. Rerun at 1, 4 and 36 nodes: the crossovers (if +// any) move with P as sum_d 2(P_d-1) vs 2(P-1) steps. +// +// SCALAR section (latency, min/mean of many reps): GlobalSum(RealD) -- which is +// the latency-hiding GlobalSumP2P (all shifts of a dimension posted at once, +// one completion per dimension) -- vs bare MPI_Allreduce of one double vs +// GlobalSumVector(1) vs the two rings at n=1. +// +// mpirun -n 8 ./Benchmark_allreduce --grid 16.16.16.32 --mpi 1.1.2.4 +// env: BENCH_MIN_KB (4) BENCH_MAX_MB (512) BENCH_REPS (5) +// BENCH_MPI_DEV_MAX_MB (4: above this MPI-dev is skipped, not attempted) +////////////////////////////////////////////////////////////////////////////// + +#include + +using namespace Grid; + +template T Fill(uint64_t i,int rank){ return T(0.5*std::sin(0.01*i+0.7*rank)+1.0e-3*rank); } + +template MPI_Datatype MpiType(void); +template<> MPI_Datatype MpiType(void){ return MPI_FLOAT; } +template<> MPI_Datatype MpiType(void){ return MPI_DOUBLE; } +template void BareAllreduce(GridCartesian *grid, T *buf, uint64_t n) +{ +#ifdef GRID_COMMS_MPI3 + int ierr = MPI_Allreduce(MPI_IN_PLACE, buf, (int)n, MpiType(), MPI_SUM, grid->communicator); + GRID_ASSERT(ierr==0); +#else + grid->GlobalSumVector(buf,(int)n); +#endif +} + +template struct Stats { double tmin=1e30, tmax=0, tsum=0; int n=0; void add(double t){ tmin=std::min(tmin,t); tmax=std::max(tmax,t); tsum+=t; n++; } }; + +template +void Run(GridCartesian *grid, const char *tname, uint64_t nbytes_lo, uint64_t nbytes_hi, int reps, uint64_t mpiDevMax) +{ + int P = grid->ProcessorCount(); + int me = grid->ThisRank(); + int Nd = grid->_ndimension; + // bytes moved per rank by each method, per byte of payload + double flatFactor = 2.0*(P-1)/P; + double cartFactor = 0.0; for(int d=0;d_processors[d]>1) cartFactor += 2.0*(grid->_processors[d]-1)/grid->_processors[d]; + + if ( me==0 ) { + std::cout << GridLogMessage << "==== " << tname << " P=" << P << " grid " << grid->_processors + << " bytes moved per payload byte: flat " << flatFactor << " cartesian " << cartFactor << std::endl; + std::cout << GridLogMessage << std::setw(10) << "bytes" << std::setw(12) << "elements" + << " | " << std::setw(9) << "MPI-host" << std::setw(9) << "MPI-bare" << std::setw(9) << "MPI-dev" << std::setw(9) << "cart-dev" << std::setw(9) << "flat-dev" + << std::setw(10) << "cart-host" << std::setw(10) << "flat-host" << " (ms, min of " << reps << ")" + << " | payload GB/s: " << "cart-dev flat-dev | wire GB/s: cart-dev flat-dev | check" << std::endl; + } + for(uint64_t bytes=nbytes_lo; bytes<=nbytes_hi; bytes*=2){ + uint64_t n = bytes/sizeof(T); + if ( n==0 ) continue; + std::vector h(n); for(uint64_t i=0;i(i,me); + std::vector ref(h); grid->GlobalSumVector(&ref[0],(int)n); // MPI-host reference (also warm-up) + std::vector out(n); + deviceVector d(n); + Stats sMh, sMb, sMd, sCd, sFd, sCh, sFh; + double worst[7]={0,0,0,0,0,0,0}; int done[7]={1,0,1,1,1,1,1}; + auto check=[&](int k, const std::vector &o){ double w=0; for(uint64_t i=0;iGlobalSumVector(&out[0],(int)n); sMh.add(usecond()-t0); } check(0,out); + // MPI-bare + out=h; { double t0=usecond(); BareAllreduce(grid,&out[0],n); sMb.add(usecond()-t0); } check(6,out); + // MPI-dev + if ( doMpiDev ) { acceleratorCopyToDevice(&h[0],&d[0],n*sizeof(T)); double t0=usecond(); grid->GlobalSumVector(&d[0],(int)n); sMd.add(usecond()-t0); + acceleratorCopyFromDevice(&d[0],&out[0],n*sizeof(T)); check(1,out); } + // cart-dev + { acceleratorCopyToDevice(&h[0],&d[0],n*sizeof(T)); double t0=usecond(); CartesianRingAllReduce(grid,&d[0],n); sCd.add(usecond()-t0); + acceleratorCopyFromDevice(&d[0],&out[0],n*sizeof(T)); check(2,out); } + // flat-dev + { acceleratorCopyToDevice(&h[0],&d[0],n*sizeof(T)); double t0=usecond(); RingAllReduce(grid,&d[0],n); sFd.add(usecond()-t0); + acceleratorCopyFromDevice(&d[0],&out[0],n*sizeof(T)); check(3,out); } + // cart-host: staged + { out=h; double t0=usecond(); acceleratorCopyToDevice(&out[0],&d[0],n*sizeof(T)); CartesianRingAllReduce(grid,&d[0],n); acceleratorCopyFromDevice(&d[0],&out[0],n*sizeof(T)); sCh.add(usecond()-t0); check(4,out); } + // flat-host: staged + { out=h; double t0=usecond(); acceleratorCopyToDevice(&out[0],&d[0],n*sizeof(T)); RingAllReduce(grid,&d[0],n); acceleratorCopyFromDevice(&d[0],&out[0],n*sizeof(T)); sFh.add(usecond()-t0); check(5,out); } + } + // rings: bitwise repeatable? + int bitwise=1; + { std::vector a(n),b(n); + acceleratorCopyToDevice(&h[0],&d[0],n*sizeof(T)); CartesianRingAllReduce(grid,&d[0],n); acceleratorCopyFromDevice(&d[0],&a[0],n*sizeof(T)); + acceleratorCopyToDevice(&h[0],&d[0],n*sizeof(T)); CartesianRingAllReduce(grid,&d[0],n); acceleratorCopyFromDevice(&d[0],&b[0],n*sizeof(T)); + if ( memcmp(&a[0],&b[0],n*sizeof(T)) ) bitwise=0; } + // worst error over ranks; tolerance scales with P and magnitude + double scale=0; for(uint64_t i=0;iGlobalMax(wmax); + RealD bw=1.0-bitwise; grid->GlobalMax(bw); + double tol = (sizeof(T)==4 ? 1.0e-5 : 1.0e-12) * std::max(scale,1.0) * P; + // times: use the slowest rank's min (the collective completes when the last rank does) + auto mx=[&](double v){ RealD x=v; grid->GlobalMax(x); return (double)x; }; + double tMh=mx(sMh.tmin), tMb=mx(sMb.tmin), tMd=doMpiDev?mx(sMd.tmin):0, tCd=mx(sCd.tmin), tFd=mx(sFd.tmin), tCh=mx(sCh.tmin), tFh=mx(sFh.tmin); + if ( me==0 ) { + std::streamsize prec=std::cout.precision(); + std::cout << GridLogMessage << std::setw(10) << bytes << std::setw(12) << n << " | " << std::fixed << std::setprecision(3) + << std::setw(9) << tMh/1000. << std::setw(9) << tMb/1000. << std::setw(9) << (doMpiDev ? tMd/1000. : -1.0) + << std::setw(9) << tCd/1000. << std::setw(9) << tFd/1000. << std::setw(10) << tCh/1000. << std::setw(10) << tFh/1000. + << " | " << std::setprecision(2) << std::setw(8) << bytes/tCd/1.0e3 << std::setw(9) << bytes/tFd/1.0e3 + << " | " << std::setw(8) << cartFactor*bytes/tCd/1.0e3 << std::setw(9) << flatFactor*bytes/tFd/1.0e3 + << " | " << (wmax0 ? " **ring not bitwise**" : "") + << (doMpiDev ? "" : " (MPI-dev skipped: > BENCH_MPI_DEV_MAX_MB)") << std::endl; + std::cout.unsetf(std::ios::fixed); std::cout.precision(prec); + } + } +} + +void Scalar(GridCartesian *grid, int reps) +{ + int me=grid->ThisRank(); + Stats sP2P, sBare, sVec, sCart, sFlat; + RealD x; deviceVector d(1); std::vector h(1); + for(int r=0;rGlobalSum(x); sP2P.add(usecond()-t0); } + h[0]=1.0+me;{ double t0=usecond(); BareAllreduce(grid,&h[0],1); sBare.add(usecond()-t0); } + h[0]=1.0+me;{ double t0=usecond(); grid->GlobalSumVector(&h[0],1); sVec.add(usecond()-t0); } + acceleratorPut(d[0],h[0]); { double t0=usecond(); CartesianRingAllReduce(grid,&d[0],1); sCart.add(usecond()-t0); } + acceleratorPut(d[0],h[0]); { double t0=usecond(); RingAllReduce(grid,&d[0],1); sFlat.add(usecond()-t0); } + } + 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 + << " 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) P=" << grid->ProcessorCount() << std::endl; + line("GlobalSum(RealD) = GlobalSumP2P", sP2P); + line("bare MPI_Allreduce(1 double)", sBare); + line("GlobalSumVector(double*,1) [MPI]", sVec); + line("CartesianRingAllReduce n=1 (device)",sCart); + line("RingAllReduce flat n=1 (device)", sFlat); +} + +int main(int argc, char **argv) +{ + Grid_init(&argc, &argv); + GridCartesian *grid = SpaceTimeGrid::makeFourDimGrid(GridDefaultLatt(), GridDefaultSimd(Nd, vComplexD::Nsimd()), GridDefaultMpi()); + Scalar(grid, getenv("BENCH_SCALAR_REPS") ? atoi(getenv("BENCH_SCALAR_REPS")) : 200); + uint64_t lo = (getenv("BENCH_MIN_KB") ? atol(getenv("BENCH_MIN_KB")) : 4) * 1024ull; + uint64_t hi = (getenv("BENCH_MAX_MB") ? atol(getenv("BENCH_MAX_MB")) : 512) * 1024ull*1024ull; + int reps = getenv("BENCH_REPS") ? atoi(getenv("BENCH_REPS")) : 5; + uint64_t mdev = (getenv("BENCH_MPI_DEV_MAX_MB") ? atol(getenv("BENCH_MPI_DEV_MAX_MB")) : 4) * 1024ull*1024ull; + std::cout << GridLogMessage << "Benchmark_allreduce: P=" << grid->ProcessorCount() << " grid " << grid->_processors + << " sizes " << lo << " .. " << hi << " bytes, reps " << reps << ", MPI-dev attempted up to " << mdev << " bytes" << std::endl; + Run(grid,"RealF (MPI_FLOAT)", lo,hi,reps,mdev); + Run(grid,"RealD (MPI_DOUBLE)",lo,hi,reps,mdev); + Grid_finalize(); +} diff --git a/systems/Frontier/allreduce_bench.job b/systems/Frontier/allreduce_bench.job new file mode 100644 index 000000000..01220b17d --- /dev/null +++ b/systems/Frontier/allreduce_bench.job @@ -0,0 +1,75 @@ +#!/bin/bash -l +#SBATCH --job-name=allreduce-bench +#SBATCH --nodes=36 +#SBATCH --ntasks-per-node=8 +#SBATCH --cpus-per-task=7 +#SBATCH --gpus-per-node=8 +#SBATCH --time=0:30:00 +#SBATCH --account=phy157_dwf +#SBATCH --gpu-bind=none +#SBATCH --exclusive +#SBATCH --mem=0 +#SBATCH -S 0 + +############################################################################## +# Benchmark_allreduce at 1, 4 and 36 nodes from one allocation: MPI_Allreduce +# (host and device buffers) vs CartesianRingAllReduce vs flat RingAllReduce +# (device, and host with explicit staging), 4 KB .. 512 MB, RealF and RealD. +# Decides whether GlobalSumVector wants a size threshold or simply the ring +# (deterministic, no size cliff) everywhere. Cray MPICH's device-buffer +# allreduce aborted above ~8 MB (MPI_FLOAT) in earlier runs: MPI-dev is +# attempted only up to BENCH_MPI_DEV_MAX_MB (4); raise it to probe the cliff. +# +# Readout: the tables (one per type per node count). Columns are ms (min of +# reps, slowest rank), payload GB/s and "wire" GB/s (bytes each ring actually +# moves / time). Compare MPI-host with cart-dev / flat-dev at each size; the +# crossovers, if any, should move with P as sum_d 2(P_d-1) vs 2(P-1) steps. +############################################################################## + +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 +export MPICH_GPU_SUPPORT_ENABLED=1 +export MPICH_SMP_SINGLE_COPY_MODE=CMA +export MPICH_OFI_NIC_POLICY=GPU +module load libfabric + +BIN=$root/benchmarks/Benchmark_allreduce +OPTS1="--accelerator-threads 8 --shm 4096 --shm-mpi 1 --device-mem 32000" +export BENCH_MIN_KB=4 +export BENCH_MAX_MB=512 +export BENCH_REPS=5 +export BENCH_MPI_DEV_MAX_MB=4 +export BENCH_SCALAR_REPS=200 + +run_cell () { + name=$1; nodes=$2; ntask=$3; geom=$4; vol=$5 + echo "----- $name : $nodes nodes, $ntask ranks, --mpi $geom -----" + fname=log.allreduce.$name + srun -N$nodes -n$ntask --kill-on-bad-exit=1 ./select_gpu $BIN --mpi $geom --grid $vol $OPTS1 > $fname 2>&1 + echo " exit $?"; sleep 20 + grep -h "====\|bytes *elements\|| .* |\| min .* us" $fname | sed 's/^Grid : Message : [0-9.]* s : //' | cut -c1-175 +} + +# name nodes ntask geom vol +run_cell N1_8gcd 1 8 1.1.2.4 16.16.32.64 +run_cell N4_32gcd 4 32 2.2.2.4 32.32.32.64 +run_cell N36_288gcd 36 288 3.6.4.4 48.48.48.96 + +echo "=========================================================" +echo "done: see log.allreduce.N1_8gcd, N4_32gcd, N36_288gcd" +echo "========================================================="