saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-11 17:48:34 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Update BFCArena logic to use backoff if cudaMalloc fails. Makes behaviour equivalent to when a CPU allocation fails. Add unit test. (#2748)

Browse source 
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.
This commit is contained in:
Scott McKay 2020-01-22 14:21:21 +10:00 committed by GitHub
parent 061f10fcd5
commit fc51473b09
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
5 changed files with 148 additions and 34 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

65
onnxruntime/core/framework/bfc_arena.cc
View file

@ -9,7 +9,9 @@ BFCArena::BFCArena(std::unique_ptr<IDeviceAllocator> 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<IDeviceAllocator> 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<BinDebugInfo, kNumBins> 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<size_t, int> 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();
}

2
onnxruntime/core/providers/cuda/cuda_call.cc
View file

@ -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;

41
onnxruntime/core/providers/cuda/cuda_common.h
View file

@ -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;

34
onnxruntime/core/providers/cuda/cuda_execution_provider.cc
View file

@ -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<CUDAAllocator>(device_id, CUDA); }, std::numeric_limits<size_t>::max()});
{OrtMemTypeDefault,
[](int device_id) {
return onnxruntime::make_unique<CUDAAllocator>(device_id, CUDA);
},
total});
InsertAllocator(CreateAllocator(default_memory_info, device_id_));
DeviceAllocatorRegistrationInfo pinned_memory_info(
{OrtMemTypeCPUOutput, [](int device_id) { return onnxruntime::make_unique<CUDAPinnedAllocator>(device_id, CUDA_PINNED); }, std::numeric_limits<size_t>::max()});
{OrtMemTypeCPUOutput,
[](int device_id) {
return onnxruntime::make_unique<CUDAPinnedAllocator>(device_id, CUDA_PINNED);
},
std::numeric_limits<size_t>::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.

40
onnxruntime/test/framework/cuda/allocator_cuda_test.cc
View file

@ -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<CUDAAllocator>(id, CUDA); },
std::numeric_limits<size_t>::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