From fc51473b094bfcee521b13a0972651e7cb97b054 Mon Sep 17 00:00:00 2001 From: Scott McKay Date: Wed, 22 Jan 2020 14:21:21 +1000 Subject: [PATCH] Update BFCArena logic to use backoff if cudaMalloc fails. Makes behaviour equivalent to when a CPU allocation fails. Add unit test. (#2748) Clear error when throwing an exception for a failed CUDA call so that there is only one error mechanism being used at a time. Minor improvements to logging to aid debugging of BFCArena behaviour. --- onnxruntime/core/framework/bfc_arena.cc | 65 ++++++++++++------- onnxruntime/core/providers/cuda/cuda_call.cc | 2 + onnxruntime/core/providers/cuda/cuda_common.h | 41 +++++++++--- .../providers/cuda/cuda_execution_provider.cc | 34 ++++++++-- .../framework/cuda/allocator_cuda_test.cc | 40 ++++++++++++ 5 files changed, 148 insertions(+), 34 deletions(-) diff --git a/onnxruntime/core/framework/bfc_arena.cc b/onnxruntime/core/framework/bfc_arena.cc index c6a569e5d6..2da45ce441 100644 --- a/onnxruntime/core/framework/bfc_arena.cc +++ b/onnxruntime/core/framework/bfc_arena.cc @@ -9,7 +9,9 @@ BFCArena::BFCArena(std::unique_ptr resource_allocator, : device_allocator_(std::move(resource_allocator)), free_chunks_list_(kInvalidChunkHandle), next_allocation_id_(1), - info_(device_allocator_->Info().name, OrtAllocatorType::OrtArenaAllocator, device_allocator_->Info().device, device_allocator_->Info().id, device_allocator_->Info().mem_type) { + info_(device_allocator_->Info().name, OrtAllocatorType::OrtArenaAllocator, + device_allocator_->Info().device, device_allocator_->Info().id, device_allocator_->Info().mem_type) { + LOGS_DEFAULT(INFO) << "Creating BFCArena for " << device_allocator_->Info().name; curr_region_allocation_bytes_ = RoundedBytes(std::min(total_memory, size_t{1048576})); // Allocate the requested amount of memory. @@ -20,10 +22,10 @@ BFCArena::BFCArena(std::unique_ptr resource_allocator, // We create bins to fit all possible ranges that cover the // memory_limit_ starting from allocations up to 256 bytes to // allocations up to (and including) the memory limit. + LOGS_DEFAULT(VERBOSE) << "Creating " << kNumBins << " bins of max chunk size " + << BinNumToSize(0) << " to " << BinNumToSize(kNumBins - 1); for (BinNum b = 0; b < kNumBins; b++) { size_t bin_size = BinNumToSize(b); - LOGS_DEFAULT(INFO) << "Creating bin of max chunk size " - << bin_size; new (BinFromIndex(b)) Bin(this, bin_size); ORT_ENFORCE(BinForSize(bin_size) == BinFromIndex(b)); ORT_ENFORCE(BinForSize(bin_size + 255) == BinFromIndex(b)); @@ -82,6 +84,12 @@ bool BFCArena::Extend(size_t rounded_bytes) { } catch (const std::bad_alloc&) { // attempted allocation can throw std::bad_alloc. we want to treat this the same as if it returned nullptr // so swallow the exception + } catch (const OnnxRuntimeException& ort_exception) { + // swallow if exception is our throw from a failed cudaMalloc call. + // re-throw otherwise. + if (std::string(ort_exception.what()).find("cudaMalloc") == std::string::npos) { + throw; + } } return new_mem; @@ -103,7 +111,7 @@ bool BFCArena::Extend(size_t rounded_bytes) { } if (mem_addr == nullptr) { - return false; + ORT_THROW("Failed to allocate memory for requested buffer of size ", rounded_bytes); } // if we didn't update already, default to growing by 2x next time @@ -210,7 +218,7 @@ size_t BFCArena::AllocatedSize(const void* ptr) { void* BFCArena::AllocateRawInternal(size_t num_bytes, bool dump_log_on_failure) { if (num_bytes == 0) { - LOGS_DEFAULT(WARNING) << "tried to allocate 0 bytes"; + LOGS_DEFAULT(VERBOSE) << "tried to allocate 0 bytes"; return nullptr; } // First, always allocate memory of at least kMinAllocationSize @@ -227,6 +235,9 @@ void* BFCArena::AllocateRawInternal(size_t num_bytes, return ptr; } + LOGS_DEFAULT(INFO) << "Extending BFCArena for " << device_allocator_->Info().name + << ". bin_num:" << bin_num << " rounded_bytes:" << rounded_bytes; + // Try to extend if (Extend(rounded_bytes)) { ptr = FindChunkPtr(bin_num, rounded_bytes, num_bytes); @@ -244,6 +255,7 @@ void* BFCArena::AllocateRawInternal(size_t num_bytes, << ". Current allocation summary follows."; DumpMemoryLog(rounded_bytes); } + return nullptr; } @@ -503,21 +515,29 @@ BFCArena::get_bin_debug_info() { void BFCArena::DumpMemoryLog(size_t num_bytes) { const std::array bin_infos = get_bin_debug_info(); + LOGS_DEFAULT(INFO) << "Allocator:" << device_allocator_->Info().name; + LOGS_DEFAULT(INFO) << "Bin size: Chunks in_use/total (if not zero). Allocated bytes in_use/total. Requested bytes."; + + size_t waste = 0; for (BinNum bin_num = 0; bin_num < kNumBins; bin_num++) { Bin* b = BinFromIndex(bin_num); const BinDebugInfo& bin_info = bin_infos[bin_num]; ORT_ENFORCE(b->free_chunks.size() == bin_info.total_chunks_in_bin - bin_info.total_chunks_in_use); - LOGS_DEFAULT(INFO) << "Bin (" << b->bin_size - << "): \tTotal Chunks: " << bin_info.total_chunks_in_bin - << ", Chunks in use: " << bin_info.total_chunks_in_use << ". " - << bin_info.total_bytes_in_bin - << " allocated for chunks. " - << bin_info.total_bytes_in_use - << " in use in bin. " - << bin_info.total_requested_bytes_in_use - << " client-requested in use in bin."; + if (bin_info.total_chunks_in_bin > 0) { + LOGS_DEFAULT(INFO) << b->bin_size + << ": Chunks " << bin_info.total_chunks_in_use << "/" << bin_info.total_chunks_in_bin + << ". Bytes " + << bin_info.total_bytes_in_use << "/" << bin_info.total_bytes_in_bin << ". " + << "Requested " << bin_info.total_requested_bytes_in_use << "."; + + waste += bin_info.total_bytes_in_use - bin_info.total_requested_bytes_in_use; + } + } + + if (waste > 0) { + LOGS_DEFAULT(INFO) << "Diff between in-use and requested bytes is " << waste; } // Find the bin that we would have liked to allocate in, so we @@ -525,16 +545,17 @@ void BFCArena::DumpMemoryLog(size_t num_bytes) { Bin* b = BinForSize(num_bytes); LOGS_DEFAULT(INFO) << "Bin for " << num_bytes - << " was " << b->bin_size + << " bytes has max bytes of " << b->bin_size << ", Chunk State: "; for (ChunkHandle h : b->free_chunks) { Chunk* c = ChunkFromHandle(h); - LOGS_DEFAULT(INFO) << c->DebugString(this, true); + LOGS_DEFAULT(INFO) << " " << c->DebugString(this, true); } // Next show the chunks that are in use, and also summarize their // number by size. + LOGS_DEFAULT(INFO) << "Overall chunks summary:"; std::map in_use_by_size; for (const auto& region : region_manager_.regions()) { ChunkHandle h = region_manager_.get_handle(region.ptr()); @@ -543,21 +564,21 @@ void BFCArena::DumpMemoryLog(size_t num_bytes) { if (c->in_use()) { in_use_by_size[c->size]++; } - LOGS_DEFAULT(INFO) << (c->in_use() ? "Chunk" : "Free ") << " at " << c->ptr + LOGS_DEFAULT(INFO) << (c->in_use() ? " Chunk" : " Free ") << " at " << c->ptr << " of size " << c->size; h = c->next; } } - LOGS_DEFAULT(INFO) << " Summary of in-use Chunks by size: "; + LOGS_DEFAULT(INFO) << "Summary of in-use chunks by size: "; size_t total_bytes = 0; for (auto& it : in_use_by_size) { - LOGS_DEFAULT(INFO) << it.second << " Chunks of size " << it.first << " totalling " - << it.first * it.second; + LOGS_DEFAULT(INFO) << " " << it.second << " chunks of size " << it.first + << ". Total " << it.first * it.second; total_bytes += (it.first * it.second); } - LOGS_DEFAULT(INFO) << "Sum Total of in-use chunks: " - << total_bytes; + + LOGS_DEFAULT(INFO) << "Sum Total of in-use chunks: " << total_bytes; LOGS_DEFAULT(INFO) << "Stats: \n" << stats_.DebugString(); } diff --git a/onnxruntime/core/providers/cuda/cuda_call.cc b/onnxruntime/core/providers/cuda/cuda_call.cc index b0f2411ef1..1c6c770f98 100644 --- a/onnxruntime/core/providers/cuda/cuda_call.cc +++ b/onnxruntime/core/providers/cuda/cuda_call.cc @@ -88,6 +88,8 @@ bool CudaCall(ERRTYPE retCode, const char* exprString, const char* libName, ERRT hostname, exprString, msg); if (THRW) { + // clear the error as we're throwing an exception with the error info + (void)cudaGetLastError(); ORT_THROW(str); } else { LOGS_DEFAULT(ERROR) << str; diff --git a/onnxruntime/core/providers/cuda/cuda_common.h b/onnxruntime/core/providers/cuda/cuda_common.h index 8807ede3c4..ba030a0a3a 100644 --- a/onnxruntime/core/providers/cuda/cuda_common.h +++ b/onnxruntime/core/providers/cuda/cuda_common.h @@ -16,12 +16,35 @@ namespace onnxruntime { namespace cuda { -#define CUDA_RETURN_IF_ERROR(expr) ORT_RETURN_IF_ERROR(CUDA_CALL(expr) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) -#define CUBLAS_RETURN_IF_ERROR(expr) ORT_RETURN_IF_ERROR(CUBLAS_CALL(expr) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) -#define CUSPARSE_RETURN_IF_ERROR(expr) ORT_RETURN_IF_ERROR(CUSPARSE_CALL(expr) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) -#define CURAND_RETURN_IF_ERROR(expr) ORT_RETURN_IF_ERROR(CURAND_CALL(expr) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) -#define CUDNN_RETURN_IF_ERROR(expr) ORT_RETURN_IF_ERROR(CUDNN_CALL(expr) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) -#define CUDNN2_RETURN_IF_ERROR(expr, m) ORT_RETURN_IF_ERROR(CUDNN_CALL2(expr, m) ? common::Status::OK() : common::Status(common::ONNXRUNTIME, common::FAIL)) +#define CUDA_RETURN_IF_ERROR(expr) \ + ORT_RETURN_IF_ERROR(CUDA_CALL(expr) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUDA error executing ", #expr)) + +#define CUBLAS_RETURN_IF_ERROR(expr) \ + ORT_RETURN_IF_ERROR(CUBLAS_CALL(expr) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUBLAS error executing ", #expr)) + +#define CUSPARSE_RETURN_IF_ERROR(expr) \ + ORT_RETURN_IF_ERROR(CUSPARSE_CALL(expr) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUSPARSE error executing ", #expr)) + +#define CURAND_RETURN_IF_ERROR(expr) \ + ORT_RETURN_IF_ERROR(CURAND_CALL(expr) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CURAND error executing ", #expr)) + +#define CUDNN_RETURN_IF_ERROR(expr) \ + ORT_RETURN_IF_ERROR(CUDNN_CALL(expr) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUDNN error executing ", #expr)) + +#define CUDNN2_RETURN_IF_ERROR(expr, m) \ + ORT_RETURN_IF_ERROR(CUDNN_CALL2(expr, m) \ + ? common::Status::OK() \ + : ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUDNN2 error executing ", #expr)) // ----------------------------------------------------------------------- // Base class for CUDA kernels @@ -41,8 +64,10 @@ class CudaKernel : public OpKernel { // __debugbreak(); if (s.IsOK()) { - // ensure no kernel launch error occurred - CUDA_RETURN_IF_ERROR(cudaGetLastError()); + auto err = cudaGetLastError(); + if (err != cudaSuccess) { + s = ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "CUDA error ", cudaGetErrorName(err), ":", cudaGetErrorString(err)); + } } return s; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 92f7ba87e6..81e49f719c 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -59,20 +59,46 @@ CUDAExecutionProvider::PerThreadContext::PerThreadContext(int device_id) { } CUDAExecutionProvider::PerThreadContext::~PerThreadContext() { - CUBLAS_CALL_THROW(cublasDestroy(cublas_handle_)); - CUDNN_CALL_THROW(cudnnDestroy(cudnn_handle_)); + // dtor shouldn't throw. if something went wrong earlier (e.g. out of CUDA memory) the handles + // here may be bad, and the destroy calls can throw. + // https://isocpp.github.io/CppCoreGuidelines/CppCoreGuidelines#Rc-dtor-noexcept + try { + CUBLAS_CALL(cublasDestroy(cublas_handle_)); + } catch (const std::exception& ex) { + LOGS_DEFAULT(ERROR) << "cublasDestroy threw:" << ex.what(); + } + + try { + CUDNN_CALL(cudnnDestroy(cudnn_handle_)); + } catch (const std::exception& ex) { + LOGS_DEFAULT(ERROR) << "cudnnDestroy threw:" << ex.what(); + } } CUDAExecutionProvider::CUDAExecutionProvider(const CUDAExecutionProviderInfo& info) : IExecutionProvider{onnxruntime::kCudaExecutionProvider}, device_id_(info.device_id) { CUDA_CALL_THROW(cudaSetDevice(device_id_)); + size_t free = 0; + size_t total = 0; + CUDA_CALL_THROW(cudaMemGetInfo(&free, &total)); + DeviceAllocatorRegistrationInfo default_memory_info( - {OrtMemTypeDefault, [](int device_id) { return onnxruntime::make_unique(device_id, CUDA); }, std::numeric_limits::max()}); + {OrtMemTypeDefault, + [](int device_id) { + return onnxruntime::make_unique(device_id, CUDA); + }, + total}); + InsertAllocator(CreateAllocator(default_memory_info, device_id_)); DeviceAllocatorRegistrationInfo pinned_memory_info( - {OrtMemTypeCPUOutput, [](int device_id) { return onnxruntime::make_unique(device_id, CUDA_PINNED); }, std::numeric_limits::max()}); + {OrtMemTypeCPUOutput, + [](int device_id) { + return onnxruntime::make_unique(device_id, CUDA_PINNED); + }, + std::numeric_limits::max()}); + InsertAllocator(CreateAllocator(pinned_memory_info, CPU_ALLOCATOR_DEVICE_ID)); // TODO: this is actually used for the cuda kernels which explicitly ask for inputs from CPU. diff --git a/onnxruntime/test/framework/cuda/allocator_cuda_test.cc b/onnxruntime/test/framework/cuda/allocator_cuda_test.cc index b0388b457a..88d8df69a6 100644 --- a/onnxruntime/test/framework/cuda/allocator_cuda_test.cc +++ b/onnxruntime/test/framework/cuda/allocator_cuda_test.cc @@ -6,6 +6,7 @@ #include "gtest/gtest.h" #include "cuda_runtime.h" #include "core/providers/cuda/cuda_allocator.h" +#include "core/providers/cuda/cuda_common.h" namespace onnxruntime { namespace test { @@ -71,5 +72,44 @@ TEST(AllocatorTest, CUDAAllocatorTest) { cuda_arena->Free(cuda_addr); pinned_allocator->Free(pinned_addr); } + +// test that we fallback to smaller allocations if the growth of the arena exceeds the available memory +TEST(AllocatorTest, CUDAAllocatorFallbackTest) { + int cuda_device_id = 0; + + size_t free = 0; + size_t total = 0; + + CUDA_CALL_THROW(cudaSetDevice(cuda_device_id)); + CUDA_CALL_THROW(cudaMemGetInfo(&free, &total)); + + // need extra test logic if this ever happens. + EXPECT_NE(free, total) << "All memory is free. Test logic does not handle this."; + + DeviceAllocatorRegistrationInfo default_memory_info( + {OrtMemTypeDefault, + [](int id) { return onnxruntime::make_unique(id, CUDA); }, + std::numeric_limits::max()}); + + auto cuda_arena = CreateAllocator(default_memory_info, cuda_device_id); + + // initial allocation that sets the growth size for the next allocation + size_t size = total / 2; + void* cuda_addr_0 = cuda_arena->Alloc(size); + EXPECT_TRUE(cuda_addr_0); + + // this should trigger an allocation equal to the current total, which should fail initially and gradually fall back + // to a smaller block. + size_t next_size = 1024; + + void* cuda_addr_1 = cuda_arena->Alloc(next_size); + EXPECT_TRUE(cuda_addr_1); + cuda_arena->Free(cuda_addr_0); + cuda_arena->Free(cuda_addr_1); + cuda_arena = nullptr; + + auto last_error = cudaGetLastError(); + EXPECT_EQ(last_error, cudaSuccess) << "Last error should be cleared if handled gracefully"; +} } // namespace test } // namespace onnxruntime