[tony-davis]

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 feat(ci): cap the gfx1250 test container's memory #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

• 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, fix 3 (container-aware memory budget, with the override) of its list.
• Depends on fix(hipblaslt): keep large pinned host buffers only while tests reuse them #14608. Next in the stack: fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610.
• feat(ci): cap the gfx1250 test container's memory #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.

[therock-pr-bot]

❌ 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?

[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                                 @@
##             users/tony-davis/AIHPBLAS-5352-host-pool-release   #14609   +/-   ##
===================================================================================
  Coverage                                                    ?   77.65%           
===================================================================================
  Files                                                       ?     2905           
  Lines                                                       ?   490744           
  Branches                                                    ?    74315           
===================================================================================
  Hits                                                        ?   381039           
  Misses                                                      ?    84684           
  Partials                                                    ?    25021           

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% <ø> (?)
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.

🚀 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]

Looks good!

[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	🟠	Every pinned host allocation fails when host memory is unknown: this PR maps 'unknown' to 0	projects/hipblaslt/clients/common/include/d_vector.hpp:94

1. 🟠 Should fix: Every pinned host allocation fails when host memory is unknown: this PR maps 'unknown' to 0

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

On Linux, get_available_host_memory now returns 0 when available_host_memory_bytes() is nullopt. That happens when /proc/meminfo has no readable MemAvailable:, no cgroup memory limit is found and HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT is unset. The h_memory and HMM d_memory constructors then fail every pinned or managed allocation and print an error each time. fast_check_memory_shortfall treats the same nullopt as unknown and lets the case run, so the two callers disagree. Returning a permissive value (e.g. SIZE_MAX, or falling back to sysinfo) would match the preflight. This only arises on kernels older than 3.14 or in sandboxes that strip MemAvailable.

Who is affected: Test and bench clients on Linux hosts whose /proc/meminfo lacks MemAvailable (pre-3.14 kernels or unusual sandboxes) with no cgroup memory limit and no override set.

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

[tony-davis]

Thanks, agreed. Fixed in 2bbfef3: when the query has no answer, get_available_host_memory falls back to the free memory sysinfo reports, as develop did, and prints the error only if that fails too. (Not unit-tested, since it reads the real host; it only changes the no-answer path.)