fix(tensilelite): share one Debug instance across the tensilelite-host boundary

projects/hipblaslt/.gitignore

@@ -49,6 +49,8 @@ CMakeCache.txt
 cmake_install.cmake
 *.make
 *.cmake
+# Keep authored CMake modules visible; the build directories above remain ignored.
+!**/cmake/**/*.cmake
 *.o.d
 
 # matrices

projects/hipblaslt/CMakeLists.txt

@@ -457,6 +457,35 @@ if(HIPBLASLT_ENABLE_HOST)
     add_subdirectory(library)
 endif()
 
+if(HIPBLASLT_ENABLE_HOST
+   AND HIPBLASLT_BUILD_SHARED_LIBS
+   AND TENSILELITE_BUILD_SHARED_LIBS
+   AND NOT WIN32
+   AND (HIPBLASLT_BUILD_TESTING OR TENSILELITE_BUILD_TESTING OR BUILD_TESTING))
+    get_filename_component(_hipblaslt_cxx_dir "${CMAKE_CXX_COMPILER}" REALPATH)
+    get_filename_component(_hipblaslt_cxx_dir "${_hipblaslt_cxx_dir}" DIRECTORY)
+    find_program(HIPBLASLT_LLVM_IFS llvm-ifs
+        HINTS "${_hipblaslt_cxx_dir}"
+        PATHS "${ROCM_PATH}/lib/llvm/bin" "$ENV{ROCM_PATH}/lib/llvm/bin" /opt/rocm/lib/llvm/bin
+    )
+    if(HIPBLASLT_LLVM_IFS)
+        enable_testing()
+        add_test(
+            NAME hipblaslt.symbol-contract.tensilelite-host
+            COMMAND "${CMAKE_COMMAND}"
+                "-DLLVM_IFS=${HIPBLASLT_LLVM_IFS}"
+                "-DCONSUMER=$<TARGET_FILE:hipblaslt>"
+                "-DPROVIDER=$<TARGET_FILE:tensilelite-host>"
+                "-DREQUIRED_IMPORTS=${CMAKE_CURRENT_SOURCE_DIR}/library/symbol-contract/tensilelite-host-imports.txt"
+                "-DNAMESPACE_PATTERN=11TensileLite"
+                "-DWORK_DIR=${CMAKE_CURRENT_BINARY_DIR}/symbol-contract"
+                -P "${CMAKE_CURRENT_SOURCE_DIR}/cmake/hipblaslt_check_symbol_contract.cmake"
+        )
+    else()
+        message(STATUS "llvm-ifs not found; hipblaslt.symbol-contract.tensilelite-host test disabled")
+    endif()
+endif()
+
 if(HIPBLASLT_ENABLE_DEVICE)
     add_subdirectory(device-library)
 endif()

projects/hipblaslt/cmake/hipblaslt_check_symbol_contract.cmake

@@ -0,0 +1,73 @@
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
+cmake_minimum_required(VERSION 3.25)
+
+foreach(_required_var LLVM_IFS CONSUMER PROVIDER REQUIRED_IMPORTS NAMESPACE_PATTERN WORK_DIR)
+    if(NOT DEFINED ${_required_var})
+        message(FATAL_ERROR "hipblaslt_check_symbol_contract: ${_required_var} is not set")
+    endif()
+endforeach()
+
+file(MAKE_DIRECTORY "${WORK_DIR}")
+
+function(_read_interface library out_defined out_undefined)
+    get_filename_component(_name "${library}" NAME)
+    set(_ifs "${WORK_DIR}/${_name}.ifs")
+    execute_process(
+        COMMAND "${LLVM_IFS}" --input-format=ELF "--output-ifs=${_ifs}" "${library}"
+        RESULT_VARIABLE _rc
+        ERROR_VARIABLE _err
+    )
+    if(NOT _rc EQUAL 0)
+        message(FATAL_ERROR "llvm-ifs failed on ${library}: ${_err}")
+    endif()
+
+    file(STRINGS "${_ifs}" _lines REGEX "^  - \\{ Name: ")
+    set(_defined "")
+    set(_undefined "")
+    foreach(_line IN LISTS _lines)
+        string(REGEX MATCH "Name: ([^, }]+)" _ "${_line}")
+        set(_symbol "${CMAKE_MATCH_1}")
+        if(_line MATCHES "Undefined: true")
+            if(NOT _line MATCHES "Weak: true")
+                list(APPEND _undefined "${_symbol}")
+            endif()
+        else()
+            list(APPEND _defined "${_symbol}")
+        endif()
+    endforeach()
+    set(${out_defined} "${_defined}" PARENT_SCOPE)
+    set(${out_undefined} "${_undefined}" PARENT_SCOPE)
+endfunction()
+
+_read_interface("${CONSUMER}" _consumer_defined _consumer_imports)
+_read_interface("${PROVIDER}" _provider_exports _provider_imports)
+
+file(STRINGS "${REQUIRED_IMPORTS}" _required_imports REGEX "[^ \t]")
+
+set(_violations "")
+
+foreach(_symbol IN LISTS _required_imports)
+    string(STRIP "${_symbol}" _symbol)
+    if(NOT _symbol IN_LIST _consumer_imports)
+        list(APPEND _violations "required import missing from ${CONSUMER}: ${_symbol}")
+    endif()
+    if(NOT _symbol IN_LIST _provider_exports)
+        list(APPEND _violations "required import not exported by ${PROVIDER}: ${_symbol}")
+    endif()
+endforeach()
+
+foreach(_symbol IN LISTS _consumer_imports)
+    if(_symbol MATCHES "${NAMESPACE_PATTERN}" AND NOT _symbol IN_LIST _provider_exports)
+        list(APPEND _violations "import not exported by ${PROVIDER}: ${_symbol}")
+    endif()
+endforeach()
+
+if(_violations)
+    list(JOIN _violations "\n  " _report)
+    message(FATAL_ERROR "Symbol contract violated:\n  ${_report}")
+endif()
+
+list(LENGTH _required_imports _required_count)
+message(STATUS "Symbol contract satisfied (${_required_count} required imports)")

