From dd4c9f8f38b1f617ae3d39e36c0805efbdd1fdcf Mon Sep 17 00:00:00 2001 From: tonde Date: Tue, 12 May 2026 17:18:52 +0200 Subject: [PATCH] cuda: add simple cudaMalloc/cudaFree allocator as opt-in (issue #2038) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Issue #2038 reports that `del model` deadlocks during cleanup on ROCm 7.2.1 + Windows + gfx1100. The trace points at the hipcub CachingDeviceAllocator's DeviceFree path — both the per-block hipEventRecord (recached branch) and the synchronous hipFree (non-recached branch) call into the ROCm runtime, and at least one of them hangs there indefinitely during a `del model`. The CudaAsyncAllocator already exists as an alternative, but it's disabled on Windows (#1072 comment) and a quick check on the local ROCm 7.2 wheels confirms hipMallocAsync still misbehaves there (invalid vector subscript during the first generate() call), so that's not a usable fallback. Add a third allocator option, "simple" / "none", that is a stateless cudaMalloc / cudaFree wrapper. It has no cache, no per-block ready events, and no per-block associated streams, so it can't trip any of the state-tracking code paths in CachingDeviceAllocator that the upstream runtime bug is sensitive to. The trade-off is that every allocation becomes a fresh cudaMalloc — fine for typical inference workloads that allocate once and reuse, more costly for workloads that allocate often. Selected via `CT2_CUDA_ALLOCATOR=simple` (or `none`). Default behaviour is unchanged. Verified locally on RX 7900 XTX (gfx1100, ROCm 7.2.0, Windows 11): - Whisper-medium load + inference + `del model` succeeds in both allocator modes (deadlock not reproducible on 7.2.0; reporter sees it on 7.2.1). - Existing flash-attention pytest suite (15 tests) passes under `CT2_CUDA_ALLOCATOR=simple` with identical results. --- src/cuda/allocator.cc | 51 +++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 51 insertions(+) diff --git a/src/cuda/allocator.cc b/src/cuda/allocator.cc index 921ed34de..0fb42cc80 100644 --- a/src/cuda/allocator.cc +++ b/src/cuda/allocator.cc @@ -13,6 +13,8 @@ #define cub hipcub #define cudaGetDevice hipGetDevice #define cudaSetDevice hipSetDevice +#define cudaFree hipFree +#define cudaMalloc hipMalloc #define cudaFreeAsync hipFreeAsync #define cudaMallocAsync hipMallocAsync #define cudaDeviceGetAttribute hipDeviceGetAttribute @@ -76,6 +78,45 @@ namespace ctranslate2 { std::unique_ptr _allocator; }; + // Direct cudaMalloc/cudaFree allocator — no caching, no event tracking. + // + // Exists primarily as an opt-in escape hatch for the deadlock on + // ROCm 7.2.1 / Windows reported in issue #2038: there, the hipcub + // CachingDeviceAllocator's per-block hipEventRecord / hipFree calls + // can hang inside the runtime during `del model` cleanup. Falling + // back to a stateless allocator (no cached blocks, no ready events, + // no per-block streams) sidesteps every code path that the upstream + // bug touches. Costs a fresh hipMalloc per allocation — fine for + // workloads that mostly run inference and rarely allocate. + // + // Enabled with `CT2_CUDA_ALLOCATOR=simple` (or `none`). + class SimpleAllocator : public Allocator { + public: + void* allocate(size_t size, int device_index) override { + int prev_device_index = -1; + if (device_index >= 0) { + CUDA_CHECK(cudaGetDevice(&prev_device_index)); + CUDA_CHECK(cudaSetDevice(device_index)); + } + void* ptr = nullptr; + CUDA_CHECK(cudaMalloc(&ptr, size)); + if (prev_device_index >= 0) + CUDA_CHECK(cudaSetDevice(prev_device_index)); + return ptr; + } + + void free(void* ptr, int device_index) override { + int prev_device_index = -1; + if (device_index >= 0) { + CUDA_CHECK(cudaGetDevice(&prev_device_index)); + CUDA_CHECK(cudaSetDevice(device_index)); + } + CUDA_CHECK(cudaFree(ptr)); + if (prev_device_index >= 0) + CUDA_CHECK(cudaSetDevice(prev_device_index)); + } + }; + class CudaAsyncAllocator : public Allocator { public: void* allocate(size_t size, int device_index) override { @@ -139,6 +180,7 @@ namespace ctranslate2 { enum class CudaAllocator { CubCaching, MallocAsync, + Simple, }; static CudaAllocator resolve_cuda_allocator() { @@ -156,6 +198,10 @@ namespace ctranslate2 { if (!cuda_malloc_async_is_supported) throw std::runtime_error("The asynchronous CUDA allocator requires CUDA >= 11.2"); allocator = CudaAllocator::MallocAsync; + } else if (allocator_name == "simple" || allocator_name == "none") { + // Stateless cudaMalloc/cudaFree — opt-in workaround for issue + // #2038 (HIP allocator free path deadlocks on ROCm 7.2.1 / Windows). + allocator = CudaAllocator::Simple; } else { throw std::invalid_argument("Invalid CUDA allocator " + allocator_name); } @@ -180,6 +226,11 @@ namespace ctranslate2 { return allocator; } + if (cuda_allocator == cuda::CudaAllocator::Simple) { + static cuda::SimpleAllocator allocator; + return allocator; + } + static cuda::CudaAsyncAllocator allocator; return allocator; }