Repository navigation
refactor(tensilelite): migrate persistent configuration names - #12540
Merged
Merged
Conversation
1 task
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
ThanHenderson
changed the title
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
jaopaulolc
added
the
rocm:gemm algos
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 25, 2026 21:47
54ed170 to
08d956d
Compare
1 task
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 25, 2026 21:52
08d956d to
2ed536f
Compare
|
🎉 All checks passed! This PR is ready for review. |
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 25, 2026 22:02
2ed536f to
1be26e9
Compare
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## users/nhenders/persistent-loop-03-native-dp #12540 +/- ##
============================================================================
Coverage 70.76% 70.76%
============================================================================
Files 2944 2944
Lines 484906 484959 +53
Branches 69944 69960 +16
============================================================================
+ Hits 343100 343156 +56
+ Misses 117935 117930 -5
- Partials 23871 23873 +2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
2 times, most recently
from
September 28, 2026 18:21
286d5d3 to
795b479
Compare
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 29, 2026 15:46
795b479 to
612aad3
Compare
ThanHenderson
marked this pull request as ready for review
September 29, 2026 16:12
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 29, 2026 17:34
612aad3 to
9838890
Compare
1 task done
caio96
reviewed
Sep 30, 2026
Contributor
|
User-facing docs for the migrated names weren't updated, and now disagree with
Worth a follow-up doc pass before the new names are the only ones anyone should reach for. |
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 30, 2026 23:37
77734eb to
e92b5a4
Compare
Contributor
Author
|
fixed in 62d1b6f: updated the environment reference, Stream-K guide, benchmark README, and changelog with the new names, aliases, precedence, and Hybrid debug override. |
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
September 30, 2026 23:41
e92b5a4 to
62d1b6f
Compare
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
October 1, 2026 02:10
62d1b6f to
dda19c0
Compare
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
October 1, 2026 02:27
dda19c0 to
55cc5d7
Compare
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
October 1, 2026 14:40
55cc5d7 to
847bbe2
Compare
jaopaulolc
approved these changes
Oct 1, 2026
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
October 1, 2026 19:28
847bbe2 to
834c085
Compare
bnemanich
pushed a commit
that referenced
this pull request
Oct 1, 2026
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](#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](#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 - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## 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
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
2 times, most recently
from
October 1, 2026 23:54
676473b to
38a2ecb
Compare
Base automatically changed from
users/nhenders/persistent-loop-03-native-dp
to
develop
October 2, 2026 00:21
bnemanich
pushed a commit
that referenced
this pull request
Oct 2, 2026
) JIRA ID : AIHPBLAS-4952 ## Motivation DataParallel assigns complete output tiles to each workgroup, yet its persistent implementation still carries StreamK iteration bookkeeping and scheduling arguments. Use a dedicated tile cursor and a smaller scheduling payload so whole-tile execution no longer depends on the split-reduction loop. ## Technical Details - Traverse whole tiles with `NextTile += PersistentGrid`, replacing StreamK's iteration cursor and tile-index division. Compute the total tile count inside the kernel. Pass `ItersPerTile` for the full-tile K loop and `PersistentGrid` in a versioned two-word scheduling payload; preserve the older payload for existing prebuilt kernels. - Track assignment reservation separately from whether prefetch loads have been issued. Reuse reserved queue work or exhaustion results without another queue pop, and skip initial loads only when prefetch loads were issued. This shared bookkeeping also changes queue/hybrid StreamK prefetch paths; existing waits and barriers still handle load completion. - Align ordinary and custom-kernel argument packing and launch geometry, including spatial clusters. Extend launch-bound regressions to both version-0 and version-1 scheduling payloads. - Update six gfx1250 saved-result cases under ADR0015: 36 kernel names reflect the new payload, which removes five scheduling arguments, adds `PersistentGrid`, and reduces argument storage by 16 bytes. Stack 3/5. Depends on [PR #12538](#12538). Followed by [PR #12540](#12540). ## Test Plan Run DataParallel assignment, prefetch, and argument-layout tests. Check complete normal/custom argument buffers and launch dimensions. Exercise zero-K, alpha-zero, batching, and edge tiles, plus static, dynamic-queue, and hybrid execution on hardware. ## Test Result Rebased onto `develop` at `7affebf2fef`, including the PR1 serialization-test cleanup and upstream StreamK partial-store fix. The full Python unit suite at stack head `795b4796ff0` passed 9,120 tests, with 47 skips, 2 expected failures, 1 unexpected pass, and all 786 snapshots passing. This includes persistent lifecycle, native DataParallel, migrated-configuration, and all 20 adaptive-GSU generation-and-assembly cases. At that stack head, the complete MessagePack host test executable built with `HIPBLASLT_ENABLE_YAML=OFF`; 138 selected tests passed, including the new upstream prediction-library tests. One existing RangeLibrary test was skipped for AIHPBLAS-3506. These are stack-head results. PR2–PR4 own patches are unchanged by the rebase. No snapshots were updated and no performance measurements were taken. Automated checks for the updated heads 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 - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level High (4/5): changes generated DataParallel tile traversal and its internal, versioned kernel payload, plus shared prefetch reservation bookkeeping used by queue/hybrid StreamK. gfx950 numerical checks and host layout/launch checks passed. The rating reflects the scope of these behavior changes and missing gfx1250 device coverage; it does not identify an unresolved correctness finding or a public hipBLASLt API change.
ThanHenderson
force-pushed
the
users/nhenders/persistent-loop-04-migration
branch
from
October 2, 2026 00:24
38a2ecb to
df134e3
Compare
bnemanich
pushed a commit
that referenced
this pull request
Oct 2, 2026
…stent tiles (#12817) JIRA ID : AIHPBLAS-3929 ## Motivation On gfx1250, TensileLite lets neighboring GPU workgroups share matrix-input loads through multicast. DataParallel previously limited these workgroups to one complete output tile each. This change lets each cluster process successive blocks of tiles while keeping its members synchronized. ## Technical Details - For `DataParallel` with `StaticGrid`, assign whole blocks of tiles matching `ClusterDim` to each cluster. Round the launch grid to complete clusters and reject padded tile counts exceeding the 32-bit scheduling limit. - Keep members beyond the matrix edge involved in shared loads by loading the nearest valid tile and suppressing their output writes. - Synchronize before reusing shared memory for another block. Keep operand and sparse-metadata multicast masks live across tiles, and balance synchronization when `K=0` or `alpha=0` skips the summation. Prefetch between output tiles (`PrefetchAcrossPersistent`) still requires one block per cluster, so clusters do not repeat in that configuration. Operand reuse (`ReuseAcrossPersistent`) and tile remapping (`SpaceFillingAlgo`) are rejected for this path. StreamK clustered reduction retains its existing scheduling. Stack 5/5, based on [PR #12540](#12540). Fixes scheduling and synchronization in the persistent DataParallel multicast support introduced by [PR #10189](#10189). ## Test Plan - Run the full `tox` suite from `tensilelite/` on gfx1250, covering Python/common tests, rocisa tests, and lint. - Run `data_parallel_cluster_multicast_boundaries.yaml` with `TENSILE_PERSISTENT_FIXED_GRID=4` and alpha values 1 and 0 to check repeated tiles, uneven batched boundaries, and skipped summations. ## Test Result Local validation on gfx1250: - Full `tox` functional suite: 10,186 passed, 0 failed, 1,758 skipped; 786 snapshots passed. - rocisa: 126 passed, 4 skipped. - Lint still fails, including 8 new unused-import diagnostics. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes GPU launch geometry, cluster synchronization, and register lifetimes, which can affect correctness and performance.
Contributor
TheRock Submodule Bump ActivityNewest first
|
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 : AIHPBLAS-4953
Motivation
Maintained configurations and runtime controls still use StreamK names for settings shared by persistent DataParallel and StreamK execution. Migrate them to named policies and persistent controls so configurations describe the selected behavior consistently.
Technical Details
HybridAssignmentPolicyclient/benchmark options. Preserve public hipBLASLt enum values and existing kernel identities. Existing serialized inputs and old runtime option spellings remain accepted at input boundaries.--global-parameters 'StreamKHybridMode=[1]'requestsDynamicWorkQueueeven when YAML uses the newer spelling. Preserve scalar/list/tuple inputs and reject conflicting aliases supplied together.Stack 4/5. Depends on PR #12539. Followed by PR #12817.
Test Plan
Compare parsed configuration content across the complete migrated corpus and emitted code for representative kernels. Test environment precedence, client parsing, and the actual CLI-to-client configuration path. Build the MessagePack host backend, hipBLASLt library, and benchmark client.
Test Result
Rebased onto
developat7affebf2fef, including the PR1 serialization-test cleanup and upstream StreamK partial-store fix. The full Python unit suite at stack head795b4796ff0passed 9,120 tests, with 47 skips, 2 expected failures, 1 unexpected pass, and all 786 snapshots passing. This includes persistent lifecycle, native DataParallel, migrated-configuration, and all 20 adaptive-GSU generation-and-assembly cases.At that stack head, the complete MessagePack host test executable built with
HIPBLASLT_ENABLE_YAML=OFF; 138 selected tests passed, including the new upstream prediction-library tests. One existing RangeLibrary test was skipped for AIHPBLAS-3506. These are stack-head results. PR2–PR4 own patches are unchanged by the rebase. No snapshots were updated and no performance measurements were taken.Automated checks for the updated heads 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): broad input and runtime-name migration with compatibility-sensitive precedence rules. Corpus equivalence, alias tests, and unchanged emitted-code checks cover the intended behavior.