From 9b0b7da96e3c22de2722e1ca614ea15869b5f8fd Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Mon, 5 Oct 2026 01:38:09 +0200 Subject: [PATCH] fix(rocm): point *_TENSILE_LIBPATH at the architecture folder (#12411) Newer ROCm releases split the rocBLAS and hipBLASLt TensileLibrary data into one folder per GPU architecture (lib/hipblaslt/library/gfx1151/...). The run.sh of the ROCm-capable backends exports ROCBLAS_TENSILE_LIBPATH and HIPBLASLT_TENSILE_LIBPATH as the parent folder, and both libraries only look directly in that folder, so they miss every kernel: rocblaslt error: Cannot read ".../hipblaslt/library/TensileLibrary_lazy_gfx1151.dat" hipModuleLoad failed: .../hipblaslt/library/Kernels.so-000-gfx1151.hsaco and fall back to slower code paths. rocm_tensile_dir keeps the folder when it has files at the top level (the older flat layout) or several architecture folders, and descends into the architecture folder when the bundle carries exactly one. Assisted-by: Claude:claude-opus-5-5 Signed-off-by: Stefan Walcz --- backend/cpp/bonsai/run.sh | 21 +++++++++++++++++++-- backend/cpp/llama-cpp/run.sh | 21 +++++++++++++++++++-- backend/cpp/turboquant/run.sh | 21 +++++++++++++++++++-- 3 files changed, 57 insertions(+), 6 deletions(-) diff --git a/backend/cpp/bonsai/run.sh b/backend/cpp/bonsai/run.sh index e59c5954e..62ebe47f6 100755 --- a/backend/cpp/bonsai/run.sh +++ b/backend/cpp/bonsai/run.sh @@ -25,6 +25,23 @@ if [ -n "$LLAMACPP_GRPC_SERVERS" ]; then fi fi +# Newer ROCm releases split the TensileLibrary data per GPU architecture +# (library/gfx1151/...). rocBLAS and hipBLASLt look for their kernels directly in +# the directory the *_TENSILE_LIBPATH variable names, so pointing at the parent +# makes them miss every kernel ("Cannot read TensileLibrary_lazy_gfx1151.dat") +# and fall back to slower paths. Point at the architecture folder when the bundle +# carries exactly one and nothing at the top level; otherwise keep the directory. +rocm_tensile_dir() { + local dir="$1" + if [ -z "$(find "$dir" -maxdepth 1 -type f -print -quit)" ]; then + set -- "$dir"/gfx*/ + if [ "$#" -eq 1 ] && [ -d "$1" ]; then + dir="${1%/}" + fi + fi + echo "$dir" +} + # Extend ld library path with the dir where this script is located/lib if [ "$(uname)" == "Darwin" ]; then export DYLD_LIBRARY_PATH="$CURDIR"/lib:$DYLD_LIBRARY_PATH @@ -32,13 +49,13 @@ else export LD_LIBRARY_PATH="$CURDIR"/lib:$LD_LIBRARY_PATH # Tell rocBLAS where to find TensileLibrary data (GPU kernel tuning files) if [ -d "$CURDIR/lib/rocblas/library" ]; then - export ROCBLAS_TENSILE_LIBPATH="$CURDIR"/lib/rocblas/library + export ROCBLAS_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/rocblas/library)" fi # Same for hipBLASLt (rocblaslt): the bundled libhipblaslt.so resolves its # TensileLibrary_lazy_gfx*.dat kernel data relative to itself, so point it at # the bundled data or it falls back to slow generic kernels (issue #10660). if [ -d "$CURDIR/lib/hipblaslt/library" ]; then - export HIPBLASLT_TENSILE_LIBPATH="$CURDIR"/lib/hipblaslt/library + export HIPBLASLT_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/hipblaslt/library)" fi # Backends built for Intel GPUs carry a copy of the Intel graphics driver, # and libze_loader is only there in those builds. Level Zero looks for a diff --git a/backend/cpp/llama-cpp/run.sh b/backend/cpp/llama-cpp/run.sh index 4c7ad19c5..cf08a65c6 100755 --- a/backend/cpp/llama-cpp/run.sh +++ b/backend/cpp/llama-cpp/run.sh @@ -27,6 +27,23 @@ if [ -n "$LLAMACPP_GRPC_SERVERS" ]; then fi fi +# Newer ROCm releases split the TensileLibrary data per GPU architecture +# (library/gfx1151/...). rocBLAS and hipBLASLt look for their kernels directly in +# the directory the *_TENSILE_LIBPATH variable names, so pointing at the parent +# makes them miss every kernel ("Cannot read TensileLibrary_lazy_gfx1151.dat") +# and fall back to slower paths. Point at the architecture folder when the bundle +# carries exactly one and nothing at the top level; otherwise keep the directory. +rocm_tensile_dir() { + local dir="$1" + if [ -z "$(find "$dir" -maxdepth 1 -type f -print -quit)" ]; then + set -- "$dir"/gfx*/ + if [ "$#" -eq 1 ] && [ -d "$1" ]; then + dir="${1%/}" + fi + fi + echo "$dir" +} + # Extend ld library path with the dir where this script is located/lib if [ "$(uname)" == "Darwin" ]; then export DYLD_LIBRARY_PATH="$CURDIR"/lib:$DYLD_LIBRARY_PATH @@ -35,13 +52,13 @@ else export LD_LIBRARY_PATH="$CURDIR"/lib:$LD_LIBRARY_PATH # Tell rocBLAS where to find TensileLibrary data (GPU kernel tuning files) if [ -d "$CURDIR/lib/rocblas/library" ]; then - export ROCBLAS_TENSILE_LIBPATH="$CURDIR"/lib/rocblas/library + export ROCBLAS_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/rocblas/library)" fi # Same for hipBLASLt (rocblaslt): the bundled libhipblaslt.so resolves its # TensileLibrary_lazy_gfx*.dat kernel data relative to itself, so point it at # the bundled data or it falls back to slow generic kernels (issue #10660). if [ -d "$CURDIR/lib/hipblaslt/library" ]; then - export HIPBLASLT_TENSILE_LIBPATH="$CURDIR"/lib/hipblaslt/library + export HIPBLASLT_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/hipblaslt/library)" fi # Backends built for Intel GPUs carry a copy of the Intel graphics driver, # and libze_loader is only there in those builds. Level Zero looks for a diff --git a/backend/cpp/turboquant/run.sh b/backend/cpp/turboquant/run.sh index 44f1f942f..327818f62 100755 --- a/backend/cpp/turboquant/run.sh +++ b/backend/cpp/turboquant/run.sh @@ -28,6 +28,23 @@ if [ -n "$LLAMACPP_GRPC_SERVERS" ]; then fi fi +# Newer ROCm releases split the TensileLibrary data per GPU architecture +# (library/gfx1151/...). rocBLAS and hipBLASLt look for their kernels directly in +# the directory the *_TENSILE_LIBPATH variable names, so pointing at the parent +# makes them miss every kernel ("Cannot read TensileLibrary_lazy_gfx1151.dat") +# and fall back to slower paths. Point at the architecture folder when the bundle +# carries exactly one and nothing at the top level; otherwise keep the directory. +rocm_tensile_dir() { + local dir="$1" + if [ -z "$(find "$dir" -maxdepth 1 -type f -print -quit)" ]; then + set -- "$dir"/gfx*/ + if [ "$#" -eq 1 ] && [ -d "$1" ]; then + dir="${1%/}" + fi + fi + echo "$dir" +} + # Extend ld library path with the dir where this script is located/lib if [ "$(uname)" == "Darwin" ]; then export DYLD_LIBRARY_PATH="$CURDIR"/lib:$DYLD_LIBRARY_PATH @@ -35,13 +52,13 @@ else export LD_LIBRARY_PATH="$CURDIR"/lib:$LD_LIBRARY_PATH # Tell rocBLAS where to find TensileLibrary data (GPU kernel tuning files) if [ -d "$CURDIR/lib/rocblas/library" ]; then - export ROCBLAS_TENSILE_LIBPATH="$CURDIR"/lib/rocblas/library + export ROCBLAS_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/rocblas/library)" fi # Same for hipBLASLt (rocblaslt): the bundled libhipblaslt.so resolves its # TensileLibrary_lazy_gfx*.dat kernel data relative to itself, so point it at # the bundled data or it falls back to slow generic kernels (issue #10660). if [ -d "$CURDIR/lib/hipblaslt/library" ]; then - export HIPBLASLT_TENSILE_LIBPATH="$CURDIR"/lib/hipblaslt/library + export HIPBLASLT_TENSILE_LIBPATH="$(rocm_tensile_dir "$CURDIR"/lib/hipblaslt/library)" fi # Backends built for Intel GPUs carry a copy of the Intel graphics driver, # and libze_loader is only there in those builds. Level Zero looks for a