Skip to content

docs(ck-tile): describe gfx1250 as CDNA5 rather than RDNA4 - #12201

Merged
ozturkosu merged 2 commits into
developfrom
users/muozturk/ck/gfx1250-cdna5-naming
Sep 17, 2026
Merged

ozturkosu merged 2 commits into
developfrom
users/muozturk/ck/gfx1250-cdna5-naming

Conversation

@ozturkosu

Copy link
Copy Markdown
Contributor

JIRA ID : AICK-2249

Summary

gfx1250 is labelled RDNA4 in several comments and one CI config string. The product generation is CDNA5. This corrects the label.

Comments and one JSON _comment only — no code, no behaviour change.

Raised in review of #12173, where the same value was corrected in arch_specs.json. Split out because renaming terminology across the tree does not belong in an LDS-capacity change.

Why the RDNA4 label was there

It was not arbitrary. gfx1250 has no MFMA units and multiplies through WMMA, which is an RDNA-family instruction, and the gfx12xx numbering points the same way. The observation about the instruction path was right; using it as a generation label was not.

arch.hpp already draws exactly this distinction:

GFX1250 is its own standalone family. Although it shares the RDNA architecture with the GFX12 family, its MMA builtins and data-type ABI differ, so it must not be treated as a GFX12-family device.

Sharing an ISA lineage is not the same as belonging to that generation. So there are two separate claims in play, and only the second one holds.

Changes

Where a generation label is meant, it now says CDNA5. Where only the instruction path matters, the generation is dropped rather than replaced — a WMMA path is a WMMA path regardless of which generation the part belongs to.

file before after
dispatcher/tests/test_gemm_utils.py gfx1250 (MI400 / RDNA4-WMMA) enablement gfx1250 (MI400 / CDNA5, WMMA) enablement
dispatcher/tests/test_batched_bridge.py same same
dispatcher/tests/test_batched_contraction_bridge.py same same
dispatcher/tests/test_multi_d_bridge.py it runs the RDNA4 WMMA path it runs the WMMA path
dispatcher/codegen/grouped_conv/grouped_config_rules_full.py such as rdna4/gfx1250 such as gfx1250
tile_engine/.../bridge_default_ci_config_gfx1250.json (MI400, RDNA4/WMMA) (MI400, CDNA5/WMMA)

Deliberately not changed

  • include/ck_tile/core/arch/arch.hpp and .../mma/scale/wmma/selector.hpp describe the ISA lineage and the selector strategy as RDNA. Both are accurate — the first is the quote above, the second documents an RDNA-style selector. Renaming these would replace a correct statement with a wrong one.
  • dispatcher/codegen/arch_specs_generated.py still reads rdna4. It is generated, and its source is the gfx1250 entry in arch_specs.json that feat(ck-tile): make dispatcher LDS capacity budget architecture-aware #12173 adds; the value is corrected there. Hand-editing a generated file is the exact trap this codebase keeps hitting, so it is not done here.

That last point means the tree is briefly inconsistent until #12173 lands. The affected field is metadata: nothing branches on its value, and its only consumer is a None check that rejects unknown architectures in strict mode.

Test plan

  • test_gemm_utils, test_batched_bridge, test_batched_contraction_bridge, test_multi_d_bridge, test_grouped_conv_codegen all pass
  • The edited CI config still parses as JSON
  • No gfx1250/MI400 reference to RDNA4 remains outside the two files listed above as intentionally kept

Note for the reviewer

The label originated in #10921, the gfx1250 bridge enablement PR, and spread from there; AICK-2082 uses the same wording. So this is a retroactive terminology correction rather than a one-off typo. If CDNA5 is not the right generation either, say so before this merges — I have taken it from review feedback on #12173 and have not independently confirmed it.

gfx1250 was labelled RDNA4 in several comments and one CI config. The
product generation is CDNA5; the RDNA association comes from the
instruction path, since the target has no MFMA units and multiplies
through WMMA, which is an RDNA-family instruction.

Those are two different claims and only the second one holds. arch.hpp
already draws the distinction: gfx1250 "is its own standalone family",
and "shares the RDNA architecture with the GFX12 family" while differing
in its MMA builtins and data-type ABI. Sharing an ISA lineage is not the
same as belonging to that generation.

Where a generation label is meant, say CDNA5. Where only the instruction
path matters, drop the generation rather than replace it: a WMMA path is
a WMMA path regardless of which generation the part belongs to.

Comments and one CI config string only. No code, no behaviour change.

Deliberately not changed:

- arch.hpp and the scale-WMMA selector describe the ISA lineage and the
  selector strategy as RDNA. Both are accurate and stay as they are.
- arch_specs_generated.py still reads rdna4. It is generated, and its
  source is the gfx1250 entry in arch_specs.json, which is added by the
  LDS capacity change. Editing a generated file by hand is what this
  codebase keeps getting bitten by, so it is corrected at the source
  there instead of here.
@github-actions github-actions Bot added project: composablekernel ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. labels Sep 16, 2026
@therock-pr-bot

therock-pr-bot Bot commented Sep 16, 2026

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass
Forbidden Files ✅ Pass
🧪 Unit Test ✅ Pass
🔎 pre-commit ✅ Pass
🚫 Draft PR 🔜 To Be Enabled
🚩 Feature Flag 🔜 To Be Enabled
📊 Code Coverage 🔜 To Be Enabled
🤖 therock-pr-bot ✅ Pass

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

Comment thread projects/composablekernel/dispatcher/tests/test_batched_bridge.py Outdated
@ozturkosu ozturkosu self-assigned this Sep 17, 2026
Review feedback on #12201: unreleased product names should not appear.
Remove the MI400 tag from the gfx1250 enablement comments and the CI
config _comment, keeping the gfx1250/CDNA5 architecture references.
@ozturkosu
ozturkosu requested a review from andriy-ca September 17, 2026 17:25
@ozturkosu
ozturkosu enabled auto-merge (squash) September 17, 2026 17:25

@andriy-ca andriy-ca 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.

LGTM!

@ozturkosu
ozturkosu merged commit 1630778 into develop Sep 17, 2026
123 checks passed
@ozturkosu
ozturkosu deleted the users/muozturk/ck/gfx1250-cdna5-naming branch September 17, 2026 20:51
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Sep 17, 2026
docs(ck-tile): describe gfx1250 as CDNA5 rather than RDNA4
 (#12201)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

JIRA ID : AICK-2249

## Summary

gfx1250 is labelled **RDNA4** in several comments and one CI config
string. The product generation is **CDNA5**. This corrects the label.

Comments and one JSON `_comment` only — no code, no behaviour change.

Raised in review of #12173, where the same value was corrected in
`arch_specs.json`. Split out because renaming terminology across the
tree does not belong in an LDS-capacity change.

## Why the RDNA4 label was there

It was not arbitrary. gfx1250 has no MFMA units and multiplies through
**WMMA**, which is an RDNA-family instruction, and the `gfx12xx`
numbering points the same way. The observation about the *instruction
path* was right; using it as a *generation* label was not.

`arch.hpp` already draws exactly this distinction:

> GFX1250 is its own standalone family. Although it **shares the RDNA
architecture** with the GFX12 family, its MMA builtins and data-type ABI
differ, so it must not be treated as a GFX12-family device.

Sharing an ISA lineage is not the same as belonging to that generation.
So there are two separate claims in play, and only the second one holds.

## Changes

Where a generation label is meant, it now says CDNA5. Where only the
instruction path matters, the generation is dropped rather than replaced
— a WMMA path is a WMMA path regardless of which generation the part
belongs to.

| file | before | after |
|---|---|---|
| `dispatcher/tests/test_gemm_utils.py` | `gfx1250 (MI400 / RDNA4-WMMA)
enablement` | `gfx1250 (MI400 / CDNA5, WMMA) enablement` |
| `dispatcher/tests/test_batched_bridge.py` | same | same |
| `dispatcher/tests/test_batched_contraction_bridge.py` | same | same |
| `dispatcher/tests/test_multi_d_bridge.py` | `it runs the RDNA4 WMMA
path` | `it runs the WMMA path` |
| `dispatcher/codegen/grouped_conv/grouped_config_rules_full.py` | `such
as rdna4/gfx1250` | `such as gfx1250` |
| `tile_engine/.../bridge_default_ci_config_gfx1250.json` | `(MI400,
RDNA4/WMMA)` | `(MI400, CDNA5/WMMA)` |

## Deliberately not changed

- **`include/ck_tile/core/arch/arch.hpp`** and
**`.../mma/scale/wmma/selector.hpp`** describe the ISA lineage and the
selector strategy as RDNA. Both are accurate — the first is the quote
above, the second documents an RDNA-style selector. Renaming these would
replace a correct statement with a wrong one.
- **`dispatcher/codegen/arch_specs_generated.py`** still reads `rdna4`.
It is generated, and its source is the gfx1250 entry in
`arch_specs.json` that #12173 adds; the value is corrected there.
Hand-editing a generated file is the exact trap this codebase keeps
hitting, so it is not done here.

That last point means the tree is briefly inconsistent until #12173
lands. The affected field is metadata: nothing branches on its value,
and its only consumer is a `None` check that rejects unknown
architectures in strict mode.

## Test plan

- [x] `test_gemm_utils`, `test_batched_bridge`,
`test_batched_contraction_bridge`, `test_multi_d_bridge`,
`test_grouped_conv_codegen` all pass
- [x] The edited CI config still parses as JSON
- [x] No `gfx1250`/`MI400` reference to RDNA4 remains outside the two
files listed above as intentionally kept

## Note for the reviewer

The label originated in #10921, the gfx1250 bridge enablement PR, and
spread from there; AICK-2082 uses the same wording. So this is a
retroactive terminology correction rather than a one-off typo. If
**CDNA5 is not the right generation either**, say so before this merges
— I have taken it from review feedback on #12173 and have not
independently confirmed it.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. organization: ROCm project: composablekernel

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants