Skip to content

feat: Build hipSPARSELt device libraries via find_package(hipblaslt); ship tensilelite-host shared #9248

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
davidd-amd wants to merge 1 commit into
base: develop
Choose a base branch
from
+91 −32
Open

feat: Build hipSPARSELt device libraries via find_package(hipblaslt); ship tensilelite-host shared #9248

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions projects/hipblaslt/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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})
Comment on lines +76 to +78

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is misleading since it makes it sound like it enforces that rocisa will be installed. But I don't think that's true.

The the python rocisa extension is gated on: (HIPBLASLT_BUNDLE_PYTHON_DEPS OR ROCISA_STANDALONE) (rocisa/CMakeLists.txt:40)

For an integrated build, the second condition is false. So we need HIPBLASLT_BUNDLE_PYTHON_DEPS to build it.

This won't error out in the build, because the install(TARGETS _rocisa ...) is gated off, but the final install tree won't have the needed rocisa files.

option(TENSILELITE_ENABLE_AUTOBUILD "Generate scripts for wrapping tensile python scripts" OFF)

if(HIPBLASLT_ENABLE_CLIENT)
Expand Down
6 changes: 5 additions & 1 deletion projects/hipblaslt/cmake/HipBLASLtCodegen.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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")
Expand Down
2 changes: 1 addition & 1 deletion projects/hipblaslt/tensilelite/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

if we now run invoke build --static we get a static hipblaslt but a shared tensilelite-host. That seems counter-intuitive for someone asking for a static build.

set(BUILD_SHARED_LIBS ${TENSILELITE_BUILD_SHARED_LIBS})
add_library(tensilelite-host)
add_library(roc::tensilelite-host ALIAS tensilelite-host)
Expand Down
20 changes: 20 additions & 0 deletions projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 $<TARGET_FILE:stinkytofu::stinkytofu>
DESTINATION "${_rocisa_install_dir}" COMPONENT runtime)
install(FILES $<TARGET_SONAME_FILE:stinkytofu::stinkytofu>
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
Expand Down
92 changes: 62 additions & 30 deletions projects/hipsparselt/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

hipblaslt only conditionally builds tensilelite-host.

The old approach forced TENSILELITE_BUILD_SHARED_LIBS to true to ensure we had tensilelite-host. This new build does not guarantee that the hipblaslt we find has tensilelite-host installed.

It's needed by hipsparselt/library/src/hcc_detail/rocsparselt/CMakeLists.txt:31


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 "$<TARGET_FILE_DIR:roc::tensilelite-host>/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)

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This looks both fragile, and linux specific.

get_filename_component(_hipsparselt_cxx_bindir "${CMAKE_CXX_COMPILER}" DIRECTORY)
set(HIPBLASLT_PYTHON_COMMAND

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

hiblaslt has a helper function for something very similar: hipblaslt_configure_bundled_python_command.

Hisparselt's cmake file is depending on implementation details from inside hipblaslt. Can this be cleaned up so there isn't as much duplication?

"${CMAKE_COMMAND}" -E env
"PYTHONPATH=$<TARGET_FILE_DIR:roc::tensilelite-host>/hipblaslt:${HIPBLASLT_CODEGEN_ROOT}"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:roc::origami>:${_hipsparselt_rocisa_dir}:$ENV{LD_LIBRARY_PATH}"

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this also needs something for libstinkytofu

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

@Alex-Vasile Alex-Vasile Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This used to call hipblaslt/device-library/CMakeLists.txt, which forwarded many build flags that are no longer being hooked up.

E.g. hipsparselt/install-host.sh:684 sets:

if [[ "${keep_build_tmp}" == true ]]; then
    tensile_opt="${tensile_opt} -DTENSILELITE_KEEP_BUILD_TMP=ON"
fi

But that goes nowhere now.

@vin-huang vin-huang Jul 11, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for pointing this out, @Alex-Vasile. I wasn't aware that some of these build flags are no longer being propagated.

One thing I'm concerned about is that we still need a mechanism to keep the generated ASM files. Previously, we relied on build flags such as TENSILELITE_KEEP_BUILD_TMP for this purpose.

@davidd-amd, is there currently any way to preserve the generated ASM files with the new build flow?

This doesn't need to be addressed in this PR, but I'd like to make sure the requirement is captured.

endblock()
endif()

add_subdirectory(library)
Expand Down
Loading