perf(tensilelite): link and bundle code objects in parallel by talumbau · Pull Request #13675 · ROCm/rocm-libraries

talumbau · 2026-10-14T20:21:09Z

JIRA ID : AIHPBLAS-5197

Motivation

After assembly, every code object's ld.lld and clang-offload-bundler steps ran one after another, about 150 s on CI for gfx942, even though each pair is independent.

Technical Details

run linking and bundling of code objects in parallel: buildAssemblyCodeObjectFiles collects (objects, raw, final, target) jobs and runs link and bundle on a ThreadPoolExecutor. The work happens in subprocesses, so threads are enough. Linker response files are already named after their output, so concurrent links don't collide. The returned code-object list keeps its order.
split the CPUs between concurrent links: ld.lld uses every hardware thread by default, so with one link per CPU in flight, each link could still start a thread per CPU. Linker now takes a threads argument, passed as -Xlinker --threads=N; each link gets an equal share of CPUThreadCount(), and the largest links start first so one of them doesn't finish last on its own.
link and bundle in a module-level function: the coverage lane runs with concurrency = multiprocessing and doesn't trace thread-pool workers, so the per-job work lives in _linkAndBundle, which tests call directly.
tests: the CPU share per link, the largest-first order with the output order kept, link and bundle with and without compression, and the linker arguments with and without a response file. From James Newling's review: two links run at once, and a link or bundle failure in a worker reaches the caller.

In the benchmarking path the common tests use, codeObjectFile is never assigned, so each step links a single code object per architecture: the pool has one worker there, and that link gets the same threads as before.

