Repository navigation
feat(TensileLite): generate deterministic assembly - #9853
Merged
mgehre-amd merged 2 commits intoAug 12, 2026
Merged
Conversation
mgehre-amd
force-pushed
the
users/mgehre/tensilelite-deterministic-labels
branch
from
July 24, 2026 13:32
8a48cd5 to
c40d41d
Compare
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
❌ 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 |
mgehre-amd
force-pushed
the
users/mgehre/tensilelite-deterministic-labels
branch
2 times, most recently
from
July 24, 2026 13:53
12cc9bc to
41f071b
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #9853 +/- ##
===========================================
- Coverage 69.61% 68.11% -1.51%
===========================================
Files 2741 2771 +30
Lines 451317 463211 +11894
Branches 66491 74315 +7824
===========================================
+ Hits 314177 315472 +1295
- Misses 116920 125355 +8435
- Partials 20220 22384 +2164
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
eble-amd
reviewed
Jul 24, 2026
TensileLite's rocisa LabelManager generated unique label names with a
random 16-char suffix (magicGenerator using rand()). Because rand() is
process-global and TensileCreateLibrary parallelizes kernel generation
across multiprocessing workers, the emitted label names differed run to
run, defeating reproducible codegen.
Replace the random suffix with a per-instance monotonically increasing
counter. The counter is carried through the pickle/deepcopy state so
copied managers keep producing unique names. Each kernel gets a fresh
LabelManager in _initKernel, so the counter restarts at 0 per kernel;
labels only need to be unique within a single kernel's assembly.
Changes:
- getUniqueName now delegates to getUniqueNamePrefix("label"); both use
the counter and keep the collision-guard loop against manually added
names.
- LabelManager gains a (data, counter) constructor and getCounter() so
__getstate__/__setstate__/__deepcopy__ preserve the counter.
mgehre-amd
force-pushed
the
users/mgehre/tensilelite-deterministic-labels
branch
from
July 27, 2026 07:23
41f071b to
42e91b1
Compare
eble-amd
approved these changes
Jul 27, 2026
KKyang
approved these changes
Jul 31, 2026
mgehre-amd
enabled auto-merge (squash)
July 31, 2026 11:28
1 task done
geotseng-amd
added a commit
that referenced
this pull request
Oct 5, 2026
test(tensilelite): normalize StinkyTofu label hashes before canonicalize_asm InsertClusterBarrierPass names its skip labels with a random 16-char [0-9a-zA-Z] hash. canonicalize_asm only renames suffixes that are all [A-Z0-9], so about one hash in 6000 took that path in one emit while its counterpart in the other emit was renamed by _stableLabels instead, making test_gfx1250v0_emits_what_gfx1250_strict_emits fail roughly once in 750 runs. Run _stableLabels on the raw source inside _emit, ahead of canonicalize_asm, so every hash is renamed by the same pass, and compare the emitted sources directly. ## Motivation `Tensile/Tests/unit/test_arch_steppings.py::test_gfx1250v0_emits_what_gfx1250_strict_emits`, added in #11777, fails intermittently in unrelated PRs' CI (for example `s_preliminary_hipblaslt_PR-12538`: 1 failed / 9822 passed). The emitted gfx1250v0 and gfx1250-strict kernels are actually identical; the failure comes from how the test normalizes random label names before comparing them. This PR removes the flake without weakening the assertion. ## Technical Details **Where the randomness comes from.** #9853 made rocisa's `LabelManager` deterministic (`prefix_<counter>`), but StinkyTofu's `InsertClusterBarrierPass::makeRandomHash()` (`shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp`) still suffixes `label_skipCBPreSignal_*` / `label_skipCBPreSignal_LCL_*` with a 16-char hash drawn from `[0-9a-zA-Z]`, seeded from `std::random_device`. This kernel (`MULTICAST_MI`, `ClusterDim=[2,1]`) carries 4 such labels per emit. **Why two normalization passes disagree.** The test normalized each emit twice, with different rules: | Pass | Matches | Renames to | |---|---|---| | `canonicalize_asm` (inside `_emit`) | `_[A-Z0-9]{16}\b` | `_LBL<n>` | | `_stableLabels` (in the test, afterwards) | `_[A-Za-z0-9]{16}(?![A-Za-z0-9])` | `_L<n>` | A hash is all `[A-Z0-9]` with probability (36/62)^16 ≈ 1/6000. When that happens in only one of the two emits, that label becomes `_LBL0` on one side and `_L3` on the other, and every later `_L<n>` shifts by one, so the comparison fails. With 8 hashes per test run, the failure rate is about 1 in 750. Fixing `_stableLabels` alone is not enough: one label is spelled `label_InitCIterWmma_label_skipCBPreSignal_<hash>_0`, where `\b` does not fire before `_`, so `canonicalize_asm` would rename an all-caps hash in some occurrences and not others. Normalization has to happen on the raw source. **Change.** - Move `_stableLabels` to module scope. - Apply it to the raw source in `_emit`, before `canonicalize_asm`. After that, `canonicalize_asm`'s label regex has nothing left to match, and it only strips `matrix_[ab]_reuse`. - `test_gfx1250v0_emits_what_gfx1250_strict_emits` now asserts `v0Src == strictSrc` directly. `test_gfx1250_emits_multicast_gfx1250_strict_does_not`, the only other `_emit` caller, is unaffected: its `!=` check holds because the multicast setup genuinely differs. The root fix is to make StinkyTofu's label naming deterministic (as #9853 did for rocisa); that is left to the StinkyTofu owners. ## Test Plan Run in a gfx1250 FFM container (ROCm 10.2 nightly), with rocisa built from this branch's tree: 1. Full file: `pytest Tensile/Tests/unit/test_arch_steppings.py`. 2. Fault injection on real emitted kernels: take raw strict and v0 sources, replace each of v0's 4 StinkyTofu hashes in turn with an all-`[A-Z0-9]` hash, and compare under the old and new normalization order. 3. Determinism check: emit 12 strict/v0 pairs (24 emits) and compare them after label normalization, against each other and across pairs, to confirm the hash is the only run-to-run variation. ## Test Result 1. `828 passed`. 2. Old order: all 4 injections mismatch (reproducing the CI failure signature `_L3` vs `_LBL0` with shifted numbering). New order: all 4 compare equal, including the hash that also appears with a `_0` suffix. 3. All 24 emits are identical after normalization; no other nondeterminism found. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. Co-authored-by: Cursor <cursoragent@cursor.com>
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.
Problem
TensileLite's rocisa
LabelManagergenerated unique branch/loop label names with arandom 16-character suffix (
magicGenerator()usingrand()), e.g.NoBranch_T8JHFHKM7BO5OHXW. Becauserand()is process-global shared state andTensileCreateLibraryparallelizes kernel generation acrossmultiprocessingworkers, the emitted label names differed run to run depending on how kernels were
distributed across workers. This is the only source of run-to-run nondeterminism in
the emitted assembly text, which defeats reproducible codegen and byte-for-byte
comparison of generated kernels.
Change
Replace the random suffix with a per-instance, monotonically increasing counter:
getUniqueNamePrefix(prefix)now emitsprefix_<counter++>, keeping thecollision-guard loop so a counter-derived name can never collide with a name that
was added manually.
getUniqueName()delegates togetUniqueNamePrefix("label").LabelManagergains a(data, counter)constructor andgetCounter()so thecounter is preserved through
__getstate__/__setstate__/__deepcopy__— acopied manager keeps producing unique names.
Each kernel receives a fresh
LabelManagerinKernelWriter._initKernel, so thecounter restarts at 0 for every kernel; labels only need to be unique within a single
kernel's assembly namespace.
Notes
getUniqueName()'s output format changes from a bare random token tolabel_<N>;it is only referenced by
rocisa/test/test_label.py. All production call sites usegetUniqueNamePrefix.Testing
rocisa/test/test_label.pypasses; verified deterministic output and that thecounter survives
pickle/deepcopyand restarts on a fresh manager.(
Tensile/Tests/unit/characterization) passes with the change, with no snapshot(
.ambr) golden updates required:_codegen(all assembly-emit tests and{basename, err}goldens): 254 passedand the remaining directories: all passed
{basename, err}only,and
codegen_harness.canonicalize_asmalready normalizes label suffixes before anystructural assertion, so counter-based labels flow through unchanged.