build-llama-rdna4: make GGML_HIP_RCCL overridable and bundle librccl

Layer split needs no collectives, so RCCL stayed OFF. The 2026-09-20 board
swap put both R9700s on CPU root ports with atomic routing, so a
tensor-parallel (row-split) runtime is now buildable: pass
GGML_HIP_RCCL=ON with a distinct BUILD_DIR beside the production one.

scripts/build-llama-rdna4.sh

@@ -38,6 +38,7 @@ JOBS="${JOBS:-$(nproc)}"
 ROCM_VERSION="${ROCM_VERSION:-7.14.0a20260624}"
 ROCM_LOCAL_PATH="${ROCM_LOCAL_PATH:-${XDG_DATA_HOME:-$HOME/.local/share}/rocm-${ROCM_VERSION}}"
 HIP_ARCH="${HIP_ARCH:-gfx1201}"
+GGML_HIP_RCCL="${GGML_HIP_RCCL:-OFF}"
 
 # ── Source ────────────────────────────────────────────────────────────
 if [[ ! -d "$LLAMA_SRC/.git" ]]; then
@@ -70,11 +71,13 @@ export ROCM_PATH="$ROCM_LOCAL_PATH"
 export HIP_PATH="$ROCM_LOCAL_PATH"
 
 # ── Configure ─────────────────────────────────────────────────────────
-# GGML_HIP_RCCL stays OFF: this host's second card cannot dispatch a RCCL
-# kernel at all (its PCIe path has no atomic routing), and layer-split serving
-# needs no collectives.  With RCCL off, llama.cpp logs "internal AllReduce init
-# failed ... falling back to meta-backend butterfly" and only --split-mode
-# tensor pays for it.
+# GGML_HIP_RCCL defaults OFF: layer-split serving needs no collectives, and
+# until 2026-09-20 this host's second card could not dispatch a RCCL kernel at
+# all (its PCIe path had no atomic routing).  With RCCL off, llama.cpp logs
+# "internal AllReduce init failed ... falling back to meta-backend butterfly"
+# and only --split-mode row/tensor pays for it.  Set GGML_HIP_RCCL=ON (with a
+# distinct BUILD_DIR) for the collective-backed runtime; the atomics
+# precondition now holds on both cards.
 cmake -S "$LLAMA_SRC" -B "$BUILD_DIR" \
     -G Ninja \
     -DCMAKE_BUILD_TYPE=Release \
@@ -83,7 +86,7 @@ cmake -S "$LLAMA_SRC" -B "$BUILD_DIR" \
     -DCMAKE_CXX_FLAGS="-I${ROCM_LOCAL_PATH}/include" \
     -DCMAKE_POSITION_INDEPENDENT_CODE=ON \
     -DGGML_HIP=ON \
-    -DGGML_HIP_RCCL=OFF \
+    -DGGML_HIP_RCCL="$GGML_HIP_RCCL" \
     -DGGML_CUDA=OFF \
     -DGGML_VULKAN=OFF \
     -DGGML_NATIVE=OFF \
@@ -116,6 +119,13 @@ if command -v patchelf >/dev/null; then
           "$ROCM_LOCAL_PATH/lib/llvm/lib"/libomp.so* \
           "$ROCM_LOCAL_PATH/lib/llvm/lib"/libomptarget.so* "$LIB_DIR/" 2>/dev/null || true
     cp -v "$ROCM_LOCAL_PATH/lib/rocm_sysdeps/lib"/librocm_sysdeps_*.so* "$LIB_DIR/" 2>/dev/null || true
+    # A RCCL build links librccl.so.1 (imported shared even with
+    # BUILD_SHARED_LIBS=OFF), and the binaries resolve libraries through
+    # $ORIGIN/../lib — so the collective library must ride beside the rest or
+    # llama-server fails to exec at all.
+    if [[ "$GGML_HIP_RCCL" == "ON" ]]; then
+        cp -v "$ROCM_LOCAL_PATH"/lib/librccl.so* "$LIB_DIR/" 2>/dev/null || true
+    fi
     # rocBLAS and hipBLASLt look for their device code at a FIXED depth under
     # the binary's lib dir: $ORIGIN/../lib/rocblas/library and
     # $ORIGIN/../lib/hipblaslt/library/<arch>. `cp -r SRC DST/` with a missing