Rebased onto develop, where code objects are built from deduplicated kernel groups and their aliases (#11202), and where the compiler target also depends on the linker, so each link job takes compilerTargetOf(name, linker.path).

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 (merged)	Compile helper kernels while generating asm kernels	−130, −104 s	−145, −230 s
#14572 (merged)	Don't return kernel source from codegen workers	−9, −28 s	−129, +41 s
#13672 (merged)	rocisa: slimmer duplicate-assignment register graph	−93, −94 s	+58, +45 s
#13673	Lighter logic loading	−82 s	−68 s
#13674 (merged)	Probe assembler capabilities for all ISAs concurrently	−10, +6 s (TensileLogic −21 s)	−193, −284 s
#13675	Link and bundle code objects in parallel	−135, −164 s	−37, +22 s

*One value per TheRock Multi-Arch CI run of the PR's current change, each against the median of other PRs' hipBLASLt builds on the same kind of runner within a day (about 1,725 s for gfx942 on 96-thread runners, 1,895 s on 80-thread runners, and 910 s for gfx110X on Windows). On one kind of Linux runner, 80% of builds are within about 20 s of the median; Windows builds vary much more, so Windows differences under about 100 s say little.

#13672 and #13674 merged on Oct 30, and #13671 and #14572 on Nov 3, so CI builds of develop and of other PRs now include them.

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 29e5b0a059e.
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

Unit and rocisa suites: no test newly failing compared with develop; the new tests pass. On the same tests, Toolchain/Assembly.py misses the same statements as on develop, so every new line is covered.
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, two runs: gfx942 hipBLASLt build −135 and −164 s, gfx950 −107 and −86 s; gfx110X on Windows −37 and +22 s, which is within Windows noise.

Submission Checklist

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

Risk level

Low. It parallelizes independent subprocess pairs; errors still propagate through executor.map.

🤖 Generated with Claude Code

therock-pr-bot · 2026-10-14T20:39:38Z

✅ 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 · 2026-10-14T20:39:40Z

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

codecov · 2026-10-14T23:40:07Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

❌ Your project check has failed because the head coverage (51.04%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #13675      +/-   ##
===========================================
+ Coverage    77.48%   77.52%   +0.04%     
===========================================
  Files         2903     2904       +1     
  Lines       490384   490460      +76     
  Branches     74217    74245      +28     
===========================================
+ Hits        379935   380187     +252     
+ Misses       85489    85346     -143     
+ Partials     24960    24927      -33

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`51.04% <ø> (+0.08%)`	⬆️
TensileLite-Unit	`84.35% <100.00%> (+0.21%)`	⬆️
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 7356d47
hipBLASLt	`68.87% <ø> (+0.09%)`	⬆️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 7356d47
hipDNN	`87.26% <ø> (ø)`		Carriedforward from 7356d47
hipFFT	`42.47% <ø> (ø)`		Carriedforward from 7356d47
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 7356d47
hipSOLVER	`69.67% <ø> (ø)`		Carriedforward from 7356d47
hipSPARSE	`87.03% <ø> (ø)`		Carriedforward from 7356d47
rocBLAS	`69.20% <ø> (ø)`		Carriedforward from 7356d47
rocFFT	`47.28% <ø> (ø)`		Carriedforward from 7356d47
rocRAND	`57.32% <ø> (ø)`		Carriedforward from 7356d47
rocSOLVER	`76.74% <ø> (ø)`		Carriedforward from 7356d47
rocSPARSE	`83.54% <ø> (ø)`		Carriedforward from 7356d47
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 7356d47

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

Files with missing lines	Coverage Δ
...aslt/tensilelite/tensilelite/Toolchain/Assembly.py	`95.74% <100.00%> (+0.74%)`	⬆️
...slt/tensilelite/tensilelite/Toolchain/Component.py	`99.43% <100.00%> (+<0.01%)`	⬆️

... and 26 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.

JIRA ID : AIHPBLAS-5194 ## Motivation Profiling assembly codegen showed about a third of rocisa's pass time going into `buildGraph` for `removeDuplicateAssignment`. It records every instruction into per-register lists for every sgpr, vgpr and mgpr, and copies every barrier into each list, but the pass only reads the sgpr lists, and the graph is discarded right afterwards. ## Technical Details - **slim down the duplicate-assignment register graph:** `buildGraph` now records only per-sgpr entries, each tagged with its barrier sequence and a precomputed item classification. Barriers advance a counter instead of being copied into every list, and resetting barriers are counted, so the pass can tell whether one fell between two entries. `removeDuplicateAssignment` walks those entries directly instead of rebuilding lists. Bounds checks, their error messages and `maxVgprSeen` are unchanged. The work is in `rocisa/include/pass.hpp` and `rocisa/src/pass/graph.cpp`. - **remove the unused NoOptItem wrapper** (from James Newling's review): NoOpt instructions are now represented by the graph's entries, so the native `NoOptItem` class had no users left. - **tests:** the pass's first direct tests (`rocisa/test/test_remove_dup_assign.py`): repeated moves, resetting and non-resetting barriers, writes to the sgpr or to a covering range, a write inside a NoOpt module between two matching moves, comments, NoOpt and nested modules, bounds checks (including ranges whose last register is out of bounds, negative indices and mgpr), and `maxVgpr` with and without vgpr references. James Newling's review added several of them. They pass on develop's rocisa and on this branch. They use only already-lowered instructions, because the pass lowers composites such as `SWaitCnt` for whichever ISA is current in rocIsa's global state. Seen while writing the tests, and not changed here: on gfx12 `SWaitCnt` lowers to `s_wait_dscnt`, which the pass doesn't treat as a reset the way it treats `s_waitcnt`. ### 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 `db2d71e0815`. - `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 - rocisa suite: 316 passed, 4 skipped. TensileLite unit suite: no test newly failing compared with develop. - 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: gfx942 hipBLASLt build 1,802 s on an 80-thread runner against that class's median of 1,896 s (−94 s), gfx950 −108 s; assembly codegen 910 → 827 s. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low to medium. It touches a codegen pass, but the output is byte-identical and the pass now has direct tests. 🤖 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>

…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>

…3671) JIRA ID : AIHPBLAS-5193 ## Motivation Device-library builds compile `Kernels.cpp` (the source helper kernels) as one single-threaded clang job, about 120 s on CI for gfx942. That job doesn't depend on the assembly kernels, but it only started after every assembly kernel had been generated, assembled, linked and bundled, so it sat on the critical path. The overlap pays off when the helper-kernel cache (`~/.tensile/helper_cache`) misses, which is every CI build because CI starts with an empty cache. With a warm cache there is little left to overlap. This PR used to also stop codegen workers from returning kernel source; that change is now #14572. ## Technical Details - **compile helper kernels while generating asm kernels:** `writeHelpers` now runs first, and `buildSourceCodeObjectFiles` is submitted to a single background thread before assembly codegen starts. It is joined after the assembly code objects are built, so a failure in the background compile still fails the build. Its outputs (`Kernels.so-*.hsaco`) and temporary paths are disjoint from the assembly path. - **tests:** the helper compile starts before assembly codegen finishes, and a failure in it is raised. Rebased onto develop: the background compile wraps the `_mapKernelsByBackend` codegen call (#13068), and assembly code objects are built from deduplicated kernel groups (#11202). ### 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, −104 s | −145, −230 s | | #14572 | Don't return kernel source from codegen workers | −9, −28 s | −129, +41 s | | #13672 (merged) | rocisa: slimmer duplicate-assignment register graph | −93, −94 s | +58, +45 s | | #13673 | Lighter logic loading | −82 s | −68 s | | #13674 (merged) | Probe assembler capabilities for all ISAs concurrently | −10, +6 s (TensileLogic −21 s) | −193, −284 s | | #13675 | Link and bundle code objects in parallel | −135, −164 s | −37, +22 s | \*One value per TheRock Multi-Arch CI run of the PR's current change, each against the median of other PRs' hipBLASLt builds on the same kind of runner within a day (about 1,725 s for gfx942 on 96-thread runners, 1,895 s on 80-thread runners, and 910 s for gfx110X on Windows). On one kind of Linux runner, 80% of builds are within about 20 s of the median; Windows builds vary much more, so Windows differences under about 100 s say little. #13672 and #13674 merged on Oct 30, so CI builds of develop and of other PRs now include them. 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 - Unit and rocisa suites: no test newly failing compared with develop; the new tests pass, and the overlap test fails on develop as expected. - 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, two runs: gfx942 hipBLASLt build −130 and −104 s, gfx950 −79 and −85 s, gfx110X on Windows −145 and −230 s (against other PRs' builds those days). - Locally, gfx942 pinned to 96 CPUs with the helper-kernel cache off, one build each: 1,449 → 1,327 s (−122 s). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. The order of build steps changes, but the outputs don't. A failure in the background compile surfaces at the join. 🤖 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>

Each code object is an independent ld.lld + clang-offload-bundler pair, but they ran one after another. Run them on a thread pool; the linker's response files are already named after their output, so concurrent links don't collide. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

ld.lld uses every hardware thread by default, so with one link per CPU running at once, each link could still start a thread per CPU. Give each link an equal share of the CPUs through a new Linker threads argument, and start the largest links first so one of them doesn't finish last on its own. Co-authored-by: Cursor <cursoragent@cursor.com>

The coverage lane runs with concurrency = multiprocessing, so it does not trace the thread-pool workers that link and bundle, and Toolchain/Assembly.py fell below its coverage floor. Moving the per-job work into _linkAndBundle lets tests call it directly. Co-authored-by: Cursor <cursoragent@cursor.com>

…14572) JIRA ID : AIHPBLAS-5193 ## Motivation TensileCreateLibrary's assembly codegen workers sent every kernel's generated source back to the parent: for gfx942, about 214k sources of about 700 KB of assembly each, compressed in memory above 10k kernels to keep that manageable. The parent never uses them: each worker already writes and assembles its own `.s` file. Even compressed, they stayed in the parent for the rest of the build, about 9 GB for gfx942. Split out of #13671 at review request, so each change can be measured on its own. ## Technical Details - **don't return kernel source from codegen workers:** the result each worker returns drops `src` and `header`, and TensileCreateLibrary no longer compresses them to ship them. This covers both codegen entry points, `processKernelSource` and the native-rocisa `processKernelSourceNativeInit`. - **mark the kernel source in codegen results as optional:** `KernelCodeGenResult.src` can now be None, as `header` already could. The compression path stays for `writeSolutionsAndKernels`, which still uses it. - **tests:** drive `writeSolutionsAndKernelsTCL` with real composed codegen on a fake kernel: the worker still writes the full source, and the result returned to the parent carries the kernel info without the source or header. ### 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, −104 s | −145, −230 s | | #14572 | Don't return kernel source from codegen workers | −9, −28 s | −129, +41 s | | #13672 (merged) | rocisa: slimmer duplicate-assignment register graph | −93, −94 s | +58, +45 s | | #13673 | Lighter logic loading | −82 s | −68 s | | #13674 (merged) | Probe assembler capabilities for all ISAs concurrently | −10, +6 s (TensileLogic −21 s) | −193, −284 s | | #13675 | Link and bundle code objects in parallel | −135, −164 s | −37, +22 s | \*One value per TheRock Multi-Arch CI run of the PR's current change, each against the median of other PRs' hipBLASLt builds on the same kind of runner within a day (about 1,725 s for gfx942 on 96-thread runners, 1,895 s on 80-thread runners, and 910 s for gfx110X on Windows). On one kind of Linux runner, 80% of builds are within about 20 s of the median; Windows builds vary much more, so Windows differences under about 100 s say little. #13672 and #13674 merged on Oct 30, so CI builds of develop and of other PRs now include them. 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 - Unit and rocisa suites: no test newly failing compared with develop; the new tests pass, and the no-source test fails on develop as expected. - 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. - Memory, full gfx942 TensileCreateLibrary run: the parent process peaks at 42.4 GB instead of 51.3 GB (−8.9 GB, −17%). - CI, two runs: gfx942 −9 and −28 s, gfx950 −21 s both times; gfx110X on Windows −129 and +41 s, which is noise. Small, but not run-to-run variation: in all four Linux runs the parallel codegen phase was faster than 76–100% of other PRs' builds on the same kind of runner. - Locally, gfx942 pinned to 96 CPUs with the helper-kernel cache off, one build each: 1,449 → 1,400 s (−49 s). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. The parent only reads the kernel info fields, which are unchanged. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: T.J. Alumbaugh <talumbau@amd.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com>

