[rocm-libraries] ROCm/rocm-libraries#12201 (commit 1630778)

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.

Author: ozturkosu

dispatcher/codegen/grouped_conv/grouped_config_rules_full.py

@@ -493,7 +493,7 @@ def get_wave_warp_pairs(
             result.append((wave, mn))
 
     # If no curated pairs survived (e.g. for architectures whose supported wave
-    # combos don't match the CDNA-derived curated strategies, such as rdna4/gfx1250),
+    # combos don't match the CDNA-derived curated strategies, such as gfx1250),
     # fall back to all tile_math-valid pairs so those arches still get kernels.
     if not result and tm_pairs:
         for wave, wt in tm_pairs:

dispatcher/tests/test_batched_bridge.py

@@ -407,7 +407,7 @@ def test_uneven_split_is_not_flagged_here(self):
         self.assertTrue(_repeat_ok(192, 192, 4, 4, 32, 32))
 
 
-# --- gfx1250 (MI400 / RDNA4-WMMA) enablement -------------------------------
+# --- gfx1250 (CDNA5, WMMA) enablement --------------------------------------
 # The batched-GEMM bridge historically allow-listed only CDNA (gfx90a/942/950,
 # MFMA). gfx1250 uses WMMA, so it needs an arch-tuple entry and a WMMA CI config
 # (warp_tile 16x16x32 -- the CDNA MFMA 32x32x16 tile does not run on gfx1250).

dispatcher/tests/test_batched_contraction_bridge.py

@@ -214,7 +214,7 @@ def test_expand_dedup_and_valid(self):
         self.assertEqual(len(cfgs), 2 * 2)  # (tile_m 2) x (pipeline 2)
 
 
-# --- gfx1250 (MI400 / RDNA4-WMMA) enablement -------------------------------
+# --- gfx1250 (CDNA5, WMMA) enablement --------------------------------------
 # The batched-contraction bridge historically allow-listed only CDNA
 # (gfx90a/942/950, MFMA). gfx1250 uses WMMA, so it needs an arch-tuple entry and
 # WMMA CI configs (warp_tile 16x16x32 -- the CDNA MFMA 32x32x64/32x32x16 tiles do

dispatcher/tests/test_gemm_utils.py

@@ -287,7 +287,7 @@ def test_multi_abd_runner_has_no_stray_multi_d_code(self):
     unittest.main()
 
 
-# --- gfx1250 (MI400 / RDNA4-WMMA) enablement -------------------------------
+# --- gfx1250 (CDNA5, WMMA) enablement --------------------------------------
 # The regular-GEMM bridge historically allow-listed only CDNA (gfx90a/942/950)
 # and carried FNUZ-only fp8 codecs. gfx1250 uses WMMA + OCP fp8, so it needs an
 # arch entry and an OCP codec path. These CPU-only tests lock that surface in.
@@ -335,7 +335,7 @@ def test_use_ocp_fp8_is_callable(self):
         self.assertIsInstance(_use_ocp_fp8(), bool)
 
 
-# --- grouped GEMM on gfx1250 (MI400) ---------------------------------------
+# --- grouped GEMM on gfx1250 -----------------------------------------------
 # The grouped bridge (#9000) shares the arch gate + WMMA warp tiles with the
 # regular bridge, so it runs on gfx1250 once gfx1250 is enabled. Its CI config
 # already uses the gfx1250-valid WMMA warp tile 16x16x32. These CPU-only checks

dispatcher/tests/test_multi_d_bridge.py

@@ -172,14 +172,14 @@ def test_configs_are_valid_sweeps(self):
 
 
 class TestMultiDGfx1250Config(unittest.TestCase):
-    """The gfx1250 (MI400) CI config must use WMMA warp tiles.
+    """The gfx1250 CI config must use WMMA warp tiles.
 
-    gfx1250 has no MFMA units: it runs the RDNA4 WMMA path whose fp16/bf16 warp
+    gfx1250 has no MFMA units: it runs the WMMA path whose fp16/bf16 warp
     tile is 16x16x32 (see arch_specs_generated.py). The merged multi_d bridge
     (#9308) shipped only an MFMA CI config (32x32x16, valid on gfx942/gfx950),
     which the kernel reports as unsupported (status -2/-1) on gfx1250. The
     gfx1250 CI config therefore pins the WMMA warp tile so the sweep produces
-    kernels that actually run on MI400.
+    kernels that actually run on gfx1250.
     """
 
     _GFX1250_CONFIG = _CONFIG_DIR / "default_ci_config_gfx1250.json"

tile_engine/ops/gemm/batched_contraction/configs/bridge_default_ci_config_gfx1250.json

@@ -1,5 +1,5 @@
 {
-    "_comment": "CI bridge sweep for batched_contraction on gfx1250 (MI400, RDNA4/WMMA): one small validated fp16 point. WMMA warp_tile 16x16x32 (the CDNA MFMA 32x32x16 does not run on gfx1250).",
+    "_comment": "CI bridge sweep for batched_contraction on gfx1250 (CDNA5/WMMA): one small validated fp16 point. WMMA warp_tile 16x16x32 (the CDNA MFMA 32x32x16 does not run on gfx1250).",
     "tile_config": {
         "tile_m": {"values": [128]},
         "tile_n": {"values": [128]},