tony-davis · Nov 6, 2026 · Oct 28, 2026 · Oct 28, 2026 · Nov 2, 2026 · Nov 2, 2026
@@ -259,7 +259,7 @@ in [`clients/tests/src/`](clients/tests/src/). The last three rows are in
 | `ulp_gtest.cpp` | `UlpMantissaBits`, `UlpAsDouble`, `UlpDistance`, `UlpAccumulate`, `UlpCheckGeneral` | Units-in-the-last-place error-measurement helpers, pure math with no GPU dependency |
 | `unit_gtest.cpp` | `UnitCheckIdentical` | Exact-storage checks used to skip element-by-element output comparisons when the addressed values have identical bytes |
 | `fast_check_gtest.cpp` | `FastCheck_pre_checkin`, `FastCheckDevice_pre_checkin` | The `fast_check` verifier: fault injection into host-computed results and the per-solution failure summary (`FastCheck_pre_checkin`, no GPU), and the device padding, poison and probe-sum kernels plus 4 GiB-boundary buffer placement on small buffers (`FastCheckDevice_pre_checkin`, no hipBLASLt kernels) |
-| `host_memory_gtest.cpp` | `HostMemory_pre_checkin` | How the client holds pinned host memory: an idle buffer of 1 GiB or more stays cached only while consecutive tests reuse it (released after a test that does not, or when a test of another shape needs new memory), smaller buffers stay cached, large buffers get no 20% reuse slack, and a failed host allocation is reported instead of used. Pins at most about 3.3 GiB; no hipBLASLt kernels |
+| `host_memory_gtest.cpp` | `HostMemory_pre_checkin` | How the client holds pinned host memory: an idle buffer of 1 GiB or more stays cached only while consecutive tests reuse it (released after a test that does not, or when a test of another shape needs new memory), a smaller one until 100 consecutive tests have not reused it, large buffers get no 20% reuse slack, and a failed host allocation is reported instead of used. The `query_*` tests point the host-memory query (`host_memory.hpp`) at temporary copies of `/proc/meminfo` and the cgroup files, and check that it takes the smaller of `MemAvailable` and the tightest container memory limit (cgroup v1 or v2, up to the root). The `limit_*` tests check that `HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` caps that budget and never raises it, and that it makes the memory preflight skip. Pins at most about 3.3 GiB; no hipBLASLt kernels |
 | `arch_revision_gtest.cpp` | `ArchRevisionSmoke` | ASIC-revision-to-library-arch mapping, a pure function reached by relative `#include` of an internal header (see point 1 below) |
 | `cotenant_gtest.cpp` | `CotenantOptions`, `CotenantLds`, `CotenantScoped` | In-process cotenant helper (Linux only): option validation and LDS-reservation sizing (host-only), plus `Scoped` start/stop and argument-error cases that need a GPU and skip without one |
 | `unit/parser_gtest.cpp` | `HostUnitInitParser`, `HostUnitActivationParser`, `HostUnitBiasSourceParser` | The client's string parsers for initialization, activation and bias source, with round-trips through their `*_to_string` helpers |
@@ -344,6 +344,12 @@ kernels on CPU and hands an artifact to a GPU stage for the run phase.
 | `full` | comprehensive + HMM (needs a managed-memory capable host) | up to 24 h | 86400 s |
 | `stress` | stress only: size-threshold and address-overflow cases needing up to about 33 GiB of device memory (the 2^32 C and D batch-stride case); for large-memory runners or a weekly run. The fast_check cases among them skip in a run with no gtest filter, such as TheRock's | ~10 min on gfx90a for the fast_check cases, longer with more solutions; the older stress cases add their host-reference time | 14400 s |
 
+The clients check host memory before pinning a host buffer, against the smaller of `MemAvailable`
+and the headroom under the container's memory limit. On a machine with no container limit, set
+`HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT` (a byte count such as `64G`) to cap that budget; it can only
+lower it. The cap applies to each check on its own: it does not limit what a process allocates in
+total, and it reserves nothing against other test processes on the same host.
+
 All tiers exclude `*known_bug*`. Multi-GPU coverage is narrow: the `*multi_gpu*` cases are built
 only with the fused GEMM plus all-to-all feature (`HIPBLASLT_ENABLE_GEMM_A2A_FUSION`, off by
 default), and they run on a two-GPU gfx950 node in Math CI's `precheckin` on pull requests.

@@ -30,6 +30,7 @@
 #include "hipblaslt_arguments.hpp"
 #include "hipblaslt_init.hpp"
 #include "hipblaslt_test.hpp"
+#include "host_memory.hpp"
 #include "singletons.hpp"
 #include <algorithm>
 #include <cinttypes>
@@ -38,7 +39,7 @@
 
 #ifdef _WIN32
 #include <windows.h>
-#else
+#elif defined(__linux__)
 #include <sys/sysinfo.h>
 #endif
 
@@ -102,25 +103,23 @@ class hip_memory
     size_t get_available_host_memory()
     {
 #ifdef __linux__
+        if(std::optional<size_t> available = available_host_memory_bytes())
+            return *available;
+        // Without /proc/meminfo or a container limit, the free memory the kernel reports.
         struct sysinfo info;
         if(sysinfo(&info) == 0)
-        {
-            // In the linux system, the host memory(pinned memory)'s capacity is the same as the free memory
-            return info.freeram;
-        }
-        else
-        {
-            hipblaslt_cerr << "Error getting available host memory" << std::endl;
-            return 0;
-        }
+            return size_t(info.freeram) * info.mem_unit;
+        hipblaslt_cerr << "Error getting available host memory" << std::endl;
+        return 0;
 #elif defined(_WIN32)
         MEMORYSTATUSEX memStatus = {};
         memStatus.dwLength = sizeof(memStatus);
         if(GlobalMemoryStatusEx(&memStatus))
         {
             // In the windows system, the host memory's capacity is the half of the physical memory
             // And previous allocation of host memory is got cleared before
-            return memStatus.ullTotalPhys / 2;
+            return std::min(size_t(memStatus.ullTotalPhys / 2),
+                            host_memory_limit_override().value_or(SIZE_MAX));
         }
         else
         {

@@ -160,9 +160,10 @@ uint64_t fast_check_probe(uint64_t seed, uint64_t index);
 
 // Returns why a case that needs device_bytes of device memory and host_bytes of host memory
 // cannot run here, or an empty string when both fit. Device memory is what hipMemGetInfo reports
-// free; host memory is MemAvailable from /proc/meminfo, which counts reclaimable page cache,
-// unlike sysinfo's freeram. Before reporting a shortfall, releases idle client-pool buffers and
-// queries again, since those buffers otherwise hide memory available to the next case.
+// free; host memory is available_host_memory_bytes() (host_memory.hpp), the smaller of
+// MemAvailable and the container's memory-limit headroom. Before reporting a shortfall, releases
+// idle client-pool buffers and queries again, since those buffers otherwise hide memory
+// available to the next case.
 std::string fast_check_memory_shortfall(size_t device_bytes, size_t host_bytes);
 
 // As above, with the free device memory read by free_device_bytes (false: unknown) instead of

@@ -0,0 +1,37 @@
+// Copyright Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <cstddef>
+#include <optional>
+#include <string>
+
+// The files the Linux host-memory query reads. Tests point them at temporary files.
+struct HostMemorySources
+{
+    std::string meminfo     = "/proc/meminfo";
+    std::string self_cgroup = "/proc/self/cgroup";
+    std::string mountinfo   = "/proc/self/mountinfo";
+    // Where the cgroup hierarchy is assumed to be when mountinfo does not say.
+    std::string cgroup_root = "/sys/fs/cgroup";
+};
+
+// Host memory this process can still allocate, in bytes: the smaller of MemAvailable and the
+// headroom under each cgroup memory limit from the process's cgroup up to the root. Inside a
+// container /proc/meminfo describes the whole host, so the cgroup limit is often the smaller.
+// On Windows, the available physical memory. Either way, HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT caps
+// the result. Returns nothing when the system cannot be read and no limit is set; 0 means that
+// no memory is left, for example in a full container.
+std::optional<size_t> available_host_memory_bytes(const HostMemorySources& sources = {});
+
+// HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT, when set, caps the free host memory that each query reports
+// to the test and bench clients: on a machine with no cgroup limit, or to make a test skip. It can
+// only lower that figure. It is not a limit on what the process allocates in total, and it
+// reserves nothing against other processes: each check compares one case or one buffer with it.
+// The value is a byte count with an optional K, M, G or T suffix in powers of 1024, such as 64G.
+// An unset variable gives nothing; an unreadable value gives nothing and a warning.
+std::optional<size_t> host_memory_limit_override();
+
+// The byte count `text` names in the format above, or nothing when it names none.
+std::optional<size_t> parse_host_memory_size(const char* text);
@@ -15,6 +15,7 @@ target_sources(hipblaslt-clients-common
         "${CMAKE_CURRENT_SOURCE_DIR}/hipblaslt_init_device.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/fast_check.cpp"
         "${CMAKE_CURRENT_SOURCE_DIR}/relaunch_compare.cpp"
+        "${CMAKE_CURRENT_SOURCE_DIR}/host_memory.cpp"
 )
 
 if(HIPBLASLT_ENABLE_BLIS)

@@ -3,27 +3,20 @@
 
 #include "fast_check.hpp"
 #include "d_vector.hpp"
+#include "host_memory.hpp"
 
 #include <hip/hip_runtime.h>
 
 #include <algorithm>
 #include <atomic>
 #include <cmath>
 #include <cstring>
-#include <fstream>
 #include <limits>
 #include <memory>
 #include <sstream>
 #include <type_traits>
 #include <vector>
 
-#ifdef _WIN32
-#ifndef NOMINMAX
-#define NOMINMAX
-#endif
-#include <windows.h>
-#endif
-
 namespace
 {
     constexpr uint64_t P = kFastCheckModulus;
@@ -1996,30 +1989,6 @@ FastCheckResult fast_check_bias_gradient(const FastCheckProblem& p,
     return result;
 }
 
-namespace
-{
-    // Host memory the operating system reports as available, or 0 when it cannot tell.
-    size_t available_host_bytes()
-    {
-#ifdef _WIN32
-        MEMORYSTATUSEX status = {};
-        status.dwLength       = sizeof(status);
-        return GlobalMemoryStatusEx(&status) ? size_t(status.ullAvailPhys) : 0;
-#else
-        std::ifstream meminfo("/proc/meminfo");
-        std::string   key;
-        size_t        kib = 0;
-        while(meminfo >> key >> kib)
-        {
-            if(key == "MemAvailable:")
-                return kib * 1024;
-            meminfo.ignore(std::numeric_limits<std::streamsize>::max(), '\n');
-        }
-        return 0;
-#endif
-    }
-}
-
 std::string fast_check_memory_shortfall(size_t device_bytes, size_t host_bytes)
 {
     return fast_check_memory_shortfall(device_bytes, host_bytes, [](size_t& free) {
@@ -2044,9 +2013,10 @@ std::string fast_check_memory_shortfall(size_t
             return "this case needs " + gib(device_bytes) + " of device memory and "
                    + gib(free_bytes) + " is free";
 
-        const size_t available = available_host_bytes();
-        if(available > 0 && host_bytes > available)
-            return "this case needs " + gib(host_bytes) + " of host memory and " + gib(available)
+        // When the system cannot report its memory, let the case run.
+        const std::optional<size_t> available = available_host_memory_bytes();
+        if(available && host_bytes > *available)
+            return "this case needs " + gib(host_bytes) + " of host memory and " + gib(*available)
                    + " is available";
         return {};
     };