Skip to content

fix(ck): Stream-K Tile Engine GPU Query Fix - #8086

Merged
arai713 merged 8 commits into
developfrom
users/arai/ck/streamk_tile_engine_gpu_fix
Jul 24, 2026
Merged

fix(ck): Stream-K Tile Engine GPU Query Fix#8086
arai713 merged 8 commits into
developfrom
users/arai/ck/streamk_tile_engine_gpu_fix

Conversation

@arai713

@arai713 arai713 commented Jun 4, 2026

Copy link
Copy Markdown
Contributor

Motivation

The Stream-K tile engine validation utilities call rocminfo at build time to detect the GPU architecture, which prevents building on CPU-only nodes. The GPU target is already known from CMake's SUPPORTED_GPU_TARGETS, so runtime hardware detection is unnecessary during code generation.

Technical Details

The gemm_streamk_validation_utils.py file used subprocess.check_output(["rocminfo"]) to query GPU hardware at CMake configure time. This call originated from get_gpu_name_by_id() and used during tile configuration validation. On CPU-only build nodes, this fails because rocminfo either doesn't exist or returns no GPU devices. The main changes are as follows:

  • Removed runtime GPU detection infrastructure: deleted get_gpu_name_by_id(), set_gpu_targets(), get_configured_gpu_targets(), the _configured_gpu_targets module variable, and the GPU_NAME_PATTERN regex.
  • Added gpu_target as an explicit parameter. is_tile_config_valid(), validate_gemm(), validate_warp_tile_combination(), validate_warp_configuration(), and validate_lds_capacity() now accept gpu_target as a required parameter instead of querying hardware internally. The corresponding changes were also made in gemm_streamk_instance_builder.py and CMakeLists.txt
  • Added WARP_SUPPORTED_COMBINATIONS for per-GPU warp config validation, and LDS_SIZE_MAP / DEFAULT_LDS_SIZE for GPU-aware LDS capacity checks.

Test Plan

The benchmarks were compiled and run on a GPU as well as a CPU only node to verify correctness.

Test Result

All tests passed

Related

JIRA ID : AICK-1635

Submission Checklist

@arai713 arai713 changed the title remove rocminfo usage and refactored to accept the GPU target as an e… [TILE_ENGINE] Stream-K Tile Engine GPU Query Fix Jun 4, 2026
@arai713
arai713 force-pushed the users/arai/ck/streamk_tile_engine_gpu_fix branch 2 times, most recently from ed06956 to 643cb67 Compare June 10, 2026 18:32
@arai713
arai713 marked this pull request as ready for review June 11, 2026 18:59
@arai713
arai713 requested a review from a team as a code owner June 11, 2026 18:59
@arai713
arai713 force-pushed the users/arai/ck/streamk_tile_engine_gpu_fix branch 2 times, most recently from cb9ab37 to d2907e1 Compare June 18, 2026 00:07
@tenpercent
tenpercent requested a review from Copilot June 18, 2026 18:59

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 removes build-time GPU hardware probing (rocminfo) from the Stream-K tile engine validation path and instead threads the configured GPU target(s) from CMake into the Python validation/generation utilities, enabling builds on CPU-only nodes.

Changes:

  • Removed runtime GPU detection and made validation functions accept an explicit gpu_target.
  • Added per-GPU validation data for warp configuration and LDS capacity budgeting.
  • Updated CMake/Python instance generation plumbing to pass --gpu_targets, and added a small pytest file to validate LDS sizing behavior.

Reviewed changes

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

File Description
projects/composablekernel/tile_engine/ops/gemm_streamk/test_lds_capacity.py Adds pytest coverage for GPU-aware LDS capacity validation.
projects/composablekernel/tile_engine/ops/gemm_streamk/gemm_streamk_validation_utils.py Removes rocminfo querying and makes validation GPU-target-driven (warp + LDS).
projects/composablekernel/tile_engine/ops/gemm_streamk/gemm_streamk_instance_builder.py Passes configured GPU targets into validation instead of probing hardware.
projects/composablekernel/tile_engine/ops/gemm_streamk/CMakeLists.txt Updates custom commands to pass --gpu_targets to the instance builder.
Comments suppressed due to low confidence (1)

