From 45161610f035ee3700f5d5f958720c69da7c7bed Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 13 Aug 2026 10:37:32 +0000 Subject: [PATCH] =?UTF-8?q?fix:=20thread=5Flocal=20reentrant=20guard=20?= =?UTF-8?q?=E2=80=94=20prevent=20cudaMalloc=20infinite=20recursion?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CUB CachingDeviceAllocator::DeviceAllocate calls cudaMalloc internally on cache miss. Without a guard, our intercepted cudaMalloc recurses into DeviceAllocate → cudaMalloc → DeviceAllocate → segfault. thread_local g_in_allocator flag detects reentrant calls and forwards them directly to the real cudaMalloc/cudaFree via dlsym(RTLD_NEXT). --- .../cccl_preload/cccl_allocator_preload.cu | 20 +++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu b/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu index 47b33448..fb3136a6 100644 --- a/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu +++ b/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu @@ -123,20 +123,32 @@ static void cccl_preload_init() { /* ======================================================================== * cudaMalloc / cudaFree intercepts + * + * CUB's DeviceAllocate internally calls cudaMalloc on cache miss. + * We must detect this reentrant call and forward to the real function, + * otherwise we get infinite recursion → segfault. * ======================================================================== */ +static thread_local bool g_in_allocator = false; + extern "C" cudaError_t cudaMalloc(void** devPtr, size_t size) { - if (!g_preload_active) { + if (!g_preload_active || g_in_allocator) { return get_real_malloc()(devPtr, size); } - return get_allocator().DeviceAllocate(devPtr, size); + g_in_allocator = true; + cudaError_t err = get_allocator().DeviceAllocate(devPtr, size); + g_in_allocator = false; + return err; } extern "C" cudaError_t cudaFree(void* devPtr) { - if (!g_preload_active || devPtr == nullptr) { + if (!g_preload_active || devPtr == nullptr || g_in_allocator) { return get_real_free()(devPtr); } - return get_allocator().DeviceFree(devPtr); + g_in_allocator = true; + cudaError_t err = get_allocator().DeviceFree(devPtr); + g_in_allocator = false; + return err; }