Skip to content

fix(rocisa): stop reading a closed pipe fd in POSIX run() - #13668

Merged
talumbau merged 1 commit into
developfrom
users/talumbau/fix-rocisa-run-closed-fd
Oct 16, 2026
Merged

talumbau merged 1 commit into
developfrom
users/talumbau/fix-rocisa-run-closed-fd

Conversation

@talumbau

Copy link
Copy Markdown
Collaborator

JIRA ID : AIHPBLAS-5173

Motivation

run() in tensilelite/rocisa/rocisa/src/helper.cpp runs the assembler for every capability probe (tryAssembler in include/hardware_caps.hpp). The POSIX version has four problems:

  1. It reads from a closed fd. The parent closes the pipe's read end p[0] and then calls read(p[0], ...). In a single-threaded process this fails with EBADF on every probe (114 times per gfx942 init on current develop). In a multithreaded process another thread can reuse that fd number between the close and the read, and the read then silently consumes up to 128 bytes of an unrelated stream.
  2. It never captures the child's output. The child dup2s its stdout/stderr onto the pipe's read end, which is not writable. The assembler's output is lost, and result (printed only in debug mode) is always empty.
  3. The pipe has no O_CLOEXEC. A child forked by another thread at the same moment can inherit the write end and keep this child's stdin from seeing EOF.
  4. With stdin closed, every capability probe succeeds (found while testing this fix). If the process starts with fd 0 closed, pipe() returns fd 0 as the read end. dup2(0, 0) does nothing, and the child's close(p[0]) then closes its own stdin. The assembler reads empty input, exits 0, so every probe reports "supported". With develop, 565 cap entries across 9 ISAs change when the process is started with <&-: 537 asm caps flip from 0 to 1, 4 asm bugs flip from False to True, and the rest are wrong counter or limit values or missing keys.

Technical Details

POSIX branch of run() only. The Windows branch is unchanged: it reads its own stdout pipe until EOF, so it does not have this bug.

  • pipe() → pipe2(p, O_CLOEXEC). The child still gets the read end as stdin because dup2 clears FD_CLOEXEC on the new descriptor. The original pipe fds are closed by execvp, so the explicit closes in the child are gone.
  • If the read end already is fd 0 (stdin was closed), dup2 would do nothing and the fd would close on exec. In that case the child clears FD_CLOEXEC with fcntl instead. This fixes problem 4.
  • When debug is false, the child's stdout/stderr go to /dev/null instead of the pipe's unwritable read end, so output is discarded on purpose rather than as a side effect of failed writes.
  • The read() from the closed fd is gone. result is now an empty string, which is what it always was in practice.
  • If fork() fails, run() closes the pipe and throws. Before, it fell into the parent branch and called waitpid(-1, ...), which can reap an unrelated child.

Scope decision: output is still discarded, not captured. result is only printed when debug is true. In debug mode the child inherits the parent's stdout/stderr, so the assembler's output already reaches the terminal. Capturing it would require a second pipe, a read-until-EOF loop before waitpid, and changes to debug-mode behavior, and it would add no information. The deadlock concern (the child filling an output pipe while the parent is still writing stdin) never arises with discarded output, and the input is a one-line asm snippet that fits well within the pipe buffer. The std::pair<int, std::string> signature stays the same, so the Windows path, which does capture output, is unaffected.

Test Plan

  • Reproduce the strace check from the bug report (ti.init((9,4,2), "/opt/rocm/llvm/bin/clang++", False) under strace -e trace=pipe,pipe2,close,read) with rocisa built from develop and from this branch.
  • Compare the full getAsmCaps(), getArchCaps(), getRegCaps(), and getAsmBugs() dicts for gfx908, gfx90a, gfx942, gfx950, gfx1030, gfx1100, gfx1200, gfx1201, and gfx1250 between the two builds. Repeat with the process started with stdin closed.
  • Compare debug=True output for a gfx942 init between the two builds.
  • New regression test rocisa/test/test_asm_caps_closed_stdin.py: probe gfx90a asm caps in a subprocess with fd 0 closed and compare against a normal probe. It needs a working amdclang++, same as the existing test_asm_caps_cwd.py.
  • Run the rocisa test suite (pytest -n 8 rocisa/test from outside the source tree) and the TensileLite unit suite (pytest -n 8 tensilelite/Tests/unit) against both builds.
  • I did not add a multithreaded test. rocIsa.init holds the GIL for the whole call, so Python threads cannot run probes concurrently, and the fd-reuse race in problem 1 cannot be triggered from Python. The O_CLOEXEC and closed-fd fixes are covered by the strace check rather than by an automated test.

