Add Stream K experimental support for RDNA by stsokolo · Pull Request #1554 · ROCm/rocm-libraries

stsokolo · 2025-09-11T13:58:31Z

This PR introduces initial changes needed for stream-k RDNA support. Some experimental tuning was conducted for a specific model, and it is working. Now we are conducting additional tuning and doing performance analysis.

Major changes:

The s_store instruction needed to be swapped with buffer_store since s_store is not supported on RDNA
Additional modifiers were added to rocisa that were missing for gfx11 and gfx12 and an "off" option was added for vgpr
Add gfx1100 and gfx1201 architectures to Origami Analytics

cliffxzx

LGTM

cmingch

LGTM, please wait for CI to pass

chiranjeevipattigidi · 2025-10-03T13:29:29Z

!verify

fjankovi · 2025-10-22T09:52:36Z

@AlexBrownAMD @cmingch @ROCm/hipblaslt-reviewers Can we get another review / approval for this PR? Thanks.

fjankovi · 2025-10-23T15:42:10Z

Can we please get a review from @ROCm/origami-reviewers? Thanks.
CC @neoblizz

neoblizz

Origami pieces LGTM!

stsokolo requested a review from a team as a code owner September 11, 2025 13:58

github-actions Bot added the project: hipblaslt label Sep 11, 2025

stsokolo changed the title ~~Initial Stream-K changes for RDNA support~~ Add Stream K support for RDNA Sep 11, 2025

assistant-librarian Bot added the organization: ROCm label Sep 11, 2025

fjankovi mentioned this pull request Sep 11, 2025

Initial Stream-K changes for Navi support #1115

Closed

stsokolo changed the title ~~Add Stream K support for RDNA~~ Add Stream K experimental support for RDNA Sep 11, 2025

aliry95amd reviewed Sep 11, 2025

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/src/analytical/Hardware.cpp Outdated

cmingch reviewed Sep 12, 2025

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/src/analytical/Hardware.cpp Outdated

Comment thread projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py Outdated

Comment thread projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py Outdated

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from 31f57ea to 86d94ee Compare September 12, 2025 09:00

cmingch requested a review from cliffxzx September 12, 2025 09:11

fjankovi requested a review from AlexBrownAMD September 12, 2025 14:19

cliffxzx reviewed Sep 15, 2025

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx12/sk_bgemm_quick.yaml

Comment thread projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp Outdated

stsokolo force-pushed the users/stsokolo/navi_sk_support branch 2 times, most recently from 2e259e9 to 814b5df Compare September 22, 2025 09:46

github-actions Bot added the shared: origami label Sep 22, 2025

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from 814b5df to 3744e4a Compare September 23, 2025 09:38

cliffxzx previously approved these changes Sep 23, 2025

View reviewed changes

stsokolo requested review from aliry95amd and cmingch September 24, 2025 13:51

stsokolo dismissed cliffxzx’s stale review via e6e67a8 September 30, 2025 10:12

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from 3744e4a to e6e67a8 Compare September 30, 2025 10:12

stsokolo requested a review from a team as a code owner September 30, 2025 10:12

cmingch reviewed Oct 1, 2025

View reviewed changes

Comment thread shared/origami/include/origami/hardware.hpp

cmingch requested a review from a team October 1, 2025 04:26

cmingch approved these changes Oct 1, 2025

View reviewed changes

jichangjichang reviewed Oct 1, 2025

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py Outdated

stsokolo force-pushed the users/stsokolo/navi_sk_support branch 2 times, most recently from ae9c6ee to 1559492 Compare October 3, 2025 10:10

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from 1559492 to 98cb6d1 Compare October 16, 2025 09:03

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from 2f2fb4b to f8383aa Compare October 16, 2025 16:18

daineAMD reviewed Oct 22, 2025

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py

daineAMD approved these changes Oct 23, 2025

View reviewed changes

neoblizz approved these changes Oct 27, 2025

View reviewed changes

stsokolo force-pushed the users/stsokolo/navi_sk_support branch from f8383aa to 93a8907 Compare October 27, 2025 09:34

stsokolo merged commit b4c7c8a into develop Oct 27, 2025
25 of 28 checks passed

stsokolo deleted the users/stsokolo/navi_sk_support branch October 27, 2025 13:46

stsokolo added 20 commits October 27, 2025 18:20

[hipBLASLt] Add initial changes for Navi SK support

ff4360b

[hipBLASLt] Update gfx1100/gfx1201 logic yaml

e943ae9

[hipBLASLt] Change Tensile version

fe4b474

[hipBLASLt] Add cache scope modifier

c9bb8eb

[hipBLASLt] Add simple hgemm and bgemm SK test configs

6a0d070

[hipBLASLt] Fix SWaitCnt

d4c7cc6

[hipBLASLt] Swap global store with buffer store instruction

89c1f07

[hipBLASLt] Fix typo

69c9868

[hipBLASLt] Minor change to MUBUFModifiers in rocisa

2058396

[hipBLASLt] Add dlc modifier to GlobalWrite

03f64b2

[hipBLASLt] Add off option to vgpr in rocisa

5a08f17

[hipBLASLt] Update container tests in rocisa

c66710c

[hipBLASLt] Remove gfx1100 and gfx1201 SK logic libraries

8a4325c

[hipBLASLt] Update SK test configs for gfx1100 and gfx1201

caed531

[hipBLASLt] Fix gfx1100 SK test config

38491a3

[hipBLASLt] Remove v_swmmac matrix instructions for gfx1201

1f50ddb

[hipBLASLt] Add s_store_b32 case in hardware caps

c044d4a

[hipBLASLt] Fix hardware caps s_store instruction

e6bf83f

[hipBLASLt] Solution.py cleanup

a3f5f02

[hipBLASLt] Fix modifiers for gfx942 and gfx950 in rocisa

93a8907

Rohan138 mentioned this pull request May 4, 2026

[Issue][hipblaslt]: enum mismatch between rocisa::DataType and origami::data_type_t in ROCm 7.2 #7056

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add Stream K experimental support for RDNA#1554

Add Stream K experimental support for RDNA#1554
stsokolo merged 20 commits into
developfrom
users/stsokolo/navi_sk_support

stsokolo commented Sep 11, 2025 •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

cliffxzx left a comment

Uh oh!

Uh oh!

cmingch left a comment

Uh oh!

Uh oh!

chiranjeevipattigidi commented Oct 3, 2025

Uh oh!

fjankovi commented Oct 22, 2025

Uh oh!

Uh oh!

fjankovi commented Oct 23, 2025 •

edited

Loading

Uh oh!

neoblizz left a comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

Conversation

stsokolo commented Sep 11, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

cliffxzx left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

cmingch left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

chiranjeevipattigidi commented Oct 3, 2025

Uh oh!

fjankovi commented Oct 22, 2025

Uh oh!

Uh oh!

fjankovi commented Oct 23, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

neoblizz left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

stsokolo commented Sep 11, 2025 •

edited

Loading

fjankovi commented Oct 23, 2025 •

edited

Loading