Repository navigation
fix(hipblaslt): enumerate solutions in index order, not heap-address order - #11954
Closed
jaopaulolc wants to merge 3 commits into
Closed
jaopaulolc wants to merge 3 commits into
jaopaulolc wants to merge 3 commits into
Conversation
…order getAllSolutions() collects candidates into a std::set<std::shared_ptr<ContractionSolution>>, which orders by raw pointer value. Under ASLR that is a different order on every run for byte-identical input. The fallback in rocblaslt_matmul_algo_get_heuristic() consumes the list positionally, taking the first entries that pass isSolutionSupported(), so the kernel it selects varied run to run. Sort the enumeration by the library-assigned solution index, which is deserialized from the library data and is stable across runs. Sorting at the consumption site rather than changing the SolutionSet comparator preserves the set's pointer-identity dedup, which the loop's duplicated_counts logic relies on. Adds a client test that pins the ordering invariant.
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (35.26%) 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 #11954 +/- ##
===========================================
+ Coverage 69.42% 70.26% +0.84%
===========================================
Files 2828 2810 -18
Lines 470113 462444 -7669
Branches 69417 68077 -1340
===========================================
- Hits 326343 324925 -1418
+ Misses 120066 113956 -6110
+ Partials 23704 23563 -141
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
AGENTS.md requires new source files to carry the two-line SPDX header rather than the legacy verbose MIT block. Match the closest neighbour, clients/tests/src/row_uniformity_gtest.cpp.
jaopaulolc
deleted the
users/jolabega/deterministic-getallsolutions-order
branch
October 6, 2026 20:30
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 : N/A
Motivation
When a caller asks for more candidate kernels than the ranked heuristic can supply, hipBLASLt falls back to
getAllSolutions(), which enumerates every kernel matching the problem's data types and collects them intoA
std::setofshared_ptrorders by raw pointer value — by heap address, which address-space layout randomization varies per process — so the enumeration comes back in a different order on every run for byte-identical input. Callers consume it by position:rocblaslt_matmul_algo_get_heuristic()appends the first entries passingisSolutionSupported(), so at the default request count of one, element zero of an address-ordered set wins.Order decides it because nothing on this path ranks. The ranked path selects using the library's tuning data — a match against tuned problem keys, or a predicted-performance model.
getAllSolutions()matches on GEMM type alone and applies no such criterion.Within one process the order is fixed, since the solution objects live in a function-local static that outlives every handle. A single process stays self-consistent; repeated runs of the same program disagree, making the selected kernel — and its runtime — irreproducible.
Technical Details
Sort the enumeration by
ContractionSolution::indexbefore consuming it.indexismapRequiredin the solution serializer, so it comes from the library data file and is identical on every run.The sort is at the point of consumption rather than in
SolutionSet's comparator, whose pointer-identity semantics the loop'sduplicated_countshandling relies on; changing it would alter deduplication for every user ofSolutionLibrary. That is a wider-blast-radius follow-up.One sort covers the GEMM, grouped-GEMM and xfloat32 fills, which share the variable.
Test Plan
clients/tests/src/get_all_algos_ordering_gtest.cpp: assertgetAllAlgos()reports strictly increasing solution indices over three data-type combinations. It checks ordering only.ldd-verified against its ownlibhipblaslt.so.1. On gfx950, under heap-layout pressure (MALLOC_TOP_PAD_), compare the enumerated list and the kernel selected on four problems whose ranked path returns fewer solutions than requested.tensilelite-testsandhipblaslt-test --gtest_filter='*RowUniformity*'.Test Result
New test. Passes 3/3 with the fix; the same binary fails 3/3 against the unpatched library on every one of 9 runs (103 to 155 adjacent-pair inversions).
Enumeration order. Over 9 runs (3 heap layouts x 3 repetitions) the unpatched build returns 7 distinct orderings of a 256-solution list, with the first divergence at exactly index 139 — where the fallback takes over. The patched build returns 1.
Selected kernel. Unpatched picks 3 to 5 distinct kernels per problem under heap-layout pressure. Patched picks one across 176 runs, the kernel the default layout already produced.
Process scope. Repeated
getAllAlgos()calls within one process return an identical index sequence on both builds, including with allocation churn or handle destroy and recreate between calls. Across separate processes the unpatched build returns a different ordering almost every time; the patched build returns the same one.Index ties. None observed (
duplicated counts: 0on both builds). Theduplicated_countshandling is kept regardless; tied entries share anindex, sogetSolutionByIndex()resolves the same kernel.Regression suites. Identical on both builds:
tensilelite-tests447 ran, 445 passed, 2 pre-existing skips, 0 failed;hipblaslt-test --gtest_filter='*RowUniformity*'62/62.Performance. At the default heap layout the selected kernel is unchanged on all four problems and its runtime matches the unpatched build. Because nothing here ranks, the entry that ends up first is not the fastest available: the deterministic pick was slower than kernels other layouts happened to put first, and on most of the four it was the slowest observed. This makes the arbitrary choice stable, not good; ranking the fallback's appended solutions is the real fix and is not attempted here.
Scope. gfx950; four problem shapes, three type combinations.
Submission Checklist
Risk level
Low. The change reorders a list that had no defined order. Every element returned before is still returned, no public signature changes, and the default-layout selection is unchanged.