Skip to content

feat(ck): enable gfx1170 and gfx1171 - #13323

Closed
stsokolo wants to merge 1 commit into
developfrom
users/stsokolo/gfx117x-ck
Closed

stsokolo wants to merge 1 commit into
developfrom
users/stsokolo/gfx117x-ck

Conversation

@stsokolo

@stsokolo stsokolo commented Oct 8, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID: ROCM-21096, AIRADSW-699, AIRADSW-588

Motivation

Enable Composable Kernel on gfx1170 and gfx1171.

Technical Details

Register gfx1170 and gfx1171 as Composable Kernel GPU targets. Use the gfx11 buffer-resource setting for those targets, and enable the WMMA path for gfx1170 and gfx1171.

Test Plan

  • Build Composable Kernel with gfx1170 and gfx1171 in the GPU target list.
  • Run the Composable Kernel tests on gfx117x.

Test Result

Pending.

Submission Checklist

@stsokolo
stsokolo requested a review from a team as a code owner October 8, 2026 16:25
@therock-pr-bot

therock-pr-bot Bot commented Oct 8, 2026 •

Copy link
Copy Markdown

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/composablekernel/include/ck/ck.hpp, projects/composablekernel/include/ck/host_utility/device_prop.hpp, projects/composablekernel/include/ck_tile/core/config.hpp; no test file found
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

📖 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 Oct 8, 2026 •

Copy link
Copy Markdown

⏳ Waiting for required checks to start or finish:

  • ⏳ pre-commit

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

The current PR isn't including support to contractions.

Besides the change I've suggested in ck.hpp, there is another required change in amd_wmma.hpp:

diff --git a/projects/composablekernel/include/ck/utility/amd_wmma.hpp b/projects/composablekernel/include/ck/utility/amd_wmma.hpp
index ed2c8a2226..8028278002 100644
--- a/projects/composablekernel/include/ck/utility/amd_wmma.hpp
+++ b/projects/composablekernel/include/ck/utility/amd_wmma.hpp
@@ -12,7 +12,8 @@ namespace ck {

 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
     defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
-    defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx11_generic__)
+    defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx1170__) || \
+    defined(__gfx1171__) || defined(__gfx11_generic__)
 #define __gfx11__
 #endif

#endif
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx11_generic__)

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.

The current implementation doesn't add gfx1170 support for the contraction target. Define the new archs as part of __gfx11__ to add the missing support:

Suggested change
defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx1170__) || \
defined(__gfx1171__) || defined(__gfx11_generic__)

@slojosic-amd slojosic-amd Oct 9, 2026 •

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.

Based on this: https://github.com/ROCm/rocm-systems/pull/11977/changes#diff-18ce3ca3324bd340f932d819048cbc60660051cac2d62705529379b3f4bbe2fbR455 gfx117X is defined as gfx11-7-generic so I am not sure that we really should add gfx117X targets under __gfx11__

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.

In this case, they should be defined gfx11-7-generic in here, and in projects/composablekernel/include/ck/utility/amd_wmma.hpp the gfx117 should properly map to the valid wmma operations.

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.

I'll remove my blocking review from this PR, but without the proper mapping for the WMMAs, the contraction kernels won't work properly.

@illsilin

illsilin commented Oct 8, 2026

Copy link
Copy Markdown
Contributor

Not sure who authorized this. I expected the gfx117 stuff to be still in the internal repo. This PR clearly won't build correctly.

@slojosic-amd

Copy link
Copy Markdown
Contributor

Not sure who authorized this. I expected the gfx117 stuff to be still in the internal repo. This PR clearly won't build correctly.

All enablement job has been done in the internal repo and we were merging develop branch into private branch regularly. After adding gfx117X support inside rocm-systems: ROCm/rocm-systems#11977 we got approval from management and BU to continue with adding gfx117X support to rocm-libraries

@stsokolo
stsokolo force-pushed the users/stsokolo/gfx117x-ck branch from 5384357 to 6c2cf93 Compare October 9, 2026 12:28

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

Removing the blocking review, but please, add the missing support so hipTensor can also be enabled for gfx1170.

@evedovelli
evedovelli dismissed their stale review October 9, 2026 13:20

Removing the blocking review, but please, add the missing support so hipTensor can also be enabled for gfx1170.

@illsilin

illsilin commented Oct 9, 2026

Copy link
Copy Markdown
Contributor

Not sure who authorized this. I expected the gfx117 stuff to be still in the internal repo. This PR clearly won't build correctly.

All enablement job has been done in the internal repo and we were merging develop branch into private branch regularly. After adding gfx117X support inside rocm-systems: ROCm/rocm-systems#11977 we got approval from management and BU to continue with adding gfx117X support to rocm-libraries

Have you actually tried building this branch for gfx117?

@illsilin

Copy link
Copy Markdown
Contributor

Created a new PR #13696, made sure it actually builds for gfx1170.

@illsilin illsilin closed this Oct 15, 2026
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.

4 participants