Skip to content

Enable gfx1250-strict in broad GPU families - #8350

Merged
marbre merged 8 commits into
mainfrom
users/marbre/gfx1250-strict-family-enablement
Oct 30, 2026
Merged

marbre merged 8 commits into
mainfrom
users/marbre/gfx1250-strict-family-enablement

Conversation

@marbre

@marbre marbre commented Sep 18, 2026 •

Copy link
Copy Markdown
Member

Motivation

Family selection must include both gfx1250 targets for builds, distributions, and artifact fetching. Default test compilation includes strict along with the other registered targets. Explicit individual target selections retain their own payload identity.

Closes #8278

Technical Details

Prerequisites before landing:

Test Plan

Fresh compiler/HIP/RAND builds and SDK, DEB/RPM, and tarball validation passed on the integration branch with temporary family membership. The exact family enablement was then integrated for a combined build with RAND tests and target-neutral hipthreads. Test target overrides were left empty to exercise the default registered-target list. These checks do not establish full component or GPU execution support.

  • python -m pytest -q -p no:cacheprovider build_tools/tests/cmake_amdgpu_targets_test.py build_tools/tests/artifact_manager_tool_test.py build_tools/github_actions/tests/configure_multi_arch_ci_test.py
  • cmake --build ../therock-build-rand-tests-gfx125X-dcgpu --target artifact-rand artifact-hipthreads
  • ../therock-build-rand-tests-gfx125X-dcgpu/compiler/amd-llvm/dist/lib/llvm/bin/llvm-objdump --offloading ../therock-build-rand-tests-gfx125X-dcgpu/math-libs/hipthreads/stage/lib/hipthreads/libhipthreads.a
  • python -m rocm_kpack.tools.verify_artifacts --artifacts-dir ../therock-build-rand-tests-gfx125X-dcgpu/artifacts --verbose --clang-offload-bundler ../therock-build-rand-tests-gfx125X-dcgpu/compiler/amd-llvm/dist/lib/llvm/bin/clang-offload-bundler

Test Result

  • 171 passed and 1 existing test skipped with umask 022.
  • rocRAND and hipRAND tests compiled for both gfx1250 targets. Each target's test kpack contains 34 entries with only its target identity.
  • Target-neutral hipthreads compiled with the default test target list. Its library contains exactly 29 distinct device images, including separate gfx1250 and gfx1250-strict images.
  • All four artifact verifier checks passed.
  • No GPU tests were executed.

Submission Checklist

Generated with Codex

Family selection must include both gfx1250 targets for builds,
distributions, and artifact fetching. Default test compilation includes
strict along with the other registered targets. Explicit individual
target selections retain their own payload identity.

Prerequisites before landing:
- #8325: shared SDK ownership.
- #8327: native package ownership.
- #8328: tarball ownership.
- A rocm-systems submodule bump containing ROCm/rocm-systems#11801.
- Established broader component readiness for family enablement.

Fresh compiler/HIP/RAND builds and SDK, DEB/RPM, and tarball validation
passed on the integration branch with temporary family membership.
The exact family enablement was then integrated for a combined build
with RAND tests and target-neutral hipthreads. Test target overrides
were left empty to exercise the default registered-target list.
These checks do not establish full component or GPU execution support.

Closes #8278

Test plan:
- python -m pytest -q -p no:cacheprovider build_tools/tests/cmake_amdgpu_targets_test.py build_tools/tests/artifact_manager_tool_test.py build_tools/github_actions/tests/configure_multi_arch_ci_test.py
- cmake --build ../therock-build-rand-tests-gfx125X-dcgpu --target artifact-rand artifact-hipthreads
- ../therock-build-rand-tests-gfx125X-dcgpu/compiler/amd-llvm/dist/lib/llvm/bin/llvm-objdump --offloading ../therock-build-rand-tests-gfx125X-dcgpu/math-libs/hipthreads/stage/lib/hipthreads/libhipthreads.a
- python -m rocm_kpack.tools.verify_artifacts --artifacts-dir ../therock-build-rand-tests-gfx125X-dcgpu/artifacts --verbose --clang-offload-bundler ../therock-build-rand-tests-gfx125X-dcgpu/compiler/amd-llvm/dist/lib/llvm/bin/clang-offload-bundler

