Skip to content

fix(tensilelite): align topology filtering with selection - #11946

Closed
davidd-amd wants to merge 3 commits into
developfrom
users/davidd-amd/tensilelite-streamk-topology-filter
Closed

davidd-amd wants to merge 3 commits into
developfrom
users/davidd-amd/tensilelite-streamk-topology-filter

Conversation

@davidd-amd

Copy link
Copy Markdown
Contributor

ISSUE ID : #11890

Motivation

Merged PR #11872 made the TensileLite benchmark client reject dynamic StreamK kernels when a device's compute-die count does not match the queue count compiled into the kernel. A post-merge review identified one exception path, a difference from library predicate ordering, and tests that did not verify the reported rejection value.

This follow-up makes the predicate return a Boolean result for unknown hardware and verifies that the benchmark client reports the same rejection reason precedence as normal library selection.

Technical Details

  • Return unsupported for SK5 on hardware without Origami analytical data before AUTO mode resolution can throw.
  • Evaluate problem and task predicates before the StreamK topology predicate, matching softwarePredicate() ordering.
  • Replace the iterator test's unused output stream with a recording reporter that verifies UNSUPPORTED_XCD_TOPOLOGY in reporting mode and no emitted value in non-reporting mode.
  • Verify that problem-assert failures take precedence over topology rejection.
  • Verify that LogReporter emits an unsupported-topology row at terse log level.

This is a follow-up to #11872 and does not change kernel generation or the hipBLASLt library interface.

Test Plan

  1. Configure and build tensilelite-tests with TENSILELITE_BUILD_TESTING=ON.
  2. Run tensilelite-tests --gtest_filter=StreamKDynamicQueueXcdGateTest.*.
  3. Run the shared hipBLASLt and TensileLite CI lanes selected for this PR.

Test Result

  • Fresh isolated C++ build against current develop: passed.
  • StreamKDynamicQueueXcdGateTest.*: 11 passed.
  • Automated GitHub and Math CI checks: pending.

Submission Checklist

Risk level

Low (2/5): the change affects benchmark-client candidate filtering and its tests. Unknown hardware now returns the documented unsupported result instead of reaching a throwing SK5 AUTO-mode calculation.

@therock-pr-bot

therock-pr-bot Bot commented Sep 10, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@codecov

codecov Bot commented Sep 10, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #11946      +/-   ##
===========================================
+ Coverage    70.54%   70.56%   +0.02%     
===========================================
  Files         2810     2810              
  Lines       462436   462508      +72     
  Branches     68077    68099      +22     
===========================================
+ Hits        326219   326347     +128     
+ Misses      112673   112610      -63     
- Partials     23544    23551       +7     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 46.69% <100.00%> (+0.24%) ⬆️
TensileLite-Unit 76.07% <ø> (+0.04%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from e71edb7
hipBLASLt 35.27% <ø> (+0.03%) ⬆️ Carriedforward from e71edb7
hipCUB 82.68% <ø> (ø) Carriedforward from e71edb7
hipDNN 87.02% <ø> (ø) Carriedforward from e71edb7
hipFFT 42.66% <ø> (ø) Carriedforward from e71edb7
hipRAND 76.12% <ø> (ø) Carriedforward from e71edb7
hipSOLVER 68.92% <ø> (ø) Carriedforward from e71edb7
hipSPARSE 86.99% <ø> (ø) Carriedforward from e71edb7
rocBLAS 48.31% <ø> (ø) Carriedforward from e71edb7
rocFFT 51.72% <ø> (ø) Carriedforward from e71edb7
rocRAND 56.90% <ø> (ø) Carriedforward from e71edb7
rocSOLVER 76.83% <ø> (ø) Carriedforward from e71edb7
rocSPARSE 74.61% <ø> (ø) Carriedforward from e71edb7
rocThrust 91.60% <ø> (ø) Carriedforward from e71edb7

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...blaslt/tensilelite/client/src/SolutionIterator.cpp 4.34% <100.00%> (+0.72%) ⬆️

... and 12 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@davidd-amd
davidd-amd requested a review from talumbau September 11, 2026 15:54
davidd-amd and others added 2 commits September 11, 2026 21:19
Follow-up to review comments left by @talumbau on #11872 after it merged,
carried by #11946.

The previous guard rejected every SK5 solution as soon as a device had no
analytical hardware, which is broader than the condition it was protecting
against. Only the origami hybrid-mode heuristic needs analytical hardware;
a mode that resolves without it was being discarded anyway.

Drop the blanket early-out and let mode resolution decide. A mode that
resolves -- explicit static, explicit dynamic, or a forced mode from
TENSILE_STREAMK5_FORCE_MODE -- is honoured: static stays selectable,
dynamic is excluded. A mode that cannot resolve (AUTO, or OFF with
sm_count_target > 0, both of which enter the heuristic) throws rather than
silently substituting a schedule the caller did not ask for. This also
makes the grouped-GEMM workspace path, which reaches mode resolution
before the topology guard runs, behave the same way as the single-GEMM
path instead of diverging from it.

This narrows the guard that #11946 added from the suggestion on
SolutionIterator.cpp:114. That suggestion rested on "the answer does not
change, because unknown hardware is already UNSUPPORTED", which does not
hold for the modes that resolve without analytical hardware: they never
reach the throwing heuristic, yet the guard rejected them too.

Unknown analytical hardware only arises when origami reports the device as
unsupported, and SK5 kernels are only built for supported architectures,
so the throwing path is defensive rather than reachable in production.

Tests: the old unknown-hardware test was vacuous whenever
TENSILE_STREAMK5_FORCE_MODE was set in the environment, because that
variable short-circuits mode resolution before the code under test runs.
Pin it with a ScopedStreamK5ForceMode guard and assert the pinned value,
then split the cases so each resolution outcome is checked on its own.
Drop the postProblem() call from the LogReporter test, which printf()s a
banner to stdout rather than to the captured stream, and add the missing
ResultReporter.hpp/<cstdint>/<cstdlib>/<stdexcept> includes.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
`TENSILELITE_BUILD_SHARED_LIBS` defaults to `ON`, and in that configuration
`tensilelite-tests` failed to link. Two separate causes.

Undefined LLVM YAML symbols. The serialization templates are instantiated in
the test translation units, not only inside `tensilelite-host`. That library
links `LLVMObjectYAML` `PRIVATE` and, when shared, is linked with
`--exclude-libs,ALL` (`tensilelite/CMakeLists.txt:40`), so it exports none of
those symbols. `RangeLibrary_test.cpp` referenced
`llvm::yaml::ScalarTraits<bool>::input`, `ScalarValueTypeInfo::Get` and
`KernelLanguageTypeInfo::Get` with nothing to resolve them against. Fixed by
linking `LLVMObjectYAML` into the test target as well, and by adding
`TENSILELITEHOST_EXPORT` to the enum-info `Get` overloads and to the
`MLFeatures` / `MLPClassification` entry points the tests call.

Duplicated `Debug` singleton. `Debug` inherits its instance accessor from
`LazySingleton`, which holds the object in a function-local static of an
inline template member. `VISIBILITY_INLINES_HIDDEN` gives that hidden
visibility, so a shared `tensilelite-host` and the test binary each construct
their own `Debug`. A test calling `reloadDebugBitsForTest()` would then mutate
an instance the library never reads. `Debug::Instance()` is now defined out of
line in `Debug.cpp` so both sides share one object.

This failure predates the Stream-K commits that precede it on this branch. A
shared build of `tensilelite-tests` at the branch point `267267524ae` emits the
same 20 undefined symbols, and the files touched here are disjoint from those
commits.

Verified in `hipblaslt-ondemand:ubuntu24-gfx950-c865f12`: `tensilelite-tests`
builds and runs in both the shared and static configurations, 416 passed /
3 skipped / 5 failed. The 5 failures are all `PciChipIdTest`, which call into
HIP and get `hipErrorNoDevice` (100) because the container has no GPU.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
@davidd-amd
davidd-amd force-pushed the users/davidd-amd/tensilelite-streamk-topology-filter branch from d0f6729 to cf40154 Compare September 11, 2026 22:37
@davidd-amd
davidd-amd requested a review from a team as a code owner September 11, 2026 22:37

@Alex-Vasile Alex-Vasile left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

CMake changes are good.

@sstamenk

sstamenk commented Sep 15, 2026 •

Copy link
Copy Markdown
Member

This PR also appears to fix the performance regression tracked in AIRADSW-923.

The issue affects LTX-2.3 on R9700S/gfx1201. Its grouped ConvTranspose1d operation is decomposed into 768 small FP32 hipBLASLt GEMMs.

Docker-level measurements:

Build Median latency
2026-09-11, before shared TensileLite host 28.249 ms
2026-09-15, with shared TensileLite host 4952.693 ms

This is approximately a 175x slowdown. Profiling showed that GPU execution remained fast; nearly all additional time was spent in hipblasLtMatmulAlgoGetHeuristic().

The underlying mechanism matches the duplicated Debug singleton described in commit cf40154:

  1. The fast solution lookup does not return a result for this small GEMM shape, so hipBLASLt enters the getAllSolutions() fallback.
  2. Before doing so, libhipblaslt.so sets excludedLibFromGetAll to avoid searching GridBased and Prediction libraries again.
  3. Because LazySingleton<Debug>::Instance() is an inline template function with hidden visibility, libhipblaslt.so and libtensilelite-host.so instantiate separate Debug objects.
  4. hipBLASLt writes the exclusion state to one instance, while ExactLogicLibrary::findAllSolutions() reads the other instance and sees an empty exclusion set.
  5. TensileLite consequently performs an exhaustive solution search for every one of the 768 GEMMs.

A predicate trace showed 176,079 candidate sections evaluated in the broken shared build versus 63 with the correctly shared state.

I validated the same out-of-line Debug::Instance() change in isolation using identical source and gfx1201 device objects:

Configuration Median latency
Static TensileLite host 30.564 ms
Broken shared TensileLite host 4887.355 ms
Shared host with exported Debug::Instance() 30.030 ms

Defining Debug::Instance() out of line in libtensilelite-host.so ensures that libhipblaslt.so resolves the same singleton. The exclusion state then reaches the search implementation, the exhaustive fallback scan disappears, and performance returns to baseline.

@sstamenk

Copy link
Copy Markdown
Member

One additional architectural question: this regression exposed that the shared-library split created two Debug singleton instances—one in libhipblaslt.so and one in libtensilelite-host.so.

The out-of-line exported Debug::Instance() correctly restores a single shared instance and fixes the regression. However, excludedLibFromGetAll is operational solution-selection state rather than purely diagnostic state.

Would maintainers prefer a follow-up refactor that passes these exclusions through an explicit per-call search context instead of mutable global Debug state? That would remove the cross-DSO singleton dependency and improve concurrency and exception safety. The current fix appears appropriate as the minimal regression correction, but reviewer guidance on the longer-term design would be appreciated.

@fjankovi

fjankovi commented Sep 23, 2026 •

Copy link
Copy Markdown
Contributor

@davidd-amd Do you plan to merge this? It addresses significant performance regressions on Radeon introduced in 10.2.

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>
@github-actions

Copy link
Copy Markdown
Contributor

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

  • Add new commits
  • Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

@github-actions github-actions Bot added the Stale PR has no activity for 25+ days label Oct 19, 2026
@github-actions

Copy link
Copy Markdown
Contributor

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

@github-actions github-actions Bot closed this Oct 25, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants