Skip to content

fix(rccl): HIPFIRE_RCCL_LIB points the loader at a librccl outside the ROCm root (nixpkgs ships it in its own store path) - #728

Open
alpineQ wants to merge 1 commit into
warpfront:masterfrom
alpineQ:fix/rccl-lib-override
Open

fix(rccl): HIPFIRE_RCCL_LIB points the loader at a librccl outside the ROCm root (nixpkgs ships it in its own store path)#728
alpineQ wants to merge 1 commit into
warpfront:masterfrom
alpineQ:fix/rccl-lib-override

Conversation

@alpineQ

@alpineQ alpineQ commented Sep 5, 2026

Copy link
Copy Markdown
Contributor

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?

  • kernelcrates/hip-bridge (dlopen candidate list in rccl.rs only; no kernel, dispatch or FFI signature change)
  • load
  • serve
  • arch crate(s)
  • crates/hipfire-quantize
  • control plane
  • docs / CI / scripts only
  • policy files

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).

Architecture-trait change?

No.

…e ROCm root (nixpkgs ships it in its own store path)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant