fix(hipblaslt): run the memory preflight for every matmul case and list memory skips by tony-davis · Pull Request #14610 · ROCm/rocm-libraries

tony-davis · 2026-10-28T20:40:22Z

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 each host buffer as large as the pool pins it (host_pool_capacity(), with the 20% extra below 1 GiB), W4A16's float copy of A, and in pointer-array mode the buffers each batch has (A, B, C, two copies of D and of the bias) rather than the strided path's. Memory held only for a while counts as the larger of two: the swizzle's scratch arrays (up to three times a swizzled operand), or the CPU reference's float or double copies of one batch of A and B, and of D whenever cblas_gemm computes in another type than the D it is given (int32 computes in double and f16 in float).
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), and cases that run on several threads at once, which share the process's peak.
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 third of four stacked PRs for AIHPBLAS-5352 and depends on #14609, which depends on #14608. #14982 builds on it to drop the host-wide lock in the gfx1250 workflow.

Test Plan

MatmulMemory_pre_checkin (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.
Copilot's review found four gaps in the estimate. I checked the two that gfx90a can run with a temporary test (not committed: the gtest data skips i8 to i32 on gfx11, gfx12 and gfx90a, and a case with no solution fails): an i8 to i32 8192 x 8192 x 16 case, and an f32 8192 x 8192 x 16 case with two pointer-array batches. W4A16 and the swizzle need a newer GPU and are from reading the code.
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 ci(gfx1250): run the large-memory hipBLASLt tests one job at a time #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 fix(hipblaslt): keep large pinned host buffers only while tests reuse them #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".
The i8 to i32 case was estimated at 1.80 GiB and used 2.40 GiB; it is now estimated at 2.30 GiB. The pointer-array case was estimated at 3.60 GiB and used 1.81 GiB; it is now estimated at 1.80 GiB.
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. 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. Her run on the rebased stack, 41717825380, was clean too. Shards still check memory one moment at a time and reserve nothing (fix 6 in the ticket), so the preflight lowers the chance of a kill rather than ruling it out. I think that's the right place to stop for this stack: the container caps keep a kill inside one job.
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 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.

Work tracking: AIHPBLAS-5352, fixes 2 (preflight for every case, one estimator) and 5 (visible memory skips) of its list.
Depends on fix(hipblaslt): make the test client's host-memory query respect container limits #14609 and fix(hipblaslt): keep large pinned host buffers only while tests reuse them #14608.
ci(gfx1250): run the large-memory hipBLASLt tests one job at a time #14545 runs the large cases after the shards behind a host-wide lock, and feat(ci): cap the gfx1250 test container's memory #14569 caps each gfx1250 container at 110 GiB. ci(gfx1250): drop the host-wide lock on the large hipBLASLt tests #14982 removes the lock; the large cases keep their own phase, after the shards.

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.

therock-pr-bot · 2026-10-28T20:46:45Z

❌ PR Check — Action Required

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	❌ Fail	Error: Check concluded with `cancelled`.
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

❌ pre-commit

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

🙋 Wish to Override Policy?

codecov · 2026-10-29T02:50:07Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@                                Coverage Diff                                @@
##           users/tony-davis/AIHPBLAS-5352-host-memory-query   #14610   +/-   ##
=================================================================================
  Coverage                                             77.65%   77.65%           
=================================================================================
  Files                                                  2905     2905           
  Lines                                                490744   490744           
  Branches                                              74315    74315           
=================================================================================
+ Hits                                                 381039   381040    +1     
+ Misses                                                84684    84683    -1     
  Partials                                              25021    25021

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

*This pull request uses carry forward flags. Click here to find out more.
see 1 file 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.

archana-ramalingam

Thanks @tony-davis!

Two gfx1250 notes (no blocker; @newling's review covers the rest):

Shards can all pass the preflight together.
- A free-memory check doesn't reserve memory against another process (newling's linked review); your description lists this as not yet measured. Each shard checks the whole container's headroom (fast_check.cpp#L2006-L2010), and gtest shards round-robin over one list, so all 4 reach the big cases together.
- On 11-02, without this stack, 9 of 12 shards in 41512870903, 41517092105 and 41519929996 died at tests 15,619–15,975 of ~17,470, around fast_check_large 16385×16383×4097 f32. With the stack, 41582794047 passed that window twice (attempts 1 and 2): 4 shards 17,421–17,427 passed each, all 6 large groups, no kill, empty [ MEMORY ].
- The device check has the same no-reservation gap. The 4 shards share one GPU, so hipMemGetInfo is device-wide; that made memory_shortfall_reclaims_idle_buffers flaky (#14177).
Lock/split removal:
- Agree with newling's review that it needs its own evidence, including no memory skips. With this stack a shortage becomes a skip, so a green run can hide them: with the 2^31/2^32 tests sharded, our 10-28 gfx1250 probe measured 84–96 GiB per shard (368 GiB total) against the 110 GiB cap. The two stack runs above kept the split.
- Please run with HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 (or check [ MEMORY ] is empty) and record per-shard peak RSS (/usr/bin/time -v, after apt-get install time in the CI image).
- Kills are bursty (none from the 10-31 host reboot until 11-02, then 3 runs in 4 h), so two clean runs show little; a few more on different days would.

tony-davis · 2026-11-03T16:31:30Z

Thanks Archana! Agreed on both points.

The shard race (fix 6 in AIHPBLAS-5352) and the device-side gap stay open. Each check is still one moment in one process, so I'd leave the large phase split out until the gfx1250 runs show whether it bites with the lock gone.
For the lock-removal PR, I'll run with HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 so a green run can't hide memory skips. I'll also record per-shard peaks (time -v, after adding time to the image) on a few different days.

The batch-offset gap from your inline comment is fixed in 23ff31ed2d7; details in the thread there.

Copilot

🟡 Changes recommended

Concurrent workers race on shared first-case and memory-measurement state, while inaccurate host-memory estimates can admit over-budget cases or unnecessarily skip valid ones.

5 open issues

What changed in this PR

This PR completes the host-memory safety stack for hipBLASLt’s test client by extending memory checks beyond fast_check cases and making memory-related skips visible. Library and kernel behavior remain unchanged.

Changes:

Runs memory preflight checks for all test-client matmul cases, including batch-offset cases, with expanded allocation estimates.
Unifies memory-shortage reporting and adds an end-of-run summary plus an option to fail on memory skips.
Adds Linux peak-memory validation, regression tests, and documentation.

File	Description
`projects/hipblaslt/TESTING.md`	Documents memory checks, skip reporting, and diagnostic flags.
`projects/hipblaslt/clients/tests/src/matmul_gtest.cpp`	Adds preflight and memory-estimate regression tests.
`projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp`	Tests memory-shortage summary formatting.
`projects/hipblaslt/clients/tests/src/hipblaslt_gtest_main.cpp`	Collects and prints memory-shortage results.
`projects/hipblaslt/clients/common/include/utility.hpp`	Removes obsolete memory-shortage message constants.
`projects/hipblaslt/clients/common/include/testing_matmul.hpp`	Broadens preflight coverage and allocation accounting.
`projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp`	Adds batch-offset memory preflight and measurement.
`projects/hipblaslt/clients/common/include/memory_estimate_check.hpp`	Adds optional Linux peak-memory validation.
`projects/hipblaslt/clients/common/include/hipblaslt_test.hpp`	Centralizes memory skips, optional failures, and summary storage.

_{Review effort: Balanced}

archana-ramalingam

Thanks @tony-davis, approving. Our feedback is addressed:

matmul_batch_offset now goes through the preflight (81234f4), with a pre_checkin test; estimate 51.55 vs 51.51 GB measured.
Small idle pool buffers are released after 100 unused tests (#14608, 8671df7).
Shard race and device-side gap stay open (fix 6, large tests kept split); lock removal gets its own PR with HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 and per-shard time -v peaks over a few days.

gfx1250 on the rebased stack (54aa7ef), 41717825380: hipBLASLt shards 17,426–17,436 passed each, all 6 large groups passed, 0 kills, no memory skips; TensileLite 216 passed / 0 failed. Clean, like the two runs on 876b837.

therock-pr-bot · 2026-11-05T16:34:22Z

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

…st memory skips The memory preflight ran only for fast_check cases, so the unit_check stress cases, which pin up to about 48 GiB of host memory on gfx90a, allocated without checking. hipblaslt-test now runs it for every matmul case; hipblaslt-bench still runs it only for fast_check. One estimate covers both: it counts the pool's 20% extra on host buffers below 1 GiB, the CPU reference's float or double copies of A and B (and of D when the output type is not the compute type), and one set of buffers per batch in pointer-array mode. Every stop for lack of memory (the preflight, a failed pinned host allocation, or a device out-of-memory error) now goes through SKIP_FOR_MEMORY: a skip whose message begins "Not enough memory:", listed under [ MEMORY ] at the end of the run. A device out-of-memory error used to be reported as a pass. HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 turns these skips into failures, and HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE=1 fails a case whose peak resident memory rose more than its estimate allows. New MatmulMemory_pre_checkin tests run the preflight and the estimate check on real cases; a HostMemory_pre_checkin test covers the summary. JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

…estimate check off Linux The MatmulMemory_pre_checkin tests called setenv and unsetenv, which Windows does not have, and the estimate test asserted a measurement that the check makes only on Linux. Use _putenv_s on Windows, and skip the estimate test where the check cannot measure. JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

…override ScopedVariable unset its variable when it left scope, so a lane that sets HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT or HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP for the whole run lost it for every later test. Restore the earlier value.

A swizzled A or B is built in a temporary pinned host buffer of the swizzled size before it is uploaded, and the estimate left it out. Count it, with the pool's 20% extra below 1 GiB, like the other host buffers. JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

matmul_batch_offset cases go through testing_matmul_batch_offset, which allocated without the preflight, so matmul_batch_offset_large (about 48 GiB of host memory) could still pin past a container's limit. Estimate its four host and device buffers, the reference D and the 128 MiB workspace bound, skip with the amounts when they do not fit, and check the estimate under HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE. JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

…reads out of the estimate check - W4A16 keeps a float copy of A, dequantized, for the CPU reference; count it. - The pointer-array path allocates A, B, C, two copies of D and two of the bias per batch, and none of the strided path's three gold buffers or E. For f32 8192 x 8192 x 16 with two batches the estimate drops from 3.60 GiB to 1.80 GiB; the case used 1.81 GiB. - The swizzle builds the layout in up to three arrays the size of the swizzled operand. That memory and the CPU reference's copies are freed before the other is made, so the larger counts. - cblas_gemm computes int32 in double and f16 in float, and copies D whenever that type differs from the D it is given. An i8 to i32 8192 x 8192 x 16 case estimated 1.80 GiB and used 2.40 GiB; it now estimates 2.30 GiB. The multithread cases run testing_matmul on several threads at once. The first-case flag of MemoryEstimateCheck is now atomic, and cases on more than one thread are not measured, since they share the process's peak. Reported by Copilot. JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

The preflight added the pool's 20% to every host buffer below kHostPoolRetainLimit. The pool now gives none when the extra would reach the limit, so the estimate uses the same host_pool_capacity(). JIRA ID : AIHPBLAS-5352 Co-authored-by: Cursor <cursoragent@cursor.com>

… 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>

…ainer limits (#14609) JIRA ID : AIHPBLAS-5352 ## Motivation Before `hipblaslt-test`, the hipBLASLt test client, pins a host buffer, it checks that enough host memory is free, so that a shortage ends in a skipped test rather than a killed process. It asked two different questions to do that: `h_memory` read `freeram` from `sysinfo()`, and the `fast_check` memory preflight read `MemAvailable` from `/proc/meminfo`. Neither looks at the memory limit of a Linux control group (cgroup), the kernel feature that containers and CI runners use to cap a job's memory, and inside a container both report the whole host. A test inside a container with a memory limit therefore sees plenty of memory, pins past the limit, and the kernel kills the process. On this PR's test host, a large `stress` case inside an 8 GiB limit is killed within 4 seconds on develop, and also with only the previous PR in this stack. ## Technical Details - New `available_host_memory_bytes()` in `clients/common/include/host_memory.hpp` and `src/host_memory.cpp`. On Linux it returns the smaller of `MemAvailable` and the room left under each cgroup memory limit, from the process's own cgroup (read from `/proc/self/cgroup`) up to the root of the hierarchy. It reads cgroup v2 (`memory.max` minus `memory.current`) and v1 (`memory.limit_in_bytes` minus `memory.usage_in_bytes`). Page cache the kernel can drop does not count as used: `active_file` and `inactive_file` in `memory.stat`, less dirty and writeback pages. A removed cgroup, which `/proc/self/cgroup` lists with a ` (deleted)` suffix, still has its limit applied. - The cgroup's path is translated through its mount in `/proc/self/mountinfo` (the cgroup2 mount, or the v1 mount that carries the memory controller, such as `cpu,memory`), and the walk never goes above that mount. Without mountinfo it falls back to `/sys/fs/cgroup`. A container may list its cgroup as `/` or by the host's path, which does not exist inside the container; the walk always checks the top it sees, which is where the container's own limit is. - `/proc/meminfo` and the cgroups each cap the answer on their own, so a full container reports 0 even when `/proc/meminfo` cannot be read. - `h_memory` (in `d_vector.hpp`) and `fast_check_memory_shortfall()` both call it on Linux. For `h_memory` this also changes the base value from `freeram` to `MemAvailable`, which counts reclaimable page cache, as the preflight already did. On this host `freeram` was 181 GiB and `MemAvailable` about 2,070 GiB, because most memory held page cache. When the query has no answer at all, `h_memory` falls back to `freeram`, as before. Inside an 8 GiB cgroup holding 6 GiB of page cache, active or inactive, a 6 GiB `hipHostMalloc` succeeded and the kernel reclaimed the cache, so counting the cache as available matches what pinning can actually get. - The managed-memory (`use_HMM`) branch of `d_memory` uses the same check, as it did with `freeram`. In these tests a managed buffer is first written on the host, so its pages are charged to the cgroup until they migrate, and the host headroom is a conservative bound. No HMM case (they are in the `full` tier) has been run under a limit yet. - Windows keeps its existing behavior at each call site: `h_memory` uses half of physical memory, and the preflight uses available physical memory. - `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT`, a new optional environment variable, caps the budget for a machine with no container limit (a development box, Windows, or bare metal), and lets a test force a memory skip. It can only lower the budget: a full container stays full, and when the system cannot be read the limit is the budget. The value is a byte count with an optional K, M, G or T suffix in powers of 1024, such as `64G`; an unreadable value is ignored with one warning. CI does not need to set it. Geo Min suggested it on #14569. - `available_host_memory_bytes()` returns `std::optional`: nothing when the system cannot be read, and 0 when no memory is left, so a full container or a limit of 0 makes the preflight skip instead of letting the case run (James Newling's commit). - The limit caps the free memory each query reports. It does not limit what a process allocates in total, and it reserves nothing against other processes; `host_memory.hpp` and `TESTING.md` say so. - The file paths the query reads are parameters (`HostMemorySources`), so tests can point it at temporary files. This is the second of four stacked PRs for AIHPBLAS-5352 and depends on #14608. With #14608, a pinned allocation that does not fit skips the test with the requested size; this PR makes the container's limit part of what "fits" means. ## Test Plan - `HostMemory_pre_checkin.query_*` (13 new tests in `clients/tests/src/host_memory_gtest.cpp`). Twelve build a temporary `/proc/meminfo`, `/proc/self/cgroup`, `/proc/self/mountinfo` and cgroup tree and check the result for: no cgroup; a v2 limit below `MemAvailable`, with page cache subtracted; a v2 cgroup with no limit (`max`); a parent cgroup with a tighter limit than the process's own; a container listing a host path, with its limit at the top, also when the mount's root is another cgroup; a mount of a subtree (`/runner`) with a tighter limit below it; a v1 memory controller mounted with cpu; cgroup v1, limited and unlimited; a full container whose usage is mostly active page cache, with and without dirty pages; a removed cgroup; a full container with no `/proc/meminfo`; and missing, unreadable or over-limit files. The thirteenth only checks that the real files on the test host are readable and that its cgroups do not raise `MemAvailable`. - `HostMemory_pre_checkin.no_host_memory_left_makes_the_preflight_skip` checks a budget of 0. The test helpers restore any `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` the run started with, and the fake-host tests clear it while they run (James Newling's commit). - `HostMemory_pre_checkin.limit_*` (4 new tests): the size format, including suffixes, overflow and malformed values; that the variable lowers the budget, never raises it, and leaves a full container at 0; that it is the budget when `/proc/meminfo` cannot be read; and that `fast_check_memory_shortfall()` reports "needs 2 GiB of host memory and 1 GiB is available" under a 1 GiB limit (skipped on a runner with less than 2 GiB free). The pool tests clear the limit while they run, since pinning past it is what they test. - To confirm the tests can fail, I built temporary variants. One stopped the cgroup walk at the process's own cgroup and counted page cache as used, and it failed 3 of the first 8 `query_*` tests. Another ignored cgroups altogether, and it failed 5; the three tests that passed under both check behavior with no limit set. A variant that ignored the environment variable failed 3 of the 4 `limit_*` tests (all but the format test), and one that let the variable replace a full container's 0 failed the cap test. The later fixes each have a test that a variant without the fix failed: no mount translation (2 tests), no v1 mount lookup, reading the cgroups only with `/proc/meminfo`, and walking above the mount. - On the real host: one `stress` case inside a transient 8 GiB cgroup (`systemd-run --user --scope -p MemoryMax=8G -p MemorySwapMax=0`), on develop, on the previous PR alone, and on this PR. - Adjacent: the rest of `HostMemory_pre_checkin`, `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin` (the preflight's suites), and the `smoke` and `quick` tiers. `hipblaslt-bench` was built but not run. ## Test Result One MI210 (gfx90a), with this branch's clients built against the 10-25 TheRock nightly hipBLASLt gfx90a device library. - `stress_matmul_size_t_index_overflow_f32_..._NN_2048_2048_2048_2_2048_2048_4_1049100_1049100_1` inside an 8 GiB limit: develop and the previous PR alone were killed (exit code 137) within 4 seconds. With this PR the test was skipped with "could not allocate 8 GiB of pinned host memory" and the process exited 0. - `HostMemory_pre_checkin`, `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin`: all passed with no limit set and with `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` at 4G, 1G and 64M for the whole run. Inside a 1.5 GiB cgroup, the preflight-limit test skips instead of failing. - On the same host, 4 GiB pinned with `hipHostMalloc` inside an 8 GiB cgroup raised its `memory.current` from 0.01 GiB to 4.03 GiB, so the container's headroom already counts the clients' pinned buffers. - `smoke` + `quick`: 9,169 passed and 2 skipped (`GroupedMatmul.smoke_Fp16NN` and `smoke_Bf16NN`, which skip on this GPU), the same as develop, in 62.3 s. - Not covered: Windows (unchanged by design), and a real cgroup v1 host. The v1 path is checked only against temporary files. - 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 Low (2/5). Every pinned host allocation in the test and bench clients now consults the new query, so a wrong answer would skip tests that should run, or let one run that should skip. Outside a memory-limited cgroup the answer is `MemAvailable`, which the preflight already used. The hipBLASLt library and device behavior are unchanged. ## Related - Work tracking: [AIHPBLAS-5352](https://amd-hub.atlassian.net/browse/AIHPBLAS-5352), fix 3 (container-aware memory budget, with the override) of its list. - Depends on #14608. Next in the stack: #14610. - #14569 caps each gfx1250 test container at 110 GiB, which is the budget this query reads there. ## 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. ## Flags / Guardrails `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` (optional, unset by default) caps the host-memory budget. It is 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>

tony-davis added this pull request to stack #14611 October 28, 2026 20:40

This was referenced Oct 28, 2026

fix(hipblaslt): keep large pinned host buffers only while tests reuse them #14608

Merged

fix(hipblaslt): make the test client's host-memory query respect container limits #14609

Merged

github-actions Bot added documentation project: hipblaslt labels Oct 28, 2026

assistant-librarian Bot added the organization: ROCm label Oct 28, 2026

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 7e45ca9 to 3e771a2 Compare October 29, 2026 15:23

tony-davis mentioned this pull request Oct 30, 2026

test(hipblaslt): make memory_shortfall_reclaims_idle_buffers deterministic #14177

Merged

1 task done

tony-davis marked this pull request as ready for review October 30, 2026 14:35

tony-davis requested review from a team as code owners October 30, 2026 14:35

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 3e771a2 to 876b837 Compare October 30, 2026 14:45

tony-davis requested review from archana-ramalingam and newling October 30, 2026 14:47

newling approved these changes Nov 2, 2026

View reviewed changes

archana-ramalingam reviewed Nov 3, 2026

View reviewed changes

Comment thread projects/hipblaslt/clients/tests/src/matmul_gtest.cpp Resolved

Copilot AI balanced review requested due to automatic review settings November 3, 2026 16:28

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 876b837 to 23ff31e Compare November 3, 2026 16:28

Copilot started reviewing on behalf of tony-davis November 3, 2026 16:29

View session

tony-davis mentioned this pull request Nov 3, 2026

ci(gfx1250): drop the host-wide lock on the large hipBLASLt tests #14982

Open

1 task done

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 23ff31e to cc4d3d5 Compare November 3, 2026 16:46

Copilot stopped reviewing on behalf of tony-davis due to an error November 3, 2026 16:49
Copilot had to stop work due to a timeout.

Copilot started reviewing on behalf of tony-davis November 3, 2026 16:50

View session

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from cc4d3d5 to 2bca6b5 Compare November 3, 2026 16:59

Copilot AI reviewed Nov 3, 2026

View reviewed changes

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 2bca6b5 to f38aba1 Compare November 3, 2026 19:46

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from f38aba1 to 54aa7ef Compare November 3, 2026 21:00

tony-davis requested a review from archana-ramalingam November 3, 2026 21:03

tony-davis enabled stacked auto-merge (squash) November 3, 2026 22:06

archana-ramalingam approved these changes Nov 3, 2026

View reviewed changes

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 54aa7ef to eab9f44 Compare November 4, 2026 16:18

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from eab9f44 to 6e5d000 Compare November 5, 2026 16:17

tony-davis and others added 7 commits November 5, 2026 13:19

tony-davis force-pushed the users/tony-davis/AIHPBLAS-5352-memory-preflight branch from 6e5d000 to c5dc0c1 Compare November 5, 2026 18:20

tony-davis merged commit 54a6037 into develop Nov 6, 2026
142 of 155 checks passed

tony-davis deleted the users/tony-davis/AIHPBLAS-5352-memory-preflight branch November 6, 2026 06:10

Navigation Menu

Repository navigation

fix(hipblaslt): run the memory preflight for every matmul case and list memory skips - #14610

fix(hipblaslt): run the memory preflight for every matmul case and list memory skips#14610
tony-davis merged 7 commits into
users/tony-davis/AIHPBLAS-5352-host-memory-queryfrom
users/tony-davis/AIHPBLAS-5352-memory-preflight

tony-davis commented Oct 28, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 28, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Oct 29, 2026 •

edited

Loading

Uh oh!

archana-ramalingam left a comment •

edited

Loading

Uh oh!

Uh oh!

tony-davis commented Nov 3, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

archana-ramalingam left a comment

Uh oh!

therock-pr-bot Bot commented Nov 5, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

tony-davis commented Oct 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Risk level

Related

Device / Architecture Coverage

Flags / Guardrails

Uh oh!

therock-pr-bot Bot commented Oct 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

❌ PR Check — Action Required

Uh oh!

codecov Bot commented Oct 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

archana-ramalingam left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

tony-davis commented Nov 3, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

archana-ramalingam left a comment

Choose a reason for hiding this comment

Uh oh!

therock-pr-bot Bot commented Nov 5, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

tony-davis commented Oct 28, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 28, 2026 •

edited

Loading

codecov Bot commented Oct 29, 2026 •

edited

Loading

archana-ramalingam left a comment •

edited

Loading