diff --git a/projects/hipblaslt/TESTING.md b/projects/hipblaslt/TESTING.md index d5e94095ca6e..28d7e6e35238 100644 --- a/projects/hipblaslt/TESTING.md +++ b/projects/hipblaslt/TESTING.md @@ -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. diff --git a/projects/hipblaslt/clients/common/include/d_vector.hpp b/projects/hipblaslt/clients/common/include/d_vector.hpp index acd9d7fb89e4..df7ad655bff0 100644 --- a/projects/hipblaslt/clients/common/include/d_vector.hpp +++ b/projects/hipblaslt/clients/common/include/d_vector.hpp @@ -30,6 +30,7 @@ #include "hipblaslt_arguments.hpp" #include "hipblaslt_init.hpp" #include "hipblaslt_test.hpp" +#include "host_memory.hpp" #include "singletons.hpp" #include #include @@ -38,7 +39,7 @@ #ifdef _WIN32 #include -#else +#elif defined(__linux__) #include #endif @@ -102,17 +103,14 @@ class hip_memory size_t get_available_host_memory() { #ifdef __linux__ + if(std::optional 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); @@ -120,7 +118,8 @@ class hip_memory { // 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 { diff --git a/projects/hipblaslt/clients/common/include/fast_check.hpp b/projects/hipblaslt/clients/common/include/fast_check.hpp index 49b33a6042a7..6e512c696bf1 100644 --- a/projects/hipblaslt/clients/common/include/fast_check.hpp +++ b/projects/hipblaslt/clients/common/include/fast_check.hpp @@ -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 diff --git a/projects/hipblaslt/clients/common/include/host_memory.hpp b/projects/hipblaslt/clients/common/include/host_memory.hpp new file mode 100644 index 000000000000..ec9e62a37e0e --- /dev/null +++ b/projects/hipblaslt/clients/common/include/host_memory.hpp @@ -0,0 +1,37 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +// 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 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 host_memory_limit_override(); + +// The byte count `text` names in the format above, or nothing when it names none. +std::optional parse_host_memory_size(const char* text); diff --git a/projects/hipblaslt/clients/common/src/CMakeLists.txt b/projects/hipblaslt/clients/common/src/CMakeLists.txt index 3d6118da86fb..055d142e5124 100755 --- a/projects/hipblaslt/clients/common/src/CMakeLists.txt +++ b/projects/hipblaslt/clients/common/src/CMakeLists.txt @@ -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) diff --git a/projects/hipblaslt/clients/common/src/fast_check.cpp b/projects/hipblaslt/clients/common/src/fast_check.cpp index 6bd9d42d00a5..9835306facb8 100644 --- a/projects/hipblaslt/clients/common/src/fast_check.cpp +++ b/projects/hipblaslt/clients/common/src/fast_check.cpp @@ -3,6 +3,7 @@ #include "fast_check.hpp" #include "d_vector.hpp" +#include "host_memory.hpp" #include @@ -10,20 +11,12 @@ #include #include #include -#include #include #include #include #include #include -#ifdef _WIN32 -#ifndef NOMINMAX -#define NOMINMAX -#endif -#include -#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::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 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 {}; }; diff --git a/projects/hipblaslt/clients/common/src/host_memory.cpp b/projects/hipblaslt/clients/common/src/host_memory.cpp new file mode 100644 index 000000000000..6da9229376ec --- /dev/null +++ b/projects/hipblaslt/clients/common/src/host_memory.cpp @@ -0,0 +1,273 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "host_memory.hpp" +#include "hipblaslt_ostream.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#ifdef _WIN32 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#include +#endif + +namespace +{ + std::optional read_number(const std::filesystem::path& file) + { + std::ifstream in(file); + std::string token; + if(!(in >> token) || token.empty() + || token.find_first_not_of("0123456789") != std::string::npos) + return std::nullopt; + try + { + return std::stoull(token); + } + catch(...) + { + return std::nullopt; + } + } + + // The value after `key` in a "key value" file such as meminfo or memory.stat. + std::optional read_field(const std::filesystem::path& file, const std::string& key) + { + std::ifstream in(file); + std::string name; + uint64_t value = 0; + while(in >> name >> value) + { + if(name == key) + return value; + in.ignore(std::numeric_limits::max(), '\n'); + } + return std::nullopt; + } + + // Memory left under one cgroup's limit, or nothing when the directory sets no limit. Page + // cache the kernel can drop does not count as used: both file lists (active_file and + // inactive_file; tmpfs and shared memory are not on them), less pages still waiting to be + // written back. A pinned allocation reclaims active and inactive cache alike. + std::optional cgroup_headroom(const std::filesystem::path& dir, bool v2) + { + // v2 writes "max" for no limit, which is not a number. v1 writes a value near 2^63, + // which MemAvailable is always below. + auto limit = read_number(dir / (v2 ? "memory.max" : "memory.limit_in_bytes")); + if(!limit) + return std::nullopt; + uint64_t used + = read_number(dir / (v2 ? "memory.current" : "memory.usage_in_bytes")).value_or(0); + auto stat = [&](const char* v2_key, const char* v1_key) { + return read_field(dir / "memory.stat", v2 ? v2_key : v1_key).value_or(0); + }; + uint64_t cache = stat("active_file", "total_active_file") + + stat("inactive_file", "total_inactive_file"); + uint64_t unwritten + = stat("file_dirty", "total_dirty") + stat("file_writeback", "total_writeback"); + uint64_t reclaimable = cache > unwritten ? cache - unwritten : 0; + used = used > reclaimable ? used - reclaimable : 0; + return *limit > used ? *limit - used : uint64_t(0); + } + + // Where a cgroup hierarchy is mounted (point), and which of its cgroups is at that point + // (root). Outside a cgroup namespace a container sees the host's paths, and its mount may + // start below the host's root. + struct CgroupMount + { + std::filesystem::path root; + std::filesystem::path point; + }; + + // mountinfo writes a space, tab, newline or backslash in a path as \ and three octal digits. + std::string unescape_mount_path(const std::string& text) + { + std::string out; + for(size_t i = 0; i < text.size(); i++) + { + if(text[i] == '\\' && i + 3 < text.size() + && std::isdigit(static_cast(text[i + 1])) + && std::isdigit(static_cast(text[i + 2])) + && std::isdigit(static_cast(text[i + 3]))) + { + out += char((text[i + 1] - '0') * 64 + (text[i + 2] - '0') * 8 + + (text[i + 3] - '0')); + i += 3; + } + else + out += text[i]; + } + return out; + } + + // The cgroup v2 mount, or the v1 mount that carries the memory controller, from + // /proc/self/mountinfo. Each line is "id parent dev root point options ... - type source + // super-options". + std::optional find_cgroup_mount(const std::string& mountinfo, bool v2) + { + std::ifstream in(mountinfo); + for(std::string line; std::getline(in, line);) + { + size_t dash = line.find(" - "); + if(dash == std::string::npos) + continue; + std::istringstream before(line.substr(0, dash)), after(line.substr(dash + 3)); + std::string id, parent, dev, root, point, type, source, options; + if(!(before >> id >> parent >> dev >> root >> point) || !(after >> type >> source)) + continue; + after >> options; + bool memory = false; + std::stringstream list(options); + for(std::string o; std::getline(list, o, ',');) + memory = memory || o == "memory"; + if(v2 ? type == "cgroup2" : type == "cgroup" && memory) + return CgroupMount{unescape_mount_path(root), unescape_mount_path(point)}; + } + return std::nullopt; + } + + // The smallest headroom from the process's cgroup up to the top of the hierarchy this + // process can see. The cgroup's path is translated through its mount when mountinfo lists + // one; otherwise the hierarchy is assumed at cgroup_root (v1: cgroup_root/memory). A + // container may list "/" or the host's path for its cgroup; either way its own limit is at + // the top it sees, so the walk always ends there. + std::optional smallest_cgroup_headroom(const HostMemorySources& s) + { + std::ifstream in(s.self_cgroup); + std::string line; + std::optional smallest; + while(std::getline(in, line)) + { + // hierarchy-id:controllers:path + size_t a = line.find(':'); + size_t b = a == std::string::npos ? a : line.find(':', a + 1); + if(b == std::string::npos) + continue; + std::string controllers = line.substr(a + 1, b - a - 1); + bool v2 = line.substr(0, a) == "0" && controllers.empty(); + bool v1 = false; + std::stringstream list(controllers); + for(std::string c; std::getline(list, c, ',');) + v1 = v1 || c == "memory"; + if(!v2 && !v1) + continue; + + // A cgroup that has been removed is listed with this suffix. + std::string path = line.substr(b + 1); + const std::string deleted = " (deleted)"; + if(path.size() >= deleted.size() + && path.compare(path.size() - deleted.size(), deleted.size(), deleted) == 0) + path.erase(path.size() - deleted.size()); + + std::filesystem::path top, rel; + if(auto mount = find_cgroup_mount(s.mountinfo, v2)) + { + // The part of the cgroup's path below the mount's root. A cgroup outside that + // root (a container listing the host's path) leaves only the top visible. + top = mount->point; + auto sub = std::filesystem::path(path).lexically_relative(mount->root); + if(!sub.empty() && *sub.begin() != ".." && sub != ".") + rel = sub; + } + else + { + top = s.cgroup_root; + if(v1) + top /= "memory"; + rel = std::filesystem::path(path).relative_path(); + } + while(true) + { + if(auto h = cgroup_headroom(top / rel, v2)) + smallest = std::min(smallest.value_or(*h), *h); + if(rel.empty()) + break; + rel = rel.parent_path(); + } + } + return smallest; + } +} // namespace + +std::optional parse_host_memory_size(const char* text) +{ + if(!text || !std::isdigit(static_cast(*text))) + return std::nullopt; + char* end = nullptr; + errno = 0; + unsigned long long value = std::strtoull(text, &end, 10); + if(errno == ERANGE) + return std::nullopt; + int shift = 0; + switch(std::toupper(static_cast(*end))) + { + case '\0': + break; + case 'K': + shift = 10; + break; + case 'M': + shift = 20; + break; + case 'G': + shift = 30; + break; + case 'T': + shift = 40; + break; + default: + return std::nullopt; + } + if(shift && *++end != '\0') + return std::nullopt; + if(value > (std::numeric_limits::max() >> shift)) + return std::nullopt; + return size_t(value) << shift; +} + +std::optional host_memory_limit_override() +{ + const char* text = std::getenv("HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT"); + if(!text || !*text) + return std::nullopt; + auto limit = parse_host_memory_size(text); + static std::atomic warned{false}; + if(!limit && !warned.exchange(true)) + hipblaslt_cerr << "warning: ignoring HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT=" << text + << ", which is not a byte count such as 64G" << std::endl; + return limit; +} + +std::optional available_host_memory_bytes(const HostMemorySources& sources) +{ + std::optional available; +#ifdef _WIN32 + (void)sources; + MEMORYSTATUSEX status = {}; + status.dwLength = sizeof(status); + if(GlobalMemoryStatusEx(&status)) + available = size_t(status.ullAvailPhys); +#else + // Each source that can be read caps the budget on its own. + if(auto kib = read_field(sources.meminfo, "MemAvailable:")) + available = size_t(*kib * 1024); + if(auto headroom = smallest_cgroup_headroom(sources)) + available = std::min(available.value_or(size_t(*headroom)), size_t(*headroom)); +#endif + // When the system cannot be read, the limit is the only budget there is. + if(auto limit = host_memory_limit_override()) + return std::min(available.value_or(*limit), *limit); + return available; +} diff --git a/projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp b/projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp index bfe6a3524413..256360a2e7ea 100644 --- a/projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/host_memory_gtest.cpp @@ -8,11 +8,22 @@ #include +#include "fast_check.hpp" #include "hipBuffer.hpp" +#include "host_memory.hpp" +#include +#include +#include +#include +#include #include #include +#ifndef _WIN32 +#include +#endif + namespace { // What the HostPoolReleaseListener in hipblaslt-test does when a test ends. @@ -32,8 +43,46 @@ namespace GTEST_SKIP() << "could not pin " #PTR " on this host"; \ } while(0) + // Sets HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT, or unsets it for nullptr, for one scope, and then + // restores the value the run started with. + class ScopedLimit + { + public: + explicit ScopedLimit(const char* value) + { + if(const char* old = std::getenv(kName)) + saved = old; + set(value); + } + ~ScopedLimit() + { + set(saved ? saved->c_str() : nullptr); + } + ScopedLimit(const ScopedLimit&) = delete; + ScopedLimit& operator=(const ScopedLimit&) = delete; + + private: + static constexpr const char* kName = "HIPBLASLT_CLIENT_HOST_MEMORY_LIMIT"; + + static void set(const char* value) + { +#ifdef _WIN32 + _putenv_s(kName, value ? value : ""); +#else + if(value) + setenv(kName, value, 1); + else + unsetenv(kName); +#endif + } + + std::optional saved; + }; + TEST(HostMemory_pre_checkin, large_buffer_released_after_a_test_that_does_not_reuse_it) { + // The pool refuses to pin beyond a host memory limit, and these sizes are the test. + ScopedLimit no_limit(nullptr); memory_pool::ReleaseCached(); { HipHostBuffer large(HIP_R_8I, kHostPoolRetainLimit); @@ -85,6 +134,7 @@ namespace { if(memory_pool::FreesBeforeEachGet()) GTEST_SKIP() << "the pool on this host frees its idle buffers before every request"; + ScopedLimit no_limit(nullptr); memory_pool::ReleaseCached(); auto first = memory_pool::Get(kHostPoolRetainLimit); SKIP_UNLESS_PINNED(first.get()); @@ -106,6 +156,7 @@ namespace TEST(HostMemory_pre_checkin, large_miss_releases_what_earlier_tests_left) { + ScopedLimit no_limit(nullptr); // A test of another shape needs a new large buffer; the previous test's ones are // released before more memory is pinned, not after this test ends. On any miss the // pool already drops the largest buffer that is too small, so two are left here. @@ -129,6 +180,7 @@ namespace TEST(HostMemory_pre_checkin, buffers_at_the_retain_limit_get_no_slack) { + ScopedLimit no_limit(nullptr); // The pool adds 20% to small requests so that a slightly larger one can reuse them. A // large buffer is reused only by the same shape, so the extra would only add pinned // memory. @@ -172,4 +224,312 @@ namespace std::string::npos) << why; } + +#ifndef _WIN32 + constexpr uint64_t kGiB = uint64_t(1) << 30; + + // A temporary stand-in for /proc/meminfo, /proc/self/cgroup and /sys/fs/cgroup. + class FakeHost + { + public: + FakeHost() + { + static int count = 0; + dir = std::filesystem::temp_directory_path() + / ("hipblaslt_host_memory_" + std::to_string(getpid()) + "_" + + std::to_string(count++)); + std::filesystem::create_directories(dir / "cgroup"); + sources.meminfo = (dir / "meminfo").string(); + sources.self_cgroup = (dir / "self_cgroup").string(); + sources.mountinfo = (dir / "mountinfo").string(); + sources.cgroup_root = (dir / "cgroup").string(); + } + ~FakeHost() + { + std::error_code ec; + std::filesystem::remove_all(dir, ec); + } + + void write(const std::filesystem::path& file, const std::string& text) const + { + std::filesystem::create_directories(file.parent_path()); + std::ofstream(file) << text; + } + void mem_available(uint64_t bytes) const + { + write(sources.meminfo, + "MemTotal: 999999999 kB\nMemFree: 1 kB\nMemAvailable: " + + std::to_string(bytes / 1024) + " kB\n"); + } + // One cgroup directory under the fake root, relative like "job/step". + void cgroup(const std::string& rel, + const std::string& limit, + uint64_t used, + uint64_t inactive_file = 0, + bool v1 = false) const + { + std::filesystem::path d = std::filesystem::path(sources.cgroup_root); + if(v1) + d /= "memory"; + d /= rel; + write(d / (v1 ? "memory.limit_in_bytes" : "memory.max"), limit + "\n"); + write(d / (v1 ? "memory.usage_in_bytes" : "memory.current"), + std::to_string(used) + "\n"); + write(d / "memory.stat", + std::string(v1 ? "total_inactive_file " : "inactive_file ") + + std::to_string(inactive_file) + "\n"); + } + + std::filesystem::path dir; + HostMemorySources sources; + + private: + // A limit set for the whole run would cap every query here. + ScopedLimit no_limit{nullptr}; + }; + + TEST(HostMemory_pre_checkin, query_reads_mem_available_without_a_cgroup) + { + FakeHost host; + host.mem_available(8 * kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + } + + TEST(HostMemory_pre_checkin, query_uses_a_cgroup_v2_limit_below_mem_available) + { + // 2 GiB limit, 1.5 GiB charged of which 0.25 GiB is page cache the kernel can drop. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", std::to_string(2 * kGiB), 3 * kGiB / 2, kGiB / 4); + EXPECT_EQ(available_host_memory_bytes(host.sources), 3 * kGiB / 4); + } + + TEST(HostMemory_pre_checkin, query_ignores_a_cgroup_v2_without_a_limit) + { + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", "max", 6 * kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + } + + TEST(HostMemory_pre_checkin, query_uses_the_tightest_limit_up_to_the_root) + { + // The job's own limit leaves 3 GiB, its parent's 0.5 GiB, and the root's 1 GiB. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/runner/job\n"); + host.cgroup("runner/job", std::to_string(4 * kGiB), kGiB); + host.cgroup("runner", std::to_string(2 * kGiB), 3 * kGiB / 2); + host.cgroup("", std::to_string(6 * kGiB), 5 * kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), kGiB / 2); + } + + TEST(HostMemory_pre_checkin, query_finds_a_container_limit_at_its_root) + { + // A container lists its cgroup by the host's path but sees only its own limit at the root. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/kubepods/pod1/ctr\n"); + host.cgroup("", std::to_string(3 * kGiB), kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), 2 * kGiB); + + // mountinfo names another cgroup at the mount, so the listed path is not below it. The + // walk stays inside the mount: a full limit beside it is not read. + const auto beside = std::filesystem::path(host.sources.cgroup_root).parent_path(); + host.write(beside / "memory.max", std::to_string(kGiB) + "\n"); + host.write(beside / "memory.current", std::to_string(kGiB) + "\n"); + host.write(host.sources.mountinfo, + "30 24 0:26 /kubepods/pod2 " + host.sources.cgroup_root + + " rw - cgroup2 cgroup2 rw\n"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 2 * kGiB); + } + + TEST(HostMemory_pre_checkin, query_reads_cgroup_v1) + { + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "7:cpu,cpuacct:/job\n5:memory:/job\n"); + host.cgroup("job", std::to_string(3 * kGiB), kGiB, 0, true); + EXPECT_EQ(available_host_memory_bytes(host.sources), 2 * kGiB); + + // v1 writes a value near 2^63 for no limit. + host.cgroup("job", "9223372036854771712", kGiB, 0, true); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + } + + TEST(HostMemory_pre_checkin, query_handles_missing_and_unexpected_files) + { + FakeHost host; + EXPECT_EQ(available_host_memory_bytes(host.sources), std::nullopt) << "no meminfo"; + + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", "not-a-number", kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB) << "garbled limit"; + + host.cgroup("job", std::to_string(kGiB), 2 * kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), 0u) << "usage above the limit"; + } + + TEST(HostMemory_pre_checkin, query_of_this_host_reads_the_real_files) + { + // The fake-tree tests cover the cgroup logic. This one only checks that the default + // sources are readable on a real host and that its cgroups can only lower MemAvailable; + // 0 is a valid answer in a full container. /proc files report a size of 0, so copy + // the text rather than the file. + FakeHost meminfo_only; + std::stringstream text; + text << std::ifstream(HostMemorySources().meminfo).rdbuf(); + meminfo_only.write(meminfo_only.sources.meminfo, text.str()); + const std::optional available = available_host_memory_bytes(); + ASSERT_TRUE(available.has_value()); + // MemAvailable moves between the two reads, so allow 1 GiB of drift. + EXPECT_LE(*available, available_host_memory_bytes(meminfo_only.sources).value_or(0) + kGiB); + } + + TEST(HostMemory_pre_checkin, query_counts_active_page_cache_as_reclaimable) + { + // A 2 GiB container that is full, mostly with page cache the kernel reclaims (active or + // inactive) before it kills anything. Pages still being written back are not counted. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", std::to_string(2 * kGiB), 2 * kGiB); + host.write(std::filesystem::path(host.sources.cgroup_root) / "job/memory.stat", + "anon 524288\nfile 1610612746\nactive_file " + std::to_string(3 * kGiB / 2) + + "\ninactive_file 10\nfile_dirty 0\nfile_writeback 0\n"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 3 * kGiB / 2 + 10); + + host.write(std::filesystem::path(host.sources.cgroup_root) / "job/memory.stat", + "active_file " + std::to_string(3 * kGiB / 2) + "\ninactive_file 10\nfile_dirty " + + std::to_string(kGiB / 4) + "\nfile_writeback " + std::to_string(kGiB / 4) + + "\n"); + EXPECT_EQ(available_host_memory_bytes(host.sources), kGiB + 10); + } + + TEST(HostMemory_pre_checkin, query_maps_the_cgroup_path_through_its_mount) + { + // The hierarchy is mounted from /runner, so /proc/self/cgroup's /runner/job is the + // mount's job directory. Its 2 GiB limit, not the mount top's 6 GiB, applies. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/runner/job\n"); + host.write(host.sources.mountinfo, + "24 1 0:22 / / rw - ext4 /dev/root rw\n30 24 0:26 /runner " + + host.sources.cgroup_root + " rw,nosuid - cgroup2 cgroup2 rw\n"); + host.cgroup("", std::to_string(6 * kGiB), kGiB); + host.cgroup("job", std::to_string(2 * kGiB), kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), kGiB); + } + + TEST(HostMemory_pre_checkin, query_finds_the_v1_memory_controller_where_it_is_mounted) + { + // The v1 memory controller shares a mount with cpu, at "cpu,memory", not "memory". + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "4:cpu,memory:/job\n"); + const std::string point = host.sources.cgroup_root + "/cpu,memory"; + host.write(host.sources.mountinfo, + "31 24 0:27 / " + point + " rw - cgroup cgroup rw,cpu,memory\n"); + host.write(point + "/job/memory.limit_in_bytes", std::to_string(3 * kGiB) + "\n"); + host.write(point + "/job/memory.usage_in_bytes", std::to_string(kGiB) + "\n"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 2 * kGiB); + } + + TEST(HostMemory_pre_checkin, query_applies_the_cgroup_limit_without_meminfo) + { + // With /proc/meminfo unreadable, a full container is still full, with or without a limit. + FakeHost host; + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", std::to_string(kGiB), kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), size_t(0)); + ScopedLimit limit("2G"); + EXPECT_EQ(available_host_memory_bytes(host.sources), size_t(0)); + } + + TEST(HostMemory_pre_checkin, query_follows_a_removed_cgroup) + { + // /proc/self/cgroup lists a cgroup that has been removed with a " (deleted)" suffix. + FakeHost host; + host.mem_available(8 * kGiB); + host.write(host.sources.self_cgroup, "0::/job (deleted)\n"); + host.cgroup("job", std::to_string(2 * kGiB), kGiB); + EXPECT_EQ(available_host_memory_bytes(host.sources), kGiB); + } + + TEST(HostMemory_pre_checkin, limit_is_a_byte_count_with_an_optional_suffix) + { + EXPECT_EQ(parse_host_memory_size("1024"), size_t(1024)); + EXPECT_EQ(parse_host_memory_size("0"), size_t(0)); + EXPECT_EQ(parse_host_memory_size("4K"), size_t(4) << 10); + EXPECT_EQ(parse_host_memory_size("512m"), size_t(512) << 20); + EXPECT_EQ(parse_host_memory_size("64G"), size_t(64) << 30); + EXPECT_EQ(parse_host_memory_size("2T"), size_t(2) << 40); + for(const char* bad : + {"", "G", "64GB", "-1", "1.5G", " 64G", "64 G", "16777216T", "99999999999999999999"}) + EXPECT_EQ(parse_host_memory_size(bad), std::nullopt) << '"' << bad << '"'; + } + + TEST(HostMemory_pre_checkin, limit_caps_the_query_and_never_raises_it) + { + FakeHost host; + host.mem_available(8 * kGiB); + { + ScopedLimit limit("2G"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 2 * kGiB); + } + { + ScopedLimit limit("16G"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + } + { + // An unreadable value is ignored, not read as 0. + ScopedLimit limit("not-a-size"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + EXPECT_EQ(host_memory_limit_override(), std::nullopt); + } + // Each limit above ended with its scope. + EXPECT_EQ(available_host_memory_bytes(host.sources), 8 * kGiB); + { + // A full container stays full. + host.write(host.sources.self_cgroup, "0::/job\n"); + host.cgroup("job", std::to_string(kGiB), kGiB); + ScopedLimit limit("2G"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 0u); + } + } + + TEST(HostMemory_pre_checkin, limit_is_the_budget_when_the_system_cannot_be_read) + { + FakeHost host; + ScopedLimit limit("3G"); + EXPECT_EQ(available_host_memory_bytes(host.sources), 3 * kGiB); + } + + TEST(HostMemory_pre_checkin, limit_makes_the_preflight_skip) + { + // Only a runner with 2 GiB free shows that the limit, not the runner, made it skip. + { + ScopedLimit no_limit(nullptr); + if(available_host_memory_bytes().value_or(0) < 2 * kGiB) + GTEST_SKIP() << "this runner has less than 2 GiB of host memory free"; + } + ScopedLimit limit("1G"); + std::string why = fast_check_memory_shortfall(0, 2 * kGiB); + EXPECT_NE(why.find("needs 2 GiB of host memory and 1 GiB is available"), std::string::npos) + << why; + } + + TEST(HostMemory_pre_checkin, no_host_memory_left_makes_the_preflight_skip) + { + // A budget of 0, from a full container or this limit, means nothing is left. It differs + // from a budget the system could not report, under which the preflight lets the case run. + ScopedLimit limit("0"); + std::string why = fast_check_memory_shortfall(0, kGiB); + EXPECT_NE(why.find("needs 1 GiB of host memory and 0 GiB is available"), std::string::npos) + << why; + } +#endif } // namespace