projects/composablekernel/tile_engine/ops/gemm_streamk/gemm_streamk_validation_utils.py:265

  • validate_warp_tile_combination looks up WARP_TILE_SUPPORTED_COMBINATIONS using the raw gpu_name. If callers pass full targets like gfx942:xnack+, the lookup misses and the function becomes permissive (and emits warnings). Normalize the GPU name similarly to validate_lds_capacity by stripping any :... suffix before the lookup/logging.
    """Validate warp tile combination against GPU-specific supported combinations."""
    # Construct the key for looking up supported combinations
    warp_tile_key = f"{a_datatype}_{b_datatype}_{c_datatype}"
    current_combination = [warp_tile_m, warp_tile_n, warp_tile_k]

    # Check if we have GPU-specific combinations
    gpu_warp_tile_combinations = WARP_TILE_SUPPORTED_COMBINATIONS.get(gpu_name, {})
    if not gpu_warp_tile_combinations:
        # If GPU not recognized, try to be permissive but log warning
        logging.warning(f"No warp tile combinations found for GPU: {gpu_name}")
        return True, ""

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread projects/composablekernel/tile_engine/ops/gemm_streamk/test_lds_capacity.py Outdated
Comment thread projects/composablekernel/tile_engine/ops/gemm_streamk/test_lds_capacity.py Outdated
Comment thread projects/composablekernel/tile_engine/ops/gemm_streamk/test_lds_capacity.py Outdated
Comment thread projects/composablekernel/tile_engine/ops/gemm_streamk/CMakeLists.txt Outdated
@arai713
arai713 force-pushed the users/arai/ck/streamk_tile_engine_gpu_fix branch from d2907e1 to 2f42c5d Compare June 22, 2026 18:43
@arai713
arai713 force-pushed the users/arai/ck/streamk_tile_engine_gpu_fix branch from 2f42c5d to eedd197 Compare June 22, 2026 20:53
@tenpercent
tenpercent requested a review from Copilot June 22, 2026 22:38

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

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

@therock-pr-bot

therock-pr-bot Bot commented Jul 11, 2026

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
🌿 Branch Name ✅ Pass
📝 PR Title/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.

@therock-pr-bot

therock-pr-bot Bot commented Jul 11, 2026

Copy link
Copy Markdown

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

@arai713 arai713 changed the title [TILE_ENGINE] Stream-K Tile Engine GPU Query Fix fix(ck): Stream-K Tile Engine GPU Query Fix Jul 11, 2026

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

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

Comment thread projects/composablekernel/tile_engine/ops/gemm_streamk/test_lds_capacity.py Outdated
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

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

comments were addressed and ci is passing. thanks!

@arai713
arai713 merged commit 33e2420 into develop Jul 24, 2026
29 checks passed
@arai713
arai713 deleted the users/arai/ck/streamk_tile_engine_gpu_fix branch July 24, 2026 17:28
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Jul 24, 2026
fix(ck): Stream-K Tile Engine GPU Query Fix

## Motivation
The Stream-K tile engine validation utilities call rocminfo at build
time to detect the GPU architecture, which prevents building on CPU-only
nodes. The GPU target is already known from CMake's
SUPPORTED_GPU_TARGETS, so runtime hardware detection is unnecessary
during code generation.

## Technical Details
The gemm_streamk_validation_utils.py file used
subprocess.check_output(["rocminfo"]) to query GPU hardware at CMake
configure time. This call originated from get_gpu_name_by_id() and used
during tile configuration validation. On CPU-only build nodes, this
fails because rocminfo either doesn't exist or returns no GPU devices.
The main changes are as follows:

- Removed runtime GPU detection infrastructure: deleted
get_gpu_name_by_id(), set_gpu_targets(), get_configured_gpu_targets(),
the _configured_gpu_targets module variable, and the GPU_NAME_PATTERN
regex.
- Added gpu_target as an explicit parameter. is_tile_config_valid(),
validate_gemm(), validate_warp_tile_combination(),
validate_warp_configuration(), and validate_lds_capacity() now accept
gpu_target as a required parameter instead of querying hardware
internally. The corresponding changes were also made in
gemm_streamk_instance_builder.py and CMakeLists.txt
- Added WARP_SUPPORTED_COMBINATIONS for per-GPU warp config validation,
and LDS_SIZE_MAP / DEFAULT_LDS_SIZE for GPU-aware LDS capacity checks.