Test results:
- 171 passed and 1 existing test skipped with umask 022.
- rocRAND and hipRAND tests compiled for both gfx1250 targets. Each
  target's test kpack contains 34 entries with only its target identity.
- Target-neutral hipthreads compiled with the default test target list.
  Its library contains exactly 29 distinct device images, including
  separate gfx1250 and gfx1250-strict images.
- All four artifact verifier checks passed.
- No GPU tests were executed.
@therock-pr-bot

therock-pr-bot Bot commented Sep 18, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 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.

self.assertEqual(
self._select("gfx125X-all;gfx1250-strict"),
self._select("gfx125X-all;gfx942"),
{

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.

gfx942 added for gfx1250 patch ????

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

This looks fine. It's a unit test, and now that gfx1250-strict is part of gfx1250X-all, gfx942 is an "explicit family" supplementing the "broad family" (unit test name).

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

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

Exactly :)

"rand_lib_gfx125X-all.tar.zst",
"rand_lib_gfx1250.tar.zst",
"rand_lib_gfx1250-strict.tar.zst",
"rand_lib_gfx942.tar.zst",

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.

same : gfx942 added for gfx1250 patch ????

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

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

Yes, this is tests for broader family enablement.

@marbre

marbre commented Oct 13, 2026

Copy link
Copy Markdown
Member Author
-- Configured AMDGPU Targets:
--   * gfx1250-strict : AMD Instinct gfx1250 strict
--   * gfx1250 : AMD Instinct MI450/MI450X/MI455X CDNA
--   Dist targets: gfx942;gfx1100;gfx1101;gfx1102;gfx1103;gfx1151;gfx1200;gfx1201;gfx1250-strict;gfx1250
--   Test targets: gfx900;gfx90c;gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201;gfx1250-strict;gfx1250
--   * Dist bundle: gfx125X-dcgpu

https://github.com/ROCm/TheRock/actions/runs/38419442156/job/115384411886?pr=8350#logs

math-libs stage for gfx125X is passing.

@marbre
marbre marked this pull request as ready for review October 13, 2026 16:51
tony-davis pushed a commit to ROCm/rocm-libraries that referenced this pull request Oct 14, 2026
JIRA ID : N/A

## Motivation
- Run hipBLASLt + TensileLite on the gfx1250-strict path (#11777) for
every PR.
- Per-PR TheRock builds don't produce gfx1250-strict yet
(ROCm/TheRock#8350, draft), so the job builds from source instead of
waiting for artifacts.
- Add Nightly schedule

## Technical Details
- Toolchain: pinned ROCm SDK nightly (`rocm_sdk_version`, default
`10.2.0a20261007`). A warn-only step flags when the SDK's compiler
differs from the one at ci-env's `therock-ref`.
- Checks that the runtime reports `gfx1250-strict`
(`HSA_DISABLE_GFX12_STRICT=0`), then builds with `install.sh -a
gfx1250-strict --skip_rocroller`.
- One runner per GPU: the job container gets the runner's
`ROCR_VISIBLE_DEVICES` (`-e ROCR_VISIBLE_DEVICES`) and fails early
unless it names exactly one GPU. Checked on the host during a CI run:
all 4 hipBLASLt shards ran on the runner's own GPU.
- hipBLASLt: loads `library/gfx1250-strict/`; fails if `hipblaslt-test`
doesn't link the library built here or its code objects aren't labelled
`gfx1250-strict`. Runs as 4 gtest shards.
- TensileLite: Uses 4 workers sharing the GPU. Unit tests dropped
(host-only, already run per PR in TheRock CI). Selection `-m "common and
not skip-gfx1250-strict"`; skipped if setup failed, still runs after a
hipBLASLt failure.
- Log scrub on build and test output: hardware details and wall-clock
times are removed before anything reaches the console or the uploaded
logs.
- Per PR: `pull_request` trigger behind a GitHub-hosted gate (non-draft,
ROCm org label, changes under `projects/hipblaslt`, its shared deps,
this workflow or `ci-env/action.yml`). One run per PR; a newer push
cancels it. Manual dispatch and the nightly are ungated. A manual run
and the nightly on the same branch don't cancel each other.
- Nightly: develop at 03:00 UTC (7 pm PST, 8 pm PDT).
- Runner state: a fresh venv per job from the image's Python 3.12 (no
`actions/setup-python`, so nothing from the runner's tool cache); the
runner's pip cache is used only by the SDK install and is read-only
except for the nightly; the job hands the root-owned workspace back to
the runner's user at the end.
- Includes #13464 (merged into this branch): `skip-gfx1250-strict` on 4
TensileLite configs and the bf8 MX gtest rows limited to gfx950, + fixed
comment saying gfx1250 has no MX kernels with bf8 inputs.
- Skips `gsu_gfx1250-strict.yaml` on gfx1250-strict for now
(`skip-gfx1250-strict`): it intermittently returns wrong values from GSU
multi-buffer kernels (see Test Result). Followed up with the config
owner; the skip comes out when that's fixed.
- Removed: the `workflow_run` trigger, artifact resolution,
commit-status posting, unused `GFX1250_HW_XFAIL`.

## Test Plan
- This PR's own `gfx1250-strict hipBLASLt + TensileLite` check on the
gfx1250 runners (the PR changes the workflow, so it runs here).

## Test Result
- CI on the gfx1250 runners ([run
38319193893](https://github.com/ROCm/rocm-libraries/actions/runs/38319193893)):
hipBLASLt 68,138 tests over 4 shards, 0 failed; TensileLite common 109
passed, 2 xfailed, 0 failed.


## Submission Checklist

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

## Risk level
- Low. Self-contained workflow. PR code runs on a self-hosted runner
only for ROCm org-labelled PRs

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude <noreply@anthropic.com>
marbre added a commit that referenced this pull request Oct 14, 2026
## Motivation

When multiple selected GPU targets share one native package owner, split
artifact manifests can retain source roots whose device images have
moved into
a shared kpack under another root. The grouped-root validation
introduced by
#8327 rejects those absent directories before the DEB and RPM backends
can
apply their established filtering.

The device-package builder then converts that exception into an empty
result.
Other variants can still be produced, causing package generation to
appear
successful and emit meta packages that depend on device packages that
were
never created. The failure is only discovered later during simulated or
native
package installation.

This was first reported in
#8327 (comment) and
reproduces
when #8350 enables `gfx1250-strict` alongside `gfx1250`.

This PR carries forward the fix and artifact evidence from #8702, then
adds RPM
coverage and makes unexpected partial package failures fatal.

Progress on #8278.

## Evidence

The original #8702 investigation used
https://github.com/ROCm/TheRock/actions/runs/36947069646. Both sparse
device
artifact manifests retained the rocSPARSE root after their device images
moved
into shared kpacks under hipSPARSELt. Each target kpack still contained
340
rocSPARSE and 3 hipSPARSELt images, while the generic artifact contained
`librocsparse.so`.

The current failure reproduces in #8350 run
https://github.com/ROCm/TheRock/actions/runs/38419442156:

- DEB package generation fails its simulated installation after the
affected
  `gfx1250` device packages are omitted:

https://github.com/ROCm/TheRock/actions/runs/38419442156/job/115420436430?pr=8350
- RPM generation completes despite omitting the affected packages:

https://github.com/ROCm/TheRock/actions/runs/38419442156/job/115420436578?pr=8350
- The downstream RHEL 8 and SLES 16 installation jobs then fail on
unresolved
  dependencies:
-
https://github.com/ROCm/TheRock/actions/runs/38419442156/job/115430597884?pr=8350
-
https://github.com/ROCm/TheRock/actions/runs/38419442156/job/115430596637?pr=8350

The omitted packages include `amdrocm-sparse`, `amdrocm-blas-test`,
`amdrocm-rand-test`, `amdrocm-ccl-test`, and `amdrocm-rccl-test` for the
shared
`gfx1250` package owner.

## Technical Details

- Ignore absent manifest roots during kpack artifact discovery, matching
the
  existing artifact-catalog behavior.
- Continue collision validation for every existing package root.
- Log skipped roots at debug level so stale manifest entries remain
  diagnosable.
- Preserve the selected `gfx1250` and `gfx1250-strict` kpack payloads in
their
  shared `gfx1250` native package.
- Propagate unexpected device-package exceptions instead of converting
them to
  empty results.
- Keep intentional empty device variants skippable.
- Return a nonzero status after the build summary whenever any base
package is
  recorded as failed.
- Add matching DEB and RPM regressions for combined-owner packages with
absent
  manifest roots.

The implementation is architecture-independent packaging logic.
Regression
fixtures use the motivating combined `gfx1250` and `gfx1250-strict`
ownership
case and verify that both selected payloads survive. No GPU execution is
required for this change.

## Test Plan

- `python build_tools/packaging/linux/tests/build_package_test.py`
- `gfx125X,gfx1250-strict` dev build:
https://github.com/ROCm/TheRock/actions/runs/38459803112

## Test Result

- 52 packaging tests passed.
- The DEB regression verifies both selected kpack filenames and bytes
survive
  staging when manifests contain absent roots.
- The RPM regression verifies both valid target roots are retained in
the
  generated spec and absent roots are excluded.
- Device-package exceptions propagate, while intentional empty variants
remain
  skippable.
- Generic recorded package failures exit with status 1 after printing
the build
  summary.

## Related

- Tracking issue: #8278
- Original fix and evidence: #8702
- Native package ownership change: #8327
- Reproducing family enablement: #8350

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/CONTRIBUTING.md.

Generated with Codex

---------

Co-authored-by: Ammar ELWazir <3832908+ammarwa@users.noreply.github.com>
@ScottTodd

Copy link
Copy Markdown
Member

PyTorch build failures: https://github.com/ROCm/TheRock/actions/runs/38587475275/job/115983081260?pr=8350#step:15:38204

/__w/TheRock/TheRock/external-builds/pytorch/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:48:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
   48 |     wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
      |                                                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:61:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
   61 |     auto flags     = CK_BUFFER_RESOURCE_3RD_DWORD;
      |                      ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:73:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
   73 |     auto flags     = CK_BUFFER_RESOURCE_3RD_DWORD;
      |                      ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
4 errors generated when compiling for gfx1250-strict.

Might need ROCm/pytorch#3639 or something similar on each pytorch release branch.

archana-ramalingam added a commit to ROCm/rocm-libraries that referenced this pull request Oct 21, 2026
It exists only because TheRock multi-arch PR builds do not produce
gfx1250-strict artifacts. Planned for removal once ROCm/TheRock#8350
lands and gfx1250-strict runners are available in TheRock CI.

Co-Authored-By: Claude <noreply@anthropic.com>
talumbau pushed a commit to ROCm/rocm-libraries that referenced this pull request Oct 22, 2026
…ite timeout; cleanup (#13771)

JIRA ID : N/A

## Motivation
Follow-ups from the #12983 and #13464 reviews, split out of #13620, plus
a fix for hipBLASLt tests running ~7x slower since per-PR runs started.

## Technical Details
- **hipBLASLt test speed:** `OMP_NUM_THREADS` and `OPENBLAS_NUM_THREADS`
= host CPUs / (4 runners x 4 shards); `BLIS_NUM_THREADS` = the same, at
most 8 (hipBLASLt's own BLIS default, which setting `OMP_NUM_THREADS`
turns off). By default the CPU reference's OpenMP and OpenBLAS pools
each start one thread per host CPU, so 4 runners x 4 shards
oversubscribe the host 16x per pool: since per-PR runs started (10-12 to
10-15), the step took ~80 min when other runners were testing vs ~15 min
alone, up to 110 min of its 120-min timeout.
- **Gate:** add `cmake/` (rocisa and stinkytofu include
`cmake/modules/default_amdclang.cmake`).
- **One-GPU check:** runs first in its step, under `set -euo pipefail`.
`ROCR_VISIBLE_DEVICES` comes from the runner's `.env`; if it is unset, N
comes from the runner name's `-gpuN` (the pool registers runner N as
`...-gpuN` with `ROCR_VISIBLE_DEVICES=N`), with a warning. It must be a
single GPU index (`^[0-9]+$`), or the job fails before the build.
  Drop `HIP_VISIBLE_DEVICES` (never forwarded into the container).
- SDK pinned to `10.2.0a20261013` (compiler `c7ceccf9a5`). Per-PR CI has
since moved to compiler `bdc35841c9` (#13835, 10-19), so the workflow's
compiler check now warns; re-pinning to `10.2.0a20261018` or later,
which carries `bdc35841c9`, is a follow-up.
- **upload-artifact:** v4.6.2 -> v7.0.2 (Node 20 deprecated). v7 needs
Actions Runner >= 2.327.1; the gfx1250 runner is 2.337.0.
- **TensileLite per-test timeout:** 600 -> 1200 s, this workflow only
(FFM and TheRock keep their own). `decouple_pgr_tdm_fuse_gfx1250.yaml`
has timed out at 600 s in most runs since #12825 grew its sweep 4.5x;
before, it passed 41 of 41.
- **Header:** the workflow is marked temporary, to be removed once
TheRock's multi-arch builds produce gfx1250-strict artifacts
(ROCm/TheRock#8350) and gfx1250-strict runners are available in TheRock
CI.

## Test Plan
- OpenMP cap: 4 jobs x 4 shards of a `*pre_checkin*` slice on a 128-core
gfx942 host, default vs capped.
- One-GPU check: 7 values under `bash -euo pipefail`.

## Test Result
- [x] Run
https://github.com/ROCm/rocm-libraries/actions/runs/38994215757:
all steps green with no warnings on SDK 10.2.0a20261013
- hipBLASLt 68,925 tests, 0 failed, in 3.9 min
(OpenMP/OpenBLAS capped, BLIS at most 8) while an uncapped run on the
same host took 52 min
- TensileLite 113 passed, including decouple_pgr_tdm_fuse_gfx1250.yaml
under the 1200 s timeout.

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

## Risk level
Low: CI-only, gfx1250 workflow.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude <noreply@anthropic.com>
@scchan

scchan commented Oct 22, 2026

Copy link
Copy Markdown

We haven't enabled the reporting of gfx1250-strict by default yet (ROCm/rocm-systems#11575). You'll have to set env var HSA_DISABLE_GFX12_STRICT=1 to enable it. I presume you'd have to set that in the environment running tests for gfx1250-strict?

@marbre

marbre commented Oct 23, 2026

Copy link
Copy Markdown
Member Author

We haven't enabled the reporting of gfx1250-strict by default yet (ROCm/rocm-systems#11575). You'll have to set env var HSA_DISABLE_GFX12_STRICT=1 to enable it. I presume you'd have to set that in the environment running tests for gfx1250-strict?

I don't think this is related to the work in this PR. This is the family enablement in builds and it does not enable running tests.

@scchan

scchan commented Oct 23, 2026

Copy link
Copy Markdown

We haven't enabled the reporting of gfx1250-strict by default yet (ROCm/rocm-systems#11575). You'll have to set env var HSA_DISABLE_GFX12_STRICT=1 to enable it. I presume you'd have to set that in the environment running tests for gfx1250-strict?

I don't think this is related to the work in this PR. This is the family enablement in builds and it does not enable running tests.

The reason I asked is because your PR description contains test results that imply that some gfx1250-strict tests have been run.

@marbre

marbre commented Oct 23, 2026

Copy link
Copy Markdown
Member Author

The reason I asked is because your PR description contains test results that imply that some gfx1250-strict tests have been run.

Please check the PR description more carefully:

  • No GPU tests were executed.

This is pure build tests.

@ScottTodd ScottTodd left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

CI results are looking good now that #9549 is merged to skip pytorch builds on gfx1250-strict.

@marbre
marbre merged commit 682b340 into main Oct 30, 2026
206 of 210 checks passed
@marbre
marbre deleted the users/marbre/gfx1250-strict-family-enablement branch October 30, 2026 05:09
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

Enable gfx1250-strict builds and shared gfx1250 packaging

4 participants