Repository navigation
test(tensilelite): run the unit tests' ParallelMap calls in-process - #14599
Merged
Alex-Vasile merged 1 commit intoNov 3, 2026
Merged
Conversation
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. |
This was referenced Oct 28, 2026
Merged
Alex-Vasile
marked this pull request as ready for review
October 28, 2026 20:37
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14599 +/- ##
===========================================
+ Coverage 77.43% 77.46% +0.03%
===========================================
Files 2902 2902
Lines 489845 489845
Branches 74142 74142
===========================================
+ Hits 379304 379453 +149
+ Misses 85573 85474 -99
+ Partials 24968 24918 -50
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
tony-davis
approved these changes
Oct 30, 2026
tony-davis
left a comment
tony-davis
left a comment
Contributor
There was a problem hiding this comment.
Hi Alex! This is a nice, well-explained first step, and the stack table in the description made the whole series easy to follow. I checked what in-process means for the production path: with CpuThreads at 1, ParallelMap uses a one-thread dummy pool and ParallelMap2 runs joblib with one job, so nothing is pickled or sent to a worker in the unit lane anymore. I think that's fine, because the real pool path still gets exercised by every device-library build (TensileCreateLibrary with the default -1), by the subprocess-based Tests/common/test_config.py runs, and by _native_rocisa_workers, which still forces 2 workers. My one worry is scope, in the inline comment below.
Non-blocking, just a thought: the unit lane used to catch things like an unpicklable payload or a worker missing parent state by accident, and now it won't. If you want to keep a cheap tripwire, one test that runs a real production map (say one kernel through ParallelMap2) with CpuThreads=2 would probably do it. characterization/Parallel/test_parallel_char.py already says the real pool paths aren't unit-driven, so a sentence there or in tensilelite/TESTING.md would also cover it.
Alex-Vasile
force-pushed
the
users/alvasile/unit-conftest-inprocess-parallelmap
branch
from
October 30, 2026 21:53
6bb9930 to
d1777c2
Compare
1 task done
CPUThreadCount() sizes its worker pool from globalParameters["CpuThreads"], whose default of -1 means every CPU the process can see, not the CI job's allocation. About 460 unit tests reach ParallelMap/ParallelMap2 with one or a few items, so each pytest and xdist process launched a machine-wide pool that did not speed those tests up and multiplied CPU use by the number of workers. Add a session-scoped autouse fixture to Tests/unit/conftest.py that sets CpuThreads to 1 in both globalParameters and defaultGlobalParameters, so restoreDefaultGlobalParameters() does not undo it. The fixture runs once in every process that runs unit tests, each xdist worker included. It is a fixture rather than a pytest_configure hook because pytest runs a nested conftest's hook as soon as it imports the conftest during collection, which would also change CpuThreads for the non-unit tests when a run collects all of Tests. Tests that need another value still set it themselves. Add one small test in characterization/Parallel that runs ParallelMap and ParallelMap2 with CpuThreads=2 and checks that the work reaches other processes with the parent's global parameters, since no other unit test uses a real pool any more. Document that in TESTING.md.
Alex-Vasile
force-pushed
the
users/alvasile/unit-conftest-inprocess-parallelmap
branch
from
November 2, 2026 20:32
d1777c2 to
9f17278
Compare
Alex-Vasile
enabled stacked auto-merge (squash)
November 2, 2026 20:32
stacked auto-merge was automatically disabled
November 2, 2026 21:14
Auto-merge was disabled for another pull request in the stack
Alex-Vasile
enabled stacked auto-merge (squash)
November 2, 2026 21:48
stacked auto-merge was automatically disabled
November 3, 2026 14:22
Auto-merge was disabled for another pull request in the stack
Alex-Vasile
deleted the
users/alvasile/unit-conftest-inprocess-parallelmap
branch
November 3, 2026 14:22
Alex-Vasile
added a commit
that referenced
this pull request
Nov 3, 2026
…14600) ISSUE ID : #14049 ## Motivation `LdsPadding.py` chooses how much padding to add to the local data share (LDS), the on-chip memory a group of GPU threads shares. It scores a candidate padding by counting bank conflicts, which happen when several threads in a wave access the same memory bank at once and their accesses have to take turns. Three functions, `_b64_wave_costs`, `_b128_wave_costs` and `_b32_wave_costs`, repeat the same loop and differ only in the bank width. A later PR in this stack adds a faster path to that loop. Putting the loop in one place first means the fast path is added once, not three times. ## Technical Details The loop moves into one `_wave_costs(rawAddrs, B, P, wOffsets, instOffs, banksPerThread)` with a helper that costs one wave. The three old names stay as thin wrappers with unchanged signatures. Behavior is meant to be identical: alignment is still checked on the summed address, a misaligned candidate still returns `None`, and an empty `instOffs` costs 0 for b64 and defaults to `(0,)` for b128 and b32. This is a production-code change in one file, with no test edits. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan Run the existing LDS padding tests, then the full CI quick-category command: `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`. ## Test Result - Wall time is unchanged within run-to-run noise: 504 s here against 511 s for the PR below, one run each, pinned to 25 CPUs on gfx942. - I did not compare the padding chosen, or the generated assembly, before and after this change. Equivalence rests on reading the diff and the existing tests; the randomized comparison against a reference arrives with the fast-path PR. - `pre-commit run --files` on the changed file passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage The change is host-side Python and needs no GPU. It sits in kernel generation code, so a behavior difference would show up in generated kernels on every architecture. Passing PR checks are sufficient only if they exercise `LdsPadding.py`, which the existing unit tests do. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): the loop is restructured in production code, with no intended change in behavior.
Contributor
TheRock Submodule Bump ActivityNewest first
|
erman-gurses
pushed a commit
that referenced
this pull request
Nov 4, 2026
…oss the unit tests (#14601) ISSUE ID : #14049 ## Motivation Many unit tests need a TensileLite `Assembler`, the object that drives the compiler toolchain, and a per-architecture capability map. The map records which instructions and features each GPU generation (an instruction set architecture, or ISA) supports. Tests built their own copies: 17 near-identical `assembler` fixtures, 9 `gfx1250_iim` fixtures, per-directory conftest copies and inline copies inside tests. The shared harness also probed all 24 supported architectures, about 40 s in a fresh process, only to obtain the assembler, which does not depend on the architecture. Every pytest process pays that cost, including each xdist worker. ## Technical Details A new `Tests/toolchain_test_helpers.py` provides a cached `shared_assembler()`, cached read-only per-ISA maps (`shared_isa_info_map` and `shared_isa_info_map_or_skip`), `copy_isa_info_map` for tests that must change a map, and `toolchain_for`. The session-scoped `assembler` and `gfx*_iim` fixtures are now defined once in `unit/conftest.py` and the copies are deleted. `test_arch_steppings.py` and `test_disableWmmaArbStall.py` change their map, so they take a copy first. `test_data_parallel_static.py` and `test_segment_interleave_policy.py` stop importing the `assembler` and `gfx1250_iim` fixtures from modules that no longer define them; without those edits the unit tree would not collect. `shared_isa_info_map_or_skip` skips only where the compiler cannot target a newer architecture (gfx1250). For gfx90a, gfx942 and gfx950 it fails, so a toolchain that drops a stable architecture is not hidden behind skips. Some of the fixtures it replaces already skipped on those architectures and others did not. A new `test_toolchain_test_helpers.py` covers that policy. The diff is 29 test files, +364 and -463 lines, and touches no product code. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan Collect the whole unit tree with `pytest --collect-only -q tensilelite/Tests/unit`, then run the CI quick-category command `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, both from `projects/hipblaslt/tensilelite`. ## Test Result - Collection (pytest's discovery of the tests) takes about 51 s here, against 52 s on develop, so it does not drop yet: the 24-architecture probe is still there. The next PR removes it. - Full-run wall time: 495 s here against 504 s for the PR below, one run each, pinned to 25 CPUs on gfx942. That gap is within noise. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage Test code only. No kernel generation, product code or GPU behavior changes, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change; product and device behavior are unchanged.
Alex-Vasile
added a commit
that referenced
this pull request
Nov 5, 2026
… parallel (#14602) ISSUE ID : #14049 ## Motivation Some unit tests need the capability map for all 24 supported GPU architectures. Each pytest process built it one architecture at a time on first use, which took about 40 s. Several copies of that builder, in conftests, the Validators test modules and `codegen_harness`, each paid the cost again, and under pytest-xdist every worker paid it while collecting tests. On develop, collection alone takes about 52 s. ## Technical Details `shared_full_isa_info_map()` is added to `toolchain_test_helpers.py`. It probes only the architectures that are not cached yet, using up to eight spawned processes, merges the results into the parent process's rocisa cache, and returns read-only entries in `SUPPORTED_ISA` order. In my runs, collection fell from about 51 s to about 12 s. The `isa_info_map` and `cxx_compiler` fixtures now live once in `unit/conftest.py`. The duplicate builders are removed, along with the LibraryIO and TensileLogic conftests that existed only to hold them, and `codegen_harness._toolchain` calls the shared helper. The two Validators modules look up one architecture when a test runs instead of building the full map when the module is imported. `test_toolchain_test_helpers.py`, added in the previous PR, gains tests of the full-map helper. A docstring notes that no test compares a spawned probe with a serial `rocIsa.init`, and that the characterization goldens guard it; I also compared the parallel result with a serial `makeIsaInfoMap` for all 24 supported ISAs and found no difference. The diff is 18 test files, +212 and -253 lines, and touches no product code. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan `pytest --collect-only -q`, then the CI quick-category command `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, from `projects/hipblaslt/tensilelite`. ## Test Result - Collection: about 12 s here against 51 s for the PR below (median of 3 runs each). - Full-run wall time: 473 s here against 495 s for the PR below, one run each, pinned to 25 CPUs on gfx942. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage Test code only. No product code or GPU behavior changes, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change; product and device behavior are unchanged.
Alex-Vasile
added a commit
that referenced
this pull request
Nov 5, 2026
…#14603) ISSUE ID : #14049 ## Motivation The true16 activation tests record how TensileLite currently generates activation code. Their module probed the gfx942 and gfx1200 architectures when it was imported, so every process that merely collected the test directory paid for those probes. The module also reset the active architecture to a hard-coded gfx942 by itself. The earlier PRs in this stack cached the per-architecture capability data these probes recompute. ## Technical Details `toolchain_test_helpers.py` gains `use_isa(arch, wavefront)`. It reads capabilities from the shared per-ISA maps, but always selects the active kernel ISA and wavefront again, because that selection cannot be cached while other tests may change the rocisa singleton (the one shared rocisa state for the process). The true16 module loses its import-time probes and its local gfx942 reset. `use_isa` and the common per-test state-preservation fixture replace them. `test_toolchain_test_helpers.py` gains four tests: reselection with a warm cache, an explicit wavefront change, state restoration, and handling of an invalid architecture. The diff is 3 test files, +83 and -47 lines, and touches no product code. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, from `projects/hipblaslt/tensilelite`. The Activation and helper tests are the ones that change. ## Test Result - Collection: about 7.8 s here against 12 s for the PR below (median of 3 runs each). - Full-run wall time: 475 s here against 473 s for the PR below, one run each, pinned to 25 CPUs on gfx942. Only collection time changes at this step. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage Test code only. No product code or GPU behavior changes, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change; product and device behavior are unchanged.
Alex-Vasile
added a commit
that referenced
this pull request
Nov 5, 2026
…ed (#14604) ISSUE ID : #14049 ## Motivation When `LdsPadding.py` scores a candidate padding for the local data share (LDS), it calls `_max_threads_per_bank` once for every wave and every instruction offset. For 18 fp8 shapes that is 1,450,696 calls that produce only 18,578 distinct answers. On the stack below it, the full unit run takes 475 s. ## Technical Details If every padded base address and every instruction offset is a multiple of `4 * banksPerThread` bytes, adding a common offset moves all threads around the memory banks together. The number of threads that collide on a bank then does not depend on the offset, so a wave needs one count, multiplied by the number of instructions. `_max_threads_per_bank` gets an opt-in `aligned` argument that counts start banks only, and `_aligned_wave_cost` is new. `_wave_costs` takes this path when the bases and offsets are aligned, and otherwise falls back to the per-instruction loop. The fallback is still needed because a misaligned base and a misaligned offset can add up to an aligned address. The shortcut is exact only when `banksPerThread` divides 64, which holds for the widths in use (1, 2 and 4), so `_wave_costs` now asserts it and the docstring says so. This changes production code in `LdsPadding.py`, plus +66 test lines in `test_LdsPadding.py`. New randomized tests compare `_wave_costs` with a reference that counts every instruction. Their inputs include bases and offsets that cancel to an aligned address, arbitrary misalignments, and empty inputs. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan `pytest tensilelite/Tests/unit/test_LdsPadding.py`, then the CI quick-category command `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`. ## Test Result - Full-run wall time: 319 s here against 475 s for the PR below, one run each, pinned to 25 CPUs on gfx942. I did not profile which tests account for the difference. - I did not compare generated assembly before and after. The equivalence check is the randomized comparison with the reference. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage The change is host-side Python and needs no GPU. It sits in kernel generation code. If the cost counts were ever wrong, padding choices would differ on every architecture. Passing PR checks are sufficient if they run the tests above. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): production padding code changes, and the new path is checked against a reference on randomized inputs.
Alex-Vasile
added a commit
that referenced
this pull request
Nov 5, 2026
…eal and eight workers (#14605) ISSUE ID : #14049 ## Motivation The CI categories that run the TensileLite unit tests use one worker (`parallel_workers: 1`), so the whole suite runs in one process. The six PRs below remove the per-process costs that made running more workers unprofitable: machine-wide process pools, repeated assembler and architecture probes, and a slow collection phase. With them in place, running the suite in parallel pays off. ## Technical Details `test_categories.yaml` sets `--dist=worksteal` and `parallel_workers: 8` for the quick and standard categories. The comprehensive and full categories also run the GPU unit tests, which take no per-GPU lock, so eight workers would all share the runner's one GPU; they keep one worker until parallel GPU runs are validated. The `hw-common`, `ffm-quick`, `emu-fast` and `emu-full` categories keep their own values. TheRock's `pytest_runner.py` passes `--numprocesses` only when a category's value is above 1, so the quick and standard categories now run under pytest-xdist. `worksteal` needs pytest-xdist 3.2 or newer, which #14798 requires. The `TESTING.md` Known Risks row that said the job runs the unit tree on one worker in every tier is updated. Under xdist, syrupy no longer prints its snapshot summary or reports unused snapshots; mismatches still fail tests, and unused snapshots were only warnings. `worksteal` lets an idle worker take tests another worker has not started yet. The default `load` scheduler hands each worker a fixed block of consecutive tests and never takes them back. About half of the suite's runtime sits in one block (`characterization/_codegen`), so some workers finished early and sat idle. The CI job has a 25-CPU allocation, and every worker repeats its own startup, so more workers stop paying for themselves. In an earlier local sample under a 25-CPU limit, one worker took 394 s, eight took 98.5 s and 25 took 102.4 s. A comment block in the file records this. Only CI configuration changes. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan Run the unit tests serially and with the new arguments, and compare outcomes: `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit` and the same command with `-n 8 --dist=worksteal`. The categories file is read by TheRock's `pytest_runner.py`, which is not in this repository, so I reproduced its arguments by hand. ## Test Result - Serial, one worker, the baseline for this PR: 321 s. - `-n 8 --dist=worksteal`: 126 s, with the same outcomes as the serial run. Both were single runs on gfx942, pinned to 25 CPUs. - The same suite took 653 s on develop. Measured cumulatively, the six PRs below bring it to 321 s and this PR brings it to 126 s. - I only ran `tensilelite/Tests/unit`, not `rocisa_tests`, and not through `pytest_runner.py`. The first CI run on this branch is the real check of how the runner reads the file. - `pre-commit run --files` on the changed file passes. Automated GitHub checks are pending. ## Related - Depends on #14798 (pytest-xdist 3.2 floor). - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage CI configuration only; no product code or GPU behavior changes. The hardware categories, and the comprehensive and full categories that include the GPU unit tests, keep their current worker counts, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): CI configuration only. If eight workers contend on a CI machine, jobs could run slower or fail intermittently; reverting the file restores one worker.
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…oss the unit tests (#14601) ISSUE ID : #14049 ## Motivation Many unit tests need a TensileLite `Assembler`, the object that drives the compiler toolchain, and a per-architecture capability map. The map records which instructions and features each GPU generation (an instruction set architecture, or ISA) supports. Tests built their own copies: 17 near-identical `assembler` fixtures, 9 `gfx1250_iim` fixtures, per-directory conftest copies and inline copies inside tests. The shared harness also probed all 24 supported architectures, about 40 s in a fresh process, only to obtain the assembler, which does not depend on the architecture. Every pytest process pays that cost, including each xdist worker. ## Technical Details A new `Tests/toolchain_test_helpers.py` provides a cached `shared_assembler()`, cached read-only per-ISA maps (`shared_isa_info_map` and `shared_isa_info_map_or_skip`), `copy_isa_info_map` for tests that must change a map, and `toolchain_for`. The session-scoped `assembler` and `gfx*_iim` fixtures are now defined once in `unit/conftest.py` and the copies are deleted. `test_arch_steppings.py` and `test_disableWmmaArbStall.py` change their map, so they take a copy first. `test_data_parallel_static.py` and `test_segment_interleave_policy.py` stop importing the `assembler` and `gfx1250_iim` fixtures from modules that no longer define them; without those edits the unit tree would not collect. `shared_isa_info_map_or_skip` skips only where the compiler cannot target a newer architecture (gfx1250). For gfx90a, gfx942 and gfx950 it fails, so a toolchain that drops a stable architecture is not hidden behind skips. Some of the fixtures it replaces already skipped on those architectures and others did not. A new `test_toolchain_test_helpers.py` covers that policy. The diff is 29 test files, +364 and -463 lines, and touches no product code. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan Collect the whole unit tree with `pytest --collect-only -q tensilelite/Tests/unit`, then run the CI quick-category command `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, both from `projects/hipblaslt/tensilelite`. ## Test Result - Collection (pytest's discovery of the tests) takes about 51 s here, against 52 s on develop, so it does not drop yet: the 24-architecture probe is still there. The next PR removes it. - Full-run wall time: 495 s here against 504 s for the PR below, one run each, pinned to 25 CPUs on gfx942. That gap is within noise. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage Test code only. No kernel generation, product code or GPU behavior changes, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change; product and device behavior are unchanged.
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
… parallel (#14602) ISSUE ID : #14049 ## Motivation Some unit tests need the capability map for all 24 supported GPU architectures. Each pytest process built it one architecture at a time on first use, which took about 40 s. Several copies of that builder, in conftests, the Validators test modules and `codegen_harness`, each paid the cost again, and under pytest-xdist every worker paid it while collecting tests. On develop, collection alone takes about 52 s. ## Technical Details `shared_full_isa_info_map()` is added to `toolchain_test_helpers.py`. It probes only the architectures that are not cached yet, using up to eight spawned processes, merges the results into the parent process's rocisa cache, and returns read-only entries in `SUPPORTED_ISA` order. In my runs, collection fell from about 51 s to about 12 s. The `isa_info_map` and `cxx_compiler` fixtures now live once in `unit/conftest.py`. The duplicate builders are removed, along with the LibraryIO and TensileLogic conftests that existed only to hold them, and `codegen_harness._toolchain` calls the shared helper. The two Validators modules look up one architecture when a test runs instead of building the full map when the module is imported. `test_toolchain_test_helpers.py`, added in the previous PR, gains tests of the full-map helper. A docstring notes that no test compares a spawned probe with a serial `rocIsa.init`, and that the characterization goldens guard it; I also compared the parallel result with a serial `makeIsaInfoMap` for all 24 supported ISAs and found no difference. The diff is 18 test files, +212 and -253 lines, and touches no product code. Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on. | Ref | Collection (s) | Run phase (s) | Total wall (s) | |---|---|---|---| | b279712 (base) | 52.3 | 600.4 | 652.7 | | #14599 | 49.4 | 461.1 | 510.5 | | #14600 | 47.0 | 456.6 | 503.6 | | #14601 | 50.9 | 443.9 | 494.8 | | #14602 | 12.1 | 461.1 | 473.2 | | #14603 | 7.8 | 466.8 | 474.6 | | #14604 | 7.6 | 311.6 | 319.2 | | #14605 | 7.7 | 313.7 | 321.4 | | #14605 (`-n 8 --dist=worksteal`) | n/a | n/a | 125.6 | ## Test Plan `pytest --collect-only -q`, then the CI quick-category command `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, from `projects/hipblaslt/tensilelite`. ## Test Result - Collection: about 12 s here against 51 s for the PR below (median of 3 runs each). - Full-run wall time: 473 s here against 495 s for the PR below, one run each, pinned to 25 CPUs on gfx942. - `pre-commit run --files` on the changed files passes. Automated GitHub checks are pending. ## Related - Work tracking: [#14049](#14049), the issue about the unit tests exceeding the 15-minute quick-tier limit on gfx94X. This PR is part of the follow-up to [#14058](#14058), which fixed the immediate cause. ## Device / Architecture Coverage Test code only. No product code or GPU behavior changes, so passing PR checks are sufficient. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change; product and device behavior are unchanged.
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.
ISSUE ID : #14049
Motivation
TensileLite's unit tests call
ParallelMap, a helper that spreads a function over a pool of worker processes. The pool size comes from theCpuThreadssetting, whose default of -1 means every CPU the process can see, not the CPUs the CI job was given. About 460 unit tests passParallelMapone item or a few, so a pool cannot speed them up. Each pytest process, and each worker started by pytest-xdist (a plugin that runs tests in several processes at once), still started a machine-wide pool.Issue #14049 reported that the TensileLite unit tests no longer fit the 15-minute quick-tier limit on gfx94X. PR #14058 fixed the immediate cause, slow parsing of custom-kernel metadata, and brought the unit-test stage from 17.5 to 14.0 minutes, which leaves little headroom against the limit. This is the first of seven PRs that push the run time down further.
Technical Details
A new session-scoped autouse fixture in
Tests/unit/conftest.pysetsCpuThreadsto 1 in bothglobalParametersanddefaultGlobalParametersand undoes that when the session ends. The second assignment matters becauserestoreDefaultGlobalParameters()copies from it and would otherwise undo the first. The fixture runs once per process that runs unit tests, including every xdist worker. Tests that need another value still set it themselves. It is a fixture rather than apytest_configurehook because pytest runs a nested conftest's hook as soon as it imports that conftest during collection, so a run that collects all oftensilelite/Tests(for exampletox -e py3 -- tensilelite/Tests -m common) would also have changedCpuThreadsfor the non-unit tests that call Tensile in-process. I checked with a probe plugin that a test underTests/commonstill sees -1 whenTests/unitis collected in the same run, and that unit tests see 1.Nothing else in the unit lane sends work to a worker process any more, so one new test in
characterization/ParallelrunsParallelMapandParallelMap2withCpuThreads=2and checks that the work runs in other processes and sees the parent's global parameters (about 0.3 s).TESTING.mdnow says where the real pool path is covered.Unit-test run time for each PR in this stack, one run each on a gfx942 host pinned to 25 CPUs, serial except the last row. Collection is the median of 3 runs. Run phase is total minus collection. The first row is the commit this stack is based on.
-n 8 --dist=worksteal)Test Plan
Run the unit tests the way the CI quick category does:
pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit, fromprojects/hipblaslt/tensilelite, with no parallelism.Test Result
pre-commit run --fileson the changed files passes.Related
Device / Architecture Coverage
The change affects how the unit tests run on the host and does not touch kernel generation or GPU code. Passing PR checks are sufficient.
Submission Checklist
Risk level
Low (1/5): test-only change; product and device behavior are unchanged.