[tony-davis]

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 ci(gfx1250): run the large-memory hipBLASLt tests one job at a time #14545), and the smoke and quick tiers, on develop and on this branch, for wall time and peak resident memory.
• Handle leaks: fix(hipblaslt): release matmul descriptors on early exits #14714 and fix(hipblaslt): release matmul timing events on early exits #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, and Archana dispatched the full stack in 41582794047, which passed hipBLASLt twice with no kills and no memory skips. Her run on the rebased stack, 41717825380, was clean too. All three kept the host-wide lock; ci(gfx1250): drop the host-wide lock on the large hipBLASLt tests #14982 covers the run without it.
• Automated GitHub checks are pending.

Submission Checklist

• 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, fixes 1 (pool retention) and 4 (null host buffer) of its list.
• Cause: test(hipblaslt): run applicable client tests on gfx1250 and gfx1250-strict #13504. CI workaround this work is meant to replace: ci(gfx1250): run the large-memory hipBLASLt tests one job at a time #14545, merged on 10-28.
• Where the pool came from: hipblaslt-bench performance improvements hipBLASLt#1202. Windows already clears the pool on smaller-memory hosts for the same reason: [hipblaslt] Fix hipblaslt-test abort issue on Windows #4793, [hipblaslt] Fix hipblaslt-test abort issue on Windows #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.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #14608      +/-   ##
===========================================
+ Coverage    77.60%   77.65%   +0.04%     
===========================================
  Files         2903     2905       +2     
  Lines       490304   490744     +440     
  Branches     74249    74315      +66     
===========================================
+ Hits        380492   381039     +547     
+ Misses       84781    84684      -97     
+ Partials     25031    25021      -10     

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	53.94% <ø> (ø)		Carriedforward from 7be597e
TensileLite-Unit	84.37% <ø> (ø)		Carriedforward from 7be597e
hipBLAS	90.62% <ø> (ø)		Carriedforward from 7be597e
hipBLASLt	70.05% <ø> (+1.08%)	⬆️
hipCUB	82.68% <ø> (ø)		Carriedforward from 7be597e
hipDNN	87.25% <ø> (ø)		Carriedforward from 7be597e
hipFFT	42.57% <ø> (ø)		Carriedforward from 7be597e
hipRAND	76.12% <ø> (ø)		Carriedforward from 7be597e
hipSOLVER	69.67% <ø> (ø)		Carriedforward from 7be597e
hipSPARSE	87.03% <ø> (ø)		Carriedforward from 7be597e
rocBLAS	69.20% <ø> (ø)		Carriedforward from 7be597e
rocFFT	46.49% <ø> (ø)		Carriedforward from 7be597e
rocRAND	57.32% <ø> (ø)		Carriedforward from 7be597e
rocSOLVER	76.74% <ø> (ø)		Carriedforward from 7be597e
rocSPARSE	83.54% <ø> (ø)		Carriedforward from 7be597e
rocThrust	93.15% <ø> (ø)		Carriedforward from 7be597e

*This pull request uses carry forward flags. Click here to find out more.
see 10 files with indirect coverage changes

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
• 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

[Alex-Vasile]

Approving CmakeLists changes

[newling]

Thanks for these PRs @tony-davis!

Overall, I’m happy with the direction of this stack. It addresses an immediate issue, and the agent review requests are small, local fixes and improvements.

Agent review: Review of the host-memory PR stack. The accompanying branches contain proposed fixes for the zero-budget handling and environment restoration.

I’m still getting familiar with the gfx1250 CI setup. I assume the workflow is configured to run automatically for eligible PRs. Just wondering if we need to do some sort of manual validation before landing.

Separately, I’d like us to revisit the complexity of host-memory allocation and caching in this part of the client, including where pinning and buffer reuse are worthwhile. That is follow-up work, not something I’m asking you to address in this stack.

[archana-ramalingam]

Thanks @tony-davis for this fix!
Approved with minor non-blocking follow-up toward AIHPBLAS-5352 (4 concurrent jobs, no OOM; a shortage skips a case instead of killing a shard), from the gfx1250 cap kills:

• Measured on a gfx1250 host before this stack (10-28, 4 shards without the 2^31/2^32 tests): shard peaks 25/25/31/53 GiB, 134 GiB if they align (cap 110). The large tests in one process held 122.6 GiB until exit.
• Their test positions align: on 11-02, without this stack, 9 of 12 shards were killed at the cap (41512870903, 41517092105, 41519929996), all at test 15,619–15,975 of ~17,470, where the 4 shards, in step, run their biggest in-shard cases (fast_check_large 16385×16383×4097 f32, complex 4096×2187×2653).
• In that probe the 4 shards together reached only ~69 GiB (host sampled every 10 s), so short spikes or CI-only usage make up the gap to 110. Shared memory was 334 of 393 GiB at the full-suite peak, likely the pinned host buffers, and the cap has no swap (relevant to @newling's pinning follow-up).

Buffers ≥ 2x a request stay in the pool (fix 1, below 1 GiB).
A miss skips them (capacity() < 2 * bytes) and evicts only the largest too-small one (#L351-L353), so they stay until a later request fits them or a shortfall clears the pool. EndTest (#L259) could age them out with the ≥ 1 GiB rule (restored_test < m_test).

gfx1250 validation (on @newling's manual-validation question):

• This PR alone passed gfx1250 once: 40934683157 (10-30, same pool code as the head): 4 shards (17,404–17,407 passed each) and the large groups, no kill. One run proves little, since pre-stack kills came in bursts.
• #14609/#14610 never triggered the workflow, so I dispatched it on users/tony-davis/AIHPBLAS-5352-memory-preflight (the full stack): 41582794047 passed hipBLASLt twice (attempts 1 and 2, 12–13 min each): 4 shards (17,421–17,427 passed each) and all 6 large groups, no kill and no memory skips (no [ MEMORY ] list). A few more runs would make it convincing.
• A gfx1250 counterpart to your large-case table: the 122.6 GiB above vs your 82 gfx90a cases at 81.9 → 49.2 GiB with this PR. Rerunning the 93 in one process with the stack would show whether the large groups can merge.
• For per-shard peaks, /usr/bin/time -v needs apt-get install time in the CI image.

Peaks with the stack, and the shared-headroom preflight, are on #14610.

[tony-davis]

Thanks @newling and @archana-ramalingam ! I took all of it, and the stack is pushed:

• Environment restore (James): cherry-picked your 5b2051dd795 into fix(hipblaslt): make the test client's host-memory query respect container limits #14609 and 3afe2cb098f into fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610, with you as author. The memory suites now pass with HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT=4G and HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 set for the whole run.
• Zero vs. unavailable (James): cherry-picked e8a05bf78ac into fix(hipblaslt): make the test client's host-memory query respect container limits #14609, so a budget of 0 now makes the preflight skip.
• Swizzle temporaries (James): checked in this PR (59065b8449c) and counted in fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610's estimate. I couldn't find a way to make just that one allocation fail in a test (the preflight trips first), so the coverage there is the shared helper's test.
• Small buffers that are 2x a request (Archana): this PR now releases a pooled buffer under 1 GiB once 100 tests in a row haven't reused it (64e1afa2974). smoke+quick on gfx90a: 59.3 s and 2.18 GiB peak, vs. 62-66 s and 2.22 GiB before.
• Wording on HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT (James): in fix(hipblaslt): make the test client's host-memory query respect container limits #14609. It caps what each query reports; it's not a total and it reserves nothing.

On gfx1250: Archana, thanks for dispatching 41582794047! I think that's our best evidence so far (two clean passes of the full stack, with the large phase still split out). I agree with you both that removing the lock should be its own PR, run with HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 and per-shard peaks over a few days, since kills come in bursts.

James, agreed on revisiting allocation and caching as follow-up. I filed AIHPBLAS-5451 for it.

[Copilot]

🟡 Changes recommended

Unresolved allocation-handling issues can cause Windows test failures, falsely successful benchmark exits, and leaked HIP resources.

3 open issues

• Skip pool-reuse test when Windows policy clears the pool · New
• Propagate host allocation failures as a nonzero benchmark exit · New
• Release HIP resources on allocation-failure skips · New

What changed in this PR

This PR reduces pinned host-memory retention in hipBLASLt’s test client and detects failed host allocations before use. It is the first step in the memory-safety stack; container-aware limits and broader memory preflight checks follow separately. The hipBLASLt library is unchanged.

Changes:

• Release unused large host buffers between tests, age out smaller buffers, and remove extra capacity from large host requests.
• Check host allocations before matmul initialization and swizzling.
• Add six regression tests and register them in the standard test tier.

File	Description
projects/​hipblaslt/​TESTING.md	Documents the host-memory regression suite.
projects/​hipblaslt/​clients/​tests/​src/​host_memory_gtest.cpp	Tests retention, reuse, allocation sizing, and failure reporting.
projects/​hipblaslt/​clients/​tests/​src/​hipblaslt_gtest_main.cpp	Registers the end-of-test pool cleanup listener.
projects/​hipblaslt/​clients/​tests/​src/​CMakeLists.txt	Includes the new tests in the test executable.
projects/​hipblaslt/​clients/​common/​include/​testing_matmul.hpp	Checks main and temporary host-buffer allocations.
projects/​hipblaslt/​clients/​common/​include/​hipBuffer.hpp	Exposes allocation status and requested-size diagnostics.
projects/​hipblaslt/​clients/​common/​include/​hipblaslt_test.hpp	Adds allocation-failure handling for tests and benchmarks.
projects/​hipblaslt/​clients/​common/​include/​d_vector.hpp	Implements host-buffer aging, reclamation, and large-request sizing.

_{Review effort: Balanced}

[peymanr]

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.)

#		Finding	Where
1	🟠	Large-host-buffer classification mismatch pins ~0.2 GiB extra: this PR split slack logic between get() and release	projects/hipblaslt/clients/common/include/d_vector.hpp:365
2	🟠	Large-miss release frees buffers the current test still reuses, causing extra pin cycles: this PR added it	projects/hipblaslt/clients/common/include/d_vector.hpp:338
3	🟠	Host OOM shows as gtest SKIPPED, not in hipBLASLt's '[ SKIPPED ] N tests' line: this PR used GTEST_SKIP	projects/hipblaslt/clients/common/include/hipblaslt_test.hpp:97
4	🟠	HostMemory_pre_checkin fails instead of skipping on a memory-starved host: this PR asserts large pinned allocations succeed	projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp:27

1. 🟠 Should fix: Large-host-buffer classification mismatch pins ~0.2 GiB extra: this PR split slack logic between get() and release

projects/hipblaslt/clients/common/include/d_vector.hpp:365

get() skips the 20% slack only when the request is >= kHostPoolRetainLimit (or a quarter of device memory), but release_from_earlier_tests classifies by capacity(). So a pinned host request of about 853 MiB to 1 GiB gets about 1.2x capacity, above 1 GiB. It is then treated as a large buffer, released after one test that does not reuse it, as every buffer of 1 GiB or more is by design. The effect is up to about 0.2 GiB of extra pinned memory and a mismatch with TESTING.md and the comments ('large buffers get no 20% reuse slack'). A missed request in this band also skips the large-miss early release of buffers left by earlier tests. Deciding the slack with bytes * 1.2 >= kHostPoolRetainLimit (or capping capacity below the limit) would make the two classifications agree.

Who is affected: hipblaslt-test runs whose host operand buffers fall between about 853 MiB and 1 GiB, on GPUs with more than about 4 GiB of device memory.

Checked against the code by an independent pass.

2. 🟠 Should fix: Large-miss release frees buffers the current test still reuses, causing extra pin cycles: this PR added it

projects/hipblaslt/clients/common/include/d_vector.hpp:338

On a host-pool miss of at least kHostPoolRetainLimit (1 GiB), get() calls release_from_earlier_tests and frees every large idle buffer earlier tests restored. That includes one a later Get in the same test would have reused (e.g. B/C/D keep the previous shape while A grows). That buffer is then pinned again with hipHostMalloc, an extra free/pin cycle of multi-GiB pinned memory. Peak pinned memory is not raised, because the buffers are freed before the new allocation. Large leftovers are also bounded: HostPoolReleaseListener::OnTestEnd drops them after any test that did not return them, so a following sub-1 GiB test holds them for at most that one test.

Who is affected: hipblaslt-test runs where consecutive tests share some multi-GiB host operands but one operand grows past the cached size.

Checked against the code by an independent pass.

3. 🟠 Should fix: Host OOM shows as gtest SKIPPED, not in hipBLASLt's '[ SKIPPED ] N tests' line: this PR used GTEST_SKIP

projects/hipblaslt/clients/common/include/hipblaslt_test.hpp:97

CHECK_HOST_ALLOCATION reports a failed pinned host allocation with GTEST_SKIP, while CHECK_DEVICE_ALLOCATION reports device OOM as SUCCEED() << LIMITED_MEMORY_STRING. The two OOM cases therefore use different channels. Host skips are not added to the custom skipped_tests counter of ConfigurableEventListener, but gtest's own printer still counts and lists them as [ SKIPPED ] with the 'could not allocate ... GiB of pinned host memory' reason. They are visible and distinguishable, just reported differently from device OOM. Unifying the two channels is optional consistency cleanup.

Who is affected: hipblaslt-test runs on hosts where a large pinned host allocation fails; the result is reported as a gtest skip rather than the custom limited-memory pass.

Checked against the code by an independent pass.

Related past reviews: #14025 on projects/hipblaslt/clients/common/include/testing_matmul.hpp, #14127 on projects/hipsparse/clients/include/hipsparse_local_dnvec.hpp, #5740 on projects/hipblaslt/clients/tests/src/hipblaslt_gtest_main.cpp, #9122 on projects/hipblaslt/clients/common/include/testing_matmul.hpp

4. 🟠 Should fix: HostMemory_pre_checkin fails instead of skipping on a memory-starved host: this PR asserts large pinned allocations succeed

projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp:27

The new HostMemory_pre_checkin tests ASSERT_NE/ASSERT_EQ that pinned allocations of kHostPoolRetainLimit (1 GiB) and up succeed; the miss test holds about 2.1 GiB at once. On Linux, h_memory refuses any request larger than sysinfo freeram, which excludes page cache, so a busy shared host fails these tests instead of skipping them. The matmul tests in this PR skip such failures with CHECK_HOST_ALLOCATION. Consider GTEST_SKIP when the large Get returns null, for consistency.

Who is affected: Linux CI hosts running *pre_checkin* while other processes leave less than about 1-2 GiB of freeram.

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 5a9c5c66; findings checked and posted by @peymanr; points already raised on this PR are left out.}

[tony-davis]

Thanks! Replies by number:

1. Agreed, fixed in d914613. A new host_pool_capacity() gives no slack when the 20% would reach the 1 GiB limit, and both get() and the preflight estimate in fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610 use it. A test pins a 0.9 GiB buffer and checks its capacity; the old rule gave it 1.08 GiB.
2. Accurate, but deliberate. Freeing before pinning is what keeps peak pinned memory down, which is the point of this PR, and you're right that it doesn't raise the peak. We measured the large cases at the same wall time as develop with this policy, so I think the extra pin cycle is worth it. No change.
3. fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610 handles this: device OOM and host OOM both go through SKIP_FOR_MEMORY, and the run lists them under [ MEMORY ]. No change here.
4. Agreed, fixed in d914613. The pool tests now skip, naming the buffer they could not pin. In a 1.5 GiB cgroup the miss test skips instead of failing.