Test Result

All results were collected on a local Linux machine with ROCm /opt/rocm amdclang++. Both builds use the same conda env contents and differ only in rocisa.

  • strace (gfx942 init, single process, no -f): read(...) = -1 EBADF occurs 114 times on develop and 0 times on this branch. All 114 pipes are now pipe2(..., O_CLOEXEC). The bug report's count of 89 came from an older develop commit, and its strace -f grep missed reads split into <unfinished ...> / resumed lines.
  • Caps equality: the caps JSON for all 9 ISAs is byte-identical between develop and this branch (cmp clean). With stdin closed, this branch still produces identical output. Develop differs in 565 entries, almost all of them spurious "supported" results.
  • Debug mode: init((9,4,2), ..., True) output is identical between builds (798 lines, 76 assembler error: lines in each).
  • rocisa tests: develop: 1 failed (the new test_asm_caps_probe_matches_with_closed_stdin, as expected), 265 passed, 4 skipped. This branch: 266 passed, 4 skipped.
  • TensileLite unit tests: both builds: 10479 passed, 384 skipped, 2 xfailed, 1 xpassed.
  • No full hipBLASLt or device-library build was run. The change is confined to rocisa's host-side probe helper.

Submission Checklist

Risk level

Low. The change is limited to the POSIX run() helper used for assembler capability probes. Capability results are byte-identical to develop for 9 ISAs in the normal case, and the only behavior change is in the closed-stdin case, where develop was wrong. pipe2 is Linux/glibc-specific, and rocisa's POSIX target is Linux.

🤖 Generated with Claude Code

run() drives every assembler capability probe (tryAssembler). On POSIX it:

- closed the pipe's read end in the parent and then read() from it. That
  fails with EBADF on every probe (114 per gfx942 init); with threads the
  fd number can be reused and the read can consume another stream's data.
- dup2'd the child's stdout/stderr onto the pipe's read end, which is not
  writable, so the "result" string was always empty.
- created the pipe without O_CLOEXEC, so children forked concurrently by
  other threads could inherit the write end and hold off this child's EOF.
- if the process started with fd 0 closed, pipe() returned fd 0 as the read
  end and the child's close(p[0]) closed its own stdin. The assembler then
  read empty input, exited 0, and every capability reported as supported.

Use pipe2(O_CLOEXEC), drop the bogus read, send the child's output to
/dev/null explicitly when not in debug mode (in debug mode it is already
inherited from the parent and printed there), clear FD_CLOEXEC when the read
end is already fd 0, and throw if fork() fails instead of calling
waitpid(-1). Capabilities are unchanged in the normal case.

Add a regression test that compares asm caps probed with stdin closed
against a normal probe.

JIRA ID : AIHPBLAS-5173

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@therock-pr-bot

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

Copy link
Copy Markdown

✅ 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

Copy link
Copy Markdown

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

@newling newling left a comment

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.

An agent review found zero issues with the code changes (the pipe vs pipe2 details are beyond me). I also got the agent to verify that this has negligable performance impact, so that's good.

For future consideration: I didn't realize this "process-launching plumbing" was how rocisa checked for instruction support, I wonder if there's a more integrated way to do this querying...

LGTM

@codecov

