Skip to content

StreamK SGPR Reduction#592

Merged
aliry95amd merged 7 commits intoROCm:developfrom
aliry95amd:StreamK_SGPR_MagicDivision
Jul 31, 2025
Merged

StreamK SGPR Reduction#592
aliry95amd merged 7 commits intoROCm:developfrom
aliry95amd:StreamK_SGPR_MagicDivision

Conversation

@aliry95amd
Copy link
Contributor

@aliry95amd aliry95amd commented Jul 10, 2025 •
edited
Loading

This PR:

  • removes magic division in StreamK and frees up 6 related SGPRs
  • packs skGrid and skTile into one SGPR and free up 1 SGPR

@aliry95amd aliry95amd force-pushed the StreamK_SGPR_MagicDivision branch from b8d8449 to 3cdf360 Compare July 10, 2025 20:58
@msujon-AMD
Copy link
Collaborator

msujon-AMD commented Jul 10, 2025

Since it changes the ABI of assembly code, we have to change the custom kernel as well, right?

@aliry95amd
Copy link
Contributor Author

aliry95amd commented Jul 10, 2025

@msujon-AMD Yeah, that's true! I have been discussing it with Brian to see what our options are.
There is also another PR #520 that requires change to custom kernels.

babakpst
babakpst requested changes Jul 10, 2025
View reviewed changes
Copy link
Contributor

@babakpst babakpst left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

until we fix custom kernels

ammallya pushed a commit that referenced this pull request Jul 22, 2025
@jsandham
Update changelog and version number (#592)
a4cc843 
* Update changelog and version number

[ROCm/hipSPARSE commit: 574e846]
@aliry95amd aliry95amd requested a review from a team as a code owner July 24, 2025 22:55
@aliry95amd aliry95amd changed the title Replace Magic Division Algorithm Used for StreamK StreamK SGPR Reduction Jul 24, 2025
@aliry95amd aliry95amd mentioned this pull request Jul 24, 2025
Closed
@aliry95amd
Copy link
Contributor Author

aliry95amd commented Jul 25, 2025 •
edited
Loading

GFX950 hipblaslt-test and tox results:

[==========] 20379 tests from 12 test suites ran. (2023285 ms total)
[  PASSED  ] 20379 tests.
hipBLASLt version: 100100
hipBLASLt git version: 63af362792
command line: ./hipblaslt-test 


=========== 87 passed, 74 skipped, 60 warnings in 6337.28s (1:45:37) ===========
___________________________________ summary ____________________________________
  py3: commands succeeded
  congratulations :)

@aliry95amd aliry95amd requested a review from babakpst July 25, 2025 18:51
AlexBrownAMD
AlexBrownAMD approved these changes Jul 28, 2025
View reviewed changes
@b-shi b-shi mentioned this pull request Jul 29, 2025
Closed
@aliry95amd aliry95amd merged commit 58d4aff into ROCm:develop Jul 31, 2025
10 of 12 checks passed
assistant-librarian bot pushed a commit to ROCm/hipBLASLt that referenced this pull request Jul 31, 2025
@aliry95amd @assistant-librarian
[rocm-libraries] ROCm/rocm-libraries#592 (commit 58d4aff)
fe4b4c1 
StreamK SGPR Reduction

This PR:
- removes magic division in StreamK and frees up 6 related SGPRs
- packs skGrid and skTile into one SGPR and free up 1 SGPR
@b-shi b-shi mentioned this pull request Aug 13, 2025
Merged
b-shi added a commit that referenced this pull request Aug 14, 2025
@b-shi
[hipblaslt] Fix typo in custom bbs tn kernel (#1194)
80c9cb1 
Bring in some SK related tile calculations from
#592. This change was missed.
Added tensilelite tests.
assistant-librarian bot pushed a commit to ROCm/hipBLASLt that referenced this pull request Aug 14, 2025
@b-shi @assistant-librarian
[rocm-libraries] ROCm/rocm-libraries#1194 (commit 80c9cb1)
813d08a 
[hipblaslt] Fix typo in custom bbs tn kernel

Bring in some SK related tile calculations from
ROCm/rocm-libraries#592. This change was missed.
Added tensilelite tests.
@aliry95amd aliry95amd deleted the StreamK_SGPR_MagicDivision branch August 26, 2025 23:05
aferoz21 pushed a commit that referenced this pull request Sep 18, 2025
@aliry95amd @b-shi @aferoz21
StreamK SGPR Reduction (#592)
7bc55ed 
This PR:
- removes magic division in StreamK and frees up 6 related SGPRs
- packs skGrid and skTile into one SGPR and free up 1 SGPR

---------

Co-authored-by: b-shi <brianshi@amd.com>
bsyrowik pushed a commit that referenced this pull request Oct 27, 2025
@dependabot @bsyrowik
Bump rocm-docs-core from 1.22.0 to 1.25.0 in /docs/sphinx (#592)
8944076 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.22.0 to 1.25.0.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/v1.25.0/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v1.22.0...v1.25.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-version: 1.25.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/rocWMMA commit: e94dab6]
bsyrowik pushed a commit that referenced this pull request Oct 27, 2025
@dependabot @bsyrowik
Bump rocm-docs-core from 1.22.0 to 1.25.0 in /docs/sphinx (#592)
a928836 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.22.0 to 1.25.0.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/v1.25.0/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v1.22.0...v1.25.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-version: 1.25.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/rocWMMA commit: e94dab6]
ammallya pushed a commit that referenced this pull request Oct 27, 2025
@dependabot
Bump rocm-docs-core from 1.22.0 to 1.25.0 in /docs/sphinx (#592)
e94dab6 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.22.0 to 1.25.0.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/v1.25.0/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v1.22.0...v1.25.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-version: 1.25.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
ammallya pushed a commit that referenced this pull request Oct 28, 2025
@dependabot
Bump rocm-docs-core from 1.22.0 to 1.25.0 in /docs/sphinx (#592)
f8d0a5d 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.22.0 to 1.25.0.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/v1.25.0/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v1.22.0...v1.25.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-version: 1.25.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/rocwmma commit: e94dab6]
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@babakpst babakpst babakpst approved these changes

@AlexBrownAMD AlexBrownAMD AlexBrownAMD approved these changes

@jichangjichang jichangjichang Awaiting requested review from jichangjichang

@KKyang KKyang Awaiting requested review from KKyang

@vin-huang vin-huang Awaiting requested review from vin-huang

@imcarsonliao imcarsonliao Awaiting requested review from imcarsonliao

@hcman2 hcman2 Awaiting requested review from hcman2

@Serge45 Serge45 Awaiting requested review from Serge45

@Jinp800125 Jinp800125 Awaiting requested review from Jinp800125

@TonyYHsieh TonyYHsieh Awaiting requested review from TonyYHsieh

@solaslin solaslin Awaiting requested review from solaslin

@aazz44ss aazz44ss Awaiting requested review from aazz44ss

@davidd-amd davidd-amd Awaiting requested review from davidd-amd

@bnemanich bnemanich Awaiting requested review from bnemanich

@msujon-AMD msujon-AMD Awaiting requested review from msujon-AMD

Assignees

No one assigned

Labels

project: hipblaslt

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

5 participants

@aliry95amd @msujon-AMD @babakpst @AlexBrownAMD @b-shi