Skip to content

fix(tensilelite): emit RPATH for tensilelite-client so the built host library wins - #12278

Merged
AlexBrownAMD merged 1 commit into
GemmFromAnywherefrom
users/talumbau/fix-tensilelite-client-runpath
Sep 17, 2026
Merged

AlexBrownAMD merged 1 commit into
GemmFromAnywherefrom
users/talumbau/fix-tensilelite-client-runpath

Conversation

@talumbau

Copy link
Copy Markdown
Collaborator

JIRA ID : AIHPBLAS-4086

Targets GemmFromAnywhere (#11155) rather than develop, so @jinchen62 can test it directly on that branch.

Motivation

The preliminary(hipblaslt) job on #11155 fails with 127 failed, 5 passed in the tensilelite common tests, while the same tests pass locally on a gfx942 box. The entire 140MB CI log contains exactly one distinct error, repeated once per failed test:

.../caches/4463c0a054b3/source/library/gfx942/TensileLibrary.yaml:55:1: error: unknown key 'solutions'
terminate called after throwing an instance of 'std::runtime_error'
  what():  Failed to load solution library
→ Aborted (core dumped), exit 134

The FileNotFoundError on 00_Final.csv that the test_use_cache tests report is purely downstream — the client aborted, so no benchmark CSV was ever written.

This is not a bug in #11155's logic. It is a pre-existing library-shadowing bug in the test harness that #11155 is simply the first PR to expose, because it is the first to change the solution schema (CustomKernelName → customKernel).

Technical Details

  1. tox.ini builds tensilelite-client into build_tmp/ and hands it to pytest via --prebuilt-client.
  2. That client links libtensilelite-host.so.1 — where the YAML deserializer (src/llvm/Loading.cpp) lives — and recorded a RUNPATH, not an RPATH.
  3. RUNPATH is searched after LD_LIBRARY_PATH. tox.ini sets LD_LIBRARY_PATH = {ROCM_PATH}/lib:..., and the Jenkins job independently exports LD_LIBRARY_PATH=/opt/rocm/lib.
  4. tensilelite/CMakeLists.txt:170 installs the tensilelite-host target, so on a CI container with hipBLASLt installed, /opt/rocm/lib/libtensilelite-host.so.1 exists.
  5. The freshly built host library is therefore silently bypassed; the client runs against the older installed one, which does not understand the new schema.

Why it cannot fail on a typical dev box: local /opt/rocm/lib has only libhipblaslt.so* and no libtensilelite-host.so, so the lookup falls through to RUNPATH and picks up the correct fresh build.

The fix links the client with -Wl,--disable-new-dtags, emitting RPATH, which is searched before LD_LIBRARY_PATH. tensilelite-client has no install() rule — it is a build/test-only binary — so there is no packaging or install-time impact, and CMAKE_INSTALL_RPATH is untouched.

Test Plan

On a gfx942 box (ROCm 7.2.4), with a stale libtensilelite-host.so.1 built from develop@merge-base placed first on LD_LIBRARY_PATH to emulate the CI container:

  1. Direct client check: linker tag, library resolution, and a real run against the exact ClientParameters.ini from a failing cache dir.
  2. End-to-end: CI's own pipeline — tox -e py3 with --cxx-flags-release=-O3, sdist install, invoke build-client — scoped to the previously failing use_cache tests.
  3. Negative control: revert the flag, relink, re-run with tox's exact LD_LIBRARY_PATH.

Test Result

Direct client check:

linker tag resolves to client exit
before RUNPATH stale lib 134, Failed to load solution library
after RPATH build-tree lib 0

End-to-end with the stale lib injected: 5 passed (test_use_cache.py::test_use_cache, test_compile_and_benchmark, test_use_cache_creates_cache_dir, test_use_cache_different_params, test_build_only.py::test_use_cache) — all previously failing in CI.

Negative control: with the flag reverted, the client picks the stale library and aborts 134 under tox's exact LD_LIBRARY_PATH. So the green run is attributable to the fix, not to the stale library never being on the path.

Caveats, stated plainly:

  • Verification used a develop-built stand-in for the stale library, not the real ROCm-installed one from CI's container (which also runs Clang 24.0.0 vs 22.0.0 locally). The mechanism is proven; the exact CI binary was not reproducible off-box. Consequently my repro reports unknown key 'customKernel' where CI reports unknown key 'solutions' — same failure mode, different key, because the CI container's library is older than develop@merge-base.
  • The real confirmation is CI itself: if this hypothesis is right, preliminary(hipblaslt) on GemmFromAnywhere + this commit should go green.
  • cpu-gemm-driver still has RUNPATH and the identical latent hazard. Deliberately left alone to keep this change minimal; it is a one-line follow-up.

Submission Checklist

Risk level

Low. Six lines in one CMakeLists.txt, affecting the link of a single non-installed test binary. The worst case is that LD_LIBRARY_PATH can no longer override the host library for tensilelite-client — which is precisely the behaviour being fixed, but worth noting if anyone deliberately relies on that override to A/B a library build.

🤖 Generated with Claude Code

… library wins

tensilelite-client links libtensilelite-host.so.1 and recorded a RUNPATH.
RUNPATH is searched after LD_LIBRARY_PATH, and tox.ini sets
LD_LIBRARY_PATH=${ROCM_PATH}/lib. On any machine where hipBLASLt is already
installed into ${ROCM_PATH}/lib -- which the CI containers are, since
tensilelite/CMakeLists.txt installs the tensilelite-host target -- the client
silently loads the installed host library instead of the one just built into
build_tmp/.

Any change to the library YAML/msgpack schema then fails in CI while passing
locally: the fresh generator writes the new schema, the stale deserializer
rejects it with "unknown key '<newkey>'" and the client aborts with SIGABRT.
The pytest failure surfaces as a misleading FileNotFoundError on 00_Final.csv,
because the aborted client never wrote the benchmark output.

Link with -Wl,--disable-new-dtags so the tag is RPATH, which is searched before
LD_LIBRARY_PATH. tensilelite-client has no install() rule, so this affects only
the build/test binary.

Co-Authored-By: Claude <noreply@anthropic.com>
@talumbau
talumbau requested a review from a team as a code owner September 17, 2026 21:56
@AlexBrownAMD
AlexBrownAMD merged commit 5b3aae7 into GemmFromAnywhere Sep 17, 2026
29 checks passed
@AlexBrownAMD
AlexBrownAMD deleted the users/talumbau/fix-tensilelite-client-runpath branch September 17, 2026 22:04
AlexBrownAMD pushed a commit that referenced this pull request Sep 18, 2026
…match the linked LLVM (#12295)

JIRA ID : AIHPBLAS-4086

Targets `GemmFromAnywhere` (#11155) rather than `develop`, so @jinchen62
can take it directly on that branch.

## Motivation

`preliminary(hipblaslt)` on #11155 fails with **127 failed, 8 passed,
1144 skipped** in the tensilelite common tests, while the same tests
pass on a local gfx942 box. The entire ~140MB CI log contains exactly
one distinct error, repeated once per failed test:

```
.../caches/<hash>/source/library/gfx942/TensileLibrary.yaml:55:1: error: unknown key 'solutions'
terminate called after throwing an instance of 'std::runtime_error'
  what():  Failed to load solution library
→ Aborted (core dumped), exit 134
```

The `FileNotFoundError` on `00_Final.csv` that `test_use_cache` reports
is a downstream symptom — the client aborted, so no benchmark CSV was
written.

**This is not a bug in the Gemm-From-Anywhere logic.** It is a
build-configuration problem that this branch is the first to trigger.

## Technical Details

`tensilelite-host` was being compiled against the **system LLVM 18**
headers (`/usr/include/llvm` → `/usr/lib/llvm-18/include/llvm`) while
linking **ROCm's LLVM 24** static libraries. LLVM 24 added `AliasMap` /
`saveAliasHNode` to `llvm::yaml::Input`, so the class layouts differ.
The compiled code touched `ValidKeys` / `Mapping` / `EC` at the wrong
offsets, which is why the failure looked impossible against the source:
`mapRequired("solutions")` runs, silently yields an empty vector, and
the key is *then* reported as never-seen.

How `-fopenmp` gets onto the host target:

1. This branch links `roc::mxDataGenerator` into `tensilelite-host`.
2. `roc::mxDataGenerator` links `OpenMP::OpenMP_CXX`
(`shared/mxdatagenerator/CMakeLists.txt:14`), whose
`INTERFACE_COMPILE_OPTIONS` carry `-fopenmp=libomp`.
3. `-fopenmp` makes clang add `${ROCM_PATH}/lib/llvm/bin/../include` as
an **implicit system** include dir in order to locate `omp.h`. That is
the same directory as the explicit `-I/opt/rocm/lib/llvm/include`, so
the explicit user `-I` is demoted, and `<llvm/...>` includes fall
through to the system LLVM 18 tree.

develop's `tensilelite-host` has no `-fopenmp` (verified: count 0 vs 1),
which is why only this branch is affected.

Evidence — preprocessed TU comparison of `src/llvm/Loading.cpp`,
identical compiler (Clang 24) for both:

| build | `/usr/include/llvm/` markers in TU | `YAMLTraits.h` resolved
to |
|---|---|---|
| develop | 0 | `/opt/rocm/lib/llvm/include/...` (LLVM 24) |
| this branch | **356** | `/usr/include/llvm/...` (**LLVM 18**) |

Flag isolation via `-E` alone:

```
PR flags (-I + -fopenmp)         -> 356
keep -fopenmp, drop explicit -I  -> 356
keep -I, drop -fopenmp           ->   0     <-- only correct configuration
drop both                        -> 356
```

**The fix:** consume mxDataGenerator's include directories explicitly,
and take the target itself as a `$<LINK_ONLY:...>` dependency so its
`INTERFACE_COMPILE_OPTIONS` never reach our compiles. The pre-swizzle
helper is correct serially, so dropping `-fopenmp` there is safe.

This branch also **reverts #12278** (`5b3aae7f8d5`). That commit was my
earlier, incorrect diagnosis of this same failure — it claimed a
`RUNPATH`/`LD_LIBRARY_PATH` shadowing of `libtensilelite-host.so`. The
CI image contains no `libtensilelite-host` and no `libhipblaslt` at all,
so there was nothing to shadow, and it did not fix anything. The RUNPATH
issue it addressed is real but latent and unrelated; it can be
re-proposed separately against `develop` on its own merits.

## Test Plan

Reproduce and verify inside the actual CI container image
(`mci-base/ubu24-stg1`, Clang 24, gfx942), which is required — the
failure cannot reproduce on a ROCm 7.2.4 / Clang 22 host.

1. Confirm `-fopenmp` is gone from `tensilelite-host` (`grep -c fopenmp
.../tensilelite-host.dir/flags.make`) while the mxdatagen include dir
and `HIPBLASLT_ENABLE_MXDATAGENERATOR` are retained.
2. Run the previously failing client invocation directly against the
generated `TensileLibrary.yaml`.
3. Clean build, full CI scope: `tox -e py3 -- -m common --durations=0`.

## Test Result

1. `-fopenmp` count on `tensilelite-host`: **0**; mxdatagen include dir
present; `HIPBLASLT_ENABLE_MXDATAGENERATOR` still defined.
2. Client exits **0** on the previously failing `TensileLibrary.yaml`
(the one containing `customKernel`).
3. Clean build, full CI scope:

```
========= 132 passed, 1144 skipped, 64 warnings in 2592.60s (0:43:12) ==========
  congratulations :)
```

`unknown key 'solutions'` occurrences: **0**. `FAILED` lines: **0**.
Same 1144 skip count as the CI run, so the same test scope — versus CI's
`127 failed, 8 passed, 1144 skipped`.

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

## Risk level

Low. One `CMakeLists.txt` hunk changing how a single target consumes
mxDataGenerator's usage requirements, plus a revert of a commit that was
verified to have no effect on this failure. No source or
kernel-generation changes.

Two follow-ups deliberately **not** included here, to keep the change
minimal:
- `tensilelite-client-common` still receives `-fopenmp`. It does on
`develop` too, and `develop` passes, so this matches known-good
behaviour — but it is worth a hardening pass.
- More broadly, **any** target on this image that picks up `-fopenmp`
will compile `<llvm/...>` from LLVM 18. A systemic fix would stop
relying on an explicit `-I` outranking clang's implicit system include
dirs.

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

---------

Co-authored-by: Claude <noreply@anthropic.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.

2 participants