newling

Not to be addressed here, but maybe for next time: I find the presentation of the observed speed-up a bit confusing, and it's not easy glancing at the description if it's real or noise. I personally like to see all the relevant values (rounded to appropriate decimal places) if fewer than 10 -- I don't think it'd take up more real estate on the PR description than the current presentation in terms of medians and percentiles. Something like:

Baseline runs from previous day: [a, b, c, d, e]. 2 runs with this PR: [f, g].

Agent review has 1 test request: https://github.com/newling/agent-public-reviews/blob/main/rocm-libraries/review_pr13675.md

but overall LGTM

jinchen62 · 2026-11-04T03:46:47Z

@newling Thanks. I pulled in your tests as is (22dab22). All 11 pass. With your probe script, the serial pool fails the overlap test and the pool that never collects results fails both error tests, while the original eight pass under both. So they cover what the original tests missed.

On the numbers, fair point; I'll lead with the raw values from now on. Here they are for this PR on gfx942, on 96-thread runners. Its three builds before the Nov 3 rebase took 1,591, 1,590 and 1,560 s. Other PRs' builds on the same kind of runner within a day of each are below:

This PR's build	Other PRs' builds	Fastest	Median	Slowest
1,591 s	524	1,655 s	1,725 s	2,559 s
1,590 s	529	1,655 s	1,726 s	2,532 s
1,560 s	509	1,614 s	1,724 s	2,426 s

