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