Skip to content

build(tensilelite): assemble gfx950 objects to the gfx950:sramecc+ target id - #14044

Merged
jaopaulolc merged 34 commits into
users/jolabega/kfa-rr-mx-1-fixesfrom
users/jolabega/kfa-rr-mx-2-target-id
Oct 30, 2026
Merged

jaopaulolc merged 34 commits into
users/jolabega/kfa-rr-mx-1-fixesfrom
users/jolabega/kfa-rr-mx-2-target-id

Conversation

@jaopaulolc

@jaopaulolc jaopaulolc commented Oct 20, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AIHPBLAS-5262

Motivation

The kernels in tensilelite/CustomKernels/rocroller/ declare a target id the build did not assemble. An AMD GPU is identified not only by its architecture (gfx950, the MI350 series) but by a target id: the architecture plus a tail of hardware feature settings, such as gfx950:sramecc+ for a part with error-correcting memory enabled. The assembler compares the whole target id and rejects a mismatch. Of the 90 checked-in kernels, 89 target gfx950 and declare .amdgcn_target "amdgcn-amd-amdhsa--gfx950:sramecc+" while the build assembled them as bare gfx950, so none of the 89 could assemble at all. The 90th targets gfx942 and was unaffected.

Technical Details

Stacked on #14043.

Every gfx950 artifact is now assembled to the target id its kernels declare. The change keeps the target id and the architecture name on separate channels, which is load-bearing rather than cosmetic:

  • The target id is produced by targetIdOf() in Python and tensilelite_offload_target() in CMake, which append the feature tail. These two feed -mcpu, --offload-arch, the .amdgcn_target directive, and the code-object bundle entry.
  • The name is produced by compilerTargetOf and baseArchName, alongside the bare ${arch}. These keep feeding directory names, file names, and the identifier accepted by the --architecture= allow-list.

The split is required because the CMake variable GPU_TARGETS is validated against an allow-list and then stripped by a regular expression into directory names. A feature-qualified value there fails configure, and no path the build produces may contain a colon.

The change covers the TensileLite assembly and HIP-source toolchains, the extop generators and their device library, the CMake path that builds the rocRoller custom kernels, and hipBLASLt's own HIP device code by way of GPU_BUILD_TARGETS.

The consequence is that gfx950 code objects are now built SRAMECC_ON with XNACK_ANY throughout, so they will not load on a gfx950 part running with sramecc disabled. That is the accepted trade. rocRoller's just-in-time compilation path already assembled to the target id the device itself reports, through the amd_comgr compiler library, so this converges the checked-in path with the just-in-time path rather than introducing a new policy.

Test Plan

  • Assemble each of the 89 checked-in gfx950 kernels twice, once against a bare gfx950 target id and once against gfx950:sramecc+, and count successes.
  • Configure and build the full tree from clean, then read the feature settings out of every gfx950 code object produced and attempt to load them on an MI355X part.
  • Check the colon-free path invariant over both the source tree and the build tree, and confirm at configure time that GPU_TARGETS and TensileCreateLibrary --architecture= receive the bare name while device code is built for the qualified target id.
  • test_arch_steppings.py: new Python tests over the architecture feature tables, including a drift guard asserting that the CMake and Python feature tables agree, a check that the feature tail never reaches a path or a file name, a check that GPU_TARGETS itself never grows a feature tail, and a check that every checked-in gfx950 kernel declares the target id the build assembles to.
  • test_custom_kernel_host.py and the toolchain component characterization suite, as regression cover for the surrounding code.

Test Result

Completed evidence:

  • Assembly, measured over the 89 gfx950 kernels present at this commit: 0 of 89 assemble against a bare gfx950 target id, 89 of 89 against gfx950:sramecc+.
  • A full build from a clean tree configures and completes. Of the 318 gfx950 code objects produced, all 318 report SRAMECC_ON with XNACK_ANY and none reports anything else. All 318 load on MI355X.
  • No file or directory in either the source tree or the build tree contains a colon. Configure shows the split directly: GPU_TARGETS and TensileCreateLibrary --architecture= receive gfx950, while device code is built for target id gfx950:sramecc+.
  • test_arch_steppings.py adds 16 tests, all passing. The kernel-declaration test asserts a relation over the shipped sources rather than a hardcoded count, so it does not need editing as kernels are added.
  • test_custom_kernel_host.py passes 54 tests and the toolchain component characterization suite passes 34.

Pending: automated GitHub checks have not reported on this branch. The MI355X load check was run locally; no shared CI lane exercises gfx950 code-object loading today.

Submission Checklist

Risk level

Medium. Every gfx950 code object the build produces changes its target id, so a gfx950 part running with sramecc disabled can no longer load them. That is a deliberate narrowing, stated above, and it matches what rocRoller's just-in-time path already did. Architecture names in paths, file names, and the --architecture= allow-list are unchanged, and the drift guard plus the colon-free path checks are there to catch a regression in that split.