## Test Plan
The benchmarks were compiled and run on a GPU as well as a CPU only node
to verify correctness.

## Test Result
All tests passed

## Related
JIRA ID : AICK-1635

## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
bghimireamd pushed a commit that referenced this pull request Jul 28, 2026
## Motivation
The Stream-K tile engine validation utilities call rocminfo at build
time to detect the GPU architecture, which prevents building on CPU-only
nodes. The GPU target is already known from CMake's
SUPPORTED_GPU_TARGETS, so runtime hardware detection is unnecessary
during code generation.

## Technical Details
The gemm_streamk_validation_utils.py file used
subprocess.check_output(["rocminfo"]) to query GPU hardware at CMake
configure time. This call originated from get_gpu_name_by_id() and used
during tile configuration validation. On CPU-only build nodes, this
fails because rocminfo either doesn't exist or returns no GPU devices.
The main changes are as follows:

- Removed runtime GPU detection infrastructure: deleted
get_gpu_name_by_id(), set_gpu_targets(), get_configured_gpu_targets(),
the _configured_gpu_targets module variable, and the GPU_NAME_PATTERN
regex.
- Added gpu_target as an explicit parameter. is_tile_config_valid(),
validate_gemm(), validate_warp_tile_combination(),
validate_warp_configuration(), and validate_lds_capacity() now accept
gpu_target as a required parameter instead of querying hardware
internally. The corresponding changes were also made in
gemm_streamk_instance_builder.py and CMakeLists.txt
- Added WARP_SUPPORTED_COMBINATIONS for per-GPU warp config validation,
and LDS_SIZE_MAP / DEFAULT_LDS_SIZE for GPU-aware LDS capacity checks.


## Test Plan
The benchmarks were compiled and run on a GPU as well as a CPU only node
to verify correctness.

## Test Result
All tests passed

## Related
JIRA ID : AICK-1635

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

---------

Co-authored-by: Maksim (Max) Podkorytov <Maksim.Podkorytov@amd.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 18, 2026
fix(ck): Stream-K Tile Engine GPU Query Fix

## Motivation
The Stream-K tile engine validation utilities call rocminfo at build
time to detect the GPU architecture, which prevents building on CPU-only
nodes. The GPU target is already known from CMake's
SUPPORTED_GPU_TARGETS, so runtime hardware detection is unnecessary
during code generation.

## Technical Details
The gemm_streamk_validation_utils.py file used
subprocess.check_output(["rocminfo"]) to query GPU hardware at CMake
configure time. This call originated from get_gpu_name_by_id() and used
during tile configuration validation. On CPU-only build nodes, this
fails because rocminfo either doesn't exist or returns no GPU devices.
The main changes are as follows:

- Removed runtime GPU detection infrastructure: deleted
get_gpu_name_by_id(), set_gpu_targets(), get_configured_gpu_targets(),
the _configured_gpu_targets module variable, and the GPU_NAME_PATTERN
regex.
- Added gpu_target as an explicit parameter. is_tile_config_valid(),
validate_gemm(), validate_warp_tile_combination(),
validate_warp_configuration(), and validate_lds_capacity() now accept
gpu_target as a required parameter instead of querying hardware
internally. The corresponding changes were also made in
gemm_streamk_instance_builder.py and CMakeLists.txt
- Added WARP_SUPPORTED_COMBINATIONS for per-GPU warp config validation,
and LDS_SIZE_MAP / DEFAULT_LDS_SIZE for GPU-aware LDS capacity checks.


## Test Plan
The benchmarks were compiled and run on a GPU as well as a CPU only node
to verify correctness.

## Test Result
All tests passed

## Related
JIRA ID : AICK-1635

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

---------

Co-authored-by: Maksim (Max) Podkorytov <Maksim.Podkorytov@amd.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@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.

4 participants