leonarski_fandClaude Fable 5.1
060a89cd7d
A long-running broker began cancelling every data collection with
Device decoding failed (CUDA (GPU) error (out of memory)), falling back to host decompression
CUDA (GPU) error (out of memory)
while nvidia-smi showed the cards less than a fifth full. Two defects, both in CUDAMemHelpers.h and
both from the pooled allocator (cudaMallocAsync) that came with rc.162.
The leak. cuda_allocation_stream() kept one stream per (thread, device) in a thread_local map of raw
cudaStream_t and never destroyed them. That was written against rugnux, where the worker threads live
as long as the process. The broker starts fresh std::async threads for every data collection - 16 or
64 of them - so every collection left that many streams behind. Measured: 0.56 MB of device memory
per leaked stream, linear to 4928 streams, never returned, with nothing on the host side growing.
The RAII wrapper (CudaStream) was there but not used at this site, and using it as-is - a stream
destroyed when its thread exits - would not have been safe: ShadowFinder builds its GPU accumulator
on a throw-away std::async thread and frees it from another thread long after, and that free is
ordered on the allocating thread's stream. So the streams are still never destroyed, but a thread
now only borrows one: CudaStream objects live in a process-wide per-device idle list, a thread takes
one on first use and hands it back when it exits. Their number is bounded by the threads that were
ever alive at once instead of by the threads ever started. The list itself is deliberately leaked, so
that nothing calls into CUDA during static destruction.
Replaying the broker's pattern against the real header, 60 collections of 64 threads:
before 3840 streams, 260 -> 2424 MB of device memory
after 64 streams, 260 -> 358 MB
The stale error. Every helper here throws a named message, yet the log carried the raw CUDA string,
so the failure came through a cuda_err() and not from an allocation. CudaDevicePtr falls back to
cudaMalloc when cudaMallocAsync fails, silently - but the failed call stays behind as the thread's
last error, and the cudaGetLastError() that follows the next kernel launch reports it. The buffers
were all allocated; the frame was lost anyway, once on the device-decode route (caught, hence the
warning) and once more on the host fallback (fatal). The pooled attempt failing, and a stream that
cannot be created, are both handled by falling back, so both now clear the error they leave.
What finite resource the production cards ran out of at under 4 GB used was not established - no
cap on the number of streams was found up to 4928 on the card this was measured on. The leak is the
only thing on this path that grows with uptime.
tests/CUDAMemHelpersTest.cpp: later threads end up on the same stream, concurrent threads on
different ones, a buffer is freed cleanly after its allocating thread has exited (and another has
borrowed its stream), and a pool that cannot serve a request leaves no error behind - the last by
capping a memory pool at 4 MB so that the pooled attempt fails and the fallback succeeds.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
104 lines
3.7 KiB
C++
104 lines
3.7 KiB
C++
// SPDX-FileCopyrightText: 2026 Filip Leonarski, Paul Scherrer Institute
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#include <catch2/catch_all.hpp>
|
|
#include "../common/CUDAWrapper.h"
|
|
|
|
#ifdef JFJOCH_USE_CUDA
|
|
|
|
#include <future>
|
|
#include <latch>
|
|
#include <set>
|
|
#include <thread>
|
|
|
|
#include "../image_analysis/indexing/CUDAMemHelpers.h"
|
|
|
|
// The broker starts fresh worker threads for every data collection. A stream per thread ever started
|
|
// is a device-memory leak that ends a long-running broker, so threads that follow one another have to
|
|
// end up on the same stream.
|
|
TEST_CASE("CudaAllocationStream_ReusedByLaterThreads", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
std::set<cudaStream_t> seen;
|
|
for (int i = 0; i < 100; i++)
|
|
std::thread([&seen] {
|
|
CudaDevicePtr<uint8_t> buffer(1 << 16);
|
|
seen.insert(cuda_allocation_stream());
|
|
}).join();
|
|
|
|
CHECK(seen.size() == 1);
|
|
CHECK(*seen.begin() != nullptr);
|
|
}
|
|
|
|
TEST_CASE("CudaAllocationStream_DistinctForConcurrentThreads", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
// Neither thread may exit - and hand its stream back - before the other has taken its own.
|
|
std::latch both_have_one(2);
|
|
auto take = [&both_have_one] {
|
|
const cudaStream_t stream = cuda_allocation_stream();
|
|
both_have_one.count_down();
|
|
both_have_one.wait();
|
|
return stream;
|
|
};
|
|
auto a = std::async(std::launch::async, take);
|
|
auto b = std::async(std::launch::async, take);
|
|
const cudaStream_t stream_a = a.get();
|
|
const cudaStream_t stream_b = b.get();
|
|
|
|
CHECK(stream_a != nullptr);
|
|
CHECK(stream_b != nullptr);
|
|
CHECK(stream_a != stream_b);
|
|
}
|
|
|
|
// The shadow finder builds its GPU accumulator on a throw-away thread and frees it from another one
|
|
// long after. The free is ordered on the allocating thread's stream, so that stream has to outlive
|
|
// the thread - including while a later thread has borrowed it.
|
|
TEST_CASE("CudaDevicePtr_FreedAfterAllocatingThreadExited", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
auto buffer = std::async(std::launch::async, [] {
|
|
return std::make_unique<CudaDevicePtr<uint8_t>>(1 << 20);
|
|
}).get();
|
|
std::thread([] { CudaDevicePtr<uint8_t> other(1 << 16); }).join();
|
|
|
|
REQUIRE(cudaMemset(buffer->get(), 0, 1 << 20) == cudaSuccess);
|
|
buffer.reset();
|
|
CHECK(cudaDeviceSynchronize() == cudaSuccess);
|
|
CHECK(cudaGetLastError() == cudaSuccess);
|
|
}
|
|
|
|
// A pool that cannot serve the request is not an error: the buffer comes from cudaMalloc instead. It
|
|
// must then not stay behind as the thread's last error, or the cudaGetLastError() after the next
|
|
// kernel launch throws "out of memory" over work that went fine.
|
|
TEST_CASE("CudaDevicePtr_HandledPoolFailureLeavesNoError", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
int device = 0;
|
|
REQUIRE(cudaGetDevice(&device) == cudaSuccess);
|
|
|
|
cudaMemPoolProps props{};
|
|
props.allocType = cudaMemAllocationTypePinned;
|
|
props.location.type = cudaMemLocationTypeDevice;
|
|
props.location.id = device;
|
|
props.maxSize = 4 << 20;
|
|
|
|
cudaMemPool_t previous = nullptr, capped = nullptr;
|
|
REQUIRE(cudaDeviceGetMemPool(&previous, device) == cudaSuccess);
|
|
REQUIRE(cudaMemPoolCreate(&capped, &props) == cudaSuccess);
|
|
REQUIRE(cudaDeviceSetMemPool(device, capped) == cudaSuccess);
|
|
{
|
|
CudaDevicePtr<uint8_t> buffer(64 << 20);
|
|
CHECK(buffer.get() != nullptr);
|
|
CHECK(cudaGetLastError() == cudaSuccess);
|
|
}
|
|
REQUIRE(cudaDeviceSetMemPool(device, previous) == cudaSuccess);
|
|
cudaMemPoolDestroy(capped);
|
|
}
|
|
|
|
#endif
|