Skip to content

fix(ck-tile): prefer amd-smi over rocm-smi in remaining TE GEMM benchmarks - #10204

Merged
geyyer merged 1 commit into
developfrom
users/muozturk/ck/smi-wrapper-benchmarks
Aug 4, 2026
Merged

geyyer merged 1 commit into
developfrom
users/muozturk/ck/smi-wrapper-benchmarks

Conversation

@ozturkosu

@ozturkosu ozturkosu commented Jul 30, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

PR #9731 centralized GPU SMI access behind tile_engine/ops/common/smi_utils.py, which prefers amd-smi and transparently falls back to rocm-smi, and migrated gemm_full_benchmark.py off direct rocm-smi calls. Three Tile Engine benchmark scripts were left calling rocm-smi --showid directly (with an amd-smi list fallback), so the migration is incomplete. Although wrapped in try/except today, these direct calls will break once rocm-smi is removed from future ROCm releases.

This completes the migration for the remaining scripts, as requested by the tech lead and tracked in JIRA ROCM-28734.

Technical Details

Migrated the identical hand-rolled detect_devices() (env → rocm-smi --showid → amd-smi list → hardcoded ["0"]) to a single smi_utils.detect_gpu_ids() call, mirroring the already-landed gemm_full_benchmark.py:

  • tile_engine/ops/gemm/block_scale_gemm/gemm_bquant/gemm_bquant_full_benchmark.py
  • tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py
  • tile_engine/ops/gemm/streamk_gemm_full_benchmark.py

Each script adds common/ to sys.path, imports detect_gpu_ids, and drops the now-unused import re. Net -67 lines.

Test Plan / Result

Validated on an MI350 (gfx950) host, exercising the exact sys.path arithmetic + smi_utils import + detect_gpu_ids() for both the shallow (gemm/) and deep (block_scale_gemm/gemm_bquant/) script locations:

  • Default (amd-smi first): detect_devices() returns the live GPU ids; count_gpus()/check_gpu_available() consistent.
  • CK_SMI_TOOL=rocm-smi (forced fallback): identical result.

Both python3 -m py_compile cleanly.

Submission Checklist

https://amd-hub.atlassian.net/jira/software/c/projects/AICK/boards/4399?selectedIssue=ROCM-28734

JIRA ID : ROCM-28734

…marks

The bquant, multi-D, and stream-K full-benchmark scripts still called
`rocm-smi --showid` directly (with an `amd-smi list` fallback) for GPU id
detection, duplicating the detection logic that PR #9731 centralized in
`tile_engine/ops/common/smi_utils.py`. This completes the SMI migration
started in #9731 so every call site prefers `amd-smi` and transparently
falls back to `rocm-smi`.

Each `detect_devices()` now delegates to `smi_utils.detect_gpu_ids()`,
matching the already-migrated `gemm_full_benchmark.py`. Added the
`common/` dir to `sys.path` and dropped the now-unused `import re`.

Validated on an MI350 (gfx950) host: both the default (amd-smi first)
and `CK_SMI_TOOL=rocm-smi` fallback paths return identical GPU ids for
both the shallow (gemm/) and deep (block_scale_gemm/gemm_bquant/) script
locations.

JIRA ID : ROCM-28734
@therock-pr-bot

therock-pr-bot Bot commented Jul 30, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

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/tile_engine/ops/gemm/block_scale_gemm/gemm_bquant/gemm_bquant_full_benchmark.py, projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py, projects/composablekernel/tile_engine/ops/gemm/streamk_gemm_full_benchmark.py; no test file found
🔎 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.

