tony-davis · Nov 6, 2026 · Oct 28, 2026 · Oct 29, 2026 · Nov 3, 2026 · Nov 3, 2026
@@ -259,6 +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 |
 | `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 |

@@ -34,6 +34,7 @@
 #include <algorithm>
 #include <cinttypes>
 #include <hipblaslt/hipblaslt.h>
+#include <type_traits>
 
 #ifdef _WIN32
 #include <windows.h>
@@ -44,6 +45,27 @@
 #define MEM_MAX_GUARD_PAD 8192
 #define MAX_DTYPE_SIZE sizeof(double)
 
+// In hipblaslt-test, an idle pinned host buffer at least this large stays cached only while
+// consecutive tests keep reusing it. The kernel cannot reclaim pinned memory, so a cached
+// multi-GiB buffer starves the other test processes on the host; smaller buffers stay cached
+// because nearly all tests reuse them.
+inline constexpr size_t kHostPoolRetainLimit = size_t(1) << 30;
+
+// The capacity the pool pins for a new host buffer of `bytes`: 20% extra so that a slightly
+// larger request can reuse it, unless that would reach kHostPoolRetainLimit. A buffer at the
+// limit is reused only by the same shape in the next test, so the extra would only add pinned
+// memory, and it would make a smaller request's buffer count as large.
+inline size_t host_pool_capacity(size_t bytes)
+{
+    const size_t padded = static_cast<size_t>(bytes * 1.2);
+    return padded < kHostPoolRetainLimit ? padded : bytes;
+}
+
+// A smaller idle pinned host buffer is released once this many consecutive tests have not
+// reused it. The pool reuses a buffer only for requests between half its size and its size, so
+// without this a run of shrinking requests would keep every earlier buffer pinned.
+inline constexpr uint64_t kHostPoolSmallIdleTests = 100;
+
 /* ============================================================================================ */
 /*! \brief  (abstract class) wrapper around a pointer to hip device or pinned host memory, including allocation size in bytes */
 class hip_memory
@@ -74,6 +96,9 @@ class hip_memory
         return capacity() < s;
     }
 
+    // The memory_pool test count when the buffer was last returned to the pool.
+    uint64_t restored_test = 0;
+
     size_t get_available_host_memory()
     {
 #ifdef __linux__
@@ -236,6 +261,20 @@ class memory_pool
         Instance().restore(dm);
     }
 
+    // True on a Windows system with not enough system memory for pool management: every request
+    // frees the idle buffers before it allocates, so nothing is reused.
+    static bool FreesBeforeEachGet()
+    {
+#ifdef _WIN32
+        MEMORYSTATUSEX memStatus = {};
+        memStatus.dwLength       = sizeof(memStatus);
+        // If the shared memory is less than 64GB(128 / 2), may not enough for the hipblaslt-test to run
+        return GlobalMemoryStatusEx(&memStatus) && memStatus.ullTotalPhys <= (128ULL << 30);
+#else
+        return false;
+#endif
+    }
+
     // Release only idle buffers returned to the pool; live allocations remain owned by callers.
     static void ReleaseCached()
     {
@@ -244,6 +283,17 @@ class memory_pool
         Instance().m_pool_managed.clear();
     }
 
+    // Ends a test: releases the idle buffers of at least `bytes` that the test did not return,
+    // so a large buffer stays cached only while consecutive tests keep reusing it, and the
+    // smaller ones that the last `small_idle_tests` tests did not return.
+    static void EndTest(size_t bytes, uint64_t small_idle_tests)
+    {
+        std::lock_guard<std::mutex> lock(Instance().m_mutex);
+        for(auto* pool : {&Instance().m_pool, &Instance().m_pool_managed})
+            Instance().release_from_earlier_tests(*pool, bytes, small_idle_tests);
+        ++Instance().m_test;
+    }
+
     // Bytes held by idle buffers, i.e. what ReleaseCached would free.
     static size_t CachedBytes()
     {
@@ -258,6 +308,24 @@ class memory_pool
 private:
     std::vector<M> m_pool, m_pool_managed;
     std::mutex m_mutex;
+    uint64_t       m_test = 0; // tests ended, counted by EndTest
+
+    // Releases the idle buffers of at least `bytes` returned before the current test and, when
+    // small_idle_tests is not 0, the smaller ones returned at least that many tests ago.
+    void release_from_earlier_tests(std::vector<M>& pool,
+                                    size_t          bytes,
+                                    uint64_t        small_idle_tests = 0)
+    {
+        pool.erase(std::remove_if(pool.begin(),
+                                  pool.end(),
+                                  [&](const M& m) {
+                                      if(m.capacity() >= bytes)
+                                          return m.restored_test < m_test;
+                                      return small_idle_tests
+                                             && m_test - m.restored_test >= small_idle_tests;
+                                  }),
+                   pool.end());
+    }
 
     static memory_pool& Instance()
     {
@@ -268,25 +336,23 @@ class memory_pool
     M get(size_t bytes, bool use_HMM = false)
     {
         auto& pool = use_HMM ? m_pool_managed : m_pool;
-
-        // For Windows system with not enough system memory, 
-        // not suitable for memory pool management when it needs another allocation
-        #ifdef _WIN32
-        MEMORYSTATUSEX memStatus = {};
-        memStatus.dwLength = sizeof(memStatus);
-        if(GlobalMemoryStatusEx(&memStatus))
+        if(FreesBeforeEachGet())
+            pool.clear();
+
+        auto  it   = std::lower_bound(pool.begin(), pool.end(), bytes);
+        bool  hit  = it != pool.end() && // found a buffer that is large enough ..
+                   it->capacity() < 2 * bytes; // but not way too large
+        if constexpr(std::is_same_v<M, h_memory>)
         {
-            // If the shared memory is less than 64GB(128 / 2), may not enough for the hipblaslt-test to run
-            if(memStatus.ullTotalPhys <= (128ULL << 30))
+            // Before pinning more host memory for a large request, release the large buffers
+            // that earlier tests left and this one has not reused.
+            if(!hit && bytes >= kHostPoolRetainLimit)
             {
-                pool.clear();
+                release_from_earlier_tests(pool, kHostPoolRetainLimit);
+                it = std::lower_bound(pool.begin(), pool.end(), bytes);
             }
         }
-        #endif
-
-        auto  it   = std::lower_bound(pool.begin(), pool.end(), bytes);
-        if(it != pool.end() && // found a buffer that is large enough ..
-           it->capacity() < 2 * bytes) // but not way too large
+        if(hit)
         {
             auto p = std::move(*it);
             p.resize(bytes);
@@ -298,9 +364,12 @@ class memory_pool
             size_t free_device_mem, total_device_mem;
             (void)hipMemGetInfo(&free_device_mem, &total_device_mem);
 
-            // Threshold for "huge" allocations — skip 20% extra allocation
+            // Threshold for "huge" allocations — skip 20% extra allocation. A host buffer gets
+            // none either when host_pool_capacity says so.
             const size_t big_thresh = total_device_mem / 4;
             bool huge_request = (bytes >= big_thresh);
+            if constexpr(std::is_same_v<M, h_memory>)
+                huge_request = huge_request || host_pool_capacity(bytes) == bytes;
 
             // remove the (largest) buffer that was too small
             if(it != pool.begin())
@@ -329,12 +398,26 @@ class memory_pool
     {
         if(!dm.get() || !dm.capacity())
             return;
-        auto& pool = dm.is_managed() ? m_pool_managed : m_pool;
+        auto& pool       = dm.is_managed() ? m_pool_managed : m_pool;
+        dm.restored_test = m_test;
         // insert in (sorted) pool
         pool.insert(std::lower_bound(pool.begin(), pool.end(), dm.capacity()), std::move(dm));
     }
 };
 
+#ifdef GOOGLE_TEST
+// After each test, releases the idle pinned host buffers at or above kHostPoolRetainLimit
+// that the test did not reuse, and the smaller ones that kHostPoolSmallIdleTests tests did not.
+class HostPoolReleaseListener : public ::testing::EmptyTestEventListener
+{
+public:
+    void OnTestEnd(const ::testing::TestInfo&) override
+    {
+        memory_pool<h_memory>::EndTest(kHostPoolRetainLimit, kHostPoolSmallIdleTests);
+    }
+};
+#endif
+
 /* ============================================================================================ */
 /*! \brief  base-class to allocate/deallocate device memory */
 template <typename T>

@@ -31,9 +31,12 @@
 #include "hip_placement.hpp"
 #include "hipblaslt_ostream.hpp"
 
+#include <initializer_list>
 #include <memory>
+#include <sstream>
 #include <string>
 #include <utility>
+#include <vector>
 
 class HipDeviceBuffer : public d_vector_type
 {
@@ -141,9 +144,9 @@ class HipHostBuffer
 {
 public:
     HipHostBuffer(hipDataType dtype, std::size_t numElements)
-        : buffer(memory_pool<h_memory>::Get(realDataTypeSize(dtype) * numElements
-                                                ? realDataTypeSize(dtype) * numElements
-                                                : realDataTypeSize(dtype)))
+        : requested(realDataTypeSize(dtype) * numElements ? realDataTypeSize(dtype) * numElements
+                                                          : realDataTypeSize(dtype))
+        , buffer(memory_pool<h_memory>::Get(requested))
     {
     }
 
@@ -193,10 +196,46 @@ class HipHostBuffer
         return reinterpret_cast<const T*>(buf());
     }
 
+    // A failed pinned allocation leaves a null buffer rather than an error.
+    hipError_t memcheck() const
+    {
+        return buffer.get() ? hipSuccess : hipErrorOutOfMemory;
+    }
+
+    std::size_t getRequestedBytes() const
+    {
+        return requested;
+    }
+
 private:
-    h_memory buffer;
+    std::size_t requested;
+    h_memory    buffer;
 };
 
+// Describes the failed allocation of `b`, or returns an empty string when it succeeded.
+inline std::string host_allocation_failure(const HipHostBuffer& b)
+{
+    if(b.memcheck() == hipSuccess)
+        return {};
+    std::ostringstream s;
+    s.precision(3);
+    s << "could not allocate " << double(b.getRequestedBytes()) / double(1 << 30)
+      << " GiB of pinned host memory";
+    return s.str();
+}
+
+// Describes the first host buffer in `groups` whose allocation failed, or returns an empty
+// string when every allocation succeeded.
+inline std::string
+    host_allocation_failure(std::initializer_list<const std::vector<HipHostBuffer>*> groups)
+{
+    for(const auto* group : groups)
+        for(const auto& b : *group)
+            if(std::string why = host_allocation_failure(b); !why.empty())
+                return why;
+    return {};
+}
+
 inline hipError_t synchronize(HipDeviceBuffer&    dBuf,
                               const HipHostBuffer& hBuf,
                               std::size_t          block_count = 1,

@@ -104,6 +104,15 @@
         }                                                     \
     } while(0)
 
+// WHY is empty when the host allocations succeeded, and otherwise says which one failed.
+#define CHECK_HOST_ALLOCATION(WHY) \
+    do                             \
+    {                              \
+        std::string why__ = (WHY); \
+        if(!why__.empty())         \
+            GTEST_SKIP() << why__; \
+    } while(0)
+
 #define CHECK_SOLUTION_FOUND(SOL_COUNT)     \
     do                                      \
     {                                       \
@@ -153,6 +162,17 @@ inline void hipblaslt_expect_status(hipblasStatus_t status, hipblasStatus_t expe
 
 #define CHECK_DEVICE_ALLOCATION CHECK_HIP_ERROR
 
+#define CHECK_HOST_ALLOCATION(WHY)                             \
+    do                                                         \
+    {                                                          \
+        std::string why__ = (WHY);                             \
+        if(!why__.empty())                                     \
+        {                                                      \
+            hipblaslt_cerr << "error: " << why__ << std::endl; \
+            exit(EXIT_FAILURE);                                \
+        }                                                      \
+    } while(0)
+
 #define EXPECT_HIPBLAS_STATUS hipblaslt_expect_status
 
 #define CHECK_SUCCESS(ERROR) \

@@ -3313,6 +3313,29 @@ void testing_matmul_with_bias(const Arguments& arg,
                 hScaleD.emplace_back(Talpha, 1);
         }
 
+        // A failed pinned allocation leaves a null buffer that the initialization below would
+        // write through.
+        CHECK_HOST_ALLOCATION(host_allocation_failure({&hA,
+                                                       &hB,
+                                                       &hC,
+                                                       &hD_gold,
+                                                       &hD_1,
+                                                       &hBias,
+                                                       &hBias_gold,
+                                                       &hE,
+                                                       &hE_gold,
+                                                       &hScaleAlphaVec,
+                                                       &hScaleA,
+                                                       &hScaleB,
+                                                       &hScaleC,
+                                                       &hScaleD,
+                                                       &hScaleE,
+                                                       &hAmaxD_gold,
+                                                       &hAmaxD,
+                                                       &hD_gold_epl,
+                                                       &hD_gold_ScaleAlpha,
+                                                       &hBias_gold_epl}));
+
         hipblaslt_seedrand();
 
         // For ULP validation, force hpl / trig_float A/B/C inputs to be
@@ -3703,6 +3726,7 @@ void testing_matmul_with_bias(const Arguments& arg,
             if(do_swizzle_a)
             {
                 HipHostBuffer tmp(TiA, size_dA[i]);
+                CHECK_HOST_ALLOCATION(host_allocation_failure(tmp));
                 swizzle_tensor_type(
                     tmp, hA[i], TiA, arg, num_batches[i], M[i], K[i], lda[i], false);
                 CHECK_HIP_ERROR(synchronize(dA[i], tmp, block_count));
@@ -3711,6 +3735,7 @@ void testing_matmul_with_bias(const Arguments& arg,
             if(do_swizzle_b)
             {
                 HipHostBuffer tmp(TiB, size_dB[i]);
+                CHECK_HOST_ALLOCATION(host_allocation_failure(tmp));
                 swizzle_tensor_type(
                     tmp, hB[i], TiB, arg, num_batches[i], N[i], K[i], ldb[i], false);
                 CHECK_HIP_ERROR(synchronize(dB[i], tmp, block_count));
@@ -4190,6 +4215,7 @@ void testing_matmul_with_bias(const Arguments& arg,
                 if(do_swizzle_a)
                 {
                     HipHostBuffer tmp(TiA, size_dA[i]);
+                    CHECK_HOST_ALLOCATION(host_allocation_failure(tmp));
                     swizzle_tensor_type(
                         tmp, hA[batchCount], TiA, arg, 1, M[i], K[i], lda[i], false);
                     CHECK_HIP_ERROR(synchronize(dA[batchCount], tmp, block_count));
@@ -4198,6 +4224,7 @@ void testing_matmul_with_bias(const Arguments& arg,
                 if(do_swizzle_b)
                 {
                     HipHostBuffer tmp(TiB, size_dB[i]);
+                    CHECK_HOST_ALLOCATION(host_allocation_failure(tmp));
                     swizzle_tensor_type(
                         tmp, hB[batchCount], TiB, arg, 1, N[i], K[i], ldb[i], false);
                     CHECK_HIP_ERROR(synchronize(dB[batchCount], tmp, block_count));

@@ -46,6 +46,7 @@ target_sources(hipblaslt-test
         ${CMAKE_CURRENT_SOURCE_DIR}/algo_errors_gtest.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/fast_check_gtest.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/relaunch_compare_gtest.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/host_memory_gtest.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/w4a16_layout_gtest.cpp
 )
 

@@ -24,6 +24,7 @@
  *
  *******************************************************************************/
 
+#include "d_vector.hpp"
 #include "hipblaslt_data.hpp"
 #include "hipblaslt_init.hpp"
 #include "hipblaslt_parse_data.hpp"
@@ -182,6 +183,7 @@ static void hipblaslt_set_listener()
     }
 
     listeners.Append(listener);
+    listeners.Append(new HostPoolReleaseListener);
 }
 
 static void hipblaslt_print_usage_warning()