Repository navigation
fix(tensilelite): deduplicate kernels by assembly identity - #11202
Merged
Merged
Conversation
ThanHenderson
changed the title
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #11202 +/- ##
===========================================
- Coverage 76.02% 73.22% -2.81%
===========================================
Files 2867 2785 -82
Lines 474741 458812 -15929
Branches 71077 68484 -2593
===========================================
- Hits 360916 335926 -24990
- Misses 89314 99467 +10153
+ Partials 24511 23419 -1092
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jaopaulolc
approved these changes
Aug 24, 2026
jaopaulolc
left a comment
jaopaulolc
left a comment
Contributor
There was a problem hiding this comment.
Nit follow-ups only; the AIHPBLAS-4388 DebugStreamK [0, 1] identity race looks closed.
The Full-vs-Min collision, empty-asm short-circuit, and shared .s/.o assemble path are addressed. Remaining notes are residual WGMXCC category coarseness, a leftover collection helper that is not on the run() path, a try/finally gap in _getName, and a missing groupKernelRecords([dsk0, dsk1]) unit assertion.
jaopaulolc
added
the
rocm:gemm algos
caio96
reviewed
Aug 25, 2026
caio96
approved these changes
Aug 25, 2026
Group compilation by emitted kernel symbol while preserving solution records and code-object placement aliases. Update naming, cache invalidation, and regression fixtures for current develop.
AlexBrownAMD
approved these changes
Oct 16, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
This was referenced Oct 28, 2026
jinchen62
added a commit
that referenced
this pull request
Nov 3, 2026
…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>
talumbau
added a commit
that referenced
this pull request
Nov 6, 2026
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](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>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AIHPBLAS-4388
Motivation
TensileLite deduplicated assembly kernels in two stages that used different identities. Early kernel collection could retain two valid StreamK solutions whose full compile identity differed, such as
DebugStreamK: [0, 1], butwriteSolutionsAndKernels()deduplicated again with a minimal basename that omittedDebugStreamK. The second kernel was then marked as a duplicate, returned an empty assembly string from the duplicate codegen path, and could race with the real kernel while both jobs wrote and assembled the same.sand.opaths. This change uses a single generated-assembly identity for codegen, keeps code-changing StreamK modes distinct, and handles code-object placement as a separate aliasing step.Technical Details
getKernelCompileKey()to describe generated-assembly identity separately fromgetKeyNoInternalArgs(), which still includes placement metadata such ascodeObjectFileandDeviceNames.kernel.duplicateflag withKernelCompileGroup, where one representative kernel is compiled and its aliases receive the same post-codegen metadata.codeObjectFileplacements.DebugStreamK,StreamKAtomic,MbskPrefetchMethod,DebugPersistentKernelLoopForever, and StreamK XCC mapping categories.Test Plan
Tensile/Tests/common/streamk/gfx1250/sk_hgemm_debug_modes.yamlTensile/Tests/common/streamk/gfx950/sk_sgemm_debug_modes.yamlTest Result
Risk level
Medium. The change is scoped to TensileLite generation and tests, but it changes kernel identity, object reuse, and emitted basenames in high-coupling generator code. Device behavior should remain unchanged for kernels whose generated assembly is identical; CI should cover the affected TensileLite unit, characterization, and architecture-specific StreamK paths.