codecov Bot commented Oct 14, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 5.00000% with 19 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...hipblaslt/tensilelite/rocisa/rocisa/src/helper.cpp 5.00% 16 Missing and 3 partials ⚠️
Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #13668      +/-   ##
===========================================
- Coverage    74.70%   74.70%   -0.01%     
===========================================
  Files         2881     2881              
  Lines       480344   480437      +93     
  Branches     71286    71294       +8     
===========================================
+ Hits        358841   358864      +23     
- Misses       97197    97265      +68     
- Partials     24306    24308       +2     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 49.66% <5.00%> (-0.04%) ⬇️
TensileLite-Unit 83.57% <ø> (ø)
hipBLAS 90.62% <ø> (ø) Carriedforward from e160811
hipBLASLt 43.29% <ø> (-0.04%) ⬇️
hipCUB 82.68% <ø> (ø) Carriedforward from e160811
hipDNN 86.80% <ø> (ø) Carriedforward from e160811
hipFFT 42.00% <ø> (ø) Carriedforward from e160811
hipRAND 76.12% <ø> (ø) Carriedforward from e160811
hipSOLVER 69.45% <ø> (ø) Carriedforward from e160811
hipSPARSE 86.99% <ø> (ø) Carriedforward from e160811
rocBLAS 68.73% <ø> (ø) Carriedforward from e160811
rocFFT 46.60% <ø> (ø) Carriedforward from e160811
rocRAND 57.32% <ø> (ø) Carriedforward from e160811
rocSOLVER 76.83% <ø> (ø) Carriedforward from e160811
rocSPARSE 77.05% <ø> (ø) Carriedforward from e160811
rocThrust 93.15% <ø> (ø) Carriedforward from e160811

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

Files with missing lines Coverage Δ
...hipblaslt/tensilelite/rocisa/rocisa/src/helper.cpp 45.16% <5.00%> (-14.45%) ⬇️

... and 3 files 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.

@Alex-Vasile Alex-Vasile left a comment

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.

Looks good

