[Alex-Vasile]

JIRA ID : AIHPBLAS-5256

Motivation

Each TensileLite benchmark step compiles Kernels.cpp, the C++ source of its helper kernels, with clang. XNACK is a GPU memory-fault replay feature; clang can build code for it on (xnack+), off (xnack-) or for either ("xnack-agnostic", the bare name such as gfx950). Benchmark and test builds compiled every helper kernel twice, once for gfx950:xnack+ and once for gfx950:xnack-, which is half of the helper compile. The helper compile was 1,008 s (39.1%) of a gfx950 per-config build aggregate. Shipped libraries already build their helpers for the bare architecture, and the runtime looks for that file first.

Technical Details

BenchmarkProblems.cmdLineArchsFor(archNames, isaInfoMap, compiler) now returns the bare architecture for every instruction set architecture (ISA) instead of its 2 xnack variants. This covers every architecture gfxVariantMap lists: gfx906, gfx908, gfx90a, gfx942 and gfx950. A step's library then holds Kernels.so-000-<arch>.hsaco instead of the -xnack+ and -xnack- files. A requested stepping still wins, as before.

When the source compiler's default arguments include -fsanitize=..., the target is <arch>:xnack+ on architectures that have that variant, as the library build does for ASAN, the address sanitizer. Tensile's own benchmark toolchain is not built that way today. Both callers pass srcToolchain.compiler.

cmdLineArchsFor also feeds the benchmark step cache key and the helper-kernel cache key. Existing --use-cache artifacts and helper-cache entries therefore miss once and are rebuilt.

Tests: new Tests/unit/test_benchmark_helper_targets.py (20 cases), and 3 test files' fake source toolchains gain a compiler with default_args.

Test Plan

Run from projects/hipblaslt/tensilelite:

• tox -e unit -- tensilelite/Tests/unit/test_benchmark_helper_targets.py tensilelite/Tests/unit/test_benchmark_steps_one_pool.py tensilelite/Tests/unit/test_tensile_backend_config.py tensilelite/Tests/unit/Ductile checks one target per architecture, the sanitizer rule, steppings, a single --offload-arch in the compile command and the changed cache key.
• pre-commit run --files on the changed files checks formatting and lint.

Test Result

Time saved: 3.26-12.26 s of build wall time per config (13.6-38.0%) on 4 gfx950 configs, and 8.27 s (39.6%) on 1 gfx942 build; the helper compile itself about halves. Measured on the original unsplit branch against the commit before it, --build-only of 1 config at a time on the gfx950 (MI355X) host, cold helper cache, 25 cores, mean of 2 runs per side (gfx942 target: 1 run per side). Per config:

config (arch)	wall	CPU (user+sys)	helper compile
gsuasb (gfx950)	10.29 to 7.03 s	36.90 to 33.31 s	6.28 to 3.15 s
gemm_bias_strided_batched_gfx940 (gfx950)	20.72 to 12.84 s	52.33 to 43.84 s	16.09 to 8.16 s
spmm_bf16_sb (gfx950)	24.10 to 20.82 s	223.91 to 216.91 s	13.82 to 7.32 s
bf16_gradient_bias (gfx950)	34.41 to 22.15 s	90.00 to 77.81 s	29.66 to 17.64 s
gemm_bias_strided_batched_gfx940 (gfx942, build only, 1 run)	20.88 to 12.61 s	51.49 to 43.84 s	15.97 to 8.05 s

• Every new helper object against the parent's -xnack+ object: every function's disassembly, the code-object metadata (except the target and build ID) and every kernel descriptor identical. Only the ELF xnack flag changes. Other artifact members were byte-identical; cache.yaml differs only in the helper file names.
• Run phase (--use-cache) of the 4 gfx950 configs: all pass, and all 27 client processes opened only Kernels.so-000-gfx950.hsaco. With the parent, the client first failed to open that file, printed hipModuleLoad failed, then loaded -xnack-.hsaco.
• The parent's gsuasb artifacts missed the cache with this change, were rebuilt and passed. The full Python unit suite: 0 failed.

Not yet run on this branch. Automated GitHub checks are pending.

Related

• Work tracking: AIHPBLAS-5256

Device / Architecture Coverage

This changes the compile targets of the benchmark helper kernels on gfx906, gfx908, gfx90a, gfx942 and gfx950. On xnack-off GPUs the benchmark client now runs xnack-agnostic helper code instead of the -xnack- object: a few more s_nop padding instructions, as shipped libraries already run. The runtime effect was not measured. The overlay requires CI coverage on the affected archs, gfx942 and gfx950 at least. An arch that is only built counts as uncovered.

Submission Checklist

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

Risk level

Medium (3/5): benchmark and test helper kernels change target and file name, and existing benchmark caches miss once. Shipped library builds are unchanged. The helper code's runtime on xnack-off hosts was not measured.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR 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.

🙋 Wish to Override Policy?

[therock-pr-bot]

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