Repository navigation
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11374 +/- ##
===========================================
+ Coverage 69.42% 69.47% +0.06%
===========================================
Files 2828 2828
Lines 470113 469724 -389
Branches 69417 69391 -26
===========================================
- Hits 326343 326337 -6
+ Misses 120066 119682 -384
- Partials 23704 23705 +1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
newling
changed the title
newling
marked this pull request as ready for review
August 28, 2026 22:02
newling
force-pushed
the
remove/hipblaslt-extop-amax
branch
from
August 31, 2026 21:10
735d8b4 to
bae10ea
Compare
Contributor
Author
newling
force-pushed
the
remove/hipblaslt-extop-amax
branch
from
September 2, 2026 15:56
bae10ea to
1780bf4
Compare
Remove the exported hipblasltExtAMax API, its generated kernels and library metadata, dedicated clients and samples, and the client-only amaxScaleA/B paths. Preserve fused GEMM AMAX_D_POINTER support and the remaining Softmax and LayerNorm extension operations.
newling
force-pushed
the
remove/hipblaslt-extop-amax
branch
from
September 9, 2026 20:10
1780bf4 to
f75587c
Compare
Contributor
Author
|
Closing in favour of the fix, #11373 |
newling
added a commit
that referenced
this pull request
Sep 15, 2026
JIRA ID: AIHPBLAS-4632
## Summary
AMax and LayerNorm reduce wave-local values through a small LDS scratch
buffer. At each tree level, upper waves publish partial results and
lower waves read those results before the same LDS slots are reused by
the next level.
The existing barrier orders the upper-wave writes before the lower-wave
reads. It does not prevent the next tree level from overwriting those
slots while the preceding reads are still in flight. That creates a
write-after-read race in both generated kernels.
This change routes every participating wave through a second barrier
after the reads complete and before the reduction loop reuses LDS. The
existing public APIs and numerical behavior are unchanged.
A focused generator test covers AMax and LayerNorm. It verifies that the
upper, lower, and inactive-wave paths converge at the LDS-reuse barrier,
and that the lower-wave read occurs between the publication and reuse
barriers. The LayerNorm generator is also installed with the TensileLite
test artifacts so this test works from the installed test tree.
## Validation
rocJITsu reproduced the write-after-read conflict before the fix:
| Operation | Target-kernel dispatches | LDS race reports |
|---|---:|---:|
| AMax | 2 | 2 |
| LayerNorm | 1 | 3 |
A representative pre-fix AMax report shows a lower wave reading the
current partial value while an upper wave begins reusing the same LDS
slot for the next tree level. Intervening address and branch
instructions are elided:
```text
RACE kernel=AMax_Ti_H_To_S_W_256_C_4
type=LDS access=write wave=1 lane=0 wg=0,0,0
Race on LDS byte 0:
==> ds_read_u16 v2, v24 ; wave 0 reads the current partial value
s_waitcnt vmcnt(63) expcnt(7) lgkmcnt(0)
v_max_f16_e32 v1, v1, v2
...
==> ds_write_b16 v24, v1 ; wave 1 reuses the slot for the next level
END_RACE
```
The second barrier ensures the read is retired before that next write
can occur.
After the fix, AMax completed 20 target-kernel dispatches and LayerNorm
completed 10 without a race report. AMax numerical comparisons remained
exact, and LayerNorm produced identical mean and inverse-variance
results across repeated runs.
The focused generator regression test passes for both operations.
Required Math CI also exercises the existing GPU numerical tests: AMax
compares twelve FP16/FP32 input-output and shape combinations against a
CPU reference, while LayerNorm compares output, mean, and inverse
variance against its CPU reference.
The rocJITsu run is retained as validation evidence rather than added as
a permanent shared sidecar workload.
## Alternative
(Decided to reject this alternative, we want to go with this fix) PR
#11374 removes the standalone AMax and LayerNorm operations and their
supporting surfaces. The two pull requests are mutually exclusive; only
one should merge.
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.
Summary
hipBLASLt currently carries standalone AMax and LayerNorm operations outside its GEMM APIs. Each operation has its own exported C entry point, assembly generator, packaged kernels, runtime model, benchmark, tests, documentation, and samples. Public usage searches found no framework consumers; the direct callers surfaced were the project-owned validation code and examples.
This removes both standalone operations as one support decision instead of retaining one side of the same maintenance tradeoff. Softmax remains as the surviving generated extension operation. Fused GEMM
HIPBLASLT_MATMUL_DESC_AMAX_D_POINTERsupport is unchanged.The build also removes stale AMax and LayerNorm generator outputs before rebuilding extension metadata. This prevents an incremental build directory from retaining metadata for kernels that are no longer linked into the extension code object.
Alternative
This is the remove-both alternative. #11373 retains both standalone AMax and LayerNorm and fixes their shared LDS-reuse race. The two pull requests are mutually exclusive and only one should merge.
Compatibility
This removes the exported
hipblasltExtAMaxandhipblasltExtLayerNormC APIs. The corresponding AMax, mixed-precision AMax, and LayerNorm examples inROCm/rocm-examplesneed companion cleanup if this direction is selected.Validation
The hipBLASLt library, test client, surviving extension clients, and gfx1151 extension library build locally. The resulting shared library exports only
hipblasltExtSoftmaxamong these generated extension operations, and the packaged extension metadata contains only Softmax.A stale-metadata incremental-build check confirms that old AMax and LayerNorm generator files are removed before packaging. The surviving Softmax benchmark runs successfully, and all six remaining extension-operation tests pass.