Each build of this PR was faster than all of them. Every build is in the tracker's CSV.

#13673) JIRA ID : AIHPBLAS-5195 ## Motivation The "Loading Logics" phase is bound by the parent process unpickling worker results on a single thread. Each worker returned a library that held a fresh `IndexSolutionLibrary` wrapper per table row (rows outnumber solutions by orders of magnitude), plus the solutions list and raw ExactLogic, which TensileCreateLibrary throws away. ## Technical Details Two commits, plus tests: - **share solution wrappers in matching tables:** `MatchingLibrary.FromOriginalState` creates one wrapper per solution index and reuses it. The wrapper only refers to its solution, and index remapping updates the solution, so sharing is safe. - **return only what TensileCreateLibrary uses:** a module-level `_parseLogicFileForLibrary` wrapper drops `solutions` and `exactLogic` in the worker. That is about 40% of the returned bytes, and parent unpickle time halves. A characterization test that stubbed `parseLibraryLogicFile` now stubs the wrapper instead. - **tests:** rows for one solution share its wrapper; the shared wrapper still follows the solution through reindexing after a pickle round trip, the way a loading worker's result comes back (James Newling's test); and the parse wrapper drops only those two fields. ### 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 (merged) | Compile helper kernels while generating asm kernels | −130, −104 s | −145, −230 s | | #14572 (merged) | Don't return kernel source from codegen workers | −9, −28 s | −129, +41 s | | #13672 (merged) | rocisa: slimmer duplicate-assignment register graph | −93, −94 s | +58, +45 s | | #13673 | Lighter logic loading | −82 s | −68 s | | #13674 (merged) | Probe assembler capabilities for all ISAs concurrently | −10, +6 s (TensileLogic −21 s) | −193, −284 s | | #13675 | Link and bundle code objects in parallel | −135, −164 s | −37, +22 s | \*One value per TheRock Multi-Arch CI run of the PR's current change, each against the median of other PRs' hipBLASLt builds on the same kind of runner within a day (about 1,725 s for gfx942 on 96-thread runners, 1,895 s on 80-thread runners, and 910 s for gfx110X on Windows). On one kind of Linux runner, 80% of builds are within about 20 s of the median; Windows builds vary much more, so Windows differences under about 100 s say little. #13672 and #13674 merged on Oct 30, and #13671 and #14572 on Nov 3, so CI builds of develop and of other PRs now include them. 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 `29e5b0a059e`. - `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 - Unit and rocisa suites: no test newly failing compared with develop; 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 against develop: gfx942 hipBLASLt build 1,727 → 1,643 s (−83 s); logic loading 155 → 102 s, writing master solution libraries 125 → 79 s. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. The same libraries are produced; the only difference is fewer duplicate objects and dropping fields that were already discarded. 🤖 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>

