Repository navigation
fix(tensilelite): export LazySingleton::Instance so the Debug singleton is not duplicated per DSO - #12465
Open
talumbau wants to merge 1 commit into
Open
fix(tensilelite): export LazySingleton::Instance so the Debug singleton is not duplicated per DSO#12465talumbau wants to merge 1 commit into
talumbau wants to merge 1 commit into
Conversation
…on 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 <tensilelitehost/export.h>
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 <noreply@anthropic.com>
talumbau
enabled auto-merge (squash)
September 23, 2026 02:00
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 check has failed because the head coverage (48.80%) 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 #12465 +/- ##
===========================================
+ Coverage 70.12% 70.23% +0.11%
===========================================
Files 2933 2933
Lines 482003 482002 -1
Branches 69493 69493
===========================================
+ Hits 337964 338517 +553
+ Misses 119775 119197 -578
- Partials 24264 24288 +24
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
jinchen62
approved these changes
Sep 23, 2026
jinchen62
left a comment
jinchen62
left a comment
Contributor
There was a problem hiding this comment.
The fix looks good. Reproduced it independently with ROCm clang 22 and the same visibility flags. Windows (dllimport) and static builds are fine too.
Description fixes before merge:
- "Restores the pre-split numbers exactly" — 16x16x16 is 25.08 ms static vs 9.93 ms fixed; the fixed build is consistently faster.
- "Cliff at 1451" — 1440 is already slow, so the cliff is in (1300, 1440].
- "
V(unique global)" —Vis weak; clang emitsWEAK, notSTB_GNU_UNIQUE. - "5 symbols, all
VorW" —debugInstanceAddress()is also a new strong export.
Test (non-blocking): with tensilelite-host static there's one link unit, so both tests pass with or without the fix — worth a GTEST_SKIP there. Also compare the probe set exactly rather than by size().
Follow-up: sharing the singleton again brings back a pre-split race. The two heuristic entry points in rocblaslt_auxiliary.cpp lock different mutexes (:2328, :2590) and get-all-algos (:2461) takes none, so it can silently drop GridBased/Prediction candidates. Not a reason to hold this PR, but worth its own ticket.
Alex-Vasile
requested changes
Sep 24, 2026
Alex-Vasile
left a comment
Alex-Vasile
left a comment
Contributor
There was a problem hiding this comment.
Also, PR description needs a bit of cleanup. E.g. The cliff appears to be between 1300 and 1440, rather than at 1451.
| // 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 |
Contributor
There was a problem hiding this comment.
Won't this test pass even without the fix if we are doing a static build? Is there a way to enforce dynamic linking for it?
| * 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(); |
Contributor
There was a problem hiding this comment.
This will show up in the public headers even though it is a private function used only for testing. Can we put it in a separate header that we compile/export only when we set TENSILELITE_BUILD_TESTING=ON?
| // 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() |
Contributor
There was a problem hiding this comment.
Also from the AI:
Applying TENSILELITEHOST_EXPORT to the generic LazySingleton::Instance() makes consumer-only specializations such as LazySingleton unresolved dllimports in Windows shared builds
Member
|
Validation on gfx1201: this PR fixes the regression. Edited: an earlier version of this comment said the PR does not recover performance, based on the result in ROCM-31245. My own A/B below shows it does. Setup
Median ms per FP32
With this PR, The out-of-line |
1 task done
Contributor
|
@talumbau I've opened #12784, which brings the symbol fixes from this PR and #11946 together. It uses the out-of-line
On a shared build, the control ( This PR has auto-merge enabled, so if you agree that #12784 supersedes it, could you disable auto-merge or close this one? I haven't changed anything here. |
davidd-amd
added a commit
that referenced
this pull request
Oct 14, 2026
…t boundary (#12784) JIRA ID : ROCM-31245 (also AIRADSW-923) ## Motivation Splitting hipBLASLt into `libhipblaslt.so` and shared `libtensilelite-host.so` created two process-local copies of `TensileLite::Debug`. `rocblaslt_auxiliary.cpp` wrote the `{GridBasedMatching, PredictionMatching}` exclusion set into the hipBLASLt copy, while `ExactLogicLibrary::findAllSolutions()` read the empty host-library copy and traversed the ranking libraries on every call. This raised `hipblasLtMatmulAlgoGetHeuristic` host 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 the `getAllSolutions()` fallback. This PR carries the symbol-related correction from #11946 and #12465. The unrelated StreamK work from #11946 remains separate. ## Technical Details ### One `Debug` instance across both shared libraries `Debug` inherited `Instance()` from the header-only `LazySingleton` template. With hidden symbol and inline visibility, each shared object owned a separate function-local static. - `Debug::Instance()` is now declared on the exported `Debug` class and defined once in `Debug.cpp`. - `Debug` no longer derives from `LazySingleton`, so no inline accessor can create another private instance. - `LazySingleton` remains unchanged. This avoids exporting a template member that would impose `dllimport` requirements on consumer-only specializations such as `LazySingleton<TestData>`. ### Complete the installed shared-library interface The hipBLASLt development package installs `roc::tensilelite-host` as a shared imported CMake target and publishes the TensileLite headers as `tensilelite_public_headers`. A consumer compiling inline or template code from those headers must be able to resolve its out-of-line calls from `libtensilelite-host.so`. This PR exports: - the complete `ScalarValueTypeInfo::Get` and `KernelLanguageTypeInfo::Get` overload families, matching the existing exported `DataTypeInfo` metadata interface; - `MLFeatures::tilesPerCU` and the stream operators for its public scale-factor types; - `MLPClassification::DenseLayer::DenseLayer`, `MLPNet::predict`, and `MLPNet::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-tests` also links `LLVMObjectYAML` directly. `tensilelite-host` links 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-host` runs `llvm-ifs` over both libraries. It verifies that: - `libhipblaslt.so` imports every symbol listed in `library/symbol-contract/tensilelite-host-imports.txt`, beginning with `TensileLite::Debug::Instance()`; - `libtensilelite-host.so` exports those required imports; - every TensileLite symbol imported by `libhipblaslt.so` is exported by its provider. The check adds no test-only function to the production library. ## Test Plan 1. Run the symbol-contract check against a pre-change TheRock artifact pair and confirm that it rejects the duplicated inline singleton. 2. Build shared `tensilelite-tests` with YAML serialization to exercise the installed-header exports and direct LLVM YAML dependency. 3. Build shared `libhipblaslt.so` with the normal msgpack serialization path and run the symbol-contract test. 4. Let the refreshed automated GitHub checks cover the shared CI lanes after merging current `develop`. ## Test Result - Pre-change control: `hipblaslt_check_symbol_contract.cmake` failed against TheRock run [37969613509](https://github.com/ROCm/TheRock/actions/runs/37969613509). `libhipblaslt.so.1.5` did not import `Debug::Instance()`, and `libtensilelite-host.so.1.0` did not export it. - Merged head `64d7532a`: shared YAML `tensilelite-tests` built successfully for gfx950 in a host-only configuration. - Merged head `64d7532a`: shared msgpack `libhipblaslt.so` built successfully; `hipblaslt.symbol-contract.tensilelite-host` passed 1/1. - Artifact audit: the scanned rocBLAS and MIOpen libraries consume hipBLASLt APIs but do not directly import the added TensileLite type/ML/MLP symbols. Their justification is the installed `roc::tensilelite-host` header interface. - Automated GitHub checks restarted after the `develop` merge and are pending. Exploratory result: combining `HIPBLASLT_ENABLE_HOST=ON` with `HIPBLASLT_ENABLE_YAML=ON` made the broad namespace check report three pre-existing msgpack imports that a YAML host library does not export (`Serialization::objectToMap`, `fileToMsgObject`, and `LoadLibraryMapping`). The supported msgpack contract above passes; this mixed serialization configuration is not counted as passing evidence for this PR. Not covered locally: - Windows shared-library behavior; automated coverage is pending. - End-to-end timing on gfx1200/gfx1201. In #12465, the same out-of-line mechanism measured 0.079 / 0.094 / 0.067 ms for 16x16x16 / 64x32x32 / 16x8x4 on gfx1201, compared with 5.16 / 5.01 / 4.99 ms before the correction. ## Follow-ups - Restoring one process-wide `Debug` object 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 through `findAllSolutions` or `SolutionLibrarySearchType`. - `EmbeddedData` and `EmbeddedLibrary` retain the same latent header-only singleton shape, but no cross-library writer/reader pair has been identified. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## 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. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
Closed
1 task done
1 task done
newling
disabled auto-merge
October 28, 2026 19:27
newling
enabled auto-merge (squash)
October 28, 2026 19:44
jaopaulolc
added a commit
that referenced
this pull request
Nov 6, 2026
… state (#14151) JIRA ID : ROCM-32224, ROCM-31982 ## Motivation hipBLASLt picks a GPU kernel (a "solution") for a matrix multiply by searching solution libraries. Every hipBLASLt handle in a process shares one static `TensileHost`, so two threads that each own a handle (one handle per stream, the pattern the header recommends) reach the same library objects. ROCM-32224 reports that this sharing causes data races. Under concurrent use they can return too few solutions, cache a wrong answer for the rest of the process, or read a half-written value. This PR started as ROCM-31982 (the first race, M3 in ROCM-32224) and now also covers two more races from the same ticket (H2 and M7) and two smaller ones found while fixing them. M3, the get-all ranking exclusion: When a heuristic query returns fewer solutions than requested, `rocblaslt_auxiliary.cpp` calls `getAllSolutions` to fill the remainder. It skips the two ranking libraries (`GridBasedMatching`, `PredictionMatching`), because their candidates were already returned. It did this by writing a set of names into a process-global member of `TensileLite::Debug`, calling `getAllSolutions`, and clearing the set afterwards. That caused three problems: 1. The setter copy-assigned a `std::set` while the getter copied it, with no common lock. This is a C++ data race. The two heuristic paths each locked a different function-local mutex, and the standalone get-all path took no lock. 2. While one call had the exclusions installed, a concurrent `getAllSolutions` on another handle also skipped GridBased/Prediction candidates and returned fewer solutions. 3. The clearing step was not RAII (automatic release when a scope ends), and both writes sat inside a `try` that rethrows `rocblaslt_status`. One exception left the exclusions installed for every later get-all in the process. Commit 091eb4d (shared `Debug` singleton, with the symbol correction from #12465) is what made the writer and the reader meet across the `libhipblaslt` / `libtensilelite-host` boundary. This work was listed as a dependency of #12465. The optimization being kept comes from #2757, per the ticket; I did not verify that independently. H2, M7, and the two smaller races: - H2: `PlaceholderLibrary` (a library shard that loads its solutions on first use) tested its `library` pointer outside the lock that guards the load, so a thread could read the pointer while another thread assigned it. - M7: `findTopSolutions()` reported whether it had returned every solution through one flag shared by all callers. Another thread could overwrite the flag before the caller read it, and `CachingLibrary` then cached the wrong answer for a problem. - `CachingLibrary` also stored the solutions and the "returned all" bit in two maps with no common lock. Two requests for the same problem with different `numSolutions` could leave 3 solutions marked as all of them, and a later request for 10 got 3. - `AMDGPU::isStandardCU()` stored 1 before it could store 0. A second thread could read the transient 1, and `ContractionSolution::isFallbackForHW` would cache the wrong answer permanently. ## Technical Details No lock is added around `hipblasLtMatmul`, `getSolutions`, or `TensileHost` dispatch. A coarse lock would serialize every dispatch, and each race above has a narrower fix. M3: The exclusion now travels as a per-call value on the existing `SolutionLibrarySearchType` argument, which `getAllSolutions` already forwards to every `findAllSolutions` override. - New enumerator `GEMM_TYPE_ONLY_NO_RANKING` (`COUNT` goes from 4 to 5). It filters like `GEMM_TYPE_ONLY`, except `ExactLogicLibrary::findAllSolutions` also skips rows whose type is `GridBasedMatching` or `PredictionMatching`, under the same `isStandardCU()` check as before. - New helper `isGemmTypeOnlySearch()`, so sites that compared `== GEMM_TYPE_ONLY` accept both values. - `getAllSolutions` (three overloads in `tensile_host.hpp`/`.cpp`) gains a trailing `searchType` parameter that defaults to `GEMM_TYPE_ONLY`. Existing callers behave as before; the two heuristic fallbacks in `rocblaslt_auxiliary.cpp` pass `GEMM_TYPE_ONLY_NO_RANKING`. - Deleted: the `m_excludedFromGetAll` member, its two accessors, both function-local mutexes, and both reset calls. H2 and M7: - `PlaceholderLibrary::loadPlaceholderLibrary()` loads the shard under `std::call_once` and returns a copy of `library`; every call site searches through that copy. The load builds what it registers in locals and merges them into the master before storing `library`. So `library` is never visible before the shard is registered, and a load that throws leaves the master unchanged for the retry. - `findTopSolutions()` keeps its `SolutionVector` return type. The shared `lastFindTopRetAll` flag and `lastFindTopAlreadyRetAll()` are removed from the eight library headers that had them. "Returned everything" means "fewer than requested", which `CachingLibrary` computes from its own call. - `CachingLibrary` keeps the solutions and the bit in one `CacheMap` entry (`TopSolutionsEntry`), so one locked `add()` publishes both and one `find()` reads both. - `AMDGPU::isStandardCUs` is now atomic. `isStandardCU()` computes the answer in a local variable and stores it once. A small copyable wrapper holds the `std::atomic<int>`, because `AMDGPU` is copied when used as a cache key. - `ContractionSolution::requiredHostWorkspaceSizePerProblem` is now `std::atomic<size_t>`. Grouped-GEMM paths fill it in lazily without a lock and every writer stores the same value. Nothing serializes or `memcpy`s the field, and the class already contains mutexes, so it is not copied. Grouped GEMM itself is not redesigned. - `hipblaslt.h` documents that distinct handles may be used concurrently from different threads and that they share process-wide caches that the library synchronizes internally. The existing handle-per-stream guidance is unchanged. No public hipBLASLt C or C++ function changes, and no device code changes. Test changes are in `ExactLogicLibrary_test.cpp`, `CachingLibrary_test.cpp`, `IndexedLibraryLoad_test.cpp`, `RangeLibrary_test.cpp`, `caching_library_gtest.cpp`, `heuristic_workspace_gtest.cpp`, `local_kernel_library_gtest.cpp`, and the new `concurrent_handles_gtest.cpp`. ## Test Plan - `tensilelite-tests`, filtered to `ExactLogicLibraryTest`: four new tests check that the new search type skips the two ranking rows on a standard-CU device, does not skip them on a non-standard-CU device (gfx942, 38 CUs), and forwards the search type to sub-libraries. - `tensilelite-tests`, filtered to `CachingLibraryTest.*:PlaceholderIndexedShardTest.ConcurrentColdLookups*`: the new `CachingLibraryTest.ConcurrentFindTopSolutionsSameProblemKeepsSolutionsAndFlagTogether` runs 4 threads on one problem (two asking for 3 solutions, two asking for 20) for 4000 rounds against a sub-library that holds 10 solutions. It fails if any result marked "returned all" holds fewer than 10 solutions, or if a final request for 10 gets fewer. `ConcurrentFindTopSolutionsDoesNotPoisonAlreadyRetAllCache` covers two different problems and stays unchanged. - `hipblaslt-test --gtest_filter='ConcurrentHandles.*'`: 4 threads running GEMM; 8 threads alternating GEMM and grouped GEMM; and the same mixed workload in a fresh process with cold lazy shards. Each thread has its own handle and stream, and every element of D is checked. - Full `tensilelite-tests`, `hipblaslt-test-hostunit`, and GPU `hipblaslt-test --gtest_filter='*quick*'` (see Test Result for which of these ran before and after the H2/M7 commit). - Before the M3 change: ThreadSanitizer on the unmodified `Debug.cpp`, and a deterministic condition-variable harness for the contamination in item 2. ## Test Result Local Release gfx950 build; automated GitHub checks have not been read for the latest push. - Negative control for the `CachingLibrary` split publish: on the code before the fix, the new same-problem test failed. Over 4000 rounds, 14 results paired the wrong "returned all" bit with their solutions, and 61 final requests for 10 solutions got fewer than 10. - After the fix: `tensilelite-tests` with `CachingLibraryTest.*:PlaceholderIndexedShardTest.ConcurrentColdLookups*` ran 7 tests, 7 passed. The new test also passed 10 of 10 repeated runs. - After the fix: `hipblaslt-test` with `ConcurrentHandles.*:HeuristicWorkspaceLimit.*:*CachingLibrary*:LocalKernel*`: 19 of 19 passed. Full `tensilelite-tests`: 845 passed, 3 skipped. - M3, before the change: ThreadSanitizer reported 20 data races between the getter's set copy (`Debug.cpp:200`) and the setter's copy-assign (`Debug.cpp:205`). The singleton was constructed before the threads started, which rules out static-initialization races. The harness showed a get-all-style reader observing the writer's two exclusions. - M3, from the earlier run on commit 95e0154: `ExactLogicLibraryTest` 11 of 11 passed. Full `tensilelite-tests` 664 passed, 4 skipped, 0 failed, of 668. `hipblaslt-test-hostunit` 14 of 14. `*quick*` on MI355X (gfx950:sramecc+:xnack-): 7740 passed, 1 failed. The failure is an FP4 subtile case (`quick_matmul_subtile_fp4_basic_..._TN_4096_4096_256_1_..._2_SAMX_32_UE8M0`, solution index 16130) with `norm_check` false at `testing_matmul.hpp:1485`. Run alone it passed 5 of 5 times, so it is intermittent. I did not run it against a build without this change and did not find a root cause. - Not rerun after the H2/M7 commits: `hipblaslt-test-hostunit` and the `*quick*` GPU category. - Grouped GEMM under concurrent handles is covered by `ConcurrentHandles.smoke_MixedGemmAndGroupedGemm` and a cold-library variant that saw 69 lazy shard loads. The cold variant also passed 30/30 against the pre-fix `PlaceholderLibrary`. The H2 tests that fail without the fix are the host tests in `IndexedLibraryLoad_test.cpp`. - Not run in the earlier M3 run: the `hipblaslt.symbol-contract.tensilelite-host` ctest is disabled at configure because `llvm-ifs` was not found. `library/symbol-contract/tensilelite-host-imports.txt` was checked by hand and references only `Debug::Instance()`. `pre-commit` fails in this environment on untouched files (`InvalidManifestError: Type tag 'textproto' is not recognized`), and `.pre-commit-config.yaml` excludes `projects/hipblaslt/.*` from its formatting hooks. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5). For a single caller, heuristic fallback and top-solution results are unchanged; they differ only where the old code lost or mis-cached candidates under concurrency. The change touches the library classes every dispatch passes through. These TensileLite host header classes change layout or vtable: `Debug` loses `m_excludedFromGetAll` and `excludedLibFromGetAll()`. `SolutionLibrary` loses the virtual `lastFindTopAlreadyRetAll()`, which changes the vtable of every library class, and eight library classes lose their `lastFindTopRetAll` member. `PlaceholderLibrary` replaces its `std::mutex` with a `std::once_flag`, returns the library from `loadPlaceholderLibrary()` instead of a `bool`, and adds the virtuals `beforeLibraryRegistered()` and `afterLibraryPublished()`. `CachingLibrary` replaces its two top-solutions caches with one. `AMDGPU::isStandardCUs` and `ContractionSolution::requiredHostWorkspaceSizePerProblem` become atomics of the same size and alignment (a `static_assert` checks `LazyFlag` against `int`), so for those two only copy behavior changes. That is safe when `libhipblaslt` and `libtensilelite-host` build together, as they do here, but a separately built consumer of the host library is not binary-compatible with this change. One unbaselined intermittent GPU failure from the earlier run remains open.
This branch has not been deployed
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
Motivation
From ROCm build
10.2.0a20260913, every small FP32 GEMM on gfx1200 pays a flat ~5 mshost-side cost per call versus ~0.025 ms on
10.2.0a20260911— a 140x–222x regression.pyt_huggingface_bartloses 91.6% of its training throughput (16.0 → 1.338 samples/s) andpytorch_utno longer completes: 42 distributed tests now exceed PyTorch's 300 s watchdog,pushing the job past the 480-minute CI limit so no results are uploaded at all. Kernel
execution is unaffected (large GEMMs do not move); the reporter's rocgdb backtrace sits in
TensileLite::PlaceholderLibrary::findAllSolutionsat 65.9% CPU, burning host cycles.The regression correlates exactly with the build that split hipBLASLt into
libhipblaslt.so.1plus a newlibtensilelite-host.so.1. Swapping only_rocm_sdk_librariesback to the good build restores full performance.Technical Details
The ticket attributes the cost to "crossing the new library boundary." That framing is
wrong — cross-DSO call overhead is nanoseconds, and the backtrace shows one crossing per
GEMM. Splitting a static archive into a DSO does not make calls slower; it changes symbol
identity.
TensileLite::Debugis aLazySingleton:Instance()is an implicitly-inline member of aheader-only class template holding a function-local static. Both
libtensilelite-hostandlibhipblasltare built with-fvisibility=hiddenand-fvisibility-inlines-hidden. WhileTensileLite was a static archive linked into
libhipblasltthere was one link unit andtherefore one instance. Once
TENSILELITE_BUILD_SHARED_LIBS=ONstarted shippinglibtensilelite-host.soas its own DSO, each shared object got its own private copy of thestatic — silently: no link error, no warning, no crash.
rocblaslt_auxiliary.cpp:2330writes the{GridBasedMatching, PredictionMatching}skip setinto libhipblaslt's copy of
Debugbefore callinggetAllSolutions.ExactLogicLibrary::findAllSolutions(ExactLogicLibrary.hpp:169) readslibtensilelite-host's copy, sees an empty set, so the
continueat line 182 never fires andthose expensive ranking sub-libraries are fully traversed. There is no memoization anywhere
on that path, so the cost is paid on every call — matching the observed steady-state
behaviour.
The fix marks
LazySingleton::Instance()TENSILELITEHOST_EXPORT, giving the functionand its local static default visibility in every object that instantiates it, so the dynamic
linker collapses them to one definition.
Singleton.hppalready included<tensilelitehost/export.h>without ever using the macro. The same change covers the otherLazySingletonderivatives (EmbeddedLibrary,EmbeddedData), which had the same latentexposure.
rocblaslt::LazySingletonis a different template used only inside libhipblasltand is untouched.
Added export surface is 5 symbols, all
V(unique global) orW(weak). libhipblasltalready exported 90
TensileLite::symbols in the shared arrangement, so this is not a newABI category.
Scope correction
The defect is architecture-independent — every 0913+ build on every arch has two
Debugsingletons. What is architecture-specific is the exposure, and it is a hard cliff rather
than a gradient. The slow path is gated behind
requestedAlgoCount > *returnAlgoCount(
rocblaslt_auxiliary.cpp:2324), so a workload is hurt iff it asks for more algos than thatarchitecture's matching tables can return for that shape. Sweeping
requestedAlgoCountongfx942 (
mm 64x32x32FP32, unfixed shared build):gfx1200/gfx1201 ship the sparsest libraries (Equality + GridBased only; 153 logic files vs
gfx942's 568+), so ordinary callers fall over the cliff on every call, and GridBased is ~73%
of what they have so losing the skip costs proportionally the most. gfx942/gfx950 have deep
enough tables that normal traffic never reaches the gate — which is why MI350X/MI350P looked
clean on the same builds. gfx1201 should be expected to be affected identically to
gfx1200 (the ticket lists it as suspected but unverified).
Test Plan
d2034be48c3) two ways, changing only-DTENSILELITE_BUILD_SHARED_LIBS, and timehipblasLtMatmulAlgoGetHeuristicfor thethree small FP32 shapes from the ticket plus
1024x1024x1024as a negative control.nmboth arrangements for the singleton symbol.&Debug::Instance()andexcludedLib.size()at both the writerand the reader.
hipblasLtMatmuldriven through thegetAllSolutionspath,checked against a CPU reference.
Test Result
All on a gfx942 host at commit
d2034be48c3. The device library was reused viaHIPBLASLT_TENSILE_LIBPATHso only the host side was rebuilt per arrangement.(1)/(4) Timing,
requestedAlgoCount=2000, ms/call:209.6x reproduces the ticket's 140x–222x band; the large-GEMM control does not move; the fix
restores the pre-split numbers exactly. The library sizes also reproduce the ticket's
fingerprint (static 9,445,024 → shared 2,860,672 + 6,848,040; ticket: 10,590,841 →
4,612,049 + 6,607,857).
(2)
nm:(3) Runtime probe:
(5) Regression test, reverting only the one-line
Singleton.hppchange:(6) Full tensilelite gtest suite, shared build: 492 tests, 485 passed, 7 skipped
(pre-existing environment skips), 0 failed.
(7) Real
hipblasLtMatmul, 16x16x16 FP32, three different returned algos each, checkedagainst a CPU reference: PASS on both shared and static, 0 mismatches. Returned-algo
counts are identical between shared+fix and static (1445 / 1451 / 1445 / 2000), i.e. the fix
changes which candidates are enumerated, not which are valid.
Not yet covered — worth a reviewer's attention:
end-to-end claim "this restores BART throughput on Navi44" has not been measured. Expect
5 ms → ~0.025 ms and 1.3 → ~13 samples/s.
HIPBLASLT_ENABLE_DEVICE=OFF, rocroller off, marker off)against a prebuilt device library. The linking arrangement that matters is identical, but
CI should cover the normal full build.
Follow-ups (deliberately not in this PR)
parameter of the query, not global debug state: it is set immediately before
getAllSolutionsand reset immediately after, guarded by astatic std::mutexthatexists solely to serialise the global. Threading it through
findAllSolutions(or intoSolutionLibrarySearchType) removes the shared mutable state, removes the lock, and isimmune to any future packaging change. This PR is the right regression fix; that is the
right design fix.
dynamic_cast/ typeinfo. The solution-library template tree carries noTENSILELITEHOST_EXPORT, so its typeinfo is now per-DSO and comparison falls back tostrcmpon very long mangled names. Not the cause of this regression — fixing thesingleton alone recovers all of the loss — but real, created by the same packaging change,
and worth its own ticket.
Submission Checklist
Risk level
Low. The functional change is a visibility attribute on one accessor; it cannot alter
computation, only how many copies of an object exist. Verified that returned-algo counts and
matmul results are bit-identical to the pre-split arrangement, and the full tensilelite suite
passes. The added ABI surface is 5 weak/unique symbols on top of the 90
TensileLite::symbols libhipblaslt already exports in this arrangement. Recovers a P2 regression that
currently blocks all gfx1200 pytorch_ut reporting.
🤖 Generated with Claude Code