@talumbau
talumbau merged commit 0477a8a into develop Oct 16, 2026
346 of 364 checks passed
@talumbau
talumbau deleted the users/talumbau/fix-rocisa-run-closed-fd branch October 16, 2026 22:17
assistant-librarian Bot pushed a commit to ROCm/hipBLASLt that referenced this pull request Oct 16, 2026
fix(rocisa): stop reading a closed pipe fd in POSIX run()
 (#13668)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

JIRA ID : AIHPBLAS-5173

## Motivation

`run()` in `tensilelite/rocisa/rocisa/src/helper.cpp` runs the assembler
for every capability probe (`tryAssembler` in
`include/hardware_caps.hpp`). The POSIX version has four problems:

1. **It reads from a closed fd.** The parent closes the pipe's read end
`p[0]` and then calls `read(p[0], ...)`. In a single-threaded process
this fails with `EBADF` on every probe (114 times per gfx942 `init` on
current develop). In a multithreaded process another thread can reuse
that fd number between the `close` and the `read`, and the `read` then
silently consumes up to 128 bytes of an unrelated stream.
2. **It never captures the child's output.** The child `dup2`s its
stdout/stderr onto the pipe's read end, which is not writable. The
assembler's output is lost, and `result` (printed only in debug mode) is
always empty.
3. **The pipe has no `O_CLOEXEC`.** A child forked by another thread at
the same moment can inherit the write end and keep this child's stdin
from seeing EOF.
4. **With stdin closed, every capability probe succeeds** (found while
testing this fix). If the process starts with fd 0 closed, `pipe()`
returns fd 0 as the read end. `dup2(0, 0)` does nothing, and the child's
`close(p[0])` then closes its own stdin. The assembler reads empty
input, exits 0, so every probe reports "supported". With develop, 565
cap entries across 9 ISAs change when the process is started with `<&-`:
537 asm caps flip from 0 to 1, 4 asm bugs flip from False to True, and
the rest are wrong counter or limit values or missing keys.

## Technical Details

POSIX branch of `run()` only. The Windows branch is unchanged: it reads
its own stdout pipe until EOF, so it does not have this bug.

- `pipe()` → `pipe2(p, O_CLOEXEC)`. The child still gets the read end as
stdin because `dup2` clears `FD_CLOEXEC` on the new descriptor. The
original pipe fds are closed by `execvp`, so the explicit `close`s in
the child are gone.
- If the read end already is fd 0 (stdin was closed), `dup2` would do
nothing and the fd would close on exec. In that case the child clears
`FD_CLOEXEC` with `fcntl` instead. This fixes problem 4.
- When `debug` is false, the child's stdout/stderr go to `/dev/null`
instead of the pipe's unwritable read end, so output is discarded on
purpose rather than as a side effect of failed writes.
- The `read()` from the closed fd is gone. `result` is now an empty
string, which is what it always was in practice.
- If `fork()` fails, `run()` closes the pipe and throws. Before, it fell
into the parent branch and called `waitpid(-1, ...)`, which can reap an
unrelated child.

**Scope decision: output is still discarded, not captured.** `result` is
only printed when `debug` is true. In debug mode the child inherits the
parent's stdout/stderr, so the assembler's output already reaches the
terminal. Capturing it would require a second pipe, a read-until-EOF
loop before `waitpid`, and changes to debug-mode behavior, and it would
add no information. The deadlock concern (the child filling an output
pipe while the parent is still writing stdin) never arises with
discarded output, and the input is a one-line asm snippet that fits well
within the pipe buffer. The `std::pair<int, std::string>` signature
stays the same, so the Windows path, which does capture output, is
unaffected.

## Test Plan

- Reproduce the strace check from the bug report (`ti.init((9,4,2),
"/opt/rocm/llvm/bin/clang++", False)` under `strace -e
trace=pipe,pipe2,close,read`) with rocisa built from develop and from
this branch.
- Compare the full `getAsmCaps()`, `getArchCaps()`, `getRegCaps()`, and
`getAsmBugs()` dicts for gfx908, gfx90a, gfx942, gfx950, gfx1030,
gfx1100, gfx1200, gfx1201, and gfx1250 between the two builds. Repeat
with the process started with stdin closed.
- Compare `debug=True` output for a gfx942 init between the two builds.
- New regression test `rocisa/test/test_asm_caps_closed_stdin.py`: probe
gfx90a asm caps in a subprocess with fd 0 closed and compare against a
normal probe. It needs a working `amdclang++`, same as the existing
`test_asm_caps_cwd.py`.
- Run the rocisa test suite (`pytest -n 8 rocisa/test` from outside the
source tree) and the TensileLite unit suite (`pytest -n 8
tensilelite/Tests/unit`) against both builds.
- I did not add a multithreaded test. `rocIsa.init` holds the GIL for
the whole call, so Python threads cannot run probes concurrently, and
the fd-reuse race in problem 1 cannot be triggered from Python. The
`O_CLOEXEC` and closed-fd fixes are covered by the strace check rather
than by an automated test.

## Test Result

All results were collected on a local Linux machine with ROCm
`/opt/rocm` amdclang++. Both builds use the same conda env contents and
differ only in rocisa.

- **strace** (gfx942 init, single process, no `-f`): `read(...) = -1
EBADF` occurs 114 times on develop and 0 times on this branch. All 114
pipes are now `pipe2(..., O_CLOEXEC)`. The bug report's count of 89 came
from an older develop commit, and its `strace -f` grep missed reads
split into `<unfinished ...>` / `resumed` lines.
- **Caps equality:** the caps JSON for all 9 ISAs is byte-identical
between develop and this branch (`cmp` clean). With stdin closed, this
branch still produces identical output. Develop differs in 565 entries,
almost all of them spurious "supported" results.
- **Debug mode:** `init((9,4,2), ..., True)` output is identical between
builds (798 lines, 76 assembler `error:` lines in each).
- **rocisa tests:** develop: 1 failed (the new
`test_asm_caps_probe_matches_with_closed_stdin`, as expected), 265
passed, 4 skipped. This branch: 266 passed, 4 skipped.
- **TensileLite unit tests:** both builds: 10479 passed, 384 skipped, 2
xfailed, 1 xpassed.
- No full hipBLASLt or device-library build was run. The change is
confined to rocisa's host-side probe helper.

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

## Risk level

Low. The change is limited to the POSIX `run()` helper used for
assembler capability probes. Capability results are byte-identical to
develop for 9 ISAs in the normal case, and the only behavior change is
in the closed-stdin case, where develop was wrong. `pipe2` is
Linux/glibc-specific, and rocisa's POSIX target is Linux.

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

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

jinchen62 pushed a commit that referenced this pull request Oct 28, 2026
makeIsaInfoMap probes each ISA with hundreds of serial assembler runs, one
ISA after another. TensileLogic asks for all 22 supported ISAs, which took
~66 s before validation could start, and TensileCreateLibrary waits on
TensileLogic.

Add rocIsa.initMany, which runs initAsmCaps for every not-yet-initialized
ISA on its own thread (GIL released), then fills the shared tables
serially. This relies on run() no longer reading a closed fd and creating
its pipe with O_CLOEXEC (#13668). Probed caps are identical for
all 22 ISAs; the probe drops to ~6 s.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jinchen62 pushed a commit that referenced this pull request Oct 29, 2026
makeIsaInfoMap probes each ISA with hundreds of serial assembler runs, one
ISA after another. TensileLogic asks for all 24 supported ISAs, which took
about 18 s on a 256-thread host before validation could start, and
TensileCreateLibrary waits on TensileLogic.

Add rocIsa.initMany, which runs initAsmCaps for every not-yet-initialized
ISA on its own thread (GIL released), then fills the shared tables
serially. On POSIX this relies on run() no longer reading a closed fd and
creating its pipe with O_CLOEXEC (#13668). Probed caps are identical for
all 24 ISAs; the probe drops to about 1 s.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jinchen62 added a commit that referenced this pull request Oct 30, 2026
…13674)

JIRA ID : AIHPBLAS-5196

## Motivation
TensileLogic validation gates TensileCreateLibrary, and it probes
assembler capabilities for every supported ISA (24 on develop) one after
another, with hundreds of assembler runs each. On Windows CI, probing is
more than half of TensileLogic's 430 s.

This PR used to also pass `--jobs` to TensileLogic from CMake and
memoize YAML scalar tag resolution. Both are dropped. TensileLogic
already runs one worker per CPU, because `ParallelMap2` ignores its
`procs` argument, so `--jobs` only changed the batch size; and neither
change gave a measurable gain on CI.

## Technical Details
- **probe assembler capabilities for all ISAs concurrently:** new
`rocIsa.initMany` runs `initAsmCaps` for each uninitialized ISA on its
own thread with the GIL released, then fills the shared tables serially.
`makeIsaInfoMap` calls it instead of `init()` for each ISA. On POSIX
this relies on #13668 (merged), which fixed `run()` reading a closed fd
and gave its pipe `O_CLOEXEC`.
- **keep concurrent Windows probes out of each other's pipes:**
`CreateProcess` hands each child every inheritable handle in the
process, and the Windows `run()` created both pipes inheritable before
clearing inheritance on the parent's ends. A probe starting in between
could pass one probe's stdin write end to another probe's assembler, and
two probes holding each other's would wait forever. The pipes are now
created non-inheritable, and only the child's ends are made inheritable.
- **failure handling** (from James Newling's review): the probe threads
are `std::jthread`s, so if starting one fails, the ones already running
are joined and the error reaches Python instead of aborting the
interpreter. If any probe fails, `initMany` raises once all of them have
finished and stores nothing.
- **smaller API points** (same review): `initMany` accepts any iterable
of ISAs, and `init` and `initMany` build an ISA's info with one shared
helper.
- **stinkytofu adaptor:** TensileLite's unit tests and gfx1250 builds
run under the stinkytofu backend, whose `rocIsa` facade had no
`initMany`; the adaptor now has one. It probes all the ISAs that
stinkytofu has no caps for with one native `initMany` call in a single
subprocess, instead of one subprocess per ISA (James Newling's commit).
- **test:** probes gfx942, gfx950 and gfx1201 once with `init()` per ISA
and once with `initMany()`, each in a fresh interpreter because rocIsa
caches capabilities per process, and requires identical asm, arch and
register caps and asm bugs.

### Series overview

This is one of six independent build-time PRs for the hipBLASLt device
library (AIHPBLAS-5192). Each one is based on develop and can be
reviewed and merged on its own.

| PR | Change | gfx942 hipBLASLt build on CI* | gfx110X on Windows* |
|---|---|---|---|
| #13671 | Compile helper kernels while generating asm kernels | −130 s
| −150 s |
| #14572 | Don't return kernel source from codegen workers | −9 s | −132
s |
| #13672 | rocisa: slimmer duplicate-assignment register graph | −94 s†
| +49 s |
| #13673 | Lighter logic loading | −82 s | −76 s |
| #13674 | Probe assembler capabilities for all ISAs concurrently | −9 s
(TensileLogic −21 s) | −198 s |
| #13675 | Link and bundle code objects in parallel | −135 s | −43 s |

\*Each PR's own TheRock Multi-Arch CI run on Oct 28, one run each,
against the median of other PRs' hipBLASLt builds on the same kind of
runner within a day: 1,725 s for gfx942 on 96-thread runners, about 915
s for gfx110X on Windows. On one kind of Linux runner, 80% of builds are
within about 20 s of the median, so smaller differences are noise;
Windows builds vary much more, so a single Windows run is only a rough
guide.

†On an 80-thread runner, against that runner class's median of 1,896 s.

In the earlier stacked version of the series, all of these together took
the gfx942 build from 1,697 to 1,221 s (−28%), hipSPARSELt down 22% and
the gfx950 hipBLASLt build down 23%.

## Test Plan
- TensileLite unit suite (`tensilelite/Tests/unit`) and rocisa suite
(`rocisa/test`), compared test by test with the same run on develop
`9d729d3280b`.
- `gemm-libs` builds for gfx942 and gfx950 of develop and of all PRs in
the series applied together, with the helper-kernel cache disabled, and
their outputs compared byte for byte.
- TheRock Multi-Arch CI for build times.

## Test Result
- Probed asm, arch and register caps and asm bugs are identical between
`init()` per ISA and `initMany()` for all 24 supported ISAs, natively
(1.2 s instead of 18.2 s on an idle 256-thread host) and through the
stinkytofu adaptor (2.1 s instead of 46 s on a busy one).
- Unit and rocisa suites: no test newly failing compared with develop,
natively and under the stinkytofu backend that CI's unit lane uses; the
adaptor's own tests too. The new tests pass.
- Byte comparison of the whole series against develop: 1,215 of 1,218
gfx942 files and 641 of 644 gfx950 files under `Tensile/library`,
`library` and `device-library` are byte-identical. The other three
differ only because the two builds come from different commits and
directories: `hipblaslt-version.h` embeds the commit, `MatchTable.yaml`
embeds the logic-file paths, and `Kernels.so-000-<arch>.hsaco` gets a
HIP CUID derived from its build path. Its code sections are identical;
only the CUID symbol and the symbol tables it reorders differ.
- CI, this PR alone: TensileLogic 120 → 99 s on gfx942, 46 → 21 s on
gfx950 and 430 → 179 s on Windows gfx110X. The whole hipBLASLt build is
9 s faster on gfx942, which is within noise, and 198 s faster on Windows
(one run each).

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

## Risk level
Medium. It adds threads to rocisa's probe code. On POSIX that relies on
#13668; on Windows on the pipe fix above, which I couldn't run locally,
but CI's rocISA Windows lane and the Windows hipBLASLt build exercise
it. Capability results are verified identical.

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

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: jinchen62 <jinchenye62@gmail.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: James Newling <james.newling@gmail.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