feat: Build hipSPARSELt device libraries via find_package(hipblaslt); ship tensilelite-host shared#9248
Conversation
…+ 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 <noreply@anthropic.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (76.92%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #9248 +/- ##
===========================================
- Coverage 64.99% 64.99% -0.00%
===========================================
Files 2685 2685
Lines 422116 422157 +41
Branches 62707 62717 +10
===========================================
+ Hits 274344 274356 +12
- Misses 127073 127102 +29
Partials 20699 20699
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
davidd-amd
changed the title
| if(NOT EXISTS "${HIPBLASLT_PATH}") | ||
| message(FATAL_ERROR "hipBLASLt required but no not found at ${HIPBLASLT_PATH}.") | ||
| endif() | ||
| find_package(hipblaslt CONFIG REQUIRED) |
There was a problem hiding this comment.
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
| set(HIPBLASLT_PYTHON_COMMAND | ||
| "${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}" |
There was a problem hiding this comment.
I think this also needs something for libstinkytofu
| $<BUILD_INTERFACE:rocisa::rocisa-cpp> | ||
| $<BUILD_INTERFACE:roc::origami> | ||
| $<BUILD_INTERFACE:hip::host> |
There was a problem hiding this comment.
Shouldn't these also be install_interface. I believe that the resulting public headers for tensillite-host will have references to all three of these dependencies. We're currently fine since all three are being important manually wherever tensilelite is used.
| 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}) |
There was a problem hiding this comment.
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.
| # 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) |
There was a problem hiding this comment.
This looks both fragile, and linux specific.
| # 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 |
There was a problem hiding this comment.
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?
| ) | ||
| endif() | ||
|
|
||
| add_subdirectory("${HIPBLASLT_PATH}" hipblaslt) |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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.
| 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) |
There was a problem hiding this comment.
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.
4 participants
Summary
This completes the TensileLite shared-library work. P1 (#7996) made origami a findable
package and P2 (#8133) turned
tensilelite-hostinto an exportable target and landed thehipblaslt_create_device_library()codegen module — but P2 shippedtensilelite-hoststatic by default and left hipSPARSELt building a private, in-tree copy of hipBLASLt via
add_subdirectory(hipblaslt).This PR flips the switch:
tensilelite-hostbuilds shared by default, and hipSPARSELtstops rebuilding hipBLASLt from source. Instead it consumes the installed hipBLASLt package
via
find_package(hipblaslt CONFIG)and generates its own sparse device libraries by callingthe exported
hipblaslt_create_device_library(). The result is a single sharedlibtensilelite-host.soacross both consumers (true structural dedup) and no duplicatehipBLASLt build inside hipSPARSELt.
What changed
Five files,
+91 / -32:projects/hipsparselt/CMakeLists.txt— the core change. Replace the inlineadd_subdirectory(hipblaslt)(which built a private HOST-off/DEVICE-on hipBLASLt inhipSPARSELt's own tree) with
find_package(hipblaslt CONFIG REQUIRED). WhenHIPSPARSELT_ENABLE_DEVICE, drive sparse codegen through the exportedhipblaslt_create_device_library()against the binary package: pointHIPBLASLT_CODEGEN_ROOTat the hipBLASLt tensilelite source (same repo, deliberatelynot installed), locate the installed
rocisafrom the importedroc::tensilelite-hosttarget, and set
PYTHONPATH/LD_LIBRARY_PATH/PATHfor the codegen subprocess. Installthe generated
library/tolib/hipsparselt.projects/hipblaslt/tensilelite/CMakeLists.txt— defaultTENSILELITE_BUILD_SHARED_LIBSto
ON, so the integrated build produces one co-exportedlibtensilelite-host.soinsteadof statically embedding a copy in each consumer (the dedup P2 deferred).
projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt— install therocisaextension(
_rocisa,__init__.py, and the bundledlibstinkytofu) tolib/hipblaslt/rocisaunderthe new option, so a downstream consumer can
import rocisawhen running codegen againstthe staged package. Guarded on
HIPBLASLT_INSTALL_ROCISA AND NOT ROCISA_STANDALONE.projects/hipblaslt/CMakeLists.txt— addoption(HIPBLASLT_INSTALL_ROCISA ...),defaulting to
HIPBLASLT_ENABLE_HOST.projects/hipblaslt/cmake/HipBLASLtCodegen.cmake— honorHIPBLASLT_CODEGEN_ROOTwhenset (falling back to the in-source path), so a separate-subproject consumer can point the
codegen at the hipBLASLt source tree.
Downstream (TheRock) coordination — not in this PR
The super-project consumer wiring lands in a separate TheRock PR (ROCm/TheRock#5650):
origami declared as a shared BLAS subproject, hipSPARSELt gaining
hipBLASLt/origami(RUNTIME) and
hipBLAS-common(BUILD) deps sofind_package(hipblaslt)and the transitivefind_dependency(hipblas-common)resolve, and the artifact manifest claiminglib/hipblaslt/rocisa/**. Because rocm-libraries -> TheRock pin bumps are automated, thatTheRock PR must be merged before the automated bump of this commit, or the composed build
breaks. This is a merge-time sequencing item, not a code dependency in this PR.
JIRA ID
Test Plan
cmake -B build -GNinja . -DTHEROCK_AMDGPU_FAMILIES=gfx90a -DTHEROCK_BUILD_TESTING=OFF -DTHEROCK_FLAG_INCLUDE_PROFILER=ONninja -C build origami hipBLASLt hipSPARSELtfind_package(hipblaslt)(noadd_subdirectory),links the single
libtensilelite-host.so,rocisais installed with hipBLASLt and foundby hipSPARSELt codegen, and sparse
TensileLibrary_*_<arch>.co/.dat.zlibaregenerated and staged.
-DTENSILELITE_BUILD_SHARED_LIBS=OFFstill builds and embeds.Test Result
Built in the TheRock manylinux container (gfx90a) with the #5650 consumer wiring
(
reland-origamitip) and the rocm-libraries submodule pinned at this commit:rc=0.find_package(hipblaslt 1.2.2 CONFIG ... PATHS .../hipBLASLt/dist/lib/cmake/hipblaslt)andbuilds host-only in ~14s with no
add_subdirectory(hipblaslt)(grep of the subprojectcache confirms it is absent).
hipblaslt-config.cmakeincludes
HipBLASLtCodegen.cmake, sohipblaslt_create_device_libraryis defined afterfind_package(hipblaslt).rocisainstalls tolib/hipblaslt/rocisa(_rocisa.abi3.so,__init__.py, bundledlibstinkytofu.so).Not yet verified locally: end-to-end sparse
.co/.datgeneration in this build. ThehipSPARSELt subproject received an empty
GPU_TARGETSin the local super-projectconfiguration, so
HIPSPARSELT_ENABLE_DEVICEwas gated off (CMakeLists.txt line ~88) and thecodegen step did not run. That is a super-project arch-propagation detail, not a defect in
this change; it is exercised in the TheRock CI pipeline once the pin is bumped. The device
codegen path should be confirmed green in the TheRock pipeline before this merges.
Related
rebased to this minimal ratchet).
roc::origami) and feat(tensilelite): stop treating tensilelite host as internal library #8133 (P2, exportabletensilelite-host+codegen module).
Submission Checklist
develop.