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

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

@@ -205,6 +205,13 @@ namespace TensileLite
         m_excludedFromGetAll = excludedSet;
     }
 
+    Debug& Debug::Instance()
+    {
+        static Debug instance;
+
+        return instance;
+    }
+
     Debug::Debug()
         : m_value(DEBUG_SM)
         , m_value2(DEBUG_SM2)
@@ -307,4 +314,11 @@ namespace TensileLite
         }
     }
 
+    TENSILELITEHOST_EXPORT const void* debugInstanceAddress();
+
+    const void* debugInstanceAddress()
+    {
+        return &Debug::Instance();
+    }
+
 } // namespace TensileLite

projects/hipblaslt/tensilelite/tests/CMakeLists.txt

@@ -28,6 +28,7 @@ target_sources(tensilelite-tests
         "${CMAKE_CURRENT_SOURCE_DIR}/DataInit_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/ExactLogicLibrary_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/HardwareMonitor_test.cpp"
+        "${CMAKE_CURRENT_SOURCE_DIR}/include/DebugInstanceProbe.hpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/include/GEMMKernelTest.hpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/include/GEMMKernelTest_impl.hpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/include/RunGEMMKernelTest.hpp"
@@ -43,6 +44,7 @@ target_sources(tensilelite-tests
         "${CMAKE_CURRENT_SOURCE_DIR}/RangeLibrary_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/ResultComparison_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/ReferenceMXFastPath_test.cpp"
+        "${CMAKE_CURRENT_SOURCE_DIR}/SingletonIdentity_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/SolutionValidation_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/StreamKLaunchSummary_test.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/TensorDescriptor_test.cpp"
@@ -88,6 +90,16 @@ endif()
 
 target_link_libraries(tensilelite-tests PUBLIC GTest::gtest)
 
+target_compile_definitions(tensilelite-tests
+    PRIVATE
+        $<$<STREQUAL:$<TARGET_PROPERTY:tensilelite-host,TYPE>,SHARED_LIBRARY>:TENSILELITE_HOST_IS_SHARED>
+)
+
+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

projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp

@@ -0,0 +1,61 @@
+// Copyright Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+
+#include <Tensile/Debug.hpp>
+
+#include "DebugInstanceProbe.hpp"
+
+namespace
+{
+    class ExcludedLibGuard
+    {
+    public:
+        ExcludedLibGuard()
+            : m_saved(TensileLite::Debug::Instance().excludedLibFromGetAll())
+        {
+        }
+
+        ~ExcludedLibGuard()
+        {
+            TensileLite::Debug::Instance().setExcludedLibFromGetAll(m_saved);
+        }
+
+    private:
+        TensileLite::StringSet m_saved;
+    };
+
+    class SingletonIdentity : public ::testing::Test
+    {
+    protected:
+        void SetUp() override
+        {
+#ifndef TENSILELITE_HOST_IS_SHARED
+            GTEST_SKIP() << "tensilelite-host is static; there is no library boundary to test.";
+#endif
+        }
+    };
+
+    TEST_F(SingletonIdentity, DebugIsSharedAcrossLinkUnits)
+    {
+        EXPECT_EQ(static_cast<const void*>(&TensileLite::Debug::Instance()),
+                  TensileLite::debugInstanceAddress())
+            << "TensileLite::Debug is duplicated across the tensilelite-host boundary. "
+               "Debug::Instance() must stay defined out of line in Debug.cpp.";
+    }
+
+    TEST_F(SingletonIdentity, DebugStatePropagatesAcrossLinkUnits)
+    {
+        ExcludedLibGuard restore;
+
+        TensileLite::StringSet probe({"GridBasedMatching", "PredictionMatching"});
+        TensileLite::Debug::Instance().setExcludedLibFromGetAll(probe);
+
+        const auto* hostDebug
+            = static_cast<const TensileLite::Debug*>(TensileLite::debugInstanceAddress());
+        EXPECT_EQ(hostDebug->excludedLibFromGetAll(), probe)
+            << "State written through Debug::Instance() here is not visible inside "
+               "tensilelite-host.";
+    }
+} // namespace

projects/hipblaslt/tensilelite/tests/include/DebugInstanceProbe.hpp

@@ -0,0 +1,11 @@
+// Copyright Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <tensilelitehost/export.h>
+
+namespace TensileLite
+{
+    TENSILELITEHOST_EXPORT const void* debugInstanceAddress();
+} // namespace TensileLite