Skip to content

fix(ck-tile): prefer amd-smi for dispatcher GPU-arch detection - #10445

Merged
ozturkosu merged 4 commits into
developfrom
users/muozturk/ck/dispatcher_amd_smi
Aug 11, 2026
Merged

ozturkosu merged 4 commits into
developfrom
users/muozturk/ck/dispatcher_amd_smi

Conversation

@ozturkosu

@ozturkosu ozturkosu commented Aug 5, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Extends the amd-smi-first migration (#9731, #10204) from the Tile Engine to the
dispatcher's GPU-architecture detection.

  • Adds detect_gpu_arch() to the shared tile_engine/ops/common/smi_utils.py
    wrapper: amd-smi static (TARGET_GRAPHICS_VERSION) first, rocm-smi
    --showproductname (GFX Version) fallback — same amd-smi-first policy as the
    rest of the wrapper.
  • dispatcher_common.detect_gpu_arch(): prefers amd-smi via the wrapper, then
    rocminfo, then the caller-supplied default. Also removes a duplicate
    definition of this function.
  • gemm_utils._get_arch(): tries amd-smi via the wrapper before rocminfo,
    preserving the raise-on-unresolved (no silent default) behavior.

Motivation

Per the directive behind #9731, device-property access should prefer
amd-smi (rocm-smi is being phased out). The Tile Engine was already migrated
(#9731, #10204); the dispatcher still detected the GPU arch by shelling out to
rocminfo directly. Tracked in JIRA ROCM-28734.

Design note

The dispatcher reaches the sibling tile_engine/ops/common/smi_utils.py through
a guarded import and falls back to rocminfo if the wrapper is unavailable, so
behavior is unchanged on hosts without amd-smi. (Open to moving the shared
wrapper to a neutral location if reviewers prefer avoiding the dispatcher →
tile_engine reach.)

Test plan (MI350X / gfx950)

  • New mock-based unit tests for detect_gpu_arch() (amd-smi preference,
    rocm-smi fallback, default) + a live check — 4/4 pass.
  • dispatcher_common.detect_gpu_arch() and gemm_utils._get_arch() both
    return gfx950 via amd-smi on the live host.
  • test_gemm_utils 21/21 still pass (no arch-detection regression).

Extends the amd-smi-first migration (PR #9731, #10204) to the dispatcher.
Adds detect_gpu_arch() to the shared smi_utils wrapper (amd-smi static
TARGET_GRAPHICS_VERSION first, rocm-smi --showproductname GFX Version
fallback) and routes the dispatcher arch detection through it:

- dispatcher_common.detect_gpu_arch(): prefer amd-smi via the wrapper, then
  rocminfo, then the caller default. Also removes a duplicate definition of
  this function.
- gemm_utils._get_arch(): try amd-smi via the wrapper before rocminfo,
  preserving the raise-on-unresolved (no silent default) behavior.

The dispatcher reaches the sibling tile_engine/ops/common/smi_utils via a
guarded import and degrades to rocminfo if it is unavailable, so behavior is
unchanged on hosts without amd-smi. Adds mock-based unit tests for the new
detect_gpu_arch() (amd-smi preference, rocm-smi fallback, default) plus a
live check. Verified on MI350X/gfx950: both dispatcher arch helpers return
gfx950 via amd-smi; test_gemm_utils 21/21 still pass.

JIRA: ROCM-28734
@ozturkosu
ozturkosu requested a review from a team as a code owner August 5, 2026 20:23
@ozturkosu ozturkosu self-assigned this Aug 5, 2026
@therock-pr-bot

therock-pr-bot Bot commented Aug 5, 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

therock-pr-bot Bot commented Aug 5, 2026

Copy link
Copy Markdown

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

Copilot AI 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.

Pull request overview

This PR extends the existing “amd-smi first, rocm-smi fallback” policy into GPU-architecture detection for the CK Tile Engine and the dispatcher, centralizing arch detection in smi_utils.detect_gpu_arch() and updating dispatcher-side detection to prefer that wrapper before falling back to rocminfo.

Changes:

  • Added detect_gpu_arch() to tile_engine/ops/common/smi_utils.py and corresponding unit/live tests.
  • Updated dispatcher_common.detect_gpu_arch() to prefer the shared wrapper (and removed a duplicate function definition).
  • Updated gemm_utils._get_arch() to try the shared wrapper before rocminfo, while preserving “raise if unresolved” behavior.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 4 comments.

File Description
projects/composablekernel/tile_engine/ops/common/test_smi_utils.py Adds unit/live tests for detect_gpu_arch() behavior and preference ordering.
projects/composablekernel/tile_engine/ops/common/smi_utils.py Introduces detect_gpu_arch() with amd-smi-first / rocm-smi fallback.
projects/composablekernel/dispatcher/python/gemm_utils.py Tries wrapper-based arch detection before rocminfo in _get_arch().
projects/composablekernel/dispatcher/python/dispatcher_common.py Prefers wrapper-based arch detection and removes duplicate detect_gpu_arch().

@ozturkosu
ozturkosu requested review from geyyer and a balanced review from Copilot August 11, 2026 00:23

This comment was marked as off-topic.

…cher

Addresses review feedback: the previous change only migrated
dispatcher_common.detect_gpu_arch() and gemm_utils._get_arch(), leaving the
busiest entry point (ctypes_utils.detect_gpu_arch(), used by every GEMM
example) and grouped_conv_utils on rocminfo. Routes both through the shared
amd-smi-first smi_utils wrapper (JIRA ROCM-28734) before rocminfo, so the
title's "dispatcher GPU-arch detection" is now true end-to-end.

DRY: the amd-smi bridge now lives once in dispatcher_common
(_detect_gpu_arch_via_amd_smi); gemm_utils/ctypes_utils/grouped_conv_utils
import it instead of re-implementing the sys.path reach.

Copilot review:
- smi_utils.detect_gpu_arch(): add Optional[str] type annotations and skip
  tools missing from PATH via shutil.which (consistent with the rest of the
  module) instead of relying on FileNotFoundError.
- TestDetectGpuArch: force CK_SMI_TOOL off and mock shutil.which so the
  amd-smi-preference assertions are deterministic regardless of ambient env
  or host GPU tooling; add a tools-absent case.
- Annotate the intentional best-effort broad excepts with noqa: BLE001.

Verified: py_compile all six files; TestDetectGpuArch 5/5 pass; live host
(gfx90a) resolves via amd-smi through all four helpers; gemm_utils reuses the
canonical bridge (identity check).
@ozturkosu

ozturkosu commented Aug 11, 2026 •

Copy link
Copy Markdown
Contributor Author

Review self-assessment + follow-up changes (commit 8085f847f3)

I re-reviewed this PR one more than and discovered : it only migrated dispatcher_common.detect_gpu_arch() and gemm_utils._get_arch(), while the busiest entry point — ctypes_utils.detect_gpu_arch(), which every GEMM example calls via default=detect_gpu_arch() — plus grouped_conv_utils were still detecting arch through rocminfo. So the title's claim ("dispatcher GPU-arch detection") wasn't true end-to-end. This commit closes that gap and addresses all Copilot feedback.

What changed

  • Completed the migration (the necessity fix). ctypes_utils.detect_gpu_arch() and grouped_conv_utils.detect_gpu_arch() now prefer amd-smi via the shared smi_utils wrapper before falling back to rocminfo, then the default. All Python arch-detection paths in the dispatcher are now amd-smi-first.
  • DRY (single source of truth). The amd-smi bridge now lives once in dispatcher_common._detect_gpu_arch_via_amd_smi(); gemm_utils, ctypes_utils, and grouped_conv_utils import it instead of re-implementing the sys.path reach. (Previously the helper was copy-pasted in two files — this PR was removing one duplicate while adding another.)
  • Copilot Initial workflows #1 — TestDetectGpuArch is now deterministic: forces CK_SMI_TOOL="" in setUp/tearDown and mocks shutil.which, plus a new tools-absent case.
  • Copilot [hipDNN] Enablement #2 — smi_utils.detect_gpu_arch() gains Optional[str] annotations and a shutil.which presence check consistent with the rest of the module.
  • Copilot Change shared libs name #3/[PR] [hipSPARSELt] Test monorepo #4 — best-effort broad excepts annotated # noqa: BLE001 with intent comments (or removed via the DRY consolidation).

Verification

  • py_compile on all six edited files — OK.
  • TestDetectGpuArch — 5/5 pass (mock-based, no GPU needed).
  • Live host (gfx90a): all four helpers (dispatcher_common, gemm_utils._get_arch, ctypes_utils, grouped_conv_utils) resolve via amd-smi; gemm_utils reuses the canonical bridge (identity check).

Deliberately left out of scope

  • Neutral-location move of smi_utils. The dispatcher still reaches the sibling tile_engine/ops/common/smi_utils via a guarded sys.path import; this is now isolated to the single dispatcher_common helper. Happy to move the wrapper to a shared/neutral module if reviewers prefer avoiding the dispatcher → tile_engine reach — flagging rather than doing it unilaterally.
  • fmha_utils ImportError fallback (only used when dispatcher_common isn't importable — a genuine standalone degrade path) and the ck-rocprof bash helper remain on rocminfo; these aren't Python dispatcher detection paths.

Note on motivation: the code being migrated used rocminfo (not rocm-smi), which isn't itself deprecated — so the honest driver here is one consistent amd-smi-first detection policy across TE + dispatcher (per #9731/#10204, JIRA ROCM-28734), plus the duplicate-definition cleanup, rather than rocminfo deprecation.

Production source should not point to internal JIRA tickets; ticket
context belongs in the PR/commit metadata instead.
@ozturkosu

Copy link
Copy Markdown
Contributor Author

Wanted to leave a note on why this change is still worth landing on top of #9731 and #10204, since at a glance it can look like more of the same amd-smi migration.

Those two PRs both live entirely in tile_engine/ and both deal with GPU id discovery (detect_gpu_ids, count_gpus, and friends). Neither of them touched the dispatcher, and neither added any way to detect the GPU architecture. On current develop, smi_utils.py has no detect_gpu_arch() at all; the only gfx material in there is the internal parse_amd_gfx / parse_rocm_gfx parsers used by the equivalence tests, which nothing external can call.

Meanwhile every arch-detection entry point in the dispatcher is still shelling out to rocminfo directly:

  • dispatcher_common.detect_gpu_arch()
  • ctypes_utils.detect_gpu_arch() (the busy one, hit by every GEMM example)
  • gemm_utils._get_arch()
  • grouped_conv_utils.detect_gpu_arch()

So this PR is doing work the earlier two never covered: it adds the arch wrapper to the shared module and routes all four dispatcher paths through it, so device-property access is amd-smi first everywhere in CK instead of the dispatcher being the one island still on rocminfo. That is the consistency the tech lead asked for under ROCM-28734.

Two other things worth calling out. First, dispatcher_common.py on develop actually defines detect_gpu_arch() twice, and the second definition silently shadows the first. This PR removes the duplicate, so it is a real correctness fix independent of the SMI story. Second, the four hand rolled detectors now share a single _detect_gpu_arch_via_amd_smi bridge rather than each re-implementing the same reach, which is a nice cleanup on its own.

One honest caveat: the "rocm-smi is going away" reasoning from #9731 does not directly apply here, since the dispatcher was on rocminfo, which is a different tool and not on the same deprecation path. rocminfo stays as a fallback, so nothing breaks on hosts without amd-smi. The value here is the end to end consistency plus the duplicate definition fix, not avoiding a future breakage.

Happy to move the shared wrapper to a neutral location if reviewers would rather not have the dispatcher reach into tile_engine/ops/common. That layering is the one open question, and it is easy to address before merge.

@geyyer geyyer 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 bf5eb0a into develop Aug 11, 2026
99 of 111 checks passed
@ozturkosu
ozturkosu deleted the users/muozturk/ck/dispatcher_amd_smi branch August 11, 2026 19:56
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 11, 2026
fix(ck-tile): prefer amd-smi for dispatcher GPU-arch
 detection (#10445)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Summary

Extends the amd-smi-first migration (#9731, #10204) from the Tile Engine
to the
dispatcher's GPU-architecture detection.

- Adds `detect_gpu_arch()` to the shared
`tile_engine/ops/common/smi_utils.py`
  wrapper: amd-smi `static` (`TARGET_GRAPHICS_VERSION`) first, rocm-smi
`--showproductname` (`GFX Version`) fallback — same amd-smi-first policy
as the
  rest of the wrapper.
- `dispatcher_common.detect_gpu_arch()`: prefers amd-smi via the
wrapper, then
rocminfo, then the caller-supplied default. Also removes a **duplicate**
  definition of this function.
- `gemm_utils._get_arch()`: tries amd-smi via the wrapper before
rocminfo,
  preserving the raise-on-unresolved (no silent default) behavior.

## Motivation

Per the  directive behind #9731, device-property access should prefer
`amd-smi` (rocm-smi is being phased out). The Tile Engine was already
migrated
(#9731, #10204); the dispatcher still detected the GPU arch by shelling
out to
`rocminfo` directly. Tracked in JIRA **ROCM-28734**.

## Design note

The dispatcher reaches the sibling `tile_engine/ops/common/smi_utils.py`
through
a guarded import and falls back to `rocminfo` if the wrapper is
unavailable, so
behavior is unchanged on hosts without amd-smi. (Open to moving the
shared
wrapper to a neutral location if reviewers prefer avoiding the
dispatcher →
tile_engine reach.)

## Test plan (MI350X / gfx950)

- [x] New mock-based unit tests for `detect_gpu_arch()` (amd-smi
preference,
      rocm-smi fallback, default) + a live check — 4/4 pass.
- [x] `dispatcher_common.detect_gpu_arch()` and `gemm_utils._get_arch()`
both
      return `gfx950` via amd-smi on the live host.
- [x] `test_gemm_utils` 21/21 still pass (no arch-detection regression).
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 18, 2026
fix(ck-tile): prefer amd-smi for dispatcher GPU-arch detection

## Summary

Extends the amd-smi-first migration (#9731, #10204) from the Tile Engine
to the
dispatcher's GPU-architecture detection.

- Adds `detect_gpu_arch()` to the shared
`tile_engine/ops/common/smi_utils.py`
  wrapper: amd-smi `static` (`TARGET_GRAPHICS_VERSION`) first, rocm-smi
`--showproductname` (`GFX Version`) fallback — same amd-smi-first policy
as the
  rest of the wrapper.
- `dispatcher_common.detect_gpu_arch()`: prefers amd-smi via the
wrapper, then
rocminfo, then the caller-supplied default. Also removes a **duplicate**
  definition of this function.
- `gemm_utils._get_arch()`: tries amd-smi via the wrapper before
rocminfo,
  preserving the raise-on-unresolved (no silent default) behavior.

## Motivation

Per the  directive behind #9731, device-property access should prefer
`amd-smi` (rocm-smi is being phased out). The Tile Engine was already
migrated
(#9731, #10204); the dispatcher still detected the GPU arch by shelling
out to
`rocminfo` directly. Tracked in JIRA **ROCM-28734**.

## Design note

The dispatcher reaches the sibling `tile_engine/ops/common/smi_utils.py`
through
a guarded import and falls back to `rocminfo` if the wrapper is
unavailable, so
behavior is unchanged on hosts without amd-smi. (Open to moving the
shared
wrapper to a neutral location if reviewers prefer avoiding the
dispatcher →
tile_engine reach.)

## Test plan (MI350X / gfx950)

- [x] New mock-based unit tests for `detect_gpu_arch()` (amd-smi
preference,
      rocm-smi fallback, default) + a live check — 4/4 pass.
- [x] `dispatcher_common.detect_gpu_arch()` and `gemm_utils._get_arch()`
both
      return `gfx950` via amd-smi on the live host.
- [x] `test_gemm_utils` 21/21 still pass (no arch-detection regression).

---------

Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
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.

3 participants