From 9df96b1eecd7bcef2424979ed990f4862bcc79bb Mon Sep 17 00:00:00 2001 From: "T.J. Alumbaugh" Date: Wed, 23 Sep 2026 00:26:52 +0000 Subject: [PATCH] fix(tensilelite): export LazySingleton::Instance so the Debug singleton is not duplicated per DSO JIRA ID : ROCM-31245 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 libhipblaslt are built with -fvisibility=hidden and -fvisibility-inlines-hidden. While TensileLite was a static archive linked into libhipblaslt there was one link unit and therefore one instance. Once TENSILELITE_BUILD_SHARED_LIBS=ON started shipping libtensilelite-host.so as its own DSO, each shared object got its own private copy of the static -- silently: no link error, no warning, no crash. The consequence is a 100x-200x host-side regression in hipblasLtMatmulAlgoGetHeuristic. rocblaslt_auxiliary.cpp writes the {GridBasedMatching, PredictionMatching} skip set into libhipblaslt's copy of Debug before calling getAllSolutions; ExactLogicLibrary::findAllSolutions reads libtensilelite-host's copy, sees an empty set, and fully traverses those expensive ranking sub-libraries on every call. There is no memoization on that path, so the cost is paid per call. Marking Instance() TENSILELITEHOST_EXPORT gives the function and its local static default visibility in both objects, so the dynamic linker collapses them to one definition. Singleton.hpp already included without using it. The same fix covers the other LazySingleton derivatives (EmbeddedLibrary, EmbeddedData), which have the same latent exposure. Adds TensileLite::debugInstanceAddress() -- a deliberately non-inline accessor in Debug.cpp -- plus SingletonIdentity_test, which compares the singleton address and state as seen from the test binary against the one inside libtensilelite-host. The test binary is a separate link unit, so it stands in for libhipblaslt. Both assertions fail without the visibility change. Co-Authored-By: Claude --- .../tensilelite/include/Tensile/Debug.hpp | 12 ++++ .../tensilelite/include/Tensile/Singleton.hpp | 9 ++- projects/hipblaslt/tensilelite/src/Debug.cpp | 5 ++ .../tensilelite/tests/CMakeLists.txt | 1 + .../tests/SingletonIdentity_test.cpp | 62 +++++++++++++++++++ 5 files changed, 88 insertions(+), 1 deletion(-) create mode 100644 projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp diff --git a/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp b/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp index 944c2eaaaa2c..edd7de23aa18 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp +++ b/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 diff --git a/projects/hipblaslt/tensilelite/include/Tensile/Singleton.hpp b/projects/hipblaslt/tensilelite/include/Tensile/Singleton.hpp index 40acc3cc049c..29d5bb04f461 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/Singleton.hpp +++ b/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; diff --git a/projects/hipblaslt/tensilelite/src/Debug.cpp b/projects/hipblaslt/tensilelite/src/Debug.cpp index 989b360d5062..91648481b1b4 100644 --- a/projects/hipblaslt/tensilelite/src/Debug.cpp +++ b/projects/hipblaslt/tensilelite/src/Debug.cpp @@ -307,4 +307,9 @@ namespace TensileLite } } + const void* debugInstanceAddress() + { + return static_cast(&Debug::Instance()); + } + } // namespace TensileLite diff --git a/projects/hipblaslt/tensilelite/tests/CMakeLists.txt b/projects/hipblaslt/tensilelite/tests/CMakeLists.txt index 7df6120fb2e5..fa13055eceff 100644 --- a/projects/hipblaslt/tensilelite/tests/CMakeLists.txt +++ b/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" diff --git a/projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp b/projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp new file mode 100644 index 000000000000..4f9110499d6c --- /dev/null +++ b/projects/hipblaslt/tensilelite/tests/SingletonIdentity_test.cpp @@ -0,0 +1,62 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include + +// 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(&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( + 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