fix(rccl): HIPFIRE_RCCL_LIB points the loader at a librccl outside the ROCm root (nixpkgs ships it in its own store path) by alpineQ · Pull Request #728 · warpfront/hipfire

alpineQ · 2026-09-05T11:44:40Z

Summary

On NixOS every multi-GPU EP load dies with failed to dlopen librccl.so, because nixpkgs ships RCCL in its own store path while rocmtoolkit-merged (the ROCm root hipfire resolves) carries HIP and HSA only; a new HIPFIRE_RCCL_LIB lets the operator point the loader at that path, and nothing changes where RCCL lives under the root.

What was wrong

RcclComms::init_all (crates/hip-bridge/src/rccl.rs) asks hipfire_config::rocm::library_candidates for librccl.so*. With a selected ROCm root that function returns paths inside the root only, by design (rocm.rs: "An explicit environment root is authoritative. Its split-tree children are considered, but unrelated installs and bare loader sonames are not"). The comment in rccl.rs still promised "bare sonames last", which the code stopped doing when that policy landed.

On nixpkgs the two never meet:

store path	libamdhip64	libhsa-runtime64	librccl
`…-rocmtoolkit-merged` (selected root, via `HIPFIRE_ROCM_PATH` or `hipcc` on `PATH`)	yes	yes	no
`…-rccl-7.2.3`			yes

So the loader tries six paths under the merged prefix and gives up. Reachable from hipfire_runtime::ep::all_reduce (ep_batch.rs), i.e. any EP/MoE model on more than one card (DeepSeek-V4, MiniMax, ornith A3B with --tp). Single-card serve and dense TP (rooted reduce-add, #662) never touch RCCL and are unaffected. Distributions that keep RCCL under the ROCm prefix (AMD apt/rpm repos, Arch /opt/rocm, Fedora and Gentoo /usr) are unaffected too, which is why nobody hit it upstream.

What changes

HIPFIRE_RCCL_LIB (read through developer_var, row in docs/env-vars.md): an explicit librccl.so path tried before the ROCm-root candidates. The root policy is untouched; no bare-soname search is added.
The stale "bare sonames last" comment is replaced by what the code does and why RCCL follows the root policy.
The dlopen error names the variable next to the existing apt / /opt/rocm hints.
crates/hip-bridge/map.md regenerated (scripts/check-crate-maps.py hip-bridge).

Which surface(s) does this touch?

Test plan

./scripts/no-gpu-ci.sh: Rust check and no-GPU unit tests pass (732 tests, 0 failures); the Python stage did not run in this shell (pytest/numpy missing and uv unavailable), and this change has no Python. scripts/check-env-docs.py and scripts/check-crate-maps.py --check print byte-identical output to a clean master checkout: the crate-map drift they report (saddle-core, rdna-compute, hipfire-dispatch, hipfire-generate, hipfire-runtime) is on master today and is not touched here
cargo build --release -p hip-bridge --features lab --example rccl_smoke clean
hip-bridge/examples/rccl_smoke on 2× RX 7900 XTX (gfx1100), nixpkgs ROCm 7.2, HIPFIRE_ROCM_PATH = rocmtoolkit-merged, HIP_VISIBLE_DEVICES=0,1, binary md5 f5b44f51bd462f6a5f6e014bac6df32d — A/B below
speed-gate: not perf-relevant

rccl_smoke A/B (same binary; only the variable differs)

A — HIPFIRE_RCCL_LIB unset, LD_LIBRARY_PATH = root only (what master does today, error text from this branch):

RCCL smoke (Rust FFI): ranks=2 iters=20

thread 'main' panicked at crates/hip-bridge/examples/rccl_smoke.rs:37:49:
init_all: RcclError { status: 0, context: "failed to dlopen librccl.so. Tried: [
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib/librccl.so\",
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib/librccl.so.1\",
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib/librccl.so.1.0\",
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib64/librccl.so\",
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib64/librccl.so.1\",
  \"/nix/store/k6dy…-rocmtoolkit-merged/lib64/librccl.so.1.0\"].
  Is RCCL installed (apt install rccl, or /opt/rocm/lib/librccl.so.1)?
  If it lives outside the ROCm root, set HIPFIRE_RCCL_LIB=/path/to/librccl.so." }

B — HIPFIRE_RCCL_LIB=/nix/store/c8dc…-rccl-7.2.3/lib/librccl.so:

RCCL smoke (Rust FFI): ranks=2 iters=20
ncclCommInitAll(n=2) OK, RCCL version = 22707

size             median µs      p10 µs      p90 µs     BW GB/s
4 KB                  72.4        71.1        81.2        0.06
32 KB                 83.0        81.9       107.2        0.39
128 KB                89.9        88.3       102.3        1.46
512 KB               124.6       115.3       130.8        4.21

rccl_smoke: PASS

Hardware validation request (optional)

{
  "routes": [
    {"mode": "battery", "tag": "qwen3.6:27b"}
  ],
  "claim": "single-card and dense-TP serve never reach this code; with RCCL under the ROCm root the candidate list is unchanged, so no hardware route changes behaviour. The fix is reachable only where the resolved root lacks librccl (nixpkgs)."
}

Related: #162, #645 (NixOS packaging).