Repository navigation
fix(tensilelite): share one Debug instance across the tensilelite-host boundary - #12784
Merged
Merged
Conversation
…t boundary
JIRA ID : ROCM-31245
Debug inherited Instance() from LazySingleton, an inline member of a
header-only template. With CXX_VISIBILITY_PRESET hidden and
VISIBILITY_INLINES_HIDDEN, a shared libtensilelite-host and libhipblaslt
each got a private copy of the function-local static. libhipblaslt wrote
the {GridBasedMatching, PredictionMatching} exclusion set into its copy
before getAllSolutions(); ExactLogicLibrary::findAllSolutions() read the
other, empty copy and traversed the ranking libraries on every call. That
is the 100x-200x host-side slowdown in hipblasLtMatmulAlgoGetHeuristic for
small GEMMs.
Define Debug::Instance() out of line in Debug.cpp so the single exported
definition owns the only instance, and drop the LazySingleton base so no
inline accessor remains that could reintroduce a second copy. The
LazySingleton template is unchanged, which avoids exporting a template
member (a dllimport hazard for consumer-only specializations on Windows).
Also make tensilelite-tests link as a shared build: export the
ScalarValueTypeInfo / KernelLanguageTypeInfo Get() overloads and the
MLFeatures / MLPClassification entry points the tests call, and link
LLVMObjectYAML into the tests in YAML builds.
SingletonIdentity_test checks instance identity and state propagation
across the library boundary through a test-only probe that is not in the
installed headers, and skips when tensilelite-host is static.
Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (49.68%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #12784 +/- ##
===========================================
- Coverage 72.37% 71.02% -1.36%
===========================================
Files 2879 2953 +74
Lines 478194 488371 +10177
Branches 70867 70671 -196
===========================================
+ Hits 346082 346828 +746
- Misses 108166 117472 +9306
- Partials 23946 24071 +125
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
1 task done
Co-authored-by: Cursor <cursoragent@cursor.com>
Alex-Vasile
reviewed
Sep 30, 2026
…tract check Drop debugInstanceAddress() and SingletonIdentity_test.cpp so no test-only code is built into libtensilelite-host. Add hipblaslt.symbol-contract.tensilelite-host, a CTest test for shared, non-Windows builds. It runs llvm-ifs on libhipblaslt and libtensilelite-host and fails if libhipblaslt does not import, or libtensilelite-host does not export, every symbol listed in library/symbol-contract/tensilelite-host-imports.txt, or if any TensileLite symbol libhipblaslt imports is not exported by libtensilelite-host. The list starts with TensileLite::Debug::Instance(); with a duplicated Debug that import disappears and the check fails. Co-authored-by: Cursor <cursoragent@cursor.com>
Alex-Vasile
approved these changes
Oct 13, 2026
talumbau
approved these changes
Oct 13, 2026
davidd-amd
enabled auto-merge (squash)
October 13, 2026 18:57
Collaborator
|
precheckin(hipblaslt) is not reporting back to the Github UI that it is done and passed. Screenshots from math-ci: 
|
talumbau
requested changes
Oct 13, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
Everything is good, but a final review shows that you are missing the file projects/hipblaslt/cmake/hipblaslt_check_symbol_contract.cmake which is added in the changes to projects/hipblaslt/CMakeLists.txt. Please add that file to the repo - thanks!
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : ROCM-31245 (also AIRADSW-923)
Motivation
Splitting hipBLASLt into
libhipblaslt.soand sharedlibtensilelite-host.socreated two process-local copies ofTensileLite::Debug.rocblaslt_auxiliary.cppwrote the{GridBasedMatching, PredictionMatching}exclusion set into the hipBLASLt copy, whileExactLogicLibrary::findAllSolutions()read the empty host-library copy and traversed the ranking libraries on every call.This raised
hipblasLtMatmulAlgoGetHeuristichost time on small GEMMs from roughly 0.03–0.07 ms to about 5 ms. The effect is most visible on gfx1200/gfx1201 shapes that reach thegetAllSolutions()fallback.This PR carries the symbol-related correction from #11946 and #12465. The unrelated StreamK work from #11946 remains separate.
Technical Details
One
Debuginstance across both shared librariesDebuginheritedInstance()from the header-onlyLazySingletontemplate. With hidden symbol and inline visibility, each shared object owned a separate function-local static.Debug::Instance()is now declared on the exportedDebugclass and defined once inDebug.cpp.Debugno longer derives fromLazySingleton, so no inline accessor can create another private instance.LazySingletonremains unchanged. This avoids exporting a template member that would imposedllimportrequirements on consumer-only specializations such asLazySingleton<TestData>.Complete the installed shared-library interface
The hipBLASLt development package installs
roc::tensilelite-hostas a shared imported CMake target and publishes the TensileLite headers astensilelite_public_headers. A consumer compiling inline or template code from those headers must be able to resolve its out-of-line calls fromlibtensilelite-host.so.This PR exports:
ScalarValueTypeInfo::GetandKernelLanguageTypeInfo::Getoverload families, matching the existing exportedDataTypeInfometadata interface;MLFeatures::tilesPerCUand the stream operators for its public scale-factor types;MLPClassification::DenseLayer::DenseLayer,MLPNet::predict, andMLPNet::valid.Six of these functions are called directly by inline/template code in shipped headers. The remaining overloads and stream operators preserve the established complete-overload and exported-value-type patterns. These are package-interface exports, not test-only hooks.
In YAML builds,
tensilelite-testsalso linksLLVMObjectYAMLdirectly.tensilelite-hostlinks that dependency privately with--exclude-libs,ALL, so the test executable cannot rely on the host library to re-export LLVM symbols.Check the binary relationship without test-only production code
For shared non-Windows test builds,
hipblaslt.symbol-contract.tensilelite-hostrunsllvm-ifsover both libraries. It verifies that:libhipblaslt.soimports every symbol listed inlibrary/symbol-contract/tensilelite-host-imports.txt, beginning withTensileLite::Debug::Instance();libtensilelite-host.soexports those required imports;libhipblaslt.sois exported by its provider.The check adds no test-only function to the production library.
Test Plan
tensilelite-testswith YAML serialization to exercise the installed-header exports and direct LLVM YAML dependency.libhipblaslt.sowith the normal msgpack serialization path and run the symbol-contract test.develop.Test Result
hipblaslt_check_symbol_contract.cmakefailed against TheRock run 37969613509.libhipblaslt.so.1.5did not importDebug::Instance(), andlibtensilelite-host.so.1.0did not export it.64d7532a: shared YAMLtensilelite-testsbuilt successfully for gfx950 in a host-only configuration.64d7532a: shared msgpacklibhipblaslt.sobuilt successfully;hipblaslt.symbol-contract.tensilelite-hostpassed 1/1.roc::tensilelite-hostheader interface.developmerge and are pending.Exploratory result: combining
HIPBLASLT_ENABLE_HOST=ONwithHIPBLASLT_ENABLE_YAML=ONmade the broad namespace check report three pre-existing msgpack imports that a YAML host library does not export (Serialization::objectToMap,fileToMsgObject, andLoadLibraryMapping). The supported msgpack contract above passes; this mixed serialization configuration is not counted as passing evidence for this PR.Not covered locally:
Follow-ups
Debugobject also restores the previous shared mutable exclusion set. The two heuristic entry points use different mutexes, and the get-all path takes none. A later design should pass exclusions per call throughfindAllSolutionsorSolutionLibrarySearchType.EmbeddedDataandEmbeddedLibraryretain the same latent header-only singleton shape, but no cross-library writer/reader pair has been identified.Submission Checklist
Risk level
Low (2/5). The runtime change moves one singleton accessor to a single exported definition. The additional exports complete an already-installed C++ target and are additive; no kernel, device-library, or dispatch algorithm changes. Windows validation and refreshed automated checks remain pending.