Skip to content

refactor(tensilelite): separate persistent execution policies - #12537

Merged
bnemanich merged 11 commits into
developfrom
users/nhenders/persistent-loop-01-policy
Oct 1, 2026
Merged

bnemanich merged 11 commits into
developfrom
users/nhenders/persistent-loop-01-policy

Conversation

@ThanHenderson

@ThanHenderson ThanHenderson commented Sep 24, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AIHPBLAS-4950

Motivation

TensileLite generates matrix-multiplication kernels and controls how GPU workgroups receive output tiles. Persistent workgroups process multiple tiles, but selecting whole-tile DataParallel work currently requires StreamK settings. Expose tile processing and work assignment separately so DataParallel can be selected directly.

Technical Details

  • Add TileProcessingStrategy and WorkAssignment to configuration, solution identity, serialization, and host launch selection. None selects ordinary execution; DataParallel selects persistent whole tiles; StreamK allows splitting tiles across workgroups.
  • Name the predicate isPersistentDataParallel to distinguish persistent whole-tile execution from ordinary matrix multiplication. Preserve both ordinary custom-kernel grid forms and the persistent launch grid.
  • Apply file-level DefaultSolution values before global defaults, with explicit solution values taking precedence. Library merging uses the same rule for policy and ordinary tuning fields.
  • Apply experimental-library --set overrides to every size group. Tests cover both policy overrides and an ordinary DepthU=32 override across multiple groups.
  • Test serialized policies through the existing exported LoadLibraryFile entry point using temporary files. Both MessagePack and LLVM YAML serialization remain supported.
  • Parse policy enums through shared C++ helpers. Remove unused scheduling-state copies and numeric mode fields from launch diagnostics; report the named policy and effective assignment.
  • Convert legacy numeric inputs at ingestion because existing production logic records still contain them. Retain the version-zero, six-word scheduling argument layout used by existing kernels. Maintained configuration names migrate in PR #12540.
  • Preserve the upstream uniform-summation rule for mapping-1 staggering using the named persistence predicate, with regression coverage for ordinary, persistent DataParallel, and StreamK kernels.
  • Bound noncluster DataParallel launches by HIP's 32-bit work-item limit. ADR0014 documents the expected name/schema changes, including the affected gfx1250 saved results.
  • Preserve the original persistence flag while generating adaptive GlobalSplitU (GSU) store branches. This fixes the undefined skBackup left by the policy migration. Add 20 generation-and-assembly regression cases across gfx90a, gfx942, gfx950, and gfx1250.

Stack 1/5, targeting develop. Followed by PR #12538.

Test Plan

Run the Python unit suite, including policy ingestion, CLI overrides, generation-and-assembly regressions, and saved-result comparisons. Run the focused experimental-library extras tests. Build the complete C++ test executable with HIPBLASLT_ENABLE_YAML=OFF and run policy, serialization, argument-layout, custom-grid, launch-bound, range, and prediction-library tests. Use MessagePack for this validation; the existing LLVM YAML shared-test link failure is outside this PR. No performance measurements.

Test Result

At 5a40ad1816b, after rebasing onto develop at 7affebf2fef, the full Python unit suite passed 8,969 tests, with 47 skips, 2 expected failures, 1 unexpected pass, and all 786 snapshots passing. This includes all 20 adaptive-GSU generation-and-assembly cases. The focused policy-ingestion, CLI, and experimental-library run passed 139 tests before the rebase.

The rebuilt MessagePack host test executable passed 118 selected tests, including the new upstream prediction-library tests. One existing RangeLibrary test was skipped for AIHPBLAS-3506. Pre-commit passed. No snapshots were updated and no performance measurements were taken.

Automated checks for the updated head are pending. Previously reported Math profiling aborts/timeouts, ASAN runner failures, the C++ coverage shortfall, and other external gfx1250 failures are not claimed fixed; gfx1250 device execution was not validated locally.

Submission Checklist

Risk level

Medium (3/5): changes policy interpretation across generation, serialization, naming, and host dispatch while retaining existing input conversion and scheduling arguments. Tests cover policy conflicts, custom grids, serialization, launch bounds, and sampled emitted instructions.

@ThanHenderson ThanHenderson changed the title refactor(hipblaslt): separate persistent execution policies refactor(tensilelite): separate persistent execution policies Sep 24, 2026
@therock-pr-bot

therock-pr-bot Bot commented Sep 24, 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

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

Copy link
Copy Markdown

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

@jaopaulolc jaopaulolc added the rocm:gemm algos Label to identify PRs of the GEMM Algorithms team. label Sep 24, 2026

@AlexBrownAMD AlexBrownAMD 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.

Nice refactor! It'd be good to also run a quick performance test just to verify modes like stream-k/persistent kernel aren't accidentally being disabled somewhere that we didn't catch in review.