@jaopaulolc jaopaulolc changed the title build(tensilelite): build gfx950 objects to gfx950:sramecc+ build(tensilelite): assemble gfx950 objects to the gfx950:sramecc+ target id Oct 20, 2026
@codecov

codecov Bot commented Oct 20, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 96.77419% with 1 line in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...ensilelite/tensilelite/TensileCreateLibrary/Run.py 83.33% 1 Missing ⚠️
Additional details and impacted files
@@                        Coverage Diff                        @@
##           users/jolabega/kfa-rr-mx-1-fixes   #14044   +/-   ##
=================================================================
  Coverage                             77.35%   77.35%           
=================================================================
  Files                                  2896     2896           
  Lines                                486053   486070   +17     
  Branches                              73475    73478    +3     
=================================================================
+ Hits                                 375957   375975   +18     
  Misses                                85264    85264           
+ Partials                              24832    24831    -1     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 50.82% <ø> (ø)
TensileLite-Unit 84.15% <96.77%> (+<0.01%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from 7157051
hipBLASLt 68.66% <ø> (+<0.01%) ⬆️
hipCUB 82.68% <ø> (ø) Carriedforward from 7157051
hipDNN 86.92% <ø> (ø) Carriedforward from 7157051
hipFFT 43.18% <ø> (ø) Carriedforward from 7157051
hipRAND 76.12% <ø> (ø) Carriedforward from 7157051
hipSOLVER 69.67% <ø> (ø) Carriedforward from 7157051
hipSPARSE 87.03% <ø> (ø) Carriedforward from 7157051
rocBLAS 68.95% <ø> (ø) Carriedforward from 7157051
rocFFT 47.36% <ø> (ø) Carriedforward from 7157051
rocRAND 57.32% <ø> (ø) Carriedforward from 7157051
rocSOLVER 76.74% <ø> (ø) Carriedforward from 7157051
rocSPARSE 83.54% <ø> (ø) Carriedforward from 7157051
rocThrust 93.15% <ø> (ø) Carriedforward from 7157051

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...lt/tensilelite/tensilelite/Common/Architectures.py 87.97% <100.00%> (+0.54%) ⬆️
...aslt/tensilelite/tensilelite/Toolchain/Assembly.py 95.00% <100.00%> (ø)
...slt/tensilelite/tensilelite/Toolchain/Component.py 99.43% <100.00%> (+<0.01%) ⬆️
...ensilelite/tensilelite/TensileCreateLibrary/Run.py 89.19% <83.33%> (ø)

... and 1 file with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@jaopaulolc
jaopaulolc added this pull request to stack #14055 October 20, 2026 20:20
@therock-pr-bot

therock-pr-bot Bot commented Oct 21, 2026 •

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 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

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

@jaopaulolc
jaopaulolc removed this pull request from stack #14055 October 21, 2026 19:52
@jaopaulolc
jaopaulolc added this pull request to stack #14148 October 21, 2026 20:25
@jaopaulolc
jaopaulolc removed this pull request from stack #14148 October 21, 2026 20:27
@jaopaulolc
jaopaulolc added this pull request to stack #14149 October 21, 2026 20:31
@therock-pr-bot

therock-pr-bot Bot commented Oct 22, 2026 •

Copy link
Copy Markdown

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

@jaopaulolc
jaopaulolc marked this pull request as ready for review October 22, 2026 18:01
@jaopaulolc
jaopaulolc requested review from a team as code owners October 22, 2026 18:01
@jaopaulolc
jaopaulolc removed this pull request from stack #14149 October 23, 2026 17:20
@jaopaulolc
jaopaulolc added this pull request to stack #14297 October 23, 2026 17:20
@jaopaulolc
jaopaulolc removed this pull request from stack #14297 October 27, 2026 15:38
@jaopaulolc
jaopaulolc added this pull request to stack #14479 October 27, 2026 16:13
@jaopaulolc jaopaulolc changed the title build(tensilelite): assemble gfx950 objects to the gfx950:sramecc+ target id feat(hipblaslt): assemble gfx950:sramecc+ and ship rocRoller MX kernels Oct 27, 2026
@jaopaulolc
jaopaulolc removed this pull request from stack #14479 October 27, 2026 17:20
@jaopaulolc
jaopaulolc added this pull request to stack #14496 October 27, 2026 17:21
@jaopaulolc
jaopaulolc removed this pull request from stack #14496 October 27, 2026 17:33
@jaopaulolc
jaopaulolc requested review from a team and newling as code owners October 30, 2026 14:13
@jaopaulolc
jaopaulolc merged commit 052f41c into develop Oct 30, 2026
3 of 4 checks passed
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.