peymanr · 2026-11-06T13:48:32Z

PR Pundit review

(Testing PR Pundit Code Review agent. Not all comments may be relevant. This is a one-time automated pass: please address what is useful, no need to request a re-review. Please reply to each point, whether it applies or not: we harvest these replies, and both kinds are enormously helpful for tuning the agent. Reach out to @peymanr for suggestions.)

#		Finding	Where
1	🟠	Link/bundle failures reported late: this PR replaced the serial loop with in-order executor.map	`projects/hipblaslt/tensilelite/tensilelite/Toolchain/Assembly.py:184`
2	🟠	CpuThreads=0/1 forces single-threaded links: this PR derives ld.lld --threads from CPUThreadCount()	`projects/hipblaslt/tensilelite/tensilelite/Toolchain/Assembly.py:178`
3	🟠	Parallel-link tests miss CpuThreads=0, threads=0 and empty/single-job edge cases	`projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_assembly_parallel_link.py:1`

1. 🟠 Should fix: Link/bundle failures reported late: this PR replaced the serial loop with in-order executor.map

projects/hipblaslt/tensilelite/tensilelite/Toolchain/Assembly.py:184

buildAssemblyCodeObjectFiles now runs links and bundles through executor.map, which returns results in submission order. A failing link or bundle is therefore reported only after every earlier-submitted job has finished. Meanwhile workers keep starting later jobs until the error is raised and the remaining queued jobs are cancelled. The build still fails with the original RuntimeError, whose Failed command includes the -o <code object> path. Using submit + as_completed with shutdown(cancel_futures=True) would fail sooner, but only the first error is ever reported, as before.

Who is affected: Developers running TensileCreateLibrary when an ld.lld or bundler invocation fails; they see the error later, after extra jobs have run.

Checked against the code by an independent pass.

2. 🟠 Should fix: CpuThreads=0/1 forces single-threaded links: this PR derives ld.lld --threads from CPUThreadCount()

projects/hipblaslt/tensilelite/tensilelite/Toolchain/Assembly.py:178

CpuThreads (-j/--jobs) now also caps ld.lld's internal threads, and this is not documented. With CpuThreads=0 (the 'no threading' mode that PythonProfile forces) CPUThreadCount() returns 0; with -j 1 it returns 1. Either way linkThreads=max(1, cpus//workers)=1, so every link runs with -Xlinker --threads=1, where before ld.lld used all hardware threads. The default (-1) is unaffected. Consider documenting that -j bounds linker threads, or deriving linkThreads from the real CPU count when CpuThreads is 0 or 1.

Who is affected: Users who run TensileCreateLibrary or benchmarks with -j 0/-j 1, CpuThreads: 0/1 in a yaml (e.g. the gfx1250 perf yamls), or PythonProfile enabled.

Checked against the code by an independent pass.

3. 🟠 Should fix: Parallel-link tests miss CpuThreads=0, threads=0 and empty/single-job edge cases

projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_assembly_parallel_link.py:1

The new tests patch CPUThreadCount only to 1, 2 or 8. They never exercise CpuThreads=0 (where CPUThreadCount() returns 0 and the code clamps to one worker with --threads=1), an empty or single-job list, or Linker(threads=0). The latter silently drops the flag because of the if threads check. The characterization _StubLinker accepts threads but does not record it; thread passing is still asserted in test_assembly_parallel_link.py. Its two unlocked appends cannot break any current assertion: the only test that reads input_types has exactly one link job. Optional: pin CPUThreadCount in the characterization tests and add a CpuThreads=0 case.

Who is affected: Only test-suite maintainers; there is no runtime impact on library builds.

Checked against the code by an independent pass.

_{Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at 22dab221; findings checked and posted by @peymanr; points already raised on this PR are left out.}

talumbau added this pull request to stack #13676 October 14, 2026 20:21

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Oct 14, 2026

assistant-librarian Bot added the organization: ROCm label Oct 14, 2026

therock-pr-bot Bot added the Not ready to Review label Oct 14, 2026

talumbau force-pushed the users/talumbau/tcl-speed-5-parallel-link branch from e5d70fb to 6abd26f Compare October 16, 2026 22:57

talumbau force-pushed the users/talumbau/tcl-speed-4-probe-validate branch from 1cf4f30 to 00aacef Compare October 16, 2026 22:57

jinchen62 removed this pull request from stack #13676 October 28, 2026 13:02

jinchen62 changed the base branch from users/talumbau/tcl-speed-4-probe-validate to develop October 28, 2026 13:02

jinchen62 force-pushed the users/talumbau/tcl-speed-5-parallel-link branch from 6abd26f to d195f2b Compare October 28, 2026 13:03

jinchen62 marked this pull request as ready for review October 28, 2026 13:05

jinchen62 requested a review from a team as a code owner October 28, 2026 13:05

jinchen62 removed the Not ready to Review label Oct 28, 2026

jinchen62 requested review from Alex-Vasile and babakpst October 28, 2026 13:06

Alex-Vasile reviewed Oct 28, 2026

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/tensilelite/Toolchain/Assembly.py Resolved

jinchen62 mentioned this pull request Oct 28, 2026

perf(tensilelite): don't return kernel source from codegen workers #14572

Merged

1 task

jinchen62 force-pushed the users/talumbau/tcl-speed-5-parallel-link branch from 88510a4 to 63548d7 Compare October 29, 2026 11:01

jinchen62 requested review from Alex-Vasile and newling November 2, 2026 03:23

Alex-Vasile approved these changes Nov 2, 2026

View reviewed changes

talumbau and others added 3 commits November 3, 2026 02:51

jinchen62 force-pushed the users/talumbau/tcl-speed-5-parallel-link branch from 63548d7 to 7356d47 Compare November 3, 2026 09:13

test(tensilelite): verify parallel link overlap and worker failures

22dab22

newling approved these changes Nov 3, 2026

View reviewed changes

talumbau merged commit d38d6b6 into develop Nov 6, 2026
496 of 521 checks passed

talumbau deleted the users/talumbau/tcl-speed-5-parallel-link branch November 6, 2026 22:41

Navigation Menu

Repository navigation

perf(tensilelite): link and bundle code objects in parallel - #13675

perf(tensilelite): link and bundle code objects in parallel#13675
talumbau merged 4 commits into
developfrom
users/talumbau/tcl-speed-5-parallel-link

talumbau commented Oct 14, 2026 •

edited by jinchen62

Loading

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

codecov Bot commented Oct 14, 2026 •

edited

Loading

Uh oh!

Uh oh!

newling left a comment •

edited

Loading

Uh oh!

jinchen62 commented Nov 4, 2026

Uh oh!

peymanr commented Nov 6, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

talumbau commented Oct 14, 2026 • edited by jinchen62 Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Series overview

Test Plan

Test Result

Submission Checklist

Risk level

Uh oh!

therock-pr-bot Bot commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Uh oh!

newling left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jinchen62 commented Nov 4, 2026

Uh oh!

peymanr commented Nov 6, 2026

PR Pundit review

1. 🟠 Should fix: Link/bundle failures reported late: this PR replaced the serial loop with in-order executor.map

2. 🟠 Should fix: CpuThreads=0/1 forces single-threaded links: this PR derives ld.lld --threads from CPUThreadCount()

3. 🟠 Should fix: Parallel-link tests miss CpuThreads=0, threads=0 and empty/single-job edge cases

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

talumbau commented Oct 14, 2026 •

edited by jinchen62

Loading

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

edited

Loading

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

edited

Loading

codecov Bot commented Oct 14, 2026 •

edited

Loading

newling left a comment •

edited

Loading