@jaopaulolc jaopaulolc 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.

LGTM so far. I am roughly 1/3 through it and will add more comments as I work through the changes.

@ThanHenderson
ThanHenderson force-pushed the users/nhenders/persistent-loop-01-policy branch 3 times, most recently from e7cfc93 to 6a3250c Compare September 25, 2026 22:01
@codecov

codecov Bot commented Sep 25, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 90.63636% with 103 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py 84.36% 15 Missing and 13 partials ⚠️
...elite/Tensile/Components/TileProcessingStrategy.py 91.46% 13 Missing and 8 partials ⚠️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py 84.40% 13 Missing and 4 partials ⚠️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py 79.66% 5 Missing and 7 partials ⚠️
...t/tensilelite/Tensile/Components/WorkAssignment.py 93.20% 6 Missing and 4 partials ⚠️
...pblaslt/tensilelite/Tensile/ExperimentalLibrary.py 91.95% 4 Missing and 3 partials ⚠️
...tensilelite/Tensile/Components/GlobalWriteBatch.py 81.25% 0 Missing and 3 partials ⚠️
...cts/hipblaslt/tensilelite/Tensile/AsmStoreState.py 71.43% 2 Missing ⚠️
...ite/Tensile/Components/Subtile/LogicalScheduler.py 66.67% 1 Missing ⚠️
...ects/hipblaslt/tensilelite/Tensile/Contractions.py 80.00% 0 Missing and 1 partial ⚠️
... and 1 more
Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #12537      +/-   ##
===========================================
+ Coverage    70.69%   70.72%   +0.03%     
===========================================
  Files         2941     2944       +3     
  Lines       484449   484876     +427     
  Branches     69855    69936      +81     