projects/hipblaslt/library/symbol-contract/tensilelite-host-imports.txt

@@ -0,0 +1 @@
+_ZN11TensileLite5Debug8InstanceEv

projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp

@@ -35,18 +35,18 @@
 #include <roctracer/roctx.h>
 #endif
 
-#include <Tensile/Singleton.hpp>
-
 namespace TensileLite
 {
     using StringSet = std::set<std::string>;
 
     /**
      * @brief Common place for defining flags which enable debug behaviour.
      */
-    class TENSILELITEHOST_EXPORT Debug : public LazySingleton<Debug>
+    class TENSILELITEHOST_EXPORT Debug
     {
     public:
+        static Debug& Instance();
+
         bool printPropertyEvaluation() const;
         bool printPredicateEvaluation() const;
         bool printPredicateEvaluationVerbose() const;
@@ -168,8 +168,6 @@ namespace TensileLite
         }
 
     private:
-        friend LazySingleton<Debug>;
-
         int         m_value;
         int         m_value2;
         bool        m_naivePropertySearch = false;

projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp

@@ -70,9 +70,9 @@ namespace TensileLite
  */
     struct KernelLanguageTypeInfo
     {
-        static KernelLanguageTypeInfo const& Get(int index);
-        static KernelLanguageTypeInfo const& Get(KernelLanguage t);
-        static KernelLanguageTypeInfo const& Get(std::string const& str);
+        TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(int index);
+        TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(KernelLanguage t);
+        TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(std::string const& str);
 
         KernelLanguage m_kernelLanguage;
         std::string    name;

projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp

@@ -59,11 +59,13 @@ namespace TensileLite
             float                     waveScale;
         };
 
-        float tilesPerCU(ContractionProblemGemm const&    problem,
-                         CUGranularityScaleFactors const& cuFactors);
+        TENSILELITEHOST_EXPORT float tilesPerCU(ContractionProblemGemm const&    problem,
+                                                CUGranularityScaleFactors const& cuFactors);
 
-        std::ostream& operator<<(std::ostream& stream, CUGranularityScaleFactors const& cugsf);
-        std::ostream& operator<<(std::ostream& stream, WaveGranularityScaleFactors const& wgsf);
+        TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream&                    stream,
+                                                        CUGranularityScaleFactors const& cugsf);
+        TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream&                      stream,
+                                                        WaveGranularityScaleFactors const& wgsf);
 
         /**
          * @brief A Property whose value is of type `float`.

projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp

@@ -73,7 +73,8 @@ namespace TensileLite
         struct DenseLayer
         {
             DenseLayer() = default;
-            DenseLayer(const std::vector<float>& weights, const std::vector<float>& bias);
+            TENSILELITEHOST_EXPORT DenseLayer(const std::vector<float>& weights,
+                                              const std::vector<float>& bias);
 
             std::vector<dtype> operator()(const std::vector<dtype>& F) const;
 
@@ -101,9 +102,10 @@ namespace TensileLite
         {
             MLPNet() = default;
 
-            std::vector<dtype> predict(std::vector<float> const& probkey) const;
+            TENSILELITEHOST_EXPORT std::vector<dtype>
+                                   predict(std::vector<float> const& probkey) const;
 
-            bool valid(bool verbose = false) const;
+            TENSILELITEHOST_EXPORT bool valid(bool verbose = false) const;
 
             std::string description() const
             {

projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp

@@ -72,9 +72,9 @@ namespace TensileLite
  */
     struct ScalarValueTypeInfo
     {
-        static ScalarValueTypeInfo const& Get(int index);
-        static ScalarValueTypeInfo const& Get(ScalarValue t);
-        static ScalarValueTypeInfo const& Get(std::string const& str);
+        TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(int index);
+        TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(ScalarValue t);
+        TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(std::string const& str);
 
         ScalarValue m_value;
         std::string name;

projects/hipblaslt/tensilelite/src/Debug.cpp

@@ -210,6 +210,13 @@ namespace TensileLite
         m_excludedFromGetAll = excludedSet;
     }
 
+    Debug& Debug::Instance()
+    {
+        static Debug instance;
+
+        return instance;
+    }
+
     Debug::Debug()
         : m_value(DEBUG_SM)
         , m_value2(DEBUG_SM2)

projects/hipblaslt/tensilelite/tests/CMakeLists.txt

@@ -100,6 +100,11 @@ endif()
 
 target_link_libraries(tensilelite-tests PUBLIC GTest::gtest)
 
+if(HIPBLASLT_ENABLE_YAML)
+    target_link_libraries(tensilelite-tests PRIVATE LLVMObjectYAML)
+    target_include_directories(tensilelite-tests PRIVATE ${LLVM_INCLUDE_DIRS})
+endif()
+
 # LaunchDimGuard_test asserts different outcomes for a cluster launch depending
 # on whether SolutionAdapter::launchKernel was compiled with the cluster path,
 # so it needs the same answer the library got. Reuses the cached result of the