From e649558197366367b4498c524a1479e7ad889b32 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Wed, 1 Jul 2026 18:11:10 +0000 Subject: [PATCH] [tensilelite] hipSPARSELt device codegen via find_package(hipblaslt) + installable rocisa hipSPARSELt consumes hipBLASLt via find_package(hipblaslt CONFIG REQUIRED) and generates its own sparse Tensile device libraries by calling the exported hipblaslt_create_device_library(), instead of add_subdirectory(hipblaslt). The Tensile codegen Python is read from the hipBLASLt source tree (same repo, not installed). The compiled rocisa extension IS a build-time requirement for a separate-subproject consumer, so it is made installable and installed with hipBLASLt (lib/hipblaslt/rocisa); hipSPARSELt derives its directory from the imported roc::tensilelite-host target's location. - tensilelite/CMakeLists.txt: default tensilelite-host to shared. - tensilelite/rocisa/CMakeLists.txt: HIPBLASLT_INSTALL_ROCISA installs rocisa (+ bundled libstinkytofu) in the integrated build. - hipblaslt/CMakeLists.txt: add HIPBLASLT_INSTALL_ROCISA option (default HOST). - cmake/HipBLASLtCodegen.cmake: honor HIPBLASLT_CODEGEN_ROOT (else source fallback). - hipsparselt/CMakeLists.txt: find_package(hipblaslt) + create_device_library, Tensile from source, rocisa via tensilelite-host target location. Co-Authored-By: Claude Opus 4.8 --- projects/hipblaslt/CMakeLists.txt | 3 + .../hipblaslt/cmake/HipBLASLtCodegen.cmake | 6 +- projects/hipblaslt/tensilelite/CMakeLists.txt | 2 +- .../tensilelite/rocisa/CMakeLists.txt | 20 ++++ projects/hipsparselt/CMakeLists.txt | 92 +++++++++++++------ 5 files changed, 91 insertions(+), 32 deletions(-) diff --git a/projects/hipblaslt/CMakeLists.txt b/projects/hipblaslt/CMakeLists.txt index 499df8a2fcde..2a56ac3e52f4 100644 --- a/projects/hipblaslt/CMakeLists.txt +++ b/projects/hipblaslt/CMakeLists.txt @@ -73,6 +73,9 @@ option(TENSILELITE_ENABLE_CLIENT "Build tensilelite client" ${TENSILELITE_BUILD_ option(HIPBLASLT_INSTALL_TENSILELITE_TEST_ARTIFACTS "Install TensileLite Python packages and test tools into the test artifact tree." OFF) +option(HIPBLASLT_INSTALL_ROCISA + "Install the rocisa Python extension with hipBLASLt so downstream consumers (e.g. hipSPARSELt) can import it to run device-library codegen against the staged hipBLASLt package." + ${HIPBLASLT_ENABLE_HOST}) option(TENSILELITE_ENABLE_AUTOBUILD "Generate scripts for wrapping tensile python scripts" OFF) if(HIPBLASLT_ENABLE_CLIENT) diff --git a/projects/hipblaslt/cmake/HipBLASLtCodegen.cmake b/projects/hipblaslt/cmake/HipBLASLtCodegen.cmake index d4651bc766a1..1124cf695a8b 100644 --- a/projects/hipblaslt/cmake/HipBLASLtCodegen.cmake +++ b/projects/hipblaslt/cmake/HipBLASLtCodegen.cmake @@ -24,7 +24,11 @@ function(hipblaslt_create_device_library) message(FATAL_ERROR "hipblaslt_create_device_library: HIPBLASLT_PYTHON_COMMAND is not set") endif() - get_filename_component(_codegen_dir "${CMAKE_CURRENT_LIST_DIR}/../tensilelite" ABSOLUTE) + if(HIPBLASLT_CODEGEN_ROOT) + set(_codegen_dir "${HIPBLASLT_CODEGEN_ROOT}") + else() + get_filename_component(_codegen_dir "${CMAKE_CURRENT_LIST_DIR}/../tensilelite" ABSOLUTE) + endif() if(NOT _cdl_TARGET) set(_cdl_TARGET "tensilelite-device-libraries") diff --git a/projects/hipblaslt/tensilelite/CMakeLists.txt b/projects/hipblaslt/tensilelite/CMakeLists.txt index 39efd988e620..a8cde63cc31d 100644 --- a/projects/hipblaslt/tensilelite/CMakeLists.txt +++ b/projects/hipblaslt/tensilelite/CMakeLists.txt @@ -14,7 +14,7 @@ endif() if(TENSILELITE_ENABLE_HOST) find_package(ZLIB REQUIRED) - option(TENSILELITE_BUILD_SHARED_LIBS "Build tensilelite-host as shared vs static" OFF) + option(TENSILELITE_BUILD_SHARED_LIBS "Build tensilelite-host as shared vs static" ON) set(BUILD_SHARED_LIBS ${TENSILELITE_BUILD_SHARED_LIBS}) add_library(tensilelite-host) add_library(roc::tensilelite-host ALIAS tensilelite-host) diff --git a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt index 2b44eb7c293a..eb28ea431576 100644 --- a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt +++ b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt @@ -250,6 +250,26 @@ if(HIPBLASLT_BUNDLE_PYTHON_DEPS OR ROCISA_STANDALONE) install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/rocisa/__init__.py" DESTINATION rocisa) endif() + # Integrated ROCm build (add_subdirectory from hipBLASLt): install rocisa so + # downstream consumers building device libraries via hipblaslt_create_device_library() + # against the staged hipBLASLt package can `import rocisa`. Installed under + # lib/hipblaslt/rocisa so PYTHONPATH=lib/hipblaslt resolves it; consumers derive + # that dir from the imported tensilelite-host target's location. + if(HIPBLASLT_INSTALL_ROCISA AND NOT ROCISA_STANDALONE) + set(_rocisa_install_dir "${CMAKE_INSTALL_LIBDIR}/hipblaslt/rocisa") + install(TARGETS _rocisa DESTINATION "${_rocisa_install_dir}" COMPONENT runtime) + install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/rocisa/__init__.py" + DESTINATION "${_rocisa_install_dir}" COMPONENT runtime) + if(_stinkytofu_from_subdirectory AND NOT WIN32) + # Bundled stinkytofu: vendor libstinkytofu next to _rocisa; _rocisa's + # INSTALL_RPATH=$ORIGIN (set above) resolves it at import time. + install(FILES $ + DESTINATION "${_rocisa_install_dir}" COMPONENT runtime) + install(FILES $ + DESTINATION "${_rocisa_install_dir}" COMPONENT runtime) + endif() + endif() + # Example plugin — the target only exists when stinkytofu is configured with # -DSTINKYTOFU_BUILD_EXAMPLES=ON (default OFF; opted in by the source-built standalone # wheel via tasks.py). Integrated/ROCm builds leave it OFF, so the target is absent and diff --git a/projects/hipsparselt/CMakeLists.txt b/projects/hipsparselt/CMakeLists.txt index 7e7f0eb0964b..741902a26eed 100644 --- a/projects/hipsparselt/CMakeLists.txt +++ b/projects/hipsparselt/CMakeLists.txt @@ -147,44 +147,76 @@ if(HIPSPARSELT_ENABLE_CUDA) endif() if(HIPSPARSELT_ENABLE_HIP) - set(HIPBLASLT_PATH "${HIPSPARSELT_HIPBLASLT_PATH}") - if(NOT EXISTS "${HIPBLASLT_PATH}") - message(FATAL_ERROR "hipBLASLt required but no not found at ${HIPBLASLT_PATH}.") - endif() + find_package(hipblaslt CONFIG REQUIRED) - block(SCOPE_FOR VARIABLES) - set(HIPBLASLT_ENABLE_HOST OFF CACHE BOOL "" FORCE) - set(HIPBLASLT_ENABLE_CLIENT OFF CACHE BOOL "" FORCE) - # Builds tensilelite's device libraries (.co/.dat) - set(HIPBLASLT_ENABLE_DEVICE ${HIPSPARSELT_ENABLE_DEVICE} CACHE BOOL "" FORCE) - set(TENSILELITE_ENABLE_HOST ON CACHE BOOL "" FORCE) - set(TENSILELITE_BUILD_SHARED_LIBS OFF CACHE BOOL "" FORCE) - set(HIPBLASLT_ENABLE_ASAN ${HIPSPARSELT_ENABLE_ASAN} CACHE BOOL "" FORCE) - set(HIPBLASLT_DEVICE_LIBRARY_PATH "${CMAKE_BINARY_DIR}/Tensile" CACHE PATH "" FORCE) - # hipSPARSELt does not use ExtOp or MatrixTransform artifacts - set(HIPBLASLT_ENABLE_EXTOPS OFF CACHE BOOL "" FORCE) - set(HIPBLASLT_ENABLE_MATRIX_TRANSFORM OFF CACHE BOOL "" FORCE) - set(HIPBLASLT_LIBLOGIC_PATH - "${CMAKE_CURRENT_SOURCE_DIR}/library/src/hcc_detail/rocsparselt/src/spmm/Tensile/Logic/asm_full" - CACHE STRING "" FORCE - ) set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md" CACHE FILEPATH "License file for package" FORCE ) - if(WIN32) - set(HIPBLASLT_TENSILE_LIBRARY_DIR "\${CPACK_PACKAGING_INSTALL_PREFIX}hipsparselt/bin" - CACHE PATH "path to tensile library" FORCE + if(HIPSPARSELT_ENABLE_DEVICE) + # hipSPARSELt ships its own sparse Tensile logic and generates its device + # libraries (.co/.dat) by calling hipBLASLt's exported + # hipblaslt_create_device_library() against the binary hipBLASLt package, + # instead of add_subdirectory(hipblaslt). + # + # The Tensile codegen Python lives in the hipBLASLt source tree (same repo, + # deliberately not installed); the compiled rocisa extension IS installed + # with hipBLASLt (lib/hipblaslt/rocisa), so its directory is derived from + # the imported roc::tensilelite-host target's location. + if(NOT HIPBLASLT_TENSILELITE_PATH) + set(HIPBLASLT_TENSILELITE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../hipblaslt/tensilelite") + endif() + if(NOT EXISTS "${HIPBLASLT_TENSILELITE_PATH}/Tensile/bin/TensileLogic") + message(FATAL_ERROR + "hipSPARSELt device codegen needs the hipBLASLt tensilelite source at " + "'${HIPBLASLT_TENSILELITE_PATH}'; set -DHIPBLASLT_TENSILELITE_PATH.") + endif() + set(HIPBLASLT_CODEGEN_ROOT "${HIPBLASLT_TENSILELITE_PATH}") + + # The installed _rocisa extension has DT_NEEDED on liborigami (linked + # PUBLIC) and libstinkytofu; its INSTALL_RPATH ($ORIGIN) only covers the + # vendored stinkytofu beside it. Since the super-project resolves each + # dependency from its own dist dir (not one merged prefix), an RPATH cannot + # span them, so put origami's and rocisa's dirs on LD_LIBRARY_PATH for the + # codegen subprocess. libamdhip64 is already resolvable from the build env. + find_package(origami CONFIG REQUIRED) + find_package(hip REQUIRED) + + find_package(Python3 COMPONENTS Interpreter REQUIRED) + set(_hipsparselt_rocisa_dir "$/hipblaslt/rocisa") + # Tensile's toolchain validator resolves hipconfig (and clang) from PATH; + # provide the hip runtime bin dir and the compiler dir so the codegen + # subprocess finds them (mirrors hipBLASLt's own bundled-python PATH). + get_filename_component(_hipsparselt_hip_bindir "${hip_DIR}/../../../bin" ABSOLUTE) + get_filename_component(_hipsparselt_cxx_bindir "${CMAKE_CXX_COMPILER}" DIRECTORY) + set(HIPBLASLT_PYTHON_COMMAND + "${CMAKE_COMMAND}" -E env + "PYTHONPATH=$/hipblaslt:${HIPBLASLT_CODEGEN_ROOT}" + "LD_LIBRARY_PATH=$:${_hipsparselt_rocisa_dir}:$ENV{LD_LIBRARY_PATH}" + "PATH=${_hipsparselt_hip_bindir}:${_hipsparselt_cxx_bindir}:$ENV{PATH}" + -- + "${Python3_EXECUTABLE}") + + set(_hipsparselt_tensile_out "${CMAKE_BINARY_DIR}/Tensile") + set(_hipsparselt_codegen_flags "") + if(HIPSPARSELT_ENABLE_ASAN) + list(APPEND _hipsparselt_codegen_flags ASAN) + endif() + hipblaslt_create_device_library( + TARGET tensilelite-device-libraries + LOGIC_PATH "${CMAKE_CURRENT_SOURCE_DIR}/library/src/hcc_detail/rocsparselt/src/spmm/Tensile/Logic/asm_full" + OUTPUT_DIR "${_hipsparselt_tensile_out}" + ARCHES ${GPU_TARGETS} + CXX_COMPILER "${CMAKE_CXX_COMPILER}" + ${_hipsparselt_codegen_flags} ) - else() - set(HIPBLASLT_TENSILE_LIBRARY_DIR - "\${CPACK_PACKAGING_INSTALL_PREFIX}${CMAKE_INSTALL_LIBDIR}/hipsparselt" - CACHE PATH "path to tensile library" FORCE + + install( + DIRECTORY "${_hipsparselt_tensile_out}/library" + DESTINATION "${CMAKE_INSTALL_LIBDIR}/hipsparselt" + COMPONENT runtime ) endif() - - add_subdirectory("${HIPBLASLT_PATH}" hipblaslt) - endblock() endif() add_subdirectory(library)