Skip to content

[Hipblaslt][CMS] CMS Support for TF32 128x192x32 NT - #3709

Closed
tywuAMD wants to merge 4 commits into
hipblaslt_common_cms_phase2from
users/tianyuwu/cms_tf32_128x192x32xNT
Closed

tywuAMD wants to merge 4 commits into
hipblaslt_common_cms_phase2from
users/tianyuwu/cms_tf32_128x192x32xNT

Conversation

@tywuAMD

@tywuAMD tywuAMD commented Jan 8, 2026

Copy link
Copy Markdown
Contributor

Motivation

Add CMS scheduling for TF32 192x256x32 NT.

Relevant PR:
TF32 192x256x32 TN: #3551
TF32 192x256x32 NN: #3677

Technical Details

Test Plan

  • Run Tensile with parameter UseCustomMainLoopSchedule: [0, 1] to compare CMS on and off.
  • Run hipblaslt-bench to compare the CMS kernel and the previous default kernel.

Test Result

Tensile

~15% uplift was observed when running Tensile

hipblaslt-bench

No gain versus the default kernel

Submission Checklist

tywuAMD and others added 4 commits January 8, 2026 10:11
Signed-off-by: tywuAMD <tianyuan.wu@amd.com>
Signed-off-by: tywuAMD <tianyuan.wu@amd.com>
@tywuAMD
tywuAMD requested a review from a team as a code owner January 8, 2026 11:05
@tywuAMD tywuAMD added the gfx950 run CI on gfx950 label Jan 8, 2026
@emezh

emezh commented Jan 9, 2026

Copy link
Copy Markdown
Contributor

test is missing in test_CustomSchedule.py

@talumbau
talumbau self-requested a review February 3, 2026 18:58

@talumbau talumbau left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The code overall looks good. A couple of changes to th test YAML are needed.

For trace/log directory this still needs:

  • hipblaslt_bench run with CMS kernels added in and --algo_method all run so we can see the kernel against all others for a representative GEMM.
  • Rocprof att trace for CMS and non-CMS cases.

- Range: [[128], [192], [1], [64, 64, 256]]
- Range: [[128], [192], [1], [1,1,64]]
- Range: [[128], [192], [1], [32, 64, 256]]
- Range: [[4096], [6144], [1], [64, 64, 256]]

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please don't add large GEMMs to the test file, especially the large M and N test values. These do not help check correctness.

- Range: [[128], [192], [1], [1,1,64]]
- Range: [[128], [192], [1], [32, 64, 256]]
- Range: [[4096], [6144], [1], [64, 64, 256]]
- Exact: [2048, 3072, 1, 8192]

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same issue here with this Exact GEMM test.

@tywuAMD tywuAMD closed this Feb 10, 2026
dileepr1 added a commit that referenced this pull request May 29, 2026
Mirrors the legacy TEST_TO_IGNORE map in
TheRock/build_tools/github_actions/test_executable_scripts/test_miopenprovider.py
which excluded the miopen_plugin_integration_tests ctest entry on
Windows+gfx110X only (TODO #3709).

- shared/ctest/parse_test_categories.py: teach the parser an optional
  `_(windows|linux)` suffix on exclude_gpu_* keys. Keys without a suffix
  (e.g. exclude_gpu_gfx950) stay OS-agnostic and behave exactly as
  before, so every existing YAML continues to parse identically. Keys
  with a suffix are skipped at YAML-parse time on the wrong OS, so the
  entry never reaches the install file for that platform. The runtime
  label stays `ex_gpu_<arch>` so test_runner.py's GPU-arch matching
  needs no changes. Same change is independently being landed via
  users/dravindr/tf_rocprim (commit ef8eee2); whichever PR merges
  first, the other will be a clean no-op.

- dnn-providers/miopen-provider/test_categories.yaml: add
  exclude_gpu_gfx110X_windows with test_patterns: ["*Integration*"].
  All TEST_P classes in integration_tests/ are named IntegrationGpu*
  and none in tests/ contain "Integration", so the filter is a no-op
  on miopen_plugin_tests and a full exclusion on
  miopen_plugin_integration_tests, matching the legacy script's
  ctest --exclude-regex "miopen_plugin_integration_tests" exactly.

Co-authored-by: Cursor <cursoragent@cursor.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