feat(tensilelite): stop treating tensilelite host as internal library

projects/hipblaslt/CMakeLists.txt

@@ -211,6 +211,11 @@ if(HIPBLASLT_ENABLE_DEVICE AND NOT WIN32)
     endif()
 endif()
 
+set(HIPBLASLT_CODEGEN_ROOT "" CACHE STRING "Root of the codegen (tensilelite) Python sources. Defaults to the in-tree 'tensilelite' directory; override to use an installed/alternate tensile.")
+if(NOT HIPBLASLT_CODEGEN_ROOT)
+    set(HIPBLASLT_CODEGEN_ROOT "${hipblaslt_SOURCE_DIR}/tensilelite")
+endif()
+
 if(HIPBLASLT_BUNDLE_PYTHON_DEPS)
     set(HIPBLASLT_PYTHON_DEPS "_rocisa")
     set(hipblaslt_python_dev Development.Module)
@@ -271,13 +276,6 @@ if(NOT ROCM_LIBS_SUPERBUILD)
     endif()
 endif()
 
-# INTERFACE library that owns the public hipblaslt API headers (the in-tree
-# `library/include` subtree plus its build-tree counterpart, where
-# `hipblaslt-export.h` and `hipblaslt-version.h` are generated). Routing the
-# include directories through a target lets consumers pick them up via the
-# build graph rather than via `target_include_directories(<tgt> BEFORE
-# PRIVATE .../library/include ...)`. `hip::host` is exposed as an INTERFACE
-# link because the in-tree headers `#include <hip/...>`.
 add_library(hipblaslt-headers INTERFACE)
 add_library(hipblaslt::headers ALIAS hipblaslt-headers)
 target_include_directories(hipblaslt-headers
@@ -336,7 +334,7 @@ if(HIPBLASLT_ENABLE_HOST)
             roc::${hipblas_target}
         PRIVATE
             hip::device
-            tensilelite::tensilelite-host
+            roc::tensilelite-host
             ${CMAKE_DL_LIBS}
             ${rocTracer}
     )
@@ -426,8 +424,12 @@ if(HIPBLASLT_ENABLE_HOST)
 
     rocm_install_targets(TARGETS hipblaslt)
 
+    set(_hipblaslt_export_targets roc::hipblaslt)
+    if(TENSILELITE_BUILD_SHARED_LIBS)
+        list(APPEND _hipblaslt_export_targets roc::tensilelite-host)
+    endif()
     rocm_export_targets(
-        TARGETS roc::hipblaslt
+        TARGETS ${_hipblaslt_export_targets}
         DEPENDS PACKAGE hip
         DEPENDS PACKAGE ${hipblas_target}
         NAMESPACE roc::
@@ -436,9 +438,17 @@ if(HIPBLASLT_ENABLE_HOST)
     configure_file(
         "${CMAKE_CURRENT_SOURCE_DIR}/cmake/hipblaslt-config.cmake.in"
         "${CMAKE_CURRENT_BINARY_DIR}/hipblaslt-config.cmake"
-        COPYONLY
+        @ONLY
     )
 
+    if(NOT WIN32)
+        install(
+            FILES "${CMAKE_CURRENT_SOURCE_DIR}/cmake/HipBLASLtCodegen.cmake"
+            DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/hipblaslt"
+            COMPONENT devel
+        )
+    endif()
+
     if( LEGACY_HIPBLAS_DIRECT )
         rocm_package_add_dependencies(DEPENDS "hipblas >= 0.50.0")
     else()

projects/hipblaslt/clients/CMakeLists.txt

@@ -26,7 +26,7 @@ target_link_libraries(hipblaslt-clients-common
         ${LAPACK_LIBRARIES}
     PRIVATE
         hip::device
-        tensilelite::tensilelite-host
+        roc::tensilelite-host
 )
 
 if(HIPBLASLT_ENABLE_MXDATAGENERATOR)
@@ -111,7 +111,7 @@ if(BUILD_TESTING OR HIPBLASLT_BUILD_TESTING)
         PRIVATE
             hip::device
             hipblaslt::hipblaslt-clients-common
-            tensilelite::tensilelite-host
+            roc::tensilelite-host
             GTest::gtest
             OpenMP::OpenMP_CXX
             ${CMAKE_DL_LIBS}

projects/hipblaslt/clients/tests/src/caching_library_gtest.cpp

@@ -59,7 +59,7 @@
 // only C++ test binary TheRock builds, ships, and runs for hipBLASLt is this client
 // `hipblaslt-test`
 // (HIPBLASLT_BUILD_TESTING), executed by test/therock/test_hipblaslt.py. Because
-// hipblaslt-test already links tensilelite::tensilelite-host (via
+// hipblaslt-test already links roc::tensilelite-host (via
 // hipblaslt-clients-common), this white-box unit test can include
 // <Tensile/CachingLibrary.hpp> directly and run with no new build dependency.
 // Placing the regression here is what makes it actually execute in CI and guard

projects/hipblaslt/cmake/HipBLASLtCodegen.cmake

@@ -0,0 +1,127 @@
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier:  MIT
+
+include_guard(GLOBAL)
+
+function(hipblaslt_create_device_library)
+    set(_opts "")
+    set(_one
+        TARGET LOGIC_PATH OUTPUT_DIR CXX_COMPILER OFFLOAD_BUNDLER JOBS LOGIC_FILTER
+        ASAN YAML_FORMAT NO_COMPRESS EXPERIMENTAL LAZY_LOAD ASM_COMMENTS KEEP_BUILD_TMP ASM_DEBUG)
+    set(_multi ARCHES)
+    cmake_parse_arguments(_cdl "${_opts}" "${_one}" "${_multi}" ${ARGN})
+
+    if(_cdl_UNPARSED_ARGUMENTS)
+        message(FATAL_ERROR "hipblaslt_create_device_library: unexpected arguments: ${_cdl_UNPARSED_ARGUMENTS}")
+    endif()
+    if(NOT _cdl_LOGIC_PATH)
+        message(FATAL_ERROR "hipblaslt_create_device_library: LOGIC_PATH is required")
+    endif()
+    if(NOT _cdl_OUTPUT_DIR)
+        message(FATAL_ERROR "hipblaslt_create_device_library: OUTPUT_DIR is required")
+    endif()
+    if(NOT HIPBLASLT_PYTHON_COMMAND)
+        message(FATAL_ERROR "hipblaslt_create_device_library: HIPBLASLT_PYTHON_COMMAND is not set")
+    endif()
+    if(NOT HIPBLASLT_CODEGEN_ROOT)
+        message(FATAL_ERROR "hipblaslt_create_device_library: HIPBLASLT_CODEGEN_ROOT is not set")
+    endif()
+
+    if(NOT _cdl_TARGET)
+        set(_cdl_TARGET "tensilelite-device-libraries")
+    endif()
+    if(NOT _cdl_ARCHES)
+        set(_cdl_ARCHES ${GPU_TARGETS})
+    endif()
+    if(NOT _cdl_ARCHES)
+        message(FATAL_ERROR "hipblaslt_create_device_library: no ARCHES given and GPU_TARGETS is empty")
+    endif()
+    if(NOT _cdl_CXX_COMPILER)
+        set(_cdl_CXX_COMPILER "${CMAKE_CXX_COMPILER}")
+    endif()
+    if(NOT DEFINED _cdl_LAZY_LOAD)
+        set(_cdl_LAZY_LOAD ON)
+    endif()
+
+    file(MAKE_DIRECTORY "${_cdl_OUTPUT_DIR}/library")
+
+    list(JOIN _cdl_ARCHES "$<SEMICOLON>" _arches_semi)
+    set(_opts_list "--architecture=${_arches_semi}" "--cxx-compiler=${_cdl_CXX_COMPILER}")
+    if(_cdl_OFFLOAD_BUNDLER)
+        list(APPEND _opts_list "--offload-bundler=${_cdl_OFFLOAD_BUNDLER}")
+    endif()
+    if(_cdl_ASAN)
+        list(APPEND _opts_list "--address-sanitizer")
+    endif()
+    if(_cdl_JOBS)
+        list(APPEND _opts_list "--jobs=${_cdl_JOBS}")
+    endif()
+    if(_cdl_KEEP_BUILD_TMP)
+        list(APPEND _opts_list "--keep-build-tmp")
+    endif()
+    if(_cdl_ASM_DEBUG)
+        list(APPEND _opts_list "--asm-debug")
+    endif()
+    if(_cdl_YAML_FORMAT)
+        list(APPEND _opts_list "--library-format=yaml")
+    endif()
+    if(_cdl_LOGIC_FILTER)
+        list(APPEND _opts_list "--logic-filter=${_cdl_LOGIC_FILTER}")
+    endif()
+    if(_cdl_NO_COMPRESS)
+        list(APPEND _opts_list "--no-compress")
+    endif()
+    if(_cdl_EXPERIMENTAL)
+        list(APPEND _opts_list "--experimental")
+    endif()
+    if(NOT _cdl_LAZY_LOAD)
+        list(APPEND _opts_list "--no-lazy-library-loading")
+    endif()
+    if(NOT _cdl_ASM_COMMENTS)
+        list(APPEND _opts_list "--disable-asm-comments")
+    endif()
+
+    set(_known_bugs "${HIPBLASLT_CODEGEN_ROOT}/Tensile/TensileLogic/known_bugs.yaml")
+    set(_logic_stamp "${CMAKE_CURRENT_BINARY_DIR}/${_cdl_TARGET}-TensileLogic.stamp")
+    add_custom_command(
+        OUTPUT "${_logic_stamp}"
+        COMMENT "Validating library logic (TensileLogic --check-all) for ${_cdl_TARGET} ..."
+        COMMAND ${HIPBLASLT_PYTHON_COMMAND}
+            "${HIPBLASLT_CODEGEN_ROOT}/Tensile/bin/TensileLogic"
+            "${_cdl_LOGIC_PATH}"
+            --known-bugs
+            "${_known_bugs}"
+            --check-all
+        COMMAND ${CMAKE_COMMAND} -E touch "${_logic_stamp}"
+        DEPENDS ${HIPBLASLT_PYTHON_DEPS} "${_known_bugs}"
+        VERBATIM
+        USES_TERMINAL
+    )
+
+    set(_output_stamp "${CMAKE_CURRENT_BINARY_DIR}/${_cdl_TARGET}.stamp")
+    set(_tcl_command
+        ${HIPBLASLT_PYTHON_COMMAND} -m Tensile.TensileCreateLibrary
+        ${_opts_list}
+        "${_cdl_LOGIC_PATH}"
+        "${_cdl_OUTPUT_DIR}"
+        HIP
+    )
+    add_custom_command(
+        OUTPUT "${_output_stamp}"
+        COMMENT "Building device libraries to ${_cdl_OUTPUT_DIR} ..."
+        COMMAND ${_tcl_command}
+        COMMAND ${CMAKE_COMMAND} -E touch "${_output_stamp}"
+        DEPENDS ${HIPBLASLT_PYTHON_DEPS} "${_logic_stamp}"
+        VERBATIM
+        USES_TERMINAL
+    )
+
+    block(SCOPE_FOR VARIABLES)
+        list(JOIN _tcl_command " " _formatted_tcl)
+        message(STATUS "Device lib build command (${_cdl_TARGET}): ${_formatted_tcl}")
+    endblock()
+
+    add_custom_target(${_cdl_TARGET} ALL
+        DEPENDS "${_output_stamp}"
+    )
+endfunction()

projects/hipblaslt/cmake/hipblaslt-config.cmake.in

@@ -18,3 +18,27 @@ block(SCOPE_FOR VARIABLES)
         message(FATAL_ERROR "Do not export targets with hip::device as an interface link library")
     endif()
 endblock()
+
+if(EXISTS "${CMAKE_CURRENT_LIST_DIR}/HipBLASLtCodegen.cmake")
+    if(NOT HIPBLASLT_PYTHON_COMMAND OR NOT HIPBLASLT_CODEGEN_ROOT)
+        find_package(Python3 QUIET COMPONENTS Interpreter)
+    endif()
+    if(NOT HIPBLASLT_PYTHON_COMMAND AND Python3_EXECUTABLE)
+        set(HIPBLASLT_PYTHON_COMMAND "${Python3_EXECUTABLE}")
+    endif()
+    if(NOT HIPBLASLT_CODEGEN_ROOT AND Python3_EXECUTABLE)
+        execute_process(
+            COMMAND "${Python3_EXECUTABLE}" -c "import os, Tensile; print(os.path.dirname(Tensile.__path__[0]))"
+            OUTPUT_VARIABLE _hipblaslt_codegen_root
+            OUTPUT_STRIP_TRAILING_WHITESPACE
+            ERROR_QUIET
+            RESULT_VARIABLE _hipblaslt_codegen_rc
+        )
+        if(_hipblaslt_codegen_rc EQUAL 0 AND _hipblaslt_codegen_root)
+            set(HIPBLASLT_CODEGEN_ROOT "${_hipblaslt_codegen_root}")
+        endif()
+        unset(_hipblaslt_codegen_root)
+        unset(_hipblaslt_codegen_rc)
+    endif()
+    include("${CMAKE_CURRENT_LIST_DIR}/HipBLASLtCodegen.cmake")
+endif()

projects/hipblaslt/cmake/hipblaslt_python.cmake

@@ -24,7 +24,7 @@ function(hipblaslt_configure_bundled_python_command python_binary_dir asan_optio
     endif()
     set(_python_path
         "${python_binary_dir}"
-        "${hipblaslt_SOURCE_DIR}/tensilelite"
+        "${HIPBLASLT_CODEGEN_ROOT}"
     )
     list(JOIN _python_path "${_ds}" _python_path)
 

projects/hipblaslt/device-library/CMakeLists.txt

@@ -1,6 +1,8 @@
 # Copyright Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier:  MIT
 
+include(HipBLASLtCodegen)
+
 set(TENSILELITE_BUILD_PARALLEL_LEVEL "" CACHE STRING "Number of CPU cores to use for building device libraries (will use nproc if unset).")
 set(TENSILELITE_KEEP_BUILD_TMP OFF CACHE STRING "Keep temporary build directory for device libraries (turning this ON  bloat the build size).")
 set(TENSILELITE_ASM_DEBUG "" CACHE STRING "Keep debug information for built code objects.")
@@ -12,104 +14,28 @@ set(TENSILELITE_OFFLOADBUNDLER "" CACHE STRING "Path to clang-offload-bundler.")
 
 set(HIPBLASLT_LIBLOGIC_PATH "" CACHE STRING "Path to library logic files (will use 'library' if unset).")
 set(HIPBLASLT_TENSILE_LIBPATH "${hipblaslt_BINARY_DIR}/Tensile" CACHE STRING "Path to output the device gemm libraries.")
-file(MAKE_DIRECTORY "${HIPBLASLT_TENSILE_LIBPATH}/library")
-
-list(JOIN GPU_TARGETS "$<SEMICOLON>" TENSILELITE_GPU_TARGETS_SEMI_ESCAPED)
-
-set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--architecture=${TENSILELITE_GPU_TARGETS_SEMI_ESCAPED}")
-set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--cxx-compiler=${CMAKE_CXX_COMPILER}")
-if(TENSILELITE_OFFLOADBUNDLER)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--offload-bundler=${TENSILELITE_OFFLOADBUNDLER}")
-endif()
-if(HIPBLASLT_ENABLE_ASAN)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--address-sanitizer")
-endif()
-if(TENSILELITE_BUILD_PARALLEL_LEVEL)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--jobs=${TENSILELITE_BUILD_PARALLEL_LEVEL}")
-endif()
-if(TENSILELITE_KEEP_BUILD_TMP)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--keep-build-tmp")
-endif()
-if(TENSILELITE_ASM_DEBUG)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--asm-debug")
-endif()
-if(HIPBLASLT_ENABLE_YAML)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--library-format=yaml")
-endif()
-if(TENSILELITE_LOGIC_FILTER)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--logic-filter=${TENSILELITE_LOGIC_FILTER}")
-endif()
-if(TENSILELITE_NO_COMPRESS)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--no-compress")
-endif()
-if(TENSILELITE_EXPERIMENTAL)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--experimental")
-endif()
-if(NOT HIPBLASLT_ENABLE_LAZY_LOAD)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--no-lazy-library-loading")
-endif()
-if(NOT TENSILELITE_ENABLE_ASM_COMMENTS)
-    set(TENSILELITE_BUILD_OPTS ${TENSILELITE_BUILD_OPTS} "--disable-asm-comments")
-endif()
 
 if(NOT HIPBLASLT_LIBLOGIC_PATH)
     set(HIPBLASLT_LIBLOGIC_PATH "${hipblaslt_SOURCE_DIR}/library")
 endif()
 
-# Pre-build gate: validate all library logic YAMLs (WorkGroup, MatrixInstruction,
-# WorkGroupMappingXCC vs CU count, etc.) before generating .dat files. Fails build
-# if any solution fails validation so bad logic is never compiled.
-set(HIPBLASLT_TENSILELOGIC_KNOWN_BUGS
-    "${hipblaslt_SOURCE_DIR}/tensilelite/Tensile/TensileLogic/known_bugs.yaml"
-)
-# Stamp DEPENDS include known_bugs.yaml but not every library logic YAML (thousands
-# of files; CONFIGURE_DEPENDS globs are costly). After editing logic only, run
-# scripts/run_tensile_logic_check.py or touch the stamp input to force re-validation.
-set(TENSILELOGIC_STAMP "${CMAKE_CURRENT_BINARY_DIR}/TensileLogic.stamp")
-add_custom_command(
-    OUTPUT "${TENSILELOGIC_STAMP}"
-    COMMENT "Validating library logic (TensileLogic --check-all) ..."
-    COMMAND ${HIPBLASLT_PYTHON_COMMAND}
-        "${hipblaslt_SOURCE_DIR}/tensilelite/Tensile/bin/TensileLogic"
-        "${HIPBLASLT_LIBLOGIC_PATH}"
-        --known-bugs
-        "${HIPBLASLT_TENSILELOGIC_KNOWN_BUGS}"
-        --check-all
-    COMMAND ${CMAKE_COMMAND} -E touch "${TENSILELOGIC_STAMP}"
-    DEPENDS ${HIPBLASLT_PYTHON_DEPS} "${HIPBLASLT_TENSILELOGIC_KNOWN_BUGS}"
-    VERBATIM
-    USES_TERMINAL
-)
-
-set(output_stamp "${CMAKE_CURRENT_BINARY_DIR}/Tensile.stamp")
-set(TENSILE_CREATE_LIBRARY_COMMAND
-    ${HIPBLASLT_PYTHON_COMMAND} -m Tensile.TensileCreateLibrary
-    ${TENSILELITE_BUILD_OPTS}
-    ${HIPBLASLT_LIBLOGIC_PATH}
-    "${HIPBLASLT_TENSILE_LIBPATH}"
-    HIP
-)
-
-add_custom_command(
-    OUTPUT "${output_stamp}"
-    COMMENT "Building device libraries to ${HIPBLASLT_TENSILE_LIBPATH} ..."
-    COMMAND ${TENSILE_CREATE_LIBRARY_COMMAND}
-    COMMAND ${CMAKE_COMMAND} -E touch "${output_stamp}"
-    DEPENDS ${HIPBLASLT_PYTHON_DEPS} "${TENSILELOGIC_STAMP}"
-    # Because the command can contain special characters
-    VERBATIM
-    # Because this can be very long running and difficult to debug deadlocks
-    # without streaming.
-    USES_TERMINAL
-)
-
-block(SCOPE_FOR VARIABLES)
-    list(JOIN TENSILE_CREATE_LIBRARY_COMMAND " " FORMATTED_TCL)
-    message(STATUS "Device lib build command: ${FORMATTED_TCL}")
-endblock()
-
-add_custom_target(tensilelite-device-libraries ALL
-    DEPENDS "${output_stamp}"
+hipblaslt_create_device_library(
+    TARGET tensilelite-device-libraries
+    LOGIC_PATH "${HIPBLASLT_LIBLOGIC_PATH}"
+    OUTPUT_DIR "${HIPBLASLT_TENSILE_LIBPATH}"
+    ARCHES ${GPU_TARGETS}
+    CXX_COMPILER "${CMAKE_CXX_COMPILER}"
+    OFFLOAD_BUNDLER "${TENSILELITE_OFFLOADBUNDLER}"
+    JOBS "${TENSILELITE_BUILD_PARALLEL_LEVEL}"
+    LOGIC_FILTER "${TENSILELITE_LOGIC_FILTER}"
+    ASAN "${HIPBLASLT_ENABLE_ASAN}"
+    YAML_FORMAT "${HIPBLASLT_ENABLE_YAML}"
+    NO_COMPRESS "${TENSILELITE_NO_COMPRESS}"
+    EXPERIMENTAL "${TENSILELITE_EXPERIMENTAL}"
+    LAZY_LOAD "${HIPBLASLT_ENABLE_LAZY_LOAD}"
+    ASM_COMMENTS "${TENSILELITE_ENABLE_ASM_COMMENTS}"
+    KEEP_BUILD_TMP "${TENSILELITE_KEEP_BUILD_TMP}"
+    ASM_DEBUG "${TENSILELITE_ASM_DEBUG}"
 )
 
 if(HIPBLASLT_ENABLE_EXTOPS)