Repository navigation
perf(tensilelite): overlap helper-kernel compile with asm codegen - #13671
Conversation
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (50.82%) 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 #13671 +/- ##
========================================
Coverage 77.30% 77.30%
========================================
Files 2893 2893
Lines 485885 485888 +3
Branches 73442 73442
========================================
+ Hits 375566 375569 +3
Misses 85553 85553
Partials 24766 24766
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
5fc2d15 to
d378fbf
Compare
d378fbf to
64c8f46
Compare
Alex-Vasile
left a comment
There was a problem hiding this comment.
- Please split this up into 2 PRs, each with its own optimization.
- Please provide time savings for each one individually, both on local machine and on CI.
It's hard to evaluate how much each is worth otherwise.
64c8f46 to
312e147
Compare
|
@Alex-Vasile Split as you asked. This PR now only overlaps the helper-kernel compile with asm codegen, and "don't return kernel source from codegen workers" is #14572. Each PR's CI run measures its own saving against develop, and I'll post those together with local numbers for each. For reference, the two together cut the gfx942 hipBLASLt CI build by 142 s (1,727 → 1,584 s). The helper compile alone takes about 122 s on CI, so most of that should be the overlap, but the split will measure it. |
Kernels.cpp holds the source helper kernels and builds as a single, single-threaded clang job (~175 s for gfx942). It does not depend on the assembly kernels, yet it only started after every assembly kernel had been generated, assembled, linked and bundled. Submit it to a background thread before assembly codegen starts and join it afterwards, taking it off the critical path. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The helper kernels must start compiling before assembly codegen finishes, and a failure in that background compile must still fail TensileCreateLibrary. Co-authored-by: Cursor <cursoragent@cursor.com>
312e147 to
d6301b0
Compare
|
@Alex-Vasile Numbers for each half, one run each:
*Device library pinned to 96 CPUs with the helper-kernel cache off: 1,449 s on develop. The CI numbers compare each PR's own run with the median of other PRs' hipBLASLt builds on the same kind of runner that day. On Linux, 80% of builds are within about 20 s of that median, so #14572's Linux gains are within noise; Windows varies a lot more. Most of the saving is the overlap. |
Are these table relative to develop, or relative to each other? |
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>
Relative to develop. Each row is that PR alone on top of the same develop commit, not one on top of the other. Locally, develop built in 1,449 s, this PR in 1,327 s and #14572 in 1,400 s. |
…embly kernels are generated What - tensilelite/TensileCreateLibrary/Run.py, writeSolutionsAndKernels (the benchmark path): Kernels.cpp/Kernels.h are written right after setup, before codegen, and buildSourceCodeObjectFiles is submitted to a one-thread ThreadPoolExecutor, so the helper compile runs while the assembly kernels are generated, checked, written, assembled and linked; it is joined after the assembly link. The executor's with block covers all of that, so any exit (including a codegen error or printExit) waits for the compile; a compile failure is re-raised by the join. Same approach as PR #13671 (27a8390) on the TensileCreateLibrary path. GenerateSourcesAndExit still writes the helpers and compiles nothing. - The compile is one more process, so it takes one slot of the pool budget (CPUThreadCount()): new benchmarkPoolSize(generateSourcesAndExit) = budget - 1 when the compile overlaps (overlapsHelperCompile(): budget > 1), else the budget. With a budget of 1 the compile finishes before codegen starts (no overlap). The codegen, invalid-kernel check (new procs parameter of removeInvalidSolutionsAndKernels) and write_assemble pools use it. - Supporting change (rule 7), tensilelite/BenchmarkProblems.py: its three solution-generation pools (fork solutions, solution pool load and construction) use the same benchmarkPoolSize (_poolSize()). joblib/loky resizes its reusable worker pool whenever n_jobs changes (stops a worker, then starts one that imports Tensile again on its first task); with solution generation at 24 and codegen at 23 that happened twice per benchmark step and made subtile_mxfp4's build 85 -> 150 s and dtl's 142 -> 151 s (measured). One size per run removes the resizes. TensileCreateLibrary/__init__.py exports benchmarkPoolSize. - Timers: the background compile keeps python_kernel_build_src_co (and its sub-timers); the main thread's wait for it is the new python_kernel_build_src_co_wait. scripts/ analyze_timing.py: the wait is in the tree under python_kernel_compilation; python_kernel_build_src_co and its sub-timers move to the overlapping totals (they run concurrently and would be double counted). In logs from before this change that time shows as untracked in python_kernel_compilation. - Tests: new tensilelite/Tests/unit/test_benchmark_helper_compile_overlap.py (12 cases: the compile is running while codegen runs, after Kernels.cpp is written; Run's and BenchmarkProblems' pools get budget - 1; benchmarkPoolSize for budgets 8, 2, 1, 0 and GenerateSourcesAndExit; budget 1 compiles before codegen; GenerateSourcesAndExit does not compile; a compile failure is raised; a codegen failure returns only after the compile finished). test_analyze_timing.py: TestHelperCompileOverlap (2). test_run_helpers_char.py: its ParallelMap2 stub accepts procs. Why - timing_runs/DECISIONS.md topic 6 (python_kernel_build_src_co), item F, plan step P4.9 (moved before the codegen restructure by the plan review); investigation timing_runs/python_kernel_build_src_co_investigation.md, "Relation to PR #13671". - The helper compile (python_kernel_build_src_co, 1,004.7 s = 39% of the gfx950 per-config build aggregate, measured with the 4c6426c tree) ran only after codegen, assembly and the assembly link, although writeHelpers only needs kernelHelperObjs, available on entry. Upper bound hidden by the overlap: 237.7 s (24%) of that aggregate (investigation, estimate; codegen and the compile share cores). Expected effect - Measured at this commit vs its parent, --build-only on 24 pinned cores (160-183, budget 24), cold helper cache per run, order old/new/new/old: fp16_tn (3 steps) 50.5, 50.6 s -> 44.4, 44.7 s (-12%) subtile_mxfp4 (130 steps) 83.2, 83.0 s -> 77.7, 78.5 s (-6%) dtl (26 steps) 142.8, 142.0 s -> 139.5, 139.6 s (-2%) Helper compile time left on the critical path (python_kernel_build_src_co, serial, vs python_kernel_build_src_co_wait): fp16_tn 12.3 -> 5.6 s, subtile 30.5 -> 26.4 s (most of its 130 steps have little codegen to hide behind), dtl 11.8 -> 7.6 s. Codegen timers unchanged within noise (subtile 20.0 -> 19.3 s, dtl 45.2 -> 45.4 s). - The full gfx950 aggregate is measured in P6 (estimate: up to the 24% bound above, less with the later helper-compile items, which shorten the compile itself). How verified - Artifacts of fp16_tn, subtile_mxfp4 and dtl, two builds per side: same member lists, non-.hsaco members byte-identical (11 / 415 / 103), helper .hsaco (6 / 260 / 52) equal except clang's random __hip_cuid_* symbol: .text byte for byte, symbols section-relative, kernel descriptors by the symbol they point at (a cuid length change can shift .rodata; two builds of the parent commit differ that way too). - Processes, sampled every 0.1 s (loky workers + 1 while a helper-compile process exists): at most 24 (= the budget) in every run of the three configs, also counting idle loky workers; the compile ran next to busy codegen/assembly workers (fp16_tn 82, subtile 42, dtl 50 samples; parent 0-2). - Real failing helper compile (Tensile_CXX_COMPILER_LAUNCHER that kills the Kernels.cpp compile after 2 s or 40 s and exits 1, fp16_tn build in its own session): the test fails with the compiler RuntimeError raised from the join in the first step (after 26.5 s / 46.7 s), and no process of the session is left 3 s later. - Solution generation with one worker less (supporting change): the parent commit with CpuThreads=23 builds subtile_mxfp4 in 83.1 s vs 84.3 / 85.1 s with 24 (measured, no loss). - TensileLite Python unit tests (tensilelite/Tests/unit, -n 32): 11299 passed, 67 skipped, 2 xfailed, 1 xpassed, 0 failed. Risks / behaviour changes - Benchmark-run pools have one process less than the budget (none when the budget is 1, or with GenerateSourcesAndExit); loky workers + compile stay within the P4.3 budget. - A failing helper compile is reported after the step's assembly work instead of right away (the step still fails with the same error). A codegen failure waits for a running compile. - Kernels.h is now always written before codegen. DisableAsmComments: True reaches the parent's rocisa output options only through in-process codegen, which then keeps it for the rest of the process. So in runs that generate in-process (pool size 1, i.e. a budget of 1 or 2) of such configs, only the first benchmark step's Kernels.h now keeps the inline-asm comments (as with parallel codegen); later steps are unchanged (no comments). C++ comments only, objects unchanged; the helper-cache key changes for that first step only. Since P4.8 no common test YAML generates serially. - analyze_timing.py: python_kernel_build_src_co is no longer in the tree (overlapping total).
There was a problem hiding this comment.
This makes sense. Hide the single threaded compilation of Kernels.cpp by overlapping it with other work. I guess it works perfectly if the other work takes longer than the compilation.
According to an agent I'm discussing this with, the Kernels.cpp (gfx942) is over 2 million lines long. One thing we could try (in the future) is finding a sensible way to split that up, so that multiple translation units can be compiled in parallel. But that can be tried later.
…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>
…embly kernels are generated What - tensilelite/TensileCreateLibrary/Run.py, writeSolutionsAndKernels (the benchmark path): Kernels.cpp/Kernels.h are written right after setup, before codegen, and buildSourceCodeObjectFiles is submitted to a one-thread ThreadPoolExecutor, so the helper compile runs while the assembly kernels are generated, checked, written, assembled and linked; it is joined after the assembly link. The executor's with block covers all of that, so any exit (including a codegen error or printExit) waits for the compile; a compile failure is re-raised by the join. Same approach as PR #13671 (27a8390) on the TensileCreateLibrary path. GenerateSourcesAndExit still writes the helpers and compiles nothing. - The compile is one more process, so it takes one slot of the pool budget (CPUThreadCount()): new benchmarkPoolSize(generateSourcesAndExit) = budget - 1 when the compile overlaps (overlapsHelperCompile(): budget > 1), else the budget. With a budget of 1 the compile finishes before codegen starts (no overlap). The codegen, invalid-kernel check (new procs parameter of removeInvalidSolutionsAndKernels) and write_assemble pools use it. - Supporting change (rule 7), tensilelite/BenchmarkProblems.py: its three solution-generation pools (fork solutions, solution pool load and construction) use the same benchmarkPoolSize (_poolSize()). joblib/loky resizes its reusable worker pool whenever n_jobs changes (stops a worker, then starts one that imports Tensile again on its first task); with solution generation at 24 and codegen at 23 that happened twice per benchmark step and made subtile_mxfp4's build 85 -> 150 s and dtl's 142 -> 151 s (measured). One size per run removes the resizes. TensileCreateLibrary/__init__.py exports benchmarkPoolSize. - Timers: the background compile keeps python_kernel_build_src_co (and its sub-timers); the main thread's wait for it is the new python_kernel_build_src_co_wait. scripts/ analyze_timing.py: the wait is in the tree under python_kernel_compilation; python_kernel_build_src_co and its sub-timers move to the overlapping totals (they run concurrently and would be double counted). In logs from before this change that time shows as untracked in python_kernel_compilation. - Tests: new tensilelite/Tests/unit/test_benchmark_helper_compile_overlap.py (12 cases: the compile is running while codegen runs, after Kernels.cpp is written; Run's and BenchmarkProblems' pools get budget - 1; benchmarkPoolSize for budgets 8, 2, 1, 0 and GenerateSourcesAndExit; budget 1 compiles before codegen; GenerateSourcesAndExit does not compile; a compile failure is raised; a codegen failure returns only after the compile finished). test_analyze_timing.py: TestHelperCompileOverlap (2). test_run_helpers_char.py: its ParallelMap2 stub accepts procs. Why - timing_runs/DECISIONS.md topic 6 (python_kernel_build_src_co), item F, plan step P4.9 (moved before the codegen restructure by the plan review); investigation timing_runs/python_kernel_build_src_co_investigation.md, "Relation to PR #13671". - The helper compile (python_kernel_build_src_co, 1,004.7 s = 39% of the gfx950 per-config build aggregate, measured with the 4c6426c tree) ran only after codegen, assembly and the assembly link, although writeHelpers only needs kernelHelperObjs, available on entry. Upper bound hidden by the overlap: 237.7 s (24%) of that aggregate (investigation, estimate; codegen and the compile share cores). Expected effect - Measured at this commit vs its parent, --build-only on 24 pinned cores (160-183, budget 24), cold helper cache per run, order old/new/new/old: fp16_tn (3 steps) 50.5, 50.6 s -> 44.4, 44.7 s (-12%) subtile_mxfp4 (130 steps) 83.2, 83.0 s -> 77.7, 78.5 s (-6%) dtl (26 steps) 142.8, 142.0 s -> 139.5, 139.6 s (-2%) Helper compile time left on the critical path (python_kernel_build_src_co, serial, vs python_kernel_build_src_co_wait): fp16_tn 12.3 -> 5.6 s, subtile 30.5 -> 26.4 s (most of its 130 steps have little codegen to hide behind), dtl 11.8 -> 7.6 s. Codegen timers unchanged within noise (subtile 20.0 -> 19.3 s, dtl 45.2 -> 45.4 s). - The full gfx950 aggregate is measured in P6 (estimate: up to the 24% bound above, less with the later helper-compile items, which shorten the compile itself). How verified - Artifacts of fp16_tn, subtile_mxfp4 and dtl, two builds per side: same member lists, non-.hsaco members byte-identical (11 / 415 / 103), helper .hsaco (6 / 260 / 52) equal except clang's random __hip_cuid_* symbol: .text byte for byte, symbols section-relative, kernel descriptors by the symbol they point at (a cuid length change can shift .rodata; two builds of the parent commit differ that way too). - Processes, sampled every 0.1 s (loky workers + 1 while a helper-compile process exists): at most 24 (= the budget) in every run of the three configs, also counting idle loky workers; the compile ran next to busy codegen/assembly workers (fp16_tn 82, subtile 42, dtl 50 samples; parent 0-2). - Real failing helper compile (Tensile_CXX_COMPILER_LAUNCHER that kills the Kernels.cpp compile after 2 s or 40 s and exits 1, fp16_tn build in its own session): the test fails with the compiler RuntimeError raised from the join in the first step (after 26.5 s / 46.7 s), and no process of the session is left 3 s later. - Solution generation with one worker less (supporting change): the parent commit with CpuThreads=23 builds subtile_mxfp4 in 83.1 s vs 84.3 / 85.1 s with 24 (measured, no loss). - TensileLite Python unit tests (tensilelite/Tests/unit, -n 32): 11299 passed, 67 skipped, 2 xfailed, 1 xpassed, 0 failed. Risks / behaviour changes - Benchmark-run pools have one process less than the budget (none when the budget is 1, or with GenerateSourcesAndExit); loky workers + compile stay within the P4.3 budget. - A failing helper compile is reported after the step's assembly work instead of right away (the step still fails with the same error). A codegen failure waits for a running compile. - Kernels.h is now always written before codegen. DisableAsmComments: True reaches the parent's rocisa output options only through in-process codegen, which then keeps it for the rest of the process. So in runs that generate in-process (pool size 1, i.e. a budget of 1 or 2) of such configs, only the first benchmark step's Kernels.h now keeps the inline-asm comments (as with parallel codegen); later steps are unchanged (no comments). C++ comments only, objects unchanged; the helper-cache key changes for that first step only. Since P4.8 no common test YAML generates serially. - analyze_timing.py: python_kernel_build_src_co is no longer in the tree (overlapping total).
TheRock Submodule Bump ActivityNewest first
|
#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>
#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>
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>
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
writeHelpersnow runs first, andbuildSourceCodeObjectFilesis 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.Rebased onto develop: the background compile wraps the
_mapKernelsByBackendcodegen 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.
*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/Tests/unit) and rocisa suite (rocisa/test), compared test by test with the same run on develop9d729d3280b.gemm-libsbuilds 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.Test Result
Tensile/library,libraryanddevice-libraryare byte-identical. The other three differ only because the two builds come from different commits and directories:hipblaslt-version.hembeds the commit,MatchTable.yamlembeds the logic-file paths, andKernels.so-000-<arch>.hsacogets a HIP CUID derived from its build path. Its code sections are identical; only the CUID symbol and the symbol tables it reorders differ.Submission Checklist
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