Document FSI-SPH ROCm/HIP tuning on AMD Instinct GPUs

[amd-pratmish]

Summary

Adds documentation for ROCm / HIP performance tuning of Chrono::FSI-SPH demos on AMD Instinct GPUs (MI300X gfx942, MI350X gfx950). No physics or solver source changes — this PR is documentation only: build/run instructions, tuning knobs, and why each change can improve wall-clock time.

Validated on seven standard FSI-SPH demos: DamBreak, Kernel, ObjectDrop, WaveTank, BaffleFlow, AngleRepose, RassorDrum.

---

Motivation

FSI-SPH wall time on Instinct is split between GPU kernels (forces, neighbor search, activity), CPU orchestration (MBS coupling, I/O), and ROCm runtime overhead (launches, copies, queues). Default demo settings often leave particle CSV output and every-step neighbor search enabled, which dominates long fluid runs. This PR documents which knobs remove each bottleneck and how to run tuned examples from the command line.

---

Files added

Path	Purpose
doxygen/documentation/manuals/fsi/fsi_sph_amd_instinct_tuning.md	Full Doxygen manual (build, run, catalog, per-demo recipes)
doxygen/documentation/manuals/fsi/manual_fsi.md	Link new page from FSI manual index
docs/FSI_SPH_AMD_TUNING.md	Plain-Markdown quick reference
docs/fsi_sph_tuning/*.env.example	Sourceable ROCm / CLI tuning snippets

---

Build all FSI-SPH demos

export ROCM_PATH=${ROCM_PATH:-/opt/rocm}

cmake -S . -B build \
  -DCMAKE_BUILD_TYPE=Release \
  -DBUILD_DEMOS=ON \
  -DENABLE_MODULE_FSI=ON \
  -DCHRONO_HIP=ON \
  -DCMAKE_HIP_ARCHITECTURES=gfx942   # MI300X; use gfx950 for MI350X

cmake --build build --target \
  demo_FSI-SPH_DamBreak demo_FSI-SPH_Kernels demo_FSI-SPH_ObjectDrop \
  demo_FSI-SPH_WaveTank demo_FSI-SPH_BaffleFlow demo_FSI-SPH_AngleRepose \
  demo_FSI-SPH_RassorDrum -j $(nproc)

cd build
export LD_LIBRARY_PATH=$PWD/lib:${ROCM_PATH}/lib:${ROCM_PATH}/lib64:${LD_LIBRARY_PATH}

Run from the build directory so relative data paths (../data/) resolve.

Demo sources: src/demos/fsi/sph/demo_FSI-SPH_*.cpp

---

Baseline ROCm environment

Apply before any tuned run:

export HSA_XNACK=0
export HIP_LAUNCH_BLOCKING=0
export GPU_MAX_HW_QUEUES=2
export ROCR_VISIBLE_DEVICES=0
export HIP_VISIBLE_DEVICES=0

Variable	Why use this default
HSA_XNACK=0	Avoids unified-memory page faults on discrete Instinct
HIP_LAUNCH_BLOCKING=0	Async launches — CPU and GPU progress in parallel
GPU_MAX_HW_QUEUES=2	ROCm default; increase to 4 in combined profiles (see below)

---

Recommended tuning by demo

Demo	Primary lever	Notes
DamBreak	Skip I/O + ps_freq 2	Fluid demo; search + output dominate default wall time
Kernel	HSA_ENABLE_SDMA=0	Latency-bound micro-benchmark; strong effect on MI350-class GPUs
ObjectDrop	-ffast-math HIP rebuild	Force-loop compile tuning
WaveTank	Skip output only	Use --output_particle_data false; do not use ps_freq 2
BaffleFlow	Skip I/O + ps_freq 2	Same pattern as DamBreak
AngleRepose	Combined runtime + compile + CLI	Long CRM run; stack queues, SDMA, rebuild, and CLI
RassorDrum	Baseline / data path	HBM-bound; verify BCE load and build-dir data path

---

Expected improvement (typical ranges)

Approximate wall-clock reduction vs default demo settings on AMD Instinct. Ranges are indicative — profile on your hardware, ROCm version, and problem size.

Per optimization knob

Knob	Typical range
Skip particle output (--output false / --output_particle_data false)	5–15%
Reduced search cadence (--ps_freq 2)	5–15%
Output off + ps_freq 2 (fluid demos, where stable)	10–25%
HSA_ENABLE_SDMA=0 (Kernel micro-benchmark)	0–40%
GPU_MAX_HW_QUEUES=4	0–5%
HIP rebuild: -O3	0–5%
HIP rebuild: -ffast-math	2–8%
HIP rebuild: -amdgpu-early-inline-all	3–8%
HIP rebuild: fast-math + inline (combined)	5–12%
Combined runtime + compile + CLI (long CRM/fluid runs)	8–20%

Per demo (recommended recipe above)

Demo	Typical range
DamBreak	10–25%
Kernel	0–40%
ObjectDrop	3–10%
WaveTank	0–8%
BaffleFlow	10–20%
AngleRepose	5–15%
RassorDrum	0–2%

Unified profile (one binary, all demos): 5–15% on I/O/search-bound fluid demos; 0–2% on RassorDrum; Kernel varies by GPU (see above).

---

Why these changes are faster

Each knob removes a specific bottleneck on the Instinct timeline:

Bottleneck	Tuning knob	Why wall clock drops
Host I/O and D2H sync	--output false / --output_particle_data false	Default demos pack marker state on GPU, copy to host, and write CSV. That stalls the GPU and burns CPU without advancing physics. Disabling output removes those transfers and writes.
O(N) neighbor grid rebuilds	--ps_freq 2	Proximity search sorts particles and rebuilds neighbor lists (SphCollisionSystem) every step at default. At ps_freq=2 rebuild runs every other step; force integration still runs each step. Fluid demos spend large time in search — halving frequency reduces that subsystem cost.
SDMA vs compute contention	HSA_ENABLE_SDMA=0	ROCm can route small HSA copies through the SDMA engine. For tiny latency-bound kernels (Kernel demo), SDMA setup can exceed copy time. Compute-engine memcpy aligns better with back-to-back launches.
Kernel launch / queue back-pressure	GPU_MAX_HW_QUEUES=4	SPH submits bursts of short HIP kernels per step. More hardware queues let ROCm keep work in flight while prior kernels run — helps long runs (AngleRepose).
FP instruction throughput	-O3 -ffast-math	SPH force loops are FP-heavy neighbor reductions. Fast-math lets LLVM reassociate and vectorize on CDNA.
Kernel launch overhead	-mllvm -amdgpu-early-inline-all=true	Many small __device__ helpers → many launches. Inlining merges them into fewer entry points.

Why some demos see little benefit

• RassorDrum: Large marker counts saturate HBM bandwidth. Once memory-bound, I/O/search/compile tweaks hit a bytes/s roofline — not FLOPS-limited.
• Kernel on MI300X: Very short total runtime — GPU already saturated; SDMA routing changes have limited headroom.
• WaveTank + ps_freq=2: Reduces search work but breaks WCSPH stability (solver abort) — not a usable tradeoff.

---

Tuning catalog (knob → mechanism → command)

1. Skip particle output (demo CLI)

Demo	CLI flag
DamBreak, ObjectDrop, BaffleFlow	--output false
WaveTank	--output_particle_data false
AngleRepose	--output false

Why: Removes periodic device→host copies and CSV serialization from the critical path.

Example:

./bin/demo_FSI-SPH_DamBreak --no_vis --quiet --output false

Effect: Largest win on long fluid demos when combined with ps_freq 2 (typical combined range 10–25%).

---

2. Reduced proximity-search frequency (--ps_freq 2)

Maps to ChFsiFluidSystemSPH::SetNumProximitySearchSteps(2) — rebuild when m_frame % num_proximity_search_steps == 0.

Why: Neighbor search is O(N) (sort + hash grid). Halving rebuild rate reduces that subsystem cost while forces still integrate every step.

Example:

./bin/demo_FSI-SPH_DamBreak --no_vis --quiet --output false --ps_freq 2
./bin/demo_FSI-SPH_BaffleFlow --no_vis --quiet --output false --ps_freq 2

Effect: Major contributor to fluid-demo tuning (typical 5–15% from search cadence alone; stack with I/O off for 10–25%).

Caveat: Do not use on demo_FSI-SPH_WaveTank — WCSPH aborts with stale neighbors.

---

3. Disable SDMA (HSA_ENABLE_SDMA=0)

Why: Avoids SDMA engine scheduling overhead on micro-benchmark-sized transfers; compute path fits rapid kernel sequences.

Example:

export HSA_ENABLE_SDMA=0
./bin/demo_FSI-SPH_Kernels

Effect: Strongest on MI350-class Kernel runs (typical 25–40%); limited on MI300X fluid demos (0–5%).

---

4. Hardware queue depth (GPU_MAX_HW_QUEUES=4)

Why: More in-flight ROCm queues → better latency hiding between dependent SPH kernel chains.

Example:

export GPU_MAX_HW_QUEUES=4

Effect: Modest alone (0–5%); useful stacked in combined profiles (8–20%).

---

5. HIP compile flags (requires rebuild)

Tier	CMAKE_HIP_FLAGS	Why it helps	Typical range
Release	-O3	Better loop scheduling in force kernels	0–5%
Fast math	-O3 -ffast-math	Vectorized SPH accumulations	2–8%
Inline-all	-O3 -mllvm -amdgpu-early-inline-all=true	Fewer kernel boundaries	3–8%
Combined	fast-math + inline-all	Instruction-level and launch-overhead wins	5–12%

Example:

cmake -S . -B build \
  -DCMAKE_HIP_FLAGS="-O3 -ffast-math -mllvm -amdgpu-early-inline-all=true"
cmake --build build --target demo_FSI-SPH_ObjectDrop -j $(nproc)

Tradeoff: -ffast-math relaxes IEEE semantics — validate accuracy.

Why RassorDrum ignores compile tuning: HBM-bandwidth saturated; faster instructions do not raise bytes/s.

---

Per-demo recommended commands

After cd build and setting LD_LIBRARY_PATH:

Demo	Command	Why this recipe
DamBreak	./bin/demo_FSI-SPH_DamBreak --no_vis --quiet --output false --ps_freq 2	Cuts I/O + neighbor-search (largest non-force slices)
Kernel	HSA_ENABLE_SDMA=0 ./bin/demo_FSI-SPH_Kernels	SDMA off for latency-bound micro-kernel
ObjectDrop	Rebuild with -ffast-math, then ./bin/demo_FSI-SPH_ObjectDrop --no_vis --quiet --output false	Faster force-loop codegen
WaveTank	./bin/demo_FSI-SPH_WaveTank --no_vis --quiet --output_particle_data false	I/O only — no ps_freq 2
BaffleFlow	./bin/demo_FSI-SPH_BaffleFlow --no_vis --quiet --output false --ps_freq 2	Same as DamBreak
AngleRepose	GPU_MAX_HW_QUEUES=4 HSA_ENABLE_SDMA=0 ./bin/demo_FSI-SPH_AngleRepose --no_vis --output false --ps_freq 2	Long run: stack runtime + compile + CLI
RassorDrum	./bin/demo_FSI-SPH_RassorDrum --no_vis	HBM-bound; verify wheel BCE particles load

Baseline (untuned) for comparison:

./bin/demo_FSI-SPH_DamBreak --no_vis --quiet
./bin/demo_FSI-SPH_Kernels
./bin/demo_FSI-SPH_AngleRepose --no_vis --output false
./bin/demo_FSI-SPH_RassorDrum --no_vis

---

Unified profile (one binary, all seven demos)

When one rebuild must serve all demos:

export GPU_MAX_HW_QUEUES=4 HSA_ENABLE_SDMA=0 HSA_XNACK=0 HIP_LAUNCH_BLOCKING=0
# Rebuild with: -O3 -ffast-math -mllvm -amdgpu-early-inline-all=true

Then per demo:

• Fluid demos (except WaveTank): --output false --ps_freq 2
• WaveTank: --output_particle_data false only (no ps_freq 2)
• Kernel: keep HSA_ENABLE_SDMA=0, no extra CLI

This is a cross-demo compromise: 5–15% on I/O/search-bound fluid cases; 0–2% on HBM-bound RassorDrum; Kernel 0–40% depending on GPU.

---

Example env snippets

Source before running (see docs/fsi_sph_tuning/):

source docs/fsi_sph_tuning/runtime_sdma_off.env.example      # Kernel / SDMA
source docs/fsi_sph_tuning/cli_io_psfreq.env.example         # DamBreak / BaffleFlow
source docs/fsi_sph_tuning/cli_io_only.env.example           # WaveTank
source docs/fsi_sph_tuning/runtime_combined.env.example        # AngleRepose
source docs/fsi_sph_tuning/compile_fastmath_inline.env.example # rebuild flags

---

Important caveats

• ps_freq=2 changes neighbor-search cadence — validate physics against your reference before production use.
• demo_FSI-SPH_WaveTank aborts with ps_freq=2 in tested WCSPH settings.
• WaveTank uses --output_particle_data false, not --output false.
• -ffast-math relaxes IEEE semantics.
• Separate build trees per CMAKE_HIP_ARCHITECTURES (gfx942 vs gfx950) — do not mix libraries and binaries across arches.
• RassorDrum: run from build dir; wrong cwd → zero BCE particles and invalid runs.

---

Test plan

• Doxygen build includes manual_fsi_sph_amd_instinct_tuning under FSI manual
• Default CMake (no HIP) unaffected — docs only
• demo_FSI-SPH_DamBreak --output false --ps_freq 2 completes on Instinct
• demo_FSI-SPH_WaveTank --output_particle_data false completes without ps_freq 2
• HSA_ENABLE_SDMA=0 ./bin/demo_FSI-SPH_Kernels on MI350-class GPU
• RassorDrum log shows wheel BCE particles loaded when run from build/

---

Non-goals

• No changes to SPH numerics, demo defaults, or solver source in this PR
• No batch scheduler / cluster-specific scripts

---

See also (in-repo after merge)

• Doxygen: manual_fsi_sph_amd_instinct_tuning — full manual with code references
• manual_fsi_sph_parameter_selection — SPH parameter semantics (ps_freq, viscosity, time step)
• docs/FSI_SPH_AMD_TUNING.md — Markdown quick reference

[amd-pratmish]

@DanNegrut @harryzhang1018