@ozturkosu ozturkosu self-assigned this Jul 30, 2026
@ozturkosu
ozturkosu marked this pull request as ready for review July 30, 2026 21:35
@ozturkosu
ozturkosu requested a review from a team as a code owner July 30, 2026 21:35
@ozturkosu
ozturkosu requested a review from geyyer July 30, 2026 21:49
@geyyer
geyyer merged commit 9116b6d into develop Aug 4, 2026
32 checks passed
@geyyer
geyyer deleted the users/muozturk/ck/smi-wrapper-benchmarks branch August 4, 2026 15:35
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 4, 2026
fix(ck-tile): prefer amd-smi over rocm-smi in remaining TE
 GEMM benchmarks (#10204)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Motivation

PR #9731 centralized GPU SMI access behind
`tile_engine/ops/common/smi_utils.py`, which prefers `amd-smi` and
transparently falls back to `rocm-smi`, and migrated
`gemm_full_benchmark.py` off direct `rocm-smi` calls. Three Tile Engine
benchmark scripts were left calling `rocm-smi --showid` directly (with
an `amd-smi list` fallback), so the migration is incomplete. Although
wrapped in try/except today, these direct calls will break once
`rocm-smi` is removed from future ROCm releases.

This completes the migration for the remaining scripts, as requested by
the tech lead and tracked in JIRA **ROCM-28734**.

## Technical Details

Migrated the identical hand-rolled `detect_devices()` (env → `rocm-smi
--showid` → `amd-smi list` → hardcoded `["0"]`) to a single
`smi_utils.detect_gpu_ids()` call, mirroring the already-landed
`gemm_full_benchmark.py`:

-
`tile_engine/ops/gemm/block_scale_gemm/gemm_bquant/gemm_bquant_full_benchmark.py`
- `tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py`
- `tile_engine/ops/gemm/streamk_gemm_full_benchmark.py`

Each script adds `common/` to `sys.path`, imports `detect_gpu_ids`, and
drops the now-unused `import re`. Net -67 lines.

## Test Plan / Result

Validated on an MI350 (gfx950) host, exercising the exact `sys.path`
arithmetic + `smi_utils` import + `detect_gpu_ids()` for both the
shallow (`gemm/`) and deep (`block_scale_gemm/gemm_bquant/`) script
locations:

- Default (amd-smi first): `detect_devices()` returns the live GPU ids;
`count_gpus()`/`check_gpu_available()` consistent.
- `CK_SMI_TOOL=rocm-smi` (forced fallback): identical result.

Both `python3 -m py_compile` cleanly.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

https://amd-hub.atlassian.net/jira/software/c/projects/AICK/boards/4399?selectedIssue=ROCM-28734

JIRA ID : ROCM-28734

Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
ozturkosu added a commit that referenced this pull request Aug 11, 2026
## 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>
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 18, 2026
fix(ck-tile): prefer amd-smi over rocm-smi in remaining TE GEMM benchmarks

## Motivation

PR #9731 centralized GPU SMI access behind
`tile_engine/ops/common/smi_utils.py`, which prefers `amd-smi` and
transparently falls back to `rocm-smi`, and migrated
`gemm_full_benchmark.py` off direct `rocm-smi` calls. Three Tile Engine
benchmark scripts were left calling `rocm-smi --showid` directly (with
an `amd-smi list` fallback), so the migration is incomplete. Although
wrapped in try/except today, these direct calls will break once
`rocm-smi` is removed from future ROCm releases.

This completes the migration for the remaining scripts, as requested by
the tech lead and tracked in JIRA **ROCM-28734**.

## Technical Details

Migrated the identical hand-rolled `detect_devices()` (env → `rocm-smi
--showid` → `amd-smi list` → hardcoded `["0"]`) to a single
`smi_utils.detect_gpu_ids()` call, mirroring the already-landed
`gemm_full_benchmark.py`:

-
`tile_engine/ops/gemm/block_scale_gemm/gemm_bquant/gemm_bquant_full_benchmark.py`
- `tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py`
- `tile_engine/ops/gemm/streamk_gemm_full_benchmark.py`

Each script adds `common/` to `sys.path`, imports `detect_gpu_ids`, and
drops the now-unused `import re`. Net -67 lines.

## Test Plan / Result

Validated on an MI350 (gfx950) host, exercising the exact `sys.path`
arithmetic + `smi_utils` import + `detect_gpu_ids()` for both the
shallow (`gemm/`) and deep (`block_scale_gemm/gemm_bquant/`) script
locations:

- Default (amd-smi first): `detect_devices()` returns the live GPU ids;
`count_gpus()`/`check_gpu_available()` consistent.
- `CK_SMI_TOOL=rocm-smi` (forced fallback): identical result.

Both `python3 -m py_compile` cleanly.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.


https://amd-hub.atlassian.net/jira/software/c/projects/AICK/boards/4399?selectedIssue=ROCM-28734

JIRA ID : ROCM-28734

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.

2 participants