diff --git a/scripts/build-llama-rdna4.sh b/scripts/build-llama-rdna4.sh index 59c7455..a1ff520 100755 --- a/scripts/build-llama-rdna4.sh +++ b/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/. `cp -r SRC DST/` with a missing