fix(tensilelite): export LazySingleton::Instance so the Debug singleton is not duplicated per DSO

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

@@ -191,5 +191,17 @@ namespace TensileLite
 
         Debug();
     };
+
+    /**
+     * @brief Address of the Debug singleton as seen from inside the TensileLite
+     *        host library itself.
+     *
+     * Deliberately non-inline (defined in Debug.cpp) so a caller in another
+     * link unit -- libhipblaslt, or a test binary -- can compare it against the
+     * address it gets from its own `Debug::Instance()` call. If the two differ,
+     * the singleton has been duplicated across the DSO boundary and every piece
+     * of state routed through it silently stops propagating. See ROCM-31245.
+     */
+    TENSILELITEHOST_EXPORT const void* debugInstanceAddress();
 } // namespace TensileLite
 

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

@@ -36,7 +36,14 @@ namespace TensileLite
     class LazySingleton
     {
     public:
-        static Class& Instance()
+        // TENSILELITEHOST_EXPORT is required, not decorative. `Instance()` is an
+        // implicitly-inline member of a header-only template, and both
+        // libhipblaslt and libtensilelite-host are built with hidden visibility
+        // plus VISIBILITY_INLINES_HIDDEN. Without an explicit default-visibility
+        // attribute each shared object gets its OWN copy of `instance`, so state
+        // written through the singleton in one library is invisible in the other
+        // -- silently, with no link error. See ROCM-31245.
+        static TENSILELITEHOST_EXPORT Class& Instance()
         {
             static Class instance;
 

projects/hipblaslt/tensilelite/src/Debug.cpp

@@ -307,4 +307,9 @@ namespace TensileLite
         }
     }
 
+    const void* debugInstanceAddress()
+    {
+        return static_cast<const void*>(&Debug::Instance());
+    }
+
 } // namespace TensileLite

projects/hipblaslt/tensilelite/tests/CMakeLists.txt

@@ -41,6 +41,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"

projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp

@@ -0,0 +1,62 @@
+// Copyright Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+
+#include <Tensile/Debug.hpp>
+
+// ROCM-31245 regression guard.
+//
+// TensileLite::Debug is a LazySingleton: `Instance()` is an implicitly-inline
+// member of a header-only class template holding a function-local static. Both
+// libtensilelite-host and its consumers (libhipblaslt, this test binary) are
+// built with -fvisibility=hidden and -fvisibility-inlines-hidden. Unless
+// `Instance()` carries an explicit default-visibility attribute, each link unit
+// gets its own private copy of the static, and there is no diagnostic of any
+// kind: no link error, no warning, no crash.
+//
+// That is exactly what happened when TensileLite was split out of libhipblaslt
+// into its own shared object. libhipblaslt wrote the
+// "skip GridBasedMatching/PredictionMatching" set into its copy of Debug;
+// ExactLogicLibrary::findAllSolutions read libtensilelite-host's copy, saw an
+// empty set, and fully traversed the expensive ranking libraries on every
+// hipblasLtMatmulAlgoGetHeuristic call -- a 100x-200x host-side regression.
+//
+// This test binary is a separate link unit from libtensilelite-host, so it
+// stands in for libhipblaslt: if the singleton can duplicate across a library
+// boundary at all, it duplicates here too.
+namespace
+{
+    TEST(SingletonIdentity, DebugSingletonIsSharedAcrossLinkUnits)
+    {
+        const void* fromThisLinkUnit
+            = static_cast<const void*>(&TensileLite::Debug::Instance());
+        const void* fromTensileLiteHost = TensileLite::debugInstanceAddress();
+
+        EXPECT_EQ(fromThisLinkUnit, fromTensileLiteHost)
+            << "TensileLite::Debug has been duplicated across a link-unit boundary. "
+               "Check that LazySingleton::Instance() still carries TENSILELITEHOST_EXPORT "
+               "(Tensile/Singleton.hpp). See ROCM-31245.";
+    }
+
+    // Address equality is necessary but not sufficient on its own to notice a
+    // regression -- verify that mutations actually propagate, which is the
+    // property callers depend on.
+    TEST(SingletonIdentity, DebugSingletonStatePropagatesAcrossLinkUnits)
+    {
+        const TensileLite::StringSet original
+            = TensileLite::Debug::Instance().excludedLibFromGetAll();
+
+        TensileLite::StringSet probe({"GridBasedMatching", "PredictionMatching"});
+        TensileLite::Debug::Instance().setExcludedLibFromGetAll(probe);
+
+        const auto* observed = static_cast<const TensileLite::Debug*>(
+            TensileLite::debugInstanceAddress());
+        EXPECT_EQ(observed->excludedLibFromGetAll().size(), 2u)
+            << "State written through Debug::Instance() in this link unit is not visible "
+               "to libtensilelite-host. See ROCM-31245.";
+
+        TensileLite::StringSet restore(original);
+        TensileLite::Debug::Instance().setExcludedLibFromGetAll(restore);
+    }
+} // namespace