From 2b0a8b8bd782ff4a23bebe61a2c6ebd4e72d2f04 Mon Sep 17 00:00:00 2001 From: dev Date: Tue, 29 Sep 2026 21:09:32 +0000 Subject: [PATCH 1/2] spec(BACKEND-CUDA-SM120): the per-allocation CUDA trace A first-forward OOM on a 24 GiB sm_120a card names one failing size, and nothing in the tree can say which allocation sequence produced the pressure. Record the instrument, its threshold rationale, the capture-sentinel test, and what it does not prove. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:opencode-go/deepseek-v4.1-flash [pi] --- .../ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md | 19 +++++ .agents/specs/cuda-alloc-trace.md | 74 +++++++++++++++++++ 2 files changed, 93 insertions(+) create mode 100644 .agents/issues/BACKEND-CUDA-SM120/ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md create mode 100644 .agents/specs/cuda-alloc-trace.md diff --git a/.agents/issues/BACKEND-CUDA-SM120/ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md b/.agents/issues/BACKEND-CUDA-SM120/ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md new file mode 100644 index 0000000000..a8f1f62f91 --- /dev/null +++ b/.agents/issues/BACKEND-CUDA-SM120/ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H +Title: A first-forward CUDA OOM reports one failing size and nothing else: there is no per-allocation trace of where device memory went +Row: BACKEND-CUDA-SM120 +State: OPEN +Kind: feature +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-29 +Updated: 2026-09-29 +Closed: - + +## Problem + +On a 24 GiB consumer-Blackwell card (sm_120a), the failure mode that matters is a first-forward cudaMalloc that does not fit, and today it produces exactly one line: `vt cuda: cudaMalloc: out of memory` naming the FAILING SIZE and nothing about the process. `VT_CUDA_ALLOC_STATS` aggregates counters, so it cannot say which allocation sequence produced the pressure; `nvidia-smi` sees the process total, not the call. The result is that an OOM is diagnosed by bisecting the forward (the Qwen3.8-27B NVFP4 arm loads ~20 GiB of packed weights into the same pool the KV cache and the repack scratch come from, so the interesting question is the ORDER and the live-bytes curve, not the total). `src/vt/cuda/cuda_backend.cu` owns Alloc/Free and `DeviceMemoryInfo`, so the instrument belongs there: `VT_CUDA_ALLOC_TRACE=1` prints every cudaMalloc of at least 16 MiB with its size, running live bytes and free device memory, every free of the same size, and the request that fails. The 16 MiB threshold is the signal-to-noise line (the pool hands out MiB-scale blocks; per-4 KiB lines would be a log, not an instrument) and a free-memory trigger below 2 GiB logs everything so the endgame is never missed. Zero-cost when unset: one process-static env read and one size compare on a path that is about to call the driver. + +## Resolution + +- diff --git a/.agents/specs/cuda-alloc-trace.md b/.agents/specs/cuda-alloc-trace.md new file mode 100644 index 0000000000..ee47a4b00f --- /dev/null +++ b/.agents/specs/cuda-alloc-trace.md @@ -0,0 +1,74 @@ +# CUDA per-allocation trace — ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H + +A first-forward `cudaMalloc` that does not fit produces one line naming the +failing size. On a 24 GiB `sm_120a` card serving a ~20 GiB NVFP4 arm out of the +same pool as the KV cache and the repack scratch, the interesting question is +which allocation sequence produced the pressure — and nothing in the tree can +answer it. + +Issue: [ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H](../issues/BACKEND-CUDA-SM120/ISSUE-LOCAL-01M3QFZD0PTVP2CT9S2RYG8M2H.md). +Owning row: `BACKEND-CUDA-SM120` ([backend-matrix.md](../backend-matrix.md)), the +consumer-Blackwell row, because that is the card this instrument exists for. + +## Premise, grounded + +| Where (line anchors at this branch's base, `b45a94273`) | What | +|---|---| +| `src/vt/cuda/cuda_backend.cu:55-59` | `Check` — the only report an OOM gets: the failing call's name and the driver string. | +| `src/vt/cuda/cuda_backend.cu:120-140` | `Alloc`/`Free`; `StatsEnabled` (`VT_CUDA_ALLOC_STATS`) aggregates counters, so it cannot attribute pressure to a sequence. | +| `DeviceMemoryInfo` | The `cudaMemGetInfo` wrapper the trace reports through, already used by the stats arm. | + +`VT_CUDA_ALLOC_STATS` is aggregate by design and stays as it is; this is a second, +opt-in instrument, not a replacement. + +## Design + +`VT_CUDA_ALLOC_TRACE=1`, read once per process: + +- Every `Alloc` of at least 16 MiB prints + `[cuda-alloc] #N size= live= free=`. The 16 MiB floor is the + signal-to-noise line: the pool hands out MiB-scale blocks, and a per-allocation + line for every 4 KiB scratch would be a log rather than an instrument. +- Below 2 GiB free, every allocation prints regardless of size — the endgame is + exactly where a size filter would hide the last few blocks. +- Every `Free` of a tracked block of at least 16 MiB prints + `[cuda-free] size=… live=… free=…`, so the live curve is readable in both + directions. The map is keyed by pointer; a free of a block the trace never + tracked is ignored. +- A failed `cudaMalloc` prints `[cuda-alloc] FAILED size=… free=… err=…` before + `Check` throws, so the failure is attributed to the live total at that moment. +- Unset: one process-static env read and one size comparison per `Alloc`. + +## Tests + +`tests/vt/test_cuda_alloc_trace.cpp` (new target, registered beside +`test_cuda_backend`): the binary enables the flag in a global initializer (the +flag is process-static, so it cannot be toggled in-process) and captures stderr +around each backend call. + +- A 32 MiB allocation prints a `[cuda-alloc]` line with `size=32.0 MiB`, and the + matching `Free` prints a `[cuda-free]` line with the same size. +- An impossible request (1 PiB) throws AND prints the `FAILED` line with `err=`, + so the next OOM names its own size and the free memory at that instant. +- The capture asserts its own sentinel, the same discipline the MoE tap's capture + uses: a capture that silently caught nothing must fail rather than read as "the + trace printed nothing". +- On a build without CUDA the case says so and exits; the flag-off path is the + same code with `AllocTraceEnabled()` false and is covered by every CUDA test in + the tree, none of which sets the variable. + +What the gate does not prove: that a real OOM on a real model is easier to read — +that is a usability claim, not a testable one. It proves the three lines exist, +carry the right size, and reach stderr. + +## Gates + +- `ctest --test-dir build -R test_cuda_alloc_trace` on a CUDA build + (`-DVLLM_CPP_CUDA_ARCHITECTURES=120a` on the local card). +- `test_cuda_backend` unchanged. +- `scripts/agent-preflight.sh --staged`. + +## Owed + +- The measurement this instrument exists for: the NVFP4 arm's live-bytes curve on + the 24 GiB `sm_120a` card, which is a row deliverable and not this PR's. From d416d6ac4d4c0a7f514942a619be8e0656a3299b Mon Sep 17 00:00:00 2001 From: dev Date: Sun, 27 Sep 2026 19:47:21 +0000 Subject: [PATCH 2/2] feat(BACKEND-CUDA-SM120): opt-in per-allocation trace VT_CUDA_ALLOC_TRACE=1 prints each cudaMalloc of at least 16 MiB with its size, live bytes and free device memory, each free of the same size, and the request that fails. A first-forward OOM becomes a size list instead of a guess. Zero-cost when unset. The threshold is the signal-to-noise line: the pool hands out MiB-scale blocks, and below 2 GiB free every allocation prints regardless of size so the endgame is never missed. The trace is keyed by pointer, so a free of an untracked block is ignored. The new test in test_cuda_alloc_trace captures stderr around a 32 MiB allocation, its free, and a request that cannot fit. It asserts its own capture with a sentinel, so a redirect that caught nothing fails instead of reading as "the trace printed nothing". The flag is process-static, so the binary enables it before main and is flag-ON by construction; the flag-off path is the same guard and is covered by every CUDA suite, none of which sets the variable. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:opencode-go/deepseek-v4.1-flash [pi] --- docs/ENVIRONMENT.md | 1 + src/vt/cuda/cuda_backend.cu | 83 +++++++++++++++++++++ tests/CMakeLists.txt | 3 + tests/vt/test_cuda_alloc_trace.cpp | 115 +++++++++++++++++++++++++++++ 4 files changed, 202 insertions(+) create mode 100644 tests/vt/test_cuda_alloc_trace.cpp diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index d514d76cda..18c2701d26 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -276,6 +276,7 @@ portable/reference path. In normal operation leave them unset. | `VT_QWEN35_STAGE_RESERVE_BYTES` | `12 GiB` (Qwen3.5/3.6 family; read once, and only by the SAFETENSORS loader) | The device memory a staged model has to leave behind for the KV cache and activations. The retag above is measurably slower than a true device copy on GB10, so a model that FITS should stage; [#1299](https://github.com/mudler/vllm.cpp/issues/1299) recorded `Qwen3.8-2.4T-A95B` exhausting a 119.631 GiB box precisely BECAUSE the CUDA arm paid for its weights twice, so a model that does not fit must not. The rule is `2 * model_weight_bytes + VT_QWEN35_STAGE_RESERVE_BYTES <= device_total_bytes`, and the factor of TWO is the substance: on a platform that does not release the host mirror after upload, a staged weight is ADDITIVE. **MEASURED on the committed gate** (`scripts/dflash2-speed-gate.sh`, Qwen3.8-27B bf16 + DFlash2 k=7, c=1, one boot, same oracle wheel): retagged **12.361** tok/s -> staged **15.029**, against vLLM's 16.363, i.e. ratio **0.7587 -> 0.9185**, **+21.6%**. **IT IS DECIDED ONCE, AT LOAD, AND THAT IS THE WHOLE POINT.** The first version of this policy asked `Backend::DeviceMemoryInfo` PER WEIGHT and compared LIVE FREE against a fraction of total; free falls as the host mirror loads, so the floor degraded into "stage the first N GiB, then stop" and captured **8.1% of the available 21.6%** (`a22030924`, gate 13.3625). The rule was already written next door: `include/vllm/platforms/interface.h` says the budget is "TOTAL rather than FREE, because `free` at load time carries the page cache and whatever else the box is doing, which would make a load-time verdict a function of contention." Only the safetensors loader latches it, so a GGUF checkpoint never stages and keeps exactly the behaviour #1299 shipped — which also makes the file total a sound proxy, since a bf16 safetensors weight occupies its file size in host RAM while a GGUF one expands on load. Parsed with `atoll`: unset, empty, unparsable and `<= 0` all fall back to 12 GiB rather than refusing, so a typo cannot silently shrink the reserve and double a model's residency. An unknown device total or an unknown model size answers NO. `VT_QWEN35_ALIAS_HOST_WEIGHTS=1` pins the retag on regardless; `=0` forces staging regardless. LOWER admits bigger models to staging; HIGHER is more conservative | | `VT_LOAD_DIRECT_UPLOAD` | on | Load a weight the device consumes VERBATIM by VIEWING the safetensors mmap (`OwnedBytes::Borrow`, keep-alive on the mapping) instead of copying it into an owned host buffer first, so the device upload reads the file mapping and the load moves those bytes ONCE rather than twice. Only whole-range same-size copies qualify — a transpose, a dtype conversion, a dequant, a concatenation or a load-time repack always takes the copy path, and the helper re-checks `numel * sizeof(dtype) == span` and fails closed to the copy on any mismatch. `0` is the same-binary A/B back to copy-then-upload. Bytes are identical either way, so tokens are unchanged. MEASURED on GB10, Qwen3.6-27B bf16 (50.098 GiB), Vulkan, same binary both arms: the weight-load phase goes **19.27 -> 12.48 s warm** (1.54x) and **52.62 -> 32.75 s cold** (1.61x), load-and-one-token **30.39 -> 22.47 s** warm and **62.98 -> 55.60 s** cold, with every ON leg beating every OFF leg. Total bytes MOVED **100.196 -> 81.260 GiB**: the host materialization pass drops **50.098 -> 31.162 GiB** while the 50.098 GiB device upload is unchanged (the model still has to be uploaded once). 37.8% of this checkpoint qualifies; the rest is merged (qkv, gate_up) or transposed (lm_head) at load and correctly still copies | | `VT_LOAD_STATS` | off | `=1` prints one line per load phase (mmap+header, weights) with its wall time, plus the bytes the load MOVED: `host_copy` (source bytes materialized into an owned host buffer), `borrowed` (source bytes viewed in place by the direct-upload path) and `device_upload` (bytes copied host to device). Diagnostic only; it changes no numerics. Issue #150 | +| `VT_CUDA_ALLOC_TRACE` | off | `=1` prints every `cudaMalloc` of at least 16 MiB with its size, live bytes and free device memory, every free of the same size, and the request that fails, so a first-forward OOM is attributable by size. Zero-cost when unset. CUDA-only | | `VT_VULKAN_ALLOC_STATS` | off | `=1` prints a device-memory line on every 1 GiB high-water crossing and a summary at exit: live buffer count, bytes REQUESTED by the caller, bytes COMMITTED by the driver (`VkMemoryRequirements::size`), peak live bytes, and the process/system context (`VmRSS`, `VmHWM`, `MemAvailable`, `Cached`) read from `/proc`. On a unified-memory device the Vulkan heap IS system RAM, so separating "the backend allocated it", "the process allocated it some other way" and "it is page cache" is the whole of a memory attribution. Diagnostic only; it changes no numerics. The counters themselves are always maintained (one relaxed atomic per allocation) and are readable from a test through `vt::vulkan::DeviceAllocStatsSnapshot()`. Vulkan-only | | `VT_VULKAN_DISPATCH_STATS` | off | `=1` traces every Vulkan compute submit to stderr (index, shader, workgroup count) BEFORE its fence wait, prints any wait over 200 ms, reports a running dispatch rate every 100 submits, and dumps a per-shader histogram at exit. Printing before the wait is what makes a HANG visible: a post-wait print never runs if the fence never signals, so the last line names the dispatch that hung. This is how the coopmat out-of-bounds load was found. Diagnostic only; it changes no numerics. Vulkan-only | | `VT_VULKAN_GEMV_UNROLL` | 4 | `=1` forces the un-unrolled decode GEMV body. Four independent accumulators keep four reads per lane in flight instead of one -- memory-level parallelism, not instruction count. It rides a specialization constant, so both arms are the same committed module and A/B in one binary. MEASURED **1.055x, 7 of 8 interleaved pairs**. Worth noting it measured 5/8 and was REVERTED earlier the same day: that test ran while the GPU was only 26% busy, where a 10% GEMV win moves e2e by 1.4% and is unresolvable against this box's noise. After the ring fix made the run GPU-bound the same code reads 7/8. A negative result is regime-dependent. Vulkan-only | diff --git a/src/vt/cuda/cuda_backend.cu b/src/vt/cuda/cuda_backend.cu index ffc0e474af..af0d0e9c08 100644 --- a/src/vt/cuda/cuda_backend.cu +++ b/src/vt/cuda/cuda_backend.cu @@ -125,7 +125,89 @@ class CudaBackend final : public Backend { static const bool e = std::getenv("VT_CUDA_ALLOC_STATS") != nullptr; return e; } + // VT_CUDA_ALLOC_TRACE: print every allocation of at least 16 MiB with its + // size, live bytes and free device memory, every free of the same size, and + // the request that fails, so a first-forward OOM is a size list rather than a + // guess. Zero-cost when unset. + static bool AllocTraceEnabled() { + static const bool e = [] { + const char* v = std::getenv("VT_CUDA_ALLOC_TRACE"); + return v != nullptr && v[0] != '0'; + }(); + return e; + } + struct AllocTrace { + std::mutex mu; + std::unordered_map sizes; + long long live = 0; + long long calls = 0; + }; + static AllocTrace& Trace() { + static AllocTrace t; + return t; + } + void TraceAlloc(void* p, size_t bytes) { + auto& t = Trace(); + std::lock_guard lk(t.mu); + t.sizes[p] = bytes; + t.live += static_cast(bytes); + const long long live = t.live; + const long long n = ++t.calls; + size_t free_b = 0, tot_b = 0; + DeviceMemoryInfo(&free_b, &tot_b); + const double mib = 1024.0 * 1024.0; + if (bytes >= (16ULL << 20) || free_b < (2ULL << 30)) { + std::fprintf(stderr, + "[cuda-alloc] #%lld size=%.1f MiB live=%.2f GiB free=%.2f GiB\n", + n, static_cast(bytes) / mib, + static_cast(live) / (1024.0 * mib), + static_cast(free_b) / (1024.0 * mib)); + std::fflush(stderr); + } + } + void TraceFree(void* p) { + auto& t = Trace(); + size_t bytes = 0; + long long live = 0; + { + std::lock_guard lk(t.mu); + auto it = t.sizes.find(p); + if (it == t.sizes.end()) return; + bytes = it->second; + t.sizes.erase(it); + t.live -= static_cast(bytes); + live = t.live; + } + size_t free_b = 0, tot_b = 0; + DeviceMemoryInfo(&free_b, &tot_b); + const double mib = 1024.0 * 1024.0; + if (bytes >= (16ULL << 20)) { + std::fprintf(stderr, + "[cuda-free] size=%.1f MiB live=%.2f GiB free=%.2f GiB\n", + static_cast(bytes) / mib, + static_cast(live) / (1024.0 * mib), + static_cast(free_b) / (1024.0 * mib)); + std::fflush(stderr); + } + } void* Alloc(size_t bytes) override { + if (AllocTraceEnabled()) { + void* p = nullptr; + cudaError_t err = cudaMalloc(&p, bytes); + if (err != cudaSuccess) { + size_t free_b = 0, tot_b = 0; + DeviceMemoryInfo(&free_b, &tot_b); + std::fprintf(stderr, + "[cuda-alloc] FAILED size=%.1f MiB free=%.2f GiB err=%s\n", + static_cast(bytes) / (1024.0 * 1024.0), + static_cast(free_b) / (1024.0 * 1024.0 * 1024.0), + cudaGetErrorString(err)); + std::fflush(stderr); + Check(err, "cudaMalloc"); + } + TraceAlloc(p, bytes); + return p; + } void* p = nullptr; Check(cudaMalloc(&p, bytes), "cudaMalloc"); if (StatsEnabled()) { @@ -135,6 +217,7 @@ class CudaBackend final : public Backend { } void Free(void* p) override { if (p == nullptr) return; + if (AllocTraceEnabled()) TraceFree(p); Check(cudaFree(p), "cudaFree"); if (StatsEnabled()) { Stats().frees.fetch_add(1, std::memory_order_relaxed); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index e2c9d553af..d4648611e8 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2831,6 +2831,9 @@ if(TARGET vllm_shared) endif() endif() vllm_cpp_add_test(test_cuda_backend vt/test_cuda_backend.cpp) +# BACKEND-CUDA-SM120: VT_CUDA_ALLOC_TRACE prints each large cudaMalloc/cudaFree +# with the live-bytes and free-memory curves, and the request that fails. +vllm_cpp_add_test(test_cuda_alloc_trace vt/test_cuda_alloc_trace.cpp) vllm_cpp_add_test(test_dropin_abi vt/test_dropin_abi.cpp) target_include_directories(test_dropin_abi PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_cuda_ops vt/test_cuda_ops.cpp) diff --git a/tests/vt/test_cuda_alloc_trace.cpp b/tests/vt/test_cuda_alloc_trace.cpp new file mode 100644 index 0000000000..840f6db298 --- /dev/null +++ b/tests/vt/test_cuda_alloc_trace.cpp @@ -0,0 +1,115 @@ +// VT_CUDA_ALLOC_TRACE: the per-allocation instrument. +// +// A first-forward OOM names one failing size and nothing else, which is not +// enough to tell which allocation sequence produced the pressure on a 24 GiB +// card. This gate pins the three lines the instrument adds: a large allocation +// with its size, the matching free, and the failing request with its error. It +// also asserts its OWN capture with a sentinel, so a redirect that silently +// caught nothing fails here instead of reading as "the trace printed nothing". +// +// The flag is read ONCE per process (a function-static in CudaBackend), so this +// binary turns it on in a global initializer, before main and thus before any +// TEST_CASE. The flag-OFF path is the same code with the guard false and is +// covered by every CUDA test in the tree, none of which sets the variable. +#include + +#include +#include +#include +#include +#include + +#if defined(__unix__) || defined(__APPLE__) +#include +#define VLLM_ALLOC_TRACE_CAPTURE 1 +#endif + +#include "vt/backend.h" +#include "vt/dtype.h" + +namespace { + +struct EnableAllocTrace { + EnableAllocTrace() { ::setenv("VT_CUDA_ALLOC_TRACE", "1", 1); } +}; +const EnableAllocTrace g_enable_alloc_trace; + +bool HasCuda() { + try { + vt::GetBackend(vt::DeviceType::kCUDA); + return true; + } catch (const std::runtime_error&) { + return false; + } +} + +#ifdef VLLM_ALLOC_TRACE_CAPTURE +// Everything written to stderr while `body` runs, as a string. +template +std::string CaptureStderr(F&& body) { + std::fflush(stderr); + int saved = ::dup(2); + char path[] = "/tmp/vllm_alloctrace_XXXXXX"; + int fd = ::mkstemp(path); + REQUIRE(saved >= 0); + REQUIRE(fd >= 0); + ::dup2(fd, 2); + body(); + std::fflush(stderr); + ::dup2(saved, 2); + ::close(saved); + ::lseek(fd, 0, SEEK_SET); + std::string out; + char buf[4096]; + ssize_t n = 0; + while ((n = ::read(fd, buf, sizeof(buf))) > 0) out.append(buf, static_cast(n)); + ::close(fd); + ::unlink(path); + return out; +} +#endif + +} // namespace + +TEST_CASE("cuda alloc trace: a large allocation, its free, and the failing request") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } + vt::Backend& gpu = vt::GetBackend(vt::DeviceType::kCUDA); + const size_t big = size_t{32} << 20; // 32 MiB, above the 16 MiB report floor + + void* p = nullptr; + std::string alloc_line; +#ifdef VLLM_ALLOC_TRACE_CAPTURE + alloc_line = CaptureStderr([&] { + std::fputs("VLLM_ALLOC_TRACE_SENTINEL\n", stderr); + p = gpu.Alloc(big); + }); + // The capture caught what we wrote. Without this the assertions below could + // read an empty string and report "the trace printed nothing" as a pass. + CHECK(alloc_line.find("VLLM_ALLOC_TRACE_SENTINEL") != std::string::npos); +#else + p = gpu.Alloc(big); +#endif + REQUIRE(p != nullptr); + CHECK(alloc_line.find("[cuda-alloc]") != std::string::npos); + CHECK(alloc_line.find("size=32.0 MiB") != std::string::npos); + +#ifdef VLLM_ALLOC_TRACE_CAPTURE + const std::string free_line = CaptureStderr([&] { gpu.Free(p); }); + CHECK(free_line.find("[cuda-free]") != std::string::npos); + CHECK(free_line.find("size=32.0 MiB") != std::string::npos); + + // A request that cannot fit anywhere: the failure names its own size and the + // free memory at that instant, and it still throws. + const std::string failed_line = CaptureStderr([&] { + CHECK_THROWS_AS(gpu.Alloc(size_t{1} << 50), std::runtime_error); + }); + CHECK(failed_line.find("FAILED") != std::string::npos); + CHECK(failed_line.find("size=") != std::string::npos); + CHECK(failed_line.find("err=") != std::string::npos); +#else + gpu.Free(p); +#endif +}