===========================================
+ Hits        342475   342916     +441     
+ Misses      118081   118056      -25     
- Partials     23893    23904      +11     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 48.82% <ø> (ø)
TensileLite-Unit 82.68% <90.64%> (+0.13%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from 69a220f
hipBLASLt 35.90% <ø> (ø) Carriedforward from 69a220f
hipCUB 82.68% <ø> (ø) Carriedforward from 69a220f
hipDNN 86.64% <ø> (ø) Carriedforward from 69a220f
hipFFT 54.32% <ø> (ø) Carriedforward from 69a220f
hipRAND 76.12% <ø> (ø) Carriedforward from 69a220f
hipSOLVER 68.96% <ø> (ø) Carriedforward from 69a220f
hipSPARSE 86.99% <ø> (ø) Carriedforward from 69a220f
rocBLAS 48.24% <ø> (ø) Carriedforward from 69a220f
rocFFT 47.11% <ø> (ø) Carriedforward from 69a220f
rocRAND 57.42% <ø> (ø) Carriedforward from 69a220f
rocSOLVER 76.83% <ø> (ø) Carriedforward from 69a220f
rocSPARSE 74.75% <ø> (ø) Carriedforward from 69a220f
rocThrust 78.45% <ø> (ø) Carriedforward from 69a220f

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

Files with missing lines Coverage Δ
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp 42.28% <ø> (ø)
...laslt/tensilelite/Tensile/AsmAddressCalculation.py 64.55% <100.00%> (+0.07%) ⬆️
.../hipblaslt/tensilelite/Tensile/BenchmarkStructs.py 89.92% <100.00%> (+0.94%) ⬆️
...slt/tensilelite/Tensile/Common/GlobalParameters.py 95.82% <ø> (ø)
...t/tensilelite/Tensile/Common/RequiredParameters.py 100.00% <ø> (ø)
.../hipblaslt/tensilelite/Tensile/Common/Utilities.py 100.00% <100.00%> (ø)
...aslt/tensilelite/Tensile/Common/ValidParameters.py 100.00% <ø> (ø)
...aslt/tensilelite/Tensile/Components/ClusterLoad.py 99.05% <100.00%> (ø)
...ts/hipblaslt/tensilelite/Tensile/Components/LSU.py 76.94% <100.00%> (ø)
...t/tensilelite/Tensile/Components/PersistentLoop.py 88.41% <100.00%> (+0.77%) ⬆️
... and 23 more

... and 1 file 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.

@jaopaulolc jaopaulolc 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.

LGTM. If a perf run does not show any unintended regressions, I am hapy to approve this.

@ThanHenderson
ThanHenderson force-pushed the users/nhenders/persistent-loop-01-policy branch from 5a40ad1 to 52c2933 Compare September 29, 2026 15:46
@ThanHenderson
ThanHenderson marked this pull request as ready for review September 29, 2026 16:11
@ThanHenderson
ThanHenderson requested review from a team as code owners September 29, 2026 16:11
@ThanHenderson
ThanHenderson force-pushed the users/nhenders/persistent-loop-01-policy branch from 52c2933 to 25fbd4e Compare September 29, 2026 17:34

@AlexBrownAMD AlexBrownAMD 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.

Thanks for the updates!

@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.

Approving CMake changes

@ThanHenderson
ThanHenderson force-pushed the users/nhenders/persistent-loop-01-policy branch 2 times, most recently from 81556c9 to 80ba953 Compare October 1, 2026 14:40
@ThanHenderson
ThanHenderson force-pushed the users/nhenders/persistent-loop-01-policy branch from 80ba953 to 46145ef Compare October 1, 2026 19:28
@bnemanich

Copy link
Copy Markdown
Contributor

All important tests passed before tests were restarted.

@bnemanich
bnemanich merged commit 812bcc6 into develop Oct 1, 2026
131 of 144 checks passed
@bnemanich
bnemanich deleted the users/nhenders/persistent-loop-01-policy branch October 1, 2026 22:56
minsukim-amd added a commit that referenced this pull request Oct 2, 2026
…20260901

Resolves three conflicts with develop's persistent execution-policy work
(#12537, #12539):

- LibraryIO.py: develop replaced parseLibraryLogicData's DefaultSolution
  merge with normalize_execution_policy_with_defaults(), which translates
  legacy StreamK selectors before Tensile's global defaults are filled in.
  This branch had moved that fill into the shared fillSolutionDefaults(), so
  give the helper develop's semantics -- normalize with the file's defaults,
  then fill defaultSolution -- and have it return the new dict, as the
  normalizer does. The converter's readers take its result too.
- TensileLibLogicToYaml.py: keep develop's normalize_execution_policy() as
  the first step of formForkParams, ahead of this branch's generic carry, so
  legacy selectors reach it as the canonical TileProcessingStrategy and
  WorkAssignment.
- DECISIONS.md: develop added its own D44 (and a second D30). Keep develop's
  entries as-is and renumber this branch's D44-D47 to D45-D48, updating the
  references in ADR 0030 and the three test files. This branch's ADRs
  (0029-0031) do not collide with develop's new 0014/0015 files.

The normalizer now always completes InternalSupportParams (it adds
PersistentLoopArgsVersion), and Tensile's reader hands that completed block
to a handwritten custom kernel, so the two custom-kernel tests expect the
normalized block rather than the raw input.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
caio96 added a commit that referenced this pull request Oct 2, 2026
…ration

SizeMapping no longer carries streamK or streamKForceDPOnly after
#12537. Gate the check on requiresPartialReduction(), which already
excludes atomic StreamK and persistent DataParallel, and migrate
sk_queue_selfclean.yaml to the named policy parameters.

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

dsaffars added a commit that referenced this pull request Oct 28, 2026
…dump (#14195)

JIRA ID: AIHPBLAS-5250

## Motivation

Dict-format library logic keeps a stable solution key order after a
Tensile merge and after geko writes a library. `SolutionIndex`,
`SolutionUID` when present, `KernelNameMin`, and `SolutionNameMin` come
first. The remaining keys are alphabetical. `InternalSupportParams` is
sorted, and `KernArgsVersion` is omitted.

The Tensile merge reorder is required after `812bcc660b`
(`refactor(tensilelite): separate persistent execution policies`,
#12537). That commit makes `reNameSolutions` rebuild every solution with
`normalize_execution_policy_with_defaults`, `sol.clear()`, and
`sol.update()`. The rebuilt insertion order is what the merge writer
saves. geko also assigns keys in place, such as `StaggerUMapping` when
`StaggerU` is 0 and `ActivationFuncCall`, and its writer emits that
insertion order, so those keys were landing at the end of the solution.

## Technical Details

- `TensileMergeLibrary.avoidRegressions` reorders every merged solution
with `LibraryIO.reorderSolutionDictForDictMerge` after
`removeDefaultInitParams` and before the YAML write.
- geko `Library.dump` applies the same layout through a local helper,
`_reorder_dict_solution`, so `libs`, `final_libs`, and `full_libs` match
merge output. Dump does not import `LibraryIO`.

## Test Plan

- `test_merge_writes_dict_solutions_in_canonical_order` merges shuffled
dict solutions and checks the lead keys, the alphabetical tail, and
sorted `InternalSupportParams` without `KernArgsVersion`.
- `test_dict_dump_orders_appended_solution_keys` checks a shuffled geko
solution that includes `SolutionUID`.
- `test_dict_library_supports_property_access_and_dump` checks the lead
order when `SolutionUID` is absent.
- Merge a real gfx950 logic file and confirm each solution starts with
the naming keys and the rest is alphabetical.

## Test Result

The new unit tests passed. A merged gfx950 library with 382 solutions
was checked separately: every solution had the naming keys first and the
remaining keys alphabetical.
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.

5 participants