Repository navigation
ci(gfx1250): run the large-memory hipBLASLt tests one job at a time - #14545
Merged
Merged
Conversation
#13504 enabled 93 hipBLASLt tests with 2^31/2^32 sizes on gfx1250 (size_t_index_overflow, size_t_output_overflow, offset_above_2_32, batch_offset_large). With 4 gtest shards every shard reaches them at the same time, and the client's pool keeps their pinned host buffers until the process exits, so a single job nearly filled the CI host's memory and concurrent jobs ran it out of memory. The shards now skip those tests. After the shards, the job runs them one group per process (each starts with an empty pool), behind a lock file shared by the host's runners (/scratch/ci-locks), so only one job at a time is in that phase. A custom hipblaslt_gtest_filter runs as given, as before. No test is dropped: the logs keep the hipblaslt_test_shard*.log names, so the report counts them. Co-Authored-By: Claude <noreply@anthropic.com>
github-actions
Bot
added
github actions
project: none
archana-ramalingam
marked this pull request as ready for review
October 28, 2026 08:31
✅ 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. |
…cesses Together they keep ~82 GiB in one process; apart the larger peaks at ~60. Co-Authored-By: Claude <noreply@anthropic.com>
tony-davis
approved these changes
Oct 28, 2026
tony-davis
left a comment
tony-davis
left a comment
Contributor
There was a problem hiding this comment.
Thanks Archana! I think this is a reasonable stopgap, and I'm okay with it going in as one. My worry is that the 93 tests from #13504 are normal tests we should be able to run, and a test shouldn't be able to take a runner down. Even when memory is short, it should skip or fail on its own and leave the host alone. I looked at the client, and I think the cause is on the harness side: the host buffer pool keeps pinned buffers until the process exits, and the memory checks are per process and point-in-time, so several processes can all pass them and then all allocate. The existing fast_check_memory_shortfall preflight also doesn't cover the stress cases. I'm taking that on as AIHPBLAS-5352 (assigned to me), so we can revert the lock later.
A few non-blocking comments on this PR.
- gtest exits 0 when a filter matches nothing, so a renamed test would make a large group silently run 0 tests. Could we fail if any
hipblaslt_test_shard_large*.logreportsRunning 0 tests? - The step timeout is 120 min and the lock wait can use 50 of it. What's the wall time of the large phase? If it's much over about 16 min, the 4th queued job could hit the wait cap.
- A custom
hipblaslt_gtest_filterstill runs the large tests in the shards with no lock, so a manual run could OOM the host again. Documenting it may be enough, but I wanted to flag it.
Also, it would be good to add a comment pointing at AIHPBLAS-5352 so the lock gets removed once the harness fix lands.
Review follow-ups: - A custom hipblaslt_gtest_filter now gets the same split: the shards add the large-memory patterns to its negatives, and the large step runs the large tests the filter selects (names from --gtest_list_tests), so a manual run can't put them back into the shards. - Fail when a large group matches no test without a custom filter (gtest passes a filter that matches nothing), or a large log runs 0 tests. - Point the stopgap at AIHPBLAS-5352 so the lock is removed with it. Co-Authored-By: Claude <noreply@anthropic.com>
Contributor
Author
Thanks Tony, agreed it's a stopgap; added AIHPBLAS-5352 to the step comment and description so the lock goes with your fix.
|
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
tony-davis
added a commit
that referenced
this pull request
Nov 6, 2026
… them (#14608) JIRA ID : AIHPBLAS-5352 ## Motivation `hipblaslt-test`, the hipBLASLt test client, keeps the host buffers that each test frees in a pool, so the next test can reuse them instead of allocating again. These buffers are pinned (allocated with `hipHostMalloc`, so the operating system can neither page them out nor reclaim them), and the pool kept every one until the process exited. On the gfx1250 hardware test host, which has about 499 GiB of memory, the 93 cases with 2^31- and 2^32-element sizes enabled in #13504 left about 123 GiB pinned in one process. With 4 CI runners each splitting the suite into 4 parallel test processes, the out-of-memory killer took the runners down on 10-23 and 10-27. Separately, a failed host allocation returns a null buffer, and the test then writes through it. I don't think this stack guarantees the gfx1250 host never runs short. What it does is make the client hold much less pinned memory, give it back sooner, and skip a case with the amounts instead of getting killed, which is the part of the problem the client can fix. If pressure comes back, the 110 GiB container caps from #14569 keep a kill inside one job, and putting the lock back (#14982 removes it) is a one-file revert. So I'm not planning a long soak test before merging. The pool comes from ROCm/hipBLASLt#1202, which reused allocations to speed up `hipblaslt-bench` across about 500,000 small matmuls. Reuse still matters for large buffers: on the test host below, `hipHostMalloc` takes about 0.18 s per GiB and `hipHostFree` about 0.06 s per GiB. ## Technical Details - A pooled host buffer of 1 GiB or more (`kHostPoolRetainLimit` in `d_vector.hpp`) now stays cached only while consecutive tests reuse it. The pool records the test count when a buffer comes back, and `HostPoolReleaseListener`, a new hook that Google Test calls after every test, frees the large buffers that the test just ended did not use. A run of same-shape cases, such as the four transpose combinations (NN, NT, TN, TT) of one size, still shares one allocation. - When a test asks for a large buffer that no cached buffer fits, the pool first frees the large buffers that earlier tests left, then pins new memory. This keeps the peak near one test's footprint when the shape changes. - A host buffer no longer gets the pool's 20% extra capacity when the extra would take it to 1 GiB or more (`host_pool_capacity()`), so a buffer under the limit stays under it. The existing cutoff for skipping the extra is a quarter of device memory, so on gfx1250 nearly every host buffer got it. - `HipHostBuffer::memcheck()` reports a failed pinned allocation, and `testing_matmul` checks every host buffer after allocating, including the four temporary buffers a swizzled A or B is built in. A failure now skips the test with the requested size ("could not allocate N GiB of pinned host memory"). - A pooled host buffer under 1 GiB is released once 100 consecutive tests have not reused it (`kHostPoolSmallIdleTests`). The pool reuses a buffer only for requests between half its size and its size, so without this a run of shrinking requests kept every earlier buffer pinned. - In `hipblaslt-bench`, a failed host allocation now exits with a failure, as a failed device allocation already did, instead of reporting success. - Device memory behaves as before. `hipblaslt-bench` does not install the listener, so it never frees large buffers early; it does see the 20% change for host buffers of 1 GiB or more. The hipBLASLt library itself is unchanged. This is the first of four stacked PRs for AIHPBLAS-5352. #14609 makes the client's host-memory query respect container memory limits, #14610 runs the memory preflight for every matmul case, and #14982 drops the host-wide lock in the gfx1250 workflow. ## Test Plan - `HostMemory_pre_checkin` (new, `clients/tests/src/host_memory_gtest.cpp`, 6 tests; the `_pre_checkin` suffix puts it in the standard tier the PR checks run). The tests check that a large buffer survives one test boundary and is freed after the next, that a same-shape test reuses it (what the first test wrote is still there), that a large miss frees two buffers earlier tests left, that a 1 GiB request gets exactly 1 GiB, that a small buffer is released after 100 tests that do not reuse it (and kept when one does), and that a failed allocation leaves a null buffer whose size is reported. A 0.9 GiB request gets exactly 0.9 GiB (the old rule gave it 1.08 GiB, over the limit). The tests skip, naming the buffer, on a host that cannot pin that much, and the reuse test skips on a Windows host where the pool frees its idle buffers before every request. A temporary variant without the small-buffer aging failed the new test. - To confirm the tests can fail, I built two temporary variants of the pool. One undid four behaviors at once (no release at the end of a test, no release on a miss, the 20% extra restored, `memcheck()` always succeeding), and it failed 4 of the 5 tests. The other freed large buffers whether or not a test used them, and it failed 3. Every test failed under at least one variant. - The 82 large-memory cases that apply to gfx90a (the five name patterns from #14545), and the `smoke` and `quick` tiers, on develop and on this branch, for wall time and peak resident memory. - Handle leaks: #14714 and #14700 on develop give the matrix layouts, matmul descriptors and HIP events owners, so an early return no longer leaks them. To check that with this stack's new early returns, a preloaded library counted creates and destroys of all three, on the full stack, over 1,892 matmul cases, once normally and once under a 1 MiB host limit so that most cases skip at the preflight. - Adjacent: the `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin` suites, including `memory_shortfall_reclaims_idle_buffers`, which uses the pool's existing `ReleaseCached()`. `hipblaslt-bench` was built, because it compiles the non-gtest form of the new check, but not run. ## Test Result One MI210 (gfx90a) on a shared host with 2.2 TB of memory. This branch's clients were built against the 10-25 TheRock nightly hipBLASLt gfx90a device library. Each pair of runs went back to back, but the host was busy (load average near 300), so the times carry some noise. | Measurement | develop | this PR | | --- | --- | --- | | 82 large-memory cases, wall time | 316.9 s | 313.8 s | | 82 large-memory cases, peak resident memory | 81.9 GiB | 49.2 GiB | | 82 large-memory cases, mean resident memory | 66.9 GiB | 16.6 GiB | | Resident memory after the last large case | 74 GiB | 2 GiB | | `smoke` + `quick` (9,171 tests), two runs | 66.6 s, 62.0 s | 63.6 s, 64.1 s | - Large-memory cases: 82 of 82 passed on both. `smoke` + `quick`: 9,169 passed and 2 skipped (`GroupedMatmul.smoke_Fp16NN` and `smoke_Bf16NN`, which skip on this GPU) on both. - `HostMemory_pre_checkin`, `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin`: 64 of 64 passed. - A first version that freed every large buffer at the end of each test took 449.6 s for the large cases against 318.2 s on develop, which is why the pool keeps a buffer while the next test reuses it. - Handle leaks: 0 layouts, descriptors or events left, both normally (13,580 created) and with 409 cases skipped at the preflight (11,235 created). - With the small-buffer aging, `smoke` + `quick` passed 9,225 tests in 59.3 s with peak resident memory 2.18 GiB. - gfx1250: this PR alone passed the hardware workflow's hipBLASLt tests in [40934683157](https://github.com/ROCm/rocm-libraries/actions/runs/40934683157), and Archana dispatched the full stack in [41582794047](https://github.com/ROCm/rocm-libraries/actions/runs/41582794047), which passed hipBLASLt twice with no kills and no memory skips. Her run on the rebased stack, [41717825380](https://github.com/ROCm/rocm-libraries/actions/runs/41717825380), was clean too. All three kept the host-wide lock; #14982 covers the run without it. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5). Every `hipblaslt-test` case passes through the new listener and the host pool, so a mistake here would affect the whole suite. The hipBLASLt library and device behavior are unchanged; only the test and bench clients change. The bench sees the 20% change near 1 GiB, and now exits with a failure when a host allocation fails. ## Related - Work tracking: [AIHPBLAS-5352](https://amd-hub.atlassian.net/browse/AIHPBLAS-5352), fixes 1 (pool retention) and 4 (null host buffer) of its list. - Cause: #13504. CI workaround this work is meant to replace: #14545, merged on 10-28. - Where the pool came from: ROCm/hipBLASLt#1202. Windows already clears the pool on smaller-memory hosts for the same reason: #4793, #6121. ## Device / Architecture Coverage Host-side test-client code only, with no kernel, library or kernel-selection change, so it does not depend on the GPU architecture and the standard PR checks are enough. Measured on gfx90a. ## Flags / Guardrails None. The 1 GiB threshold is the constant `kHostPoolRetainLimit`. [AIHPBLAS-5352]: https://amd-hub.atlassian.net/browse/AIHPBLAS-5352?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
added a commit
that referenced
this pull request
Nov 6, 2026
…st memory skips (#14610) JIRA ID : AIHPBLAS-5352 ## Motivation Before allocating, `hipblaslt-test`, the hipBLASLt test client, estimates how much device and host memory a matmul case needs, and skips the case with both amounts when either does not fit. That memory preflight ran only for `fast_check` cases. The `stress` cases checked with `unit_check`, such as `matmul_size_t_index_overflow` and its `_hpa` variant, pin up to about 48 GiB of host memory each on gfx90a and allocated without any check, so on a host or container without room they were killed rather than skipped. Memory skips were also easy to miss: a device out-of-memory error was reported as a passing test, and nothing summarized which cases did not run. ## Technical Details - Every `hipblaslt-test` matmul case in `testing_matmul` now runs the preflight; `hipblaslt-bench` still runs it only for `fast_check` cases. One estimate in `testing_matmul.hpp` serves both kinds of case. Besides the buffers it already counted, it now adds the pool's 20% extra capacity on host buffers below 1 GiB, the CPU reference's float or double copies of one batch of A and B (and of D when the output type is not the compute type), the temporary buffers a swizzled A or B is built in, and one set of buffers per batch in pointer-array mode. - `matmul_batch_offset` cases go through their own function, `testing_matmul_batch_offset`, which now runs the same preflight. Its estimate counts the four host and device buffers, the reference D and the 128 MiB workspace bound. - Every stop for lack of memory goes through one macro, `SKIP_FOR_MEMORY`: the preflight, a failed pinned host allocation, and a device out-of-memory error. Each one skips the test with a message that begins "Not enough memory:", and the end of the run lists those tests under `[ MEMORY ]`, one line each with the amounts. The unused `LIMITED_MEMORY_STRING` messages are removed. - `HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1` turns those skips into failures, so a lane on a large-memory host can require that every case ran. - `HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE=1` (Linux) checks the estimate on real cases. For a case estimated at 256 MiB or more, it resets the process's peak resident memory (`VmHWM`, which counts pinned and pageable memory alike), and fails the case when the peak rose more than the estimate plus 10% and 256 MiB. It leaves out the first case in a process, which also loads the kernel library (about 0.6 GiB here). - Behavior change for reviewers to weigh: a device out-of-memory error at allocation (`CHECK_DEVICE_ALLOCATION`) used to call `SUCCEED()`, which counts as a pass. It is now a skip. On a GPU with enough memory nothing changes; on a smaller one those cases move from "passed" to "skipped", which is what they were. This is the last of three stacked PRs for AIHPBLAS-5352 and depends on #14609, which depends on #14608. ## Test Plan - `MatmulMemory_pre_checkin` (2 new tests in `clients/tests/src/matmul_gtest.cpp`). One runs a 4096 x 4096 x 16 f32 `unit_check` case under a 64 MiB `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` with `HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1`, and expects the preflight's message. The other runs two cases under `HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE=1` and bounds the estimate from both sides: f32 8192 x 8192 x 16, where six 256 MiB host buffers with their 20% extra make up the estimate, and f16 16384 x 16 x 16384, where the CPU reference's 1 GiB float copy of A is most of it. Together they take about 3 s and pin up to about 1.9 GiB. - `MatmulMemory_pre_checkin.preflight_stops_a_batch_offset_case_that_does_not_fit` does the same for a `matmul_batch_offset` case. The test helpers restore any environment variable the run started with (James Newling's commit). - `HostMemory_pre_checkin.memory_shortages_are_listed_with_their_reasons` checks the `[ MEMORY ]` summary. - To confirm the tests can fail, I built three temporary variants. Dropping the 20% extra from the estimate failed the f32 case (estimate 1.5 GiB, used 1.93 GiB), dropping the reference copies failed the f16 case (0.606 GiB against 1.67 GiB), and running the preflight only for `fast_check` again failed both tests. - The standard tier (`*smoke*:*quick*:*pre_checkin*`) on develop, on this PR, and on this PR with `HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE=1`. - The 82 large-memory cases that apply to gfx90a (the five name patterns from #14545): normally, with the estimate check, and inside a transient 8 GiB cgroup (`systemd-run --user --scope -p MemoryMax=8G -p MemorySwapMax=0`). ## Test Result One MI210 (gfx90a), with this branch's clients built against the 10-25 TheRock nightly hipBLASLt gfx90a device library. - Standard tier: develop passed 19,695 and skipped 26; this PR passed 19,715 and skipped the same 26. The 20 extra passes are the new tests across the three PRs, so the preflight added no skips there. With the estimate check on, nothing failed. Peak resident memory for the tier fell from 30.1 GiB to 21.5 GiB, mostly from #14608. - In the estimate tests the f32 case was estimated at 1.94 GB and used 1.94 GB, and the f16 case 1.73 GB and 1.71 GB. - Large-memory cases: 82 of 82 passed (327.5 s, peak resident memory 49.2 GiB). With the estimate check all 82 passed again; for the 35 `stress` cases it measured (the 36th was the first case in the process), the rise in peak resident memory was between 0.996 and 1.003 times the estimate, which reached 48.1 GiB. The 46 `fast_check` cases estimate less than 256 MiB of host memory each, so the check does not apply to them. - Inside the 8 GiB cgroup: develop was killed (exit code 137) within 4 seconds, at the first `stress` case. This PR exited 0 after 84 s: the 46 `fast_check` cases passed, and the 36 `stress` cases skipped and were listed, for example "this case needs 16.2 GiB of host memory and 7.41 GiB is available". - `matmul_batch_offset_large`'s shape (2^32-element offsets, run by hand on gfx90a because the YAML case is limited to gfx942, gfx950 and gfx1250): the estimate was 51.55 GB and the measured rise 51.51 GB. - `smoke`, the memory and `fast_check` suites and the `batch_offset` cases: 2,802 passed and 9 `GroupedMatmul` cases skipped, both with nothing set and with `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT=4G` and `HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1` set for the whole run. No memory skips. - gfx1250: Archana dispatched the full stack in [41582794047](https://github.com/ROCm/rocm-libraries/actions/runs/41582794047). It passed hipBLASLt twice (4 shards and all 6 large groups), with no kills and an empty `[ MEMORY ]` list, with the large-case phase still separate from the shards. Shards still check memory one moment at a time and reserve nothing; whether that matters with the lock removed needs its own runs. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5). Every `hipblaslt-test` matmul and batch-offset case now runs the preflight, so an estimate that is too high would skip cases that would have run; the standard tier shows the same skips as develop, and the estimate tests bound overestimates. Device out-of-memory errors now report as skips instead of passes. The hipBLASLt library and device behavior are unchanged. ## Related - Work tracking: [AIHPBLAS-5352](https://amd-hub.atlassian.net/browse/AIHPBLAS-5352), fixes 2 (preflight for every case, one estimator) and 5 (visible memory skips) of its list. - Depends on #14609 and #14608. - #14545 runs the large cases after the shards behind a host-wide lock, and #14569 caps each gfx1250 container at 110 GiB. With this stack and the cap in place, a follow-up change to the gfx1250 workflow can remove the lock; the large cases keep their own phase until the question of several shards in one container (fix 6 in the ticket) is settled. ## Device / Architecture Coverage Host-side test-client code only, with no kernel, library or kernel-selection change, so the standard PR checks are enough for correctness. The gfx1250 hardware run in the ticket's acceptance checks the result where the failure happened. ## Flags / Guardrails `HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP` and `HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE`, both off by default and documented in `TESTING.md`. [AIHPBLAS-5352]: https://amd-hub.atlassian.net/browse/AIHPBLAS-5352?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : N/A
Motivation
The gfx1250 HW host ran out of memory (10-23, 10-27); the workflow is disabled. Cause: #13504 enabled 93 hipBLASLt tests with 2^31/2^32 sizes on gfx1250. All 4 shards reach them at once, and the client's pool keeps their pinned host buffers, so one job nearly fills the host.
Technical Details
flock(/scratch/ci-locks), one job at a time; 50-min wait cap.hipblaslt_gtest_filtergets the same split: the shards add the 93 to its negatives, and the large step runs the large tests the filter selects (from--gtest_list_tests).hipblaslt_test_shard_large*.log.Test Plan
hipblaslt-test.Test Result
Submission Checklist
Risk level
Low. gfx1250 HW workflow only; same tests, different order.
🤖 Generated with Claude Code