| name | ai4science-perf-analysis |
| description | Runs the AMD AI agents bottleneck-analysis workflow on a HydraGNN or ORBIT-2 training job using TraceLens + Omnistat user-mode + paired analyst/verifier subagents, or drives the iterative perf-optimizer-loop (accept/revert on throughput or epoch time). Use when the user wants to analyze a perf run, analyze a 2-node HydraGNN run, find bottlenecks with TraceLens/Omnistat, run the AMD AI agents on a job, run the perf-optimizer-loop, tune ORBIT-2 throughput, or optimize HydraGNN training performance. |
AI4Science perf-analysis (multi-subagent bottleneck workflow)
When this skill applies
The user wants an automated bottleneck analysis of a multi-node training/inference run using AMD's open-source observability tooling (TraceLens, Omnistat). This is distinct from the ai4science-run-models skill — that one launches a model; this one diagnoses a model after it runs.
Default target: HydraGNN on AMD MI355X. ORBIT-2 training uses the same perf-analysis pattern; see earth_science/models/ORBIT-2/recipes/perf-analysis/. ORBIT-2 iterative sysopt (throughput-primary FOM) lives in earth_science/models/ORBIT-2/recipes/perf-optimizer-loop/ (run_optimizer_loop.sh, lever_catalog.yaml).
Repository entry points
HydraGNN
ORBIT-2
Orchestration loop (what the main agent does)
- Read the recipe README to refresh on the run topology and artifact layout.
- Dispatch the launcher subagent (Task tool,
generalPurpose or shell) with the prompt at agents/launcher.md. Wait for the job to complete and the manifest to be written.
- Dispatch the two analyst subagents in parallel (one Task call per subagent in the same message):
tracelens_analyst — prompt at agents/tracelens_analyst.md
omnistat_analyst — prompt at agents/omnistat_analyst.md
- Dispatch the two verifier subagents in parallel once analysts return:
tracelens_verifier
omnistat_verifier
- Dispatch the synthesizer subagent, hand it both
verified_claims.json files.
- Print the combined report to the user with a short executive summary in chat.
Subagent contract (every subagent must)
- Read only the files listed in its
## Inputs section.
- Produce only the file(s) listed in its
## Outputs section.
- Write a single line
STATUS=ok|partial|fail; reason=<short> to stdout as the last line.
- Never escalate to multi-node
srun; verifiers may use at most one 1-node srun -N1 --time=00:05:00 interactive probe (partition/account from .cluster-config.yaml).
- Never edit files outside
$AI4S_SHARED_DIR/models/HydraGNN/perf-runs/<jobid>/.
Cluster constraints (AMD Instinct MI355X)
- Partition / account: read from .cluster-config.yaml.
- Central Prometheus may be unreachable from compute nodes (check your cluster's network policy). Always use the user-mode VictoriaMetrics that the launcher started.
- System-mode Omnistat may be running on every compute node at port 8001. Don't touch it; we run our own user-mode collector alongside.
- The login node typically has no GPU and no MPI. All analysis runs after the job — no live profiling on the login node.
- Node-specific mount faults are real and silent. Any single compute node can land an allocation with a broken per-user autofs/NFS mount of
/home/$USER or $AI4S_SHARED_DIR, while sibling nodes in the same allocation are fine. Symptom in slurmd logs on the broken node: Home Directory for <user> Not Found ... Please contact system admin and lstat ...: no such file or directory; the surviving nodes wedge in a pmix_coll_ring collective fence until SLURM kills the job at the wall.
- Confirmed bad node patterns we hit (MI355X, May 2026): intermittent dual-NUMA STREAM degradation on one node, broken container bind-mounts (
$HOME/$AI4S_SHARED_DIR) on another, PMIx ring timeout on a third. Node names are cluster-specific; see the untracked microbench outputs for hostnames.
- Diagnostic rule: never blacklist a node class based on a single failure. Always probe the specific allocated nodes and exclude only those that fail by name. Probe is a one-task-per-node srun that checks
/home/$USER, $AI4S_SHARED_DIR, and the SIF path; see the Per-node mount-health probe section of sbatch_train_perf_amd.sh (exit 42 on fail). The perf-optimizer-loop orchestrator handles exit 42 by re-submitting with --exclude=<bad-node> and adding the node to loop-<uuid>/known_bad_nodes.txt; the lever is NOT penalized.
When something goes wrong
- If the launcher times out waiting for
sacct to report a terminal state, it must scontrol show job <id> and write state=TIMEOUT_WAITING in the manifest; the orchestrator should surface this to the user and stop.
- If TraceLens or omnistat-inspect can't be installed (network, py-version), the launcher's install step writes a clear error to stderr and exits non-zero — the orchestrator must NOT fall back to "skip analysis", it must surface the error.
- If a verifier refutes an analyst's top claim, it stays in the report (with
verdict=refuted) so future iterations don't re-derive the same wrong conclusion.
Lessons captured (smoke-test on JOB 6762)
The first 2-node end-to-end run exposed five real bugs that are now all worked around in sbatch_train_perf_amd.sh and the agent prompts:
%(SLURM_JOB_ID)s in omnistat config is broken — configparser doesn't interpolate os.environ. Use @JOB_DIR@ placeholder and sed at submit time.
/tmp/omni_rmsjobinfo permission collision with system-mode Omnistat — override job_detection_file to a per-job path under $AI4S_SHARED_DIR/....
- VictoriaMetrics on login node needs
-fs.disableMmap to load even a 1.6 MB DB (cgroup mmap restriction).
- HydraGNN
Profile block must be under NeuralNetwork, not at the top level — train_validate_test() is called with config["NeuralNetwork"] so that's the scope its Profiler reads.
- PromQL with
jobid="..." requires joining via rmsjob_info — rocm_* metrics don't carry a jobid label directly. Verifier subagents must use:
metric * on (instance) group_left() (max by (instance) (rmsjob_info{jobid="..."}))
Lessons captured (perf-optimizer-loop pilot, loop-43b33ec1, 2026-05-22)
- Diagnose broken nodes individually, not by class (see Cluster constraints above for full detail). The optimizer-loop's iter-1 wrongly concluded "a-class nodes lack /home" from a single failed alloc — only one specific node was broken. The fix is a per-job mount-health probe (now in
sbatch_train_perf_amd.sh, exit code 42) plus orchestrator step 2e-bis: on exit 42, parse node_health_probe.txt, append the bad node(s) to loop-<uuid>/known_bad_nodes.txt, and re-submit with --exclude=<list>. Lever is untouched.
- HydraGNN
torch.compile hooks must wrap the model object, not a module path. import hydragnn.train.train (the path the iter-4 hook tried) does not exist in the installed package. Wrap the constructed model directly: model = torch.compile(model, backend="inductor", mode="reduce-overhead", fullgraph=False) at the rank-script level, by patching the entrypoint script (gfm_mlip_all_mpnn.py) immediately after model construction. (Now superseded by Lesson #6 below — torch.compile is BLOCKED for HydraGNN MLIP regardless of where it wraps.)
parse_convergence.py over-counts epochs by N_ranks because each rank emits a tqdm=100% line and the parser counts every one. Always read epoch wall-time from rank-0 tqdm directly (s/it × max_num_batch), not from the parser's epoch count. Fixed in fom_extractor; long-term fix is to filter on rank=0 lines in the parser.
- MFMA TFLOPS methodology must be declared explicitly. Burst-rate via PromQL
rate([10s])×4 (iter-0/3 values ≈ 0.006-0.012 TFLOPS) and time-averaged total_accumulated_ops / duration (iter-4 value ≈ 0.75 TFLOPS) differ by ~60× for this workload and must not be compared. fom_extractor should standardize on time-averaged across the full job window, with the burst-rate as a separate *_peak_burst field.
epoch_time_s is the correct primary FOM for HydraGNN-style latency-bound GNN workloads. Throughput (samples/s) misleads when batch size is a lever, because epoch wall time scales with batches-per-epoch, not just per-sample work.
Lessons captured (loop-c3e4df1c + compile-investigation-aa480554, 2026-05-22..23)
torch.compile is BLOCKED for HydraGNN GFM-MLIP MACE under PyTorch 2.10 / ROCm 7.2.2. Six hooks across three hypotheses (inner-model compile, outer + dynamo.allow_in_graph(autograd.grad), single-submodule compile) all fail with the same verbatim PyTorch error: RuntimeError: torch.compile with aot_autograd does not currently support double backward (torch/_functorch/_aot_autograd/runtime_wrappers.py:2356). Root cause: HydraGNN's energy_force_loss (hydragnn/models/create.py:718) calls torch.autograd.grad(graph_energy_pred, data.pos, create_graph=True) to compute forces as derivatives of energy w.r.t. positions; with create_graph=True, training requires DOUBLE BACKWARD (loss.backward differentiates through forces). Every torch.compile mode (default, reduce-overhead, max-autotune) goes through AOT Autograd which pre-compiles the backward as a joint program and cannot itself be differentiated. To unblock requires upstream HydraGNN source change: move energy_force_loss INTO EnhancedModelWrapper.forward() so autograd.grad is inside the compiled forward (the canonical MACE/Allegro/NequIP pattern). ~30-50 LOC across create.py:622-758 + train_validate_test.py:725-735, config-gated. The lever torch_compile_e3nn is marked status: blocked in lever_catalog.yaml with full evidence; the lever_picker drops blocked levers from its candidate set.
dynamo.allow_in_graph(autograd.grad) is weaker than MACE upstream docs suggest. It only prevents dynamo from TRACING THROUGH the call; it does NOT prevent dynamo's FakeTensor symbolic-execution pass from running autograd.grad against fake tensors (which fails with the allow_unused=True error because fake tensors have no real autograd graph). The MACE prepare() pattern works only when autograd.grad lives inside the compiled function's forward() — not when it's in a separate method called from outside.
- TunableOp live tuning (
PYTORCH_TUNABLEOP_TUNING=1) is unsafe on MI355X / ROCm 7.2.2 with HydraGNN. All 16 GPUs hit Memory access fault by GPU node-N during the hipBLASLt kernel autotuning phase (loop-c3e4df1c iter-3 / job 7188). Use a 2-phase pattern instead: dedicated warmup sbatch with PYTORCH_TUNABLEOP_TUNING=1 to generate tunableop_results_<N>.csv files, then production runs with PYTORCH_TUNABLEOP_ENABLED=1 only (no _TUNING=1). The tunable_op_warmup_then_use lever in the catalog encodes this. The original tunable_op_live lever is now marked status: blocked.
- Lever payoff can be run-to-run noise, not node-class signal.
num_workers_12 improved one run (-12.4%) and regressed another (+10.7%), both on the same node pair on different days. The lever picker should treat num_workers_* results as high-variance until replicated within ±5% on the same node pair.
- Always read
nodes_list from iter-0-baseline.json before drawing any class conclusion from a FOM delta. Two baselines that appear to "use different node classes" can in fact be the same physical nodes scheduled by SLURM on different days — especially when one iter-1 failed and the successful iter-0 landed wherever SLURM placed it.
hydragnn/train/__init__.py re-export shadow breaks naive import hydragnn.train.train_validate_test as tvt. The line from .train_validate_test import train_validate_test, train, ... makes hydragnn.train.train_validate_test resolve to the function train_validate_test (re-exported as attribute of the package), not the module. Use importlib.import_module("hydragnn.train.train_validate_test") instead. Applies to any rank hook that needs to monkey-patch train_validate_test.train.
sbatch --export=ALL,... does NOT include shell vars that are not in --export's explicit list. When using --export=ALL,K1=V1,K2=V2, vars set in the calling shell but not listed here are propagated only if they're in the user's persistent env (login shell init). Always explicitly list cluster-pathing vars like AI4S_SHARED_DIR in --export.
Lessons captured (pathc-pathb-a3b8f3f0 + loop-e3fac2af, 2026-05-23)
- fp32-vs-fp64 epoch-time comparison is the canonical low-cost dispatch-vs-compute discriminator. On gfx950 (MI355X), the fp32 MFMA peak is 2× the fp64 MFMA peak. A compute-bound workload should show ≥30% speedup when switching fp64→fp32. HydraGNN GFM-MLIP MACE showed ~0% speedup (78→79 s, +1.3% within noise, fp32 verified by
precision-diagnostic rank-0 logs model_param_dtype=torch.float32 first_batch_float_dtype=torch.float32). This proves the workload is dispatch-bound. Run this comparison ONCE per workload-class as a single sbatch; the verdict shapes the entire subsequent lever strategy:
- fp32 ≈ fp64 → dispatch-bound → compute-side levers (
batch_*, torch_compile_*, tunable_op_*, precision_fp32) are dead ends. Focus on kernel-trace analysis to identify the dispatch culprit, then either upstream kernel fusion or reduce launch count (larger blocks, mega-kernels).
- fp32 ≪ fp64 → compute-bound → standard compute-side levers apply. fp32 itself becomes a real lever (within numerical tolerance).
- fp32 > fp64 → memory-bound → bandwidth-targeting levers (batch packing, fewer kernel launches over more data) apply.
torch.jit.script is ALSO blocked for HydraGNN MLIP MACE (in addition to torch.compile). Root cause: MACEStack.py:397 uses **conv_args keyword-arg expansion when calling the interaction block; TorchScript JIT does not support **kwargs expansion (NotSupportedError: keyword-arg expansion is not supported). Despite the upstream @compile_mode("script") decorator from e3nn, the actual call site is not scriptable. Both PyTorch compiler paths are now definitively dead for HydraGNN MLIP MACE as shipped — unblocking either requires upstream source changes:
torch.compile path: move energy_force_loss into EnhancedModelWrapper.forward() (~30-50 LOC; addresses AOT-double-backward; see lesson #6)
torch.jit.script path: unpack **conv_args at MACEStack.py:397 (small diff but in the read-only /opt/hydragnn-pkgs overlay)
- TunableOp
_TUNING=1 is unsafe regardless of phase on MI355X / ROCm 7.2.2 / PyTorch 2.10. Both the live-tuning lever (tunable_op_live, loop-c3e4df1c iter-3) and the warmup-then-use 2-phase pattern (tunable_op_warmup_then_use, loop-e3fac2af iter-2) hit the same Memory access fault by GPU node-N on all 16 GPUs during the hipBLASLt tuning routine. The 2-phase split doesn't help — the fault is in the tuning routine itself, not in the production consumption. Both levers are marked status: blocked in the catalog. Re-test only after a ROCm stack update with a documented hipBLASLt fix.
- The "dispatch ceiling" conclusion is solid, but do NOT frame it as node-class-specific. All three loop baselines ran on the same physical node pair; every compute-side lever tested (batch_800, num_workers_12, precision_fp32) regressed or was neutral. That confirms dispatch-bound character for this workload — but makes no statement about node classes. The microbench (lesson #18) later showed all tested nodes have identical HIP launch latency, so the ceiling is workload-driven, not hardware-class-driven.
- Cross-loop best is composed of a free finding + a small lever: current best (75.0 s/epoch) = baseline-on-b1-b2 (76.0 s) +
rccl_high_priority (-1.3%). The single accepted lever across 3 loops + 1 investigation is rccl_high_priority (TORCH_NCCL_HIGH_PRIORITY=1 GPU_MAX_HW_QUEUES=2); it should be baked into the baseline contract for HydraGNN GFM-MLIP going forward (now marked status: accepted-baked-in in the catalog).
Lessons captured (a-vs-b parity microbench, 2026-05-26)
- A "node-class performance gap" narrative was debunked by a controlled microbench. A 4-test microbench sweep (script:
examples/microbench_node_health.sh; raw outputs in untracked $AI4S_SHARED_DIR/microbench-a-vs-b/) found:
- All tested nodes had identical HW at the host level (same CPU model, NUMA layout, GPUs, NICs, firmware).
- HIP empty-kernel launch latency 3.5–3.8 µs/launch on every GPU of every node; NCCL single-node all-reduce and STREAM single-NUMA TRIAD were statistically equivalent across all nodes tested (one outlier node had degraded dual-NUMA STREAM COPY, flagged for sysadmin).
- A controlled 2-node sanity HydraGNN run using nodes from each "class" showed < 2% wall-time difference, well inside noise.
- Root cause of the original gap narrative: all three baseline loop runs landed on the same physical node pair. The 92.5 s → 76 s spread was run-to-run variance on the same pair on different days. The misattribution arose from labeling one loop as "class A" because its iter-1 failed on a broken-mount node — but the successful baseline iter-0 landed on a different pair via SLURM scheduling.
- Action for future loops: always
jq .nodes_list $LOOP_DIR/iter-0-baseline.json before drawing any class conclusion. The launcher and synthesizer subagents should log the assigned nodelist front-and-center in their output.
- Per-node health regressions are real and worth a 30-second prologue probe. In 2 weeks we hit three different per-node failure modes: container bind-mount EIO on one node, PMIx/NCCL ring timeout on another, dual-NUMA bandwidth degradation on a third. Each manifested as a confusing application-level failure hours into a job. All three would have been caught at job start by a 30-s SLURM Prolog running
microbench_node_health.sh (mount-write probe + STREAM dual-NUMA + NCCL single-node + HIP launch latency). Recommend adding this as a SLURM Prolog to your cluster admin.
Lessons captured (attribution pass on job 7187, 2026-05-27)
- Attribution pass tooling: Use
examples/run_fom_extractor.py + TraceLens on $AI4S_SHARED_DIR/models/HydraGNN/perf-runs/<jobid>/. Backfill manifest.json if missing (loop jobs before launcher wrote it). VictoriaMetrics PromQL on perf-run DBs requires time=<unix> at the job window — instant queries on the login node return empty series even when omnistat_hardware_counter data exists. kernel_correlation_summary.attribution_quality=poor on a ~6 s profile-epoch kineto slice is often a windowing artifact (busy_frac <0.5 in every 1 s bucket); trust TraceLens ops_summary.csv (aten::mm + aten::bmm ≈80% kernel time) before escalating to OMNISTAT_KERNEL_TRACE=1. See recipes/perf-optimizer-loop/dispatch-attribution.md.
Lessons captured (ORBIT-2 bf16 enablement, 2026-06-11)
- ROCm Flash SDPA has a 65535 batch-grid cap — fatal for attention that inflates the batch dim. ORBIT-2 Bayes-CAST EDM crashed on the first
training_step in bf16 with HIP error: invalid argument (hipErrorInvalidValue), apparently at var_agg's proj Linear (attention.py:280). The Linear was a red herring (async error misattribution): a standalone bf16 Linear at that shape passes even at M=663552. The real culprit is var_agg/temporal_agg cross-attention, which flattens (B, History, L) into the SDPA batch dim, so B = batch·tokens (e.g. batch·648). PyTorch SDPA's default backend picks ROCm Flash, which maps that batch dim to a HIP grid dimension capped at 65535 → fails once batch·648 > 65535 (i.e. batch ≥ ~128; batch 64→PASS, 128→FAIL). fp32 "worked" only because it dispatched to the math backend.
- Isolate fast with a faithful op-replica on 1 GPU (no data, ~15 s):
crossattn mode reproduces; math/efficient/eager backends all PASS, default FAILs at batch 128. This pattern (faithful op replica on one GPU) beats full-node data-loading jobs for kernel bugs.
- Fix: wrap the offending module's SDPA in
with torch.nn.attention.sdpa_kernel([SDPBackend.EFFICIENT_ATTENTION, SDPBackend.MATH]): to exclude Flash (MATH fallback-only). Validated end-to-end: ORBIT-2 job 9238 (bf16 b256) COMPLETES with decreasing loss. Prefer applying it globally (all self- + cross-attention SDPA): the --orbit-selfattn benchmark shows EFFICIENT within ±3–5% of Flash for normal self-attention (q=kv=648) across batch 64/256/1024, so unifying on EFFICIENT removes the foot-gun at ~noise-level cost. Do NOT make MATH primary — it's 10–15× slower (≈15 vs ≈290 TFLOPS); it just never gets selected when EFFICIENT supports the shape. Use --orbit-varagg to confirm the Flash cap and --orbit-selfattn to confirm the perf parity before deciding scope.
- Generalize: any model that folds extra dims into the attention batch (per-variable, per-timestep, per-patch cross-attention) can hit this in bf16 at large batch on ROCm. Suspect Flash's grid cap before blaming GEMM/BLAS;
HIP_LAUNCH_BLOCKING=1 mis-points to the next CPU-sync op, so trust an op-replica probe over the traceback line.
- FOM parser (RESOLVED 2026-06-11): this EDM build logs per-step loss as
epoch: N batch_idx M ... loss tensor(L) and per-step wall time as my rank 0. tic4-tic1 in X seconds — it has no Batch N: … seconds / Epoch N completed line. parse_training_log.py now reads both this EDM format and the older intermediate_downscaling format (pairs each tic4-tic1 to the preceding step; synthesizes per-epoch loss from per-step losses). Validated on jobs 9234 (1236 s/s) & 9257 (4097 s/s). Crashed jobs (no batch_idx lines) correctly yield epoch_records=0 + null throughput — that's the "never reached steady state" signal, not a parser bug. Do not rewrite the parser.
Lessons captured (ORBIT-2 bf16 optimizer loop, loop-bf16-loop-001, 2026-06-12)
- ORBIT-2 EDM (1×8 MI355X) is COMPUTE-bound, not dispatch-/dataloader-bound — the opposite of HydraGNN MLIP. The fp32-vs-bf16 discriminator (job 9264 vs 9261, same batch 1024) showed bf16 = 2.38× fp32 (4089 vs 1718 samples/s). A ≥1.4× precision speedup ⇒ MFMA/GEMM-bound. Consequences confirmed in-loop:
torch_compile_edm was neutral (-0.3%, within noise — Inductor finds little to fuse on an already-GEMM-heavy ViT/EDM), and num_workers_8 regressed -5.6% because 8 workers × 8 ranks = 64 dataloader procs oversubscribe the 56 schedulable CPUs (cpus-per-task=7). For compute-bound ORBIT-2 the only real throughput levers are bigger effective batch (needs more staged data) or GEMM tuning (TunableOp — later shown NOT blocked, see #28). Baseline (bf16 + EFFICIENT SDPA + batch 1024) remained the best config; no catalog lever beat it at 1 node. Keep num_workers ≤ 4 for ORBIT-2 to avoid CPU oversubscription.
num_workers silently changes the EFFECTIVE per-step batch in the Bayes-CAST IterableDataset → phantom throughput. Loop overnight-3x-001 iter-1 set num_workers=4 and run_fom_extractor.py reported +129% (7215 vs 3150 s/s) — a pure measurement artifact. Throughput = nominal global_batch_size ÷ steady step time, but with more workers the real per-step batch shrank ~2.4× (HBM 236→98 GB, max_batches_per_epoch 3→4, epoch loss 2.86→3.50). Guard: foms.json now carries hbm_reserved_GB, hbm_reserved_pct_288, max_batches_per_epoch; reject any cross-run throughput delta where HBM% or batches/epoch deviates from baseline. The Opus-4.8 overnight orchestrator caught this unaided via the HBM+batch cross-check — an unguarded loop would have crowned a phantom 2.3× best and poisoned every downstream iteration. For dataloader levers, compare only runs with matching HBM%/batches-per-epoch.
- 3× ERA5 staging breaks the 1704-sample cap and enables true HBM saturation.
stage_era5_3x_symlink.sh symlinks the one staged year (1979, 20 shards) as 1980_/1981_ (60 shards; the loader globs train/*.npz with no year filter — itermodule.py:133). On 3× data, bf16 batch 4096 reaches 82% HBM (236/288 GB) at 3152 s/s, and HBM scales ~linearly with batch (1024→20.6%, 2048→40.5%, 4096→82.0%). Perf/throughput study only — data repeats, not valid for science. The overnight loop on this config (4 iters) found no env lever beats baseline (num_workers artifact rejected; fsdp_prefetch +0.99%, sdpa_efficient +1.03% — both noise-floor), reconfirming compute-bound. Real wins need multi-node scale-out, real multi-year data, or GEMM tuning (TunableOp — now experimental, not blocked; see #28).
- Max-VRAM ≠ max-throughput at a small data cap. At the 1704-sample staging cap, batch 1024 (~21% HBM) gave the best throughput (4098 s/s) while batch 1704 (~34% HBM) was 27% slower (2975 s/s) — batch > sample-count just repeats/partial-fills a step. Lock the optimizer baseline at the throughput-optimal batch, and only chase HBM% after staging more ERA5 (raises the cap so larger batches do real work). See BASELINE_LOCKIN.md.
- Multi-node ORBIT-2 on MI355X needs GID-consistent nodes + a non-restrictive
LD_LIBRARY_PATH (RoCE/RCCL). First 2-node ORBIT-2 attempts (jobs 10583/10584) aborted with ionic_comp_msn: cqe with error 12 / NET/IB ... status=12 vendor err 11 RETRY_EXC. Two distinct causes, both fixed in sbatch_train_perf_amd.sh: (a) it pinned --env LD_LIBRARY_PATH=…/torch/lib only, so the host ANP plugin (/opt/rocm/lib/librccl-anp.so) couldn't resolve its /opt/rocm/lib deps — widened to include /opt/rocm/lib:/usr/lib/x86_64-linux-gnu (HydraGNN never set this var, which is why it worked). (b) RoCE GID-index asymmetry: NCCL_IB_GID_INDEX=1 assumes the fabric ULA (fd93:16d3:59b6:014f/RoCEv2) sits at GID idx1 on every node, but a subset of nodes carry a stray global IPv6 (2001:19f0:…) at idx1**, shifting the fabric ULA to idx2 → localGid(ULA)/remoteGid(global) mismatch → no route → retry-exceeded abort. The other nodes were uniform (idx1=ULA); pinning to GID-consistent nodes (--exclude the bad range) makes 2-node work with 0 cqe errors (validated job 10587 on a4-a5, then loop jobs 10588+). Enumerate GIDs with cat /sys/class/infiniband/ionic_3/ports/1/gids/{1,2} and …/gid_attrs/types/*. NCCL_IB_GID_INDEX is now overridable in the sbatch. Proper long-term fix is admin-side (remove stray global IPv6 from b-nodes' RoCE NICs, or subnet-based GID selection). Driver run_2node_scaleout_loop.sh can auto-exclude a known-bad range via O2_EXCLUDE_PATTERN (default: none). Always verify GID-table consistency before blaming the model for a multi-node hang. 2-node result (loop-2node-scaleout-001): baseline 6046 s/s vs 1-node 2987 → 2.02× (101% strong-scaling efficiency) — FSDP all-gather over RoCE is fully hidden behind compute (consistent with compute-bound). All comm/compute levers neutral (IB QPS 2/4, NCCL_MIN_NCHANNELS=112, torch.compile: ±1.5%); none beat baseline. GPU_MAX_HW_QUEUES=4 aborts at init (rc=134, "terminate called recursively", before any step) on MI355X — keep the validated =2.
- TunableOp is NOT platform-blocked on MI355X / ROCm 7.2.2 / PyTorch 2.10 — the HydraGNN "blocked" verdict (#8/#15) does NOT generalize. A dedicated isolated probe (single GPU,
perf-runs/tunableop-probe/, jobs 10595/10596) ran PYTORCH_TUNABLEOP_ENABLED=1 PYTORCH_TUNABLEOP_TUNING=1 over 17 GEMMs — including the ORBIT-2 large-M bf16 var_agg shape class (mm NN and F.linear TN-with-bias: M up to 262144, e.g. tn_1024_131072_1024) that the sbatch comment claims throws hipErrorInvalidValue under hipBLASLt. All 14/14 matmuls tuned cleanly, 0 memory faults, a valid tunableop_results.csv was written (hipBLASLt was actually selected for every large-M bias GEMM; rocBLAS won a few NN cases). Prereqs are all present: gfx950 Tensile logic libs exist in /opt/rocm/lib/hipblaslt/library (439 gfx950 files) and /opt/rocm/lib/rocblas/library (156) — nothing is missing from the container. Validators: HIPBLASLT_VERSION=100202-dabb6df2b9, ROCBLAS_VERSION=5.2.0.dabb6df2b9, GCN_ARCH_NAME=gfx950:sramecc+:xnack-. Interpretation: the HydraGNN "Memory access fault on all 16 GPUs during tuning" (#8/#15) is workload/concurrency-specific (full distributed run, MACE double-backward shapes), NOT a hipBLASLt-on-gfx950 prerequisite gap. Action: ORBIT-2's tunable_op_warmup_then_use is now status: experimental (not blocked); the remaining unknown is a full FSDP/all-rank warmup run (the HydraGNN fault only manifested under 16-GPU distributed tuning). Probe API note: torch.cuda.tunable.write_file() does not exist in 2.10 — the cache auto-writes on exit; use set_filename() + get_results(). Lesson: do not copy a blocked verdict across workloads/models without re-probing — the cheap isolated GEMM probe (~2 min, 1 GPU) settles it.
- TunableOp on ORBIT-2's 8-rank FSDP path does NOT fault (HydraGNN block disproved again), BUT live tuning is impractically slow for its giant patch-token GEMMs. Sanity jobs 10597/10598 ran the full 8×MI355X FSDP EDM with
ORBIT2_TUNABLEOP_MODE=tune (PYTORCH_TUNABLEOP_TUNING=1): 0 memory faults, per-rank caches (tunableop_results_rank<procid><dev>.csv) written cleanly — the distributed path HydraGNN crashed on (#8/#15) is fine for ORBIT-2. However, the dominant GEMM is tn_256_5308416_256 (M=5.3M flattened patch tokens; the backward is a K=5.3M reduction), and TunableOp spent ~3.5 min on the first giant op and >12 min on the second before it was cancelled — a full tune would be many hours. Root cause: per-candidate kernel instantiation (hipBLASLt+rocBLAS enumerate hundreds of algos; each is JIT-instantiated on these unusual huge shapes). This is NOT fixable by env knobs: PYTORCH_TUNABLEOP_ROTATING_BUFFER_SIZE=0 (verified to deactivate rotation per the API docstring — "equal to zero means deactivate"), MAX_TUNING_ITERATIONS, and MAX_TUNING_DURATION_MS only bound per-solution timing, not the candidate count. Practical guidance: for giant-GEMM models, either (a) time-box the tune and accept a partial cache (the 2-3 giant ops dominate runtime, so even a partial cache captures most of any uplift), or (b) skip TunableOp and spend the node-hours on 2-node scale-out (known 2.02×). Wiring added to sbatch_train_perf_amd.sh: ORBIT2_TUNABLEOP_MODE=off|tune|use + ORBIT2_TUNABLEOP_DIR (rw-bound shared cache; per-rank files keyed on SLURM_PROCID so 8 ranks don't clobber one results0.csv).
- END-TO-END VERDICT: TunableOp does NOT help ORBIT-2 and is a correctness hazard — do not adopt. Live tuning of the giant patch-token GEMMs is impractically slow (#29), so it was driven offline (per-shape GPU-pinned subprocesses; all 102 GEMM shapes incl. M=31.8M/15.8M/5.3M giants tuned). Result: 1-node tuned DIVERGES TO NaN (reproduced 2×: epoch-0 loss 5.5 vs baseline 0.87, then NaN); 2-node tuned = +0.44% (noise). Two lessons: (a) TunableOp has no numerical-correctness gate, so a fast-but-wrong bf16 kernel can win and get cached — and a tuned cache is layout-specific (bit 1-node
fsdp=8 shapes, not 2-node fsdp=2/ddp=8), so it must be re-validated per parallelism config. (b) Even with all giants tuned there is zero perf upside — the default rocBLAS/hipBLASLt path is already near-optimal, confirming ORBIT-2 is bottlenecked by the GEMMs themselves, not kernel selection. Offline-tuning traps if ever revisited: pin workers with only ROCR_VISIBLE_DEVICES (setting HIP_VISIBLE_DEVICES too double-filters → hipErrorNoDevice); key cache-completeness on the core (op, transAB_m_n_k) (TunableOp rewrites the _ld_* suffix). Real throughput levers remain: multi-node scale-out (linear) and bigger effective batch — not TunableOp. If revisited, add a per-kernel correctness gate before caching.
- Perf/throughput runs must NOT persist model checkpoints — the Bayes-CAST EDM trainer saves one every epoch by default.
train_edm.py:save_checkpoint() is called unconditionally at the end of every epoch (line ~1390) and writes a ~71 MB interm_epoch_<N>.ckpt to the hardcoded cp_save_path = "checkpoints/bayes-cast" (relative to launch/, i.e. $ORBIT2_ROOT/launch/checkpoints/bayes-cast/). The config's trainer.checkpoint: field is empty but that only controls loading/resume, not saving — saving is unconditional. Files are overwritten per run (fixed names) so they don't grow per-job, but a 6-epoch run still leaves ~340 MB of unwanted state. Fix: added an env gate ORBIT2_DISABLE_CKPT=1 → save_checkpoint() early-returns on ALL ranks together (no unbalanced FSDP state_dict() gather / barrier). sbatch_train_perf_amd.sh now defaults it to 1 (perf script → never save); sbatch_train_amd.sh defaults it to 0 (general training keeps checkpoints) and run_scaling_study.sh passes =1. All perf submitters (sweep_*, submit_perf_baseline_*, the optimizer/2-node/tunableop loops) route through sbatch_train_perf_amd.sh so they inherit the skip. The initial_*.pth save path (train_edm.py:188) only fires when tensor_par_size>1 (we use 1) so it is not a concern. Lesson: when standing up a throughput harness around an upstream trainer, audit its default checkpoint/save behavior — it will silently write state every epoch unless gated.
num_workers is NOT a throughput lever for ORBIT-2 — confirmed by a clean controlled test, and the FOM extractor is now hardened against the artifact. The Bayes-CAST IterableDataset shards files by per_worker = n_files // (num_workers × data_par_size) (data_par_size = fsdp × simple_ddp = 8), and each worker batches its own shard independently — so more workers means smaller/fragmented batches. Staging 5× ERA5 (100 shards) so per_worker = 3 lets full batches refill to 4096 at num_workers=4, but each worker still emits a partial 1016-sample trailing batch (4096 + 1016 = 3 files × ~1704). Jobs 10504 (nw=1) vs 10505 (nw=4), batch 4096, identical otherwise: nominal throughput said +80% (2993→5379 s/s); real-per-step throughput showed +2.7% (2987→3068 s/s) = noise, and the full-4096 step time was ~equal (11.0 vs 10.8 s). Fix shipped: parse_training_log.py now reads the real per-step batch from the EDM y.shape line and aggregate_foms computes throughput_method=real_per_step_batch (Σ real_samples × ranks ÷ Σ step_time), plus partial_step_fraction and steady_realized_batch_dims as integrity surfaces. Still reject cross-run deltas where steady_realized_batch_dims/partial_step_fraction differ from baseline. Reconfirms ORBIT-2 is compute-bound; keep num_workers ≤ 4 for I/O overlap only.
The full list (with cross-cutting context) is at the end of .cursor/skills/ai4science-studio/SKILL.md. When fixing a bug here, check there first and propagate the lesson.