Repository navigation
test(tensilelite): defer true16 ISA probes and share target selection - #14603
Merged
Alex-Vasile merged 1 commit intoNov 5, 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 #14603 +/- ##
========================================
Coverage 77.71% 77.72%
========================================
Files 2904 2904
Lines 490964 490964
Branches 74412 74412
========================================
+ Hits 381539 381553 +14
+ Misses 84432 84424 -8
+ Partials 24993 24987 -6
*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 one looks good to me. I was a little nervous about dropping _restore_legacy_isa, since its comment says sibling Activation modules assume gfx942. But _isolate_rocisa_state in unit/conftest.py is the outer fixture, so it already restored the pre-test kernel after _restore_legacy_isa ran. I think the old fixture was effectively a no-op. Taking the probes out of import time is also a nice determinism win for #14605: what rocisa has selected after collection no longer depends on which modules a worker happened to collect.
The four new use_isa tests cover the cases I'd want (warm reselection, wavefront change, restore, bad arch). No inline comments from me.
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
October 30, 2026 21:53
4431901 to
7d2aa93
Compare
Contributor
PR Pundit review(Testing PR Pundit Code Review agent. Not all comments may be relevant. This is a one-time automated pass: please address what is useful, no need to request a re-review. A quick reply per point (useful or not, and why) helps us calibrate the agent. Reach out to @peymanr for suggestions and feedback.)
1. 🟠 Should fix: Plain tuple ISA args produce tuple-typed keys in shared lru_cache, currently harmless
Who is affected: Unit-test workers where a tuple-argument test builds the gfx1200/gfx942 map before a gfx-name caller; no test observes a difference today. Checked against the code by an independent pass. Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at |
1 task done
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 2, 2026 20:32
7d2aa93 to
bb530c1
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
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 2, 2026 21:44
bb530c1 to
f1de304
Compare
Contributor
Author
Thanks, fixed. |
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
added a commit
that referenced
this pull request
Nov 3, 2026
…14599) 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 the `CpuThreads` setting, whose default of -1 means every CPU the process can see, not the CPUs the CI job was given. About 460 unit tests pass `ParallelMap` one 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.py` sets `CpuThreads` to 1 in both `globalParameters` and `defaultGlobalParameters` and undoes that when the session ends. The second assignment matters because `restoreDefaultGlobalParameters()` 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 a `pytest_configure` hook because pytest runs a nested conftest's hook as soon as it imports that conftest during collection, so a run that collects all of `tensilelite/Tests` (for example `tox -e py3 -- tensilelite/Tests -m common`) would also have changed `CpuThreads` for the non-unit tests that call Tensile in-process. I checked with a probe plugin that a test under `Tests/common` still sees -1 when `Tests/unit` is 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/Parallel` runs `ParallelMap` and `ParallelMap2` with `CpuThreads=2` and checks that the work runs in other processes and sees the parent's global parameters (about 0.3 s). `TESTING.md` now 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. | 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 the way the CI quick category does: `pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit`, from `projects/hipblaslt/tensilelite`, with no parallelism. ## Test Result - Wall time for the full run, one run each on a gfx942 host with the process pinned to 25 CPUs: develop 653 s, this PR 511 s. - `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 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 - [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
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 3, 2026 14:22
f1de304 to
06ae8a7
Compare
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 3, 2026 14:42
06ae8a7 to
b2a181a
Compare
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.
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 3, 2026 18:58
b2a181a to
a157e58
Compare
Contributor
PR Pundit review(Testing PR Pundit Code Review agent. Not all comments may be relevant. This is a one-time automated pass: please address what is useful, no need to request a re-review. Please reply to each point, whether it applies or not: we harvest these replies, and both kinds are enormously helpful for tuning the agent. Reach out to @peymanr for suggestions.)
1. 🟠 Should fix: Test loads rocisa_test_state twice as a second module copy: this PR used a bare import
The new test imports Who is affected: Only readers and maintainers of Checked against the code by an independent pass. 2. 🟠 Should fix: true16 activation tests error with FileNotFoundError without amdclang++: this PR switched to compiler-probing use_isa
Who is affected: Developers running only this test module locally without Checked against the code by an independent pass. Related past reviews: #9711 on Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at |
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.
erman-gurses
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 4, 2026 17:46
a157e58 to
f77792f
Compare
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 4, 2026 19:38
f77792f to
5291a2d
Compare
Contributor
Author
|
Thanks @peymanr. 1. Bare 2. true16 tests failing without
I checked this by making
So without a compiler, the change turns an import-time crash into a clear per-test failure. The compiler lookup now also checks |
Alex-Vasile
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 4, 2026 20:21
5291a2d to
dcebc4a
Compare
Base automatically changed from
users/alvasile/tests-parallel-full-isa-map
to
develop
November 5, 2026 03:40
Add use_isa(arch, wavefront) to toolchain_test_helpers. It initializes capabilities through the shared per-ISA maps but always reselects the active kernel ISA and wavefront, because that selection cannot be cached while other tests may change the rocisa singleton. Remove the true16 Activation characterization module's import-time gfx942 and gfx1200 probes and its local hardcoded gfx942 reset; use_isa and the common per-test state-preservation fixture replace them. Add contract tests for warm-cache reselection, explicit wavefront changes, state restoration and invalid architecture handling.
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
force-pushed
the
users/alvasile/tests-lazy-true16-isa-selection
branch
from
November 5, 2026 03:40
dcebc4a to
bf43e3a
Compare
Alex-Vasile
deleted the
users/alvasile/tests-lazy-true16-isa-selection
branch
November 5, 2026 14:13
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.
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
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.pygainsuse_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_isaand the common per-test state-preservation fixture replace them.test_toolchain_test_helpers.pygains 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.
-n 8 --dist=worksteal)Test Plan
pytest -m "not gpu" --snapshot-warn-unused tensilelite/Tests/unit, fromprojects/hipblaslt/tensilelite. The Activation and helper tests are the ones that change.Test Result
pre-commit run --fileson the changed files passes. Automated GitHub checks are pending.Related
Device / Architecture Coverage
Test code only. No product code or GPU behavior changes, so passing PR checks are sufficient.
Submission Checklist
Risk level
Low (1/5): test-only change; product and device behavior are unchanged.