From c444b9d76a9fda3775026ed3a4120ae1a6caaeb3 Mon Sep 17 00:00:00 2001 From: KeDengMS Date: Mon, 12 Oct 2020 22:12:05 -0700 Subject: [PATCH] Add CUDA option to run copy in default stream (#5445) * Add CUDA option to run copy in default stream This change fixes #4829. Thanks @maherzog for providing the repro! The bug is caused by memory reuse in BFC arena, where copy and compute stream in CUDA has a racing condition. BFC arena is an arena allocator on top of cudaMalloc/Free to reduce the cost in syncing CPU and GPU when alloc/free. It means when CPU alloc/free the memory, GPU might not finished previous work on the memory, so that CPU and GPU could run asynchronously. This is OK if there's only one stream, where the execution order in CPU and GPU are consistent. For example, if we have two kernels A and B, CPU runs allocA->computeA->freeA->allocB->computeB->freeB, A and B could shares the same memory since computeA and computeB will not have racing as long as they run in the same GPU compute stream. However, if CPU runs allocA->CopyA->freeA->allocB->computeB->freeB, the order of execution in GPU could have copyA happen after computeB, if copy and compute happens in different GPU streams. This change makes copy to run in default compute stream, while adding an option to fall back to previous behavior if there's perf hit. This is a short term fix before BFC arena could support multiple streams. User may use following options to revert to previous behavior: C API: struct OrtCUDAProviderOptions cudaProviderOpt; cudaProviderOpt.do_copy_in_default_stream = false; C++ API: CUDAExecutionProviderInfo cudaEPInfo; cudaEPInfo.do_copy_in_default_stream = false; C# API: pending... Python: import onnxruntime onnxruntime.capi._pybind_state.set_do_copy_in_default_stream(False) * Confirmed the test failes in CI when doing copy in separate stream Revert the test to get CI pass now * Fix Windows test * Address CR --- .../core/session/onnxruntime_c_api.h | 23 +++++++-------- .../providers/cuda/cuda_execution_provider.cc | 5 ++-- .../providers/cuda/cuda_execution_provider.h | 2 ++ .../providers/cuda/cuda_provider_factory.cc | 16 +++++++---- .../core/providers/cuda/gpu_data_transfer.cc | 27 +++++++++++++----- .../core/providers/cuda/gpu_data_transfer.h | 2 +- .../python/onnxruntime_pybind_state.cc | 14 ++++++--- onnxruntime/test/onnx/main.cc | 1 + .../test/perftest/command_args_parser.cc | 6 +++- onnxruntime/test/perftest/ort_test_session.cc | 3 +- .../test/perftest/test_configuration.h | 1 + .../test/python/onnxruntime_test_python.py | 18 ++++++++++++ onnxruntime/test/testdata/issue4829.onnx | Bin 0 -> 873 bytes onnxruntime/test/util/default_providers.cc | 3 +- orttraining/orttraining/models/bert/main.cc | 3 +- orttraining/orttraining/models/gpt2/main.cc | 3 +- orttraining/orttraining/models/mnist/main.cc | 3 +- 17 files changed, 93 insertions(+), 37 deletions(-) create mode 100755 onnxruntime/test/testdata/issue4829.onnx diff --git a/include/onnxruntime/core/session/onnxruntime_c_api.h b/include/onnxruntime/core/session/onnxruntime_c_api.h index c02f6cb300..9e89923c58 100644 --- a/include/onnxruntime/core/session/onnxruntime_c_api.h +++ b/include/onnxruntime/core/session/onnxruntime_c_api.h @@ -7,7 +7,6 @@ #include #include "onnxruntime_session_options_config_keys.h" - // This value is used in structures passed to ORT so that a newer version of ORT will still work with them #define ORT_API_VERSION 5 @@ -261,16 +260,17 @@ typedef enum OrtMemType { } OrtMemType; typedef enum OrtCudnnConvAlgoSearch { - EXHAUSTIVE, // expensive exhaustive benchmarking using cudnnFindConvolutionForwardAlgorithmEx - HEURISTIC, // lightweight heuristic based search using cudnnGetConvolutionForwardAlgorithm_v7 - DEFAULT, // default algorithm using CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM + EXHAUSTIVE, // expensive exhaustive benchmarking using cudnnFindConvolutionForwardAlgorithmEx + HEURISTIC, // lightweight heuristic based search using cudnnGetConvolutionForwardAlgorithm_v7 + DEFAULT, // default algorithm using CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM } OrtCudnnConvAlgoSearch; typedef struct OrtCUDAProviderOptions { - int device_id; // cuda device with id=0 as default device. + int device_id; // cuda device with id=0 as default device. OrtCudnnConvAlgoSearch cudnn_conv_algo_search; // cudnn conv algo search option - size_t cuda_mem_limit; // default cuda memory limitation to maximum finite value of size_t. - int arena_extend_strategy; // default area extend strategy to KNextPowerOfTwo. + size_t cuda_mem_limit; // default cuda memory limitation to maximum finite value of size_t. + int arena_extend_strategy; // default area extend strategy to KNextPowerOfTwo. + int do_copy_in_default_stream; } OrtCUDAProviderOptions; struct OrtApi; @@ -1046,8 +1046,7 @@ struct OrtApi { */ ORT_API2_STATUS(SessionGetProfilingStartTimeNs, _In_ const OrtSession* sess, _Outptr_ uint64_t* out); - -/** + /** * Use this API to configure the global thread pool options to be used in the call to CreateEnvWithGlobalThreadPools. * A value of 0 means ORT will pick the default. * A value of 1 means the invoking thread will be used; no threads will be created in the thread pool. @@ -1077,7 +1076,7 @@ struct OrtApi { */ ORT_API2_STATUS(AddInitializer, _Inout_ OrtSessionOptions* options, _In_z_ const char* name, _In_ const OrtValue* val); - + /** * Creates a custom environment with global threadpools and logger that will be shared across sessions. * Use this in conjunction with DisablePerSessionThreads API or else the session will use @@ -1088,8 +1087,8 @@ struct OrtApi { ORT_API2_STATUS(CreateEnvWithCustomLoggerAndGlobalThreadPools, OrtLoggingFunction logging_function, _In_opt_ void* logger_param, OrtLoggingLevel logging_level, _In_ const char* logid, _In_ const struct OrtThreadingOptions* tp_options, _Outptr_ OrtEnv** out); - #ifdef USE_CUDA - /** +#ifdef USE_CUDA + /** * Append CUDA execution provider */ ORT_API2_STATUS(OrtSessionOptionsAppendExecutionProvider_CUDA, diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 2a8a92b9be..8c9d554984 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -121,7 +121,8 @@ CUDAExecutionProvider::CUDAExecutionProvider(const CUDAExecutionProviderInfo& in device_id_(info.device_id), cuda_mem_limit_(info.cuda_mem_limit), arena_extend_strategy_(info.arena_extend_strategy), - cudnn_conv_algo_(info.cudnn_conv_algo) { + cudnn_conv_algo_(info.cudnn_conv_algo), + do_copy_in_default_stream_(info.do_copy_in_default_stream) { CUDA_CALL_THROW(cudaSetDevice(device_id_)); // must wait GPU idle, otherwise cudaGetDeviceProperties might fail @@ -1811,7 +1812,7 @@ static bool CastNeedFallbackToCPU(const onnxruntime::Node& node) { } std::unique_ptr CUDAExecutionProvider::GetDataTransfer() const { - return onnxruntime::make_unique(); + return onnxruntime::make_unique(do_copy_in_default_stream_); } std::vector> diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.h b/onnxruntime/core/providers/cuda/cuda_execution_provider.h index 139902c731..2f915477af 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.h @@ -25,6 +25,7 @@ struct CUDAExecutionProviderInfo { size_t cuda_mem_limit{std::numeric_limits::max()}; ArenaExtendStrategy arena_extend_strategy{ArenaExtendStrategy::kNextPowerOfTwo}; OrtCudnnConvAlgoSearch cudnn_conv_algo{OrtCudnnConvAlgoSearch::EXHAUSTIVE}; + bool do_copy_in_default_stream{true}; }; // Logical device representation. @@ -87,6 +88,7 @@ private: size_t cuda_mem_limit_; ArenaExtendStrategy arena_extend_strategy_; int cudnn_conv_algo_; + bool do_copy_in_default_stream_; struct DeferredReleaseCPUPtrs { bool recorded = false; diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 7518ed09dd..c12ae89cf0 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -18,11 +18,13 @@ struct CUDAProviderFactory : IExecutionProviderFactory { CUDAProviderFactory(OrtDevice::DeviceId device_id, size_t cuda_mem_limit = std::numeric_limits::max(), ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, - OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE) + OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE, + bool do_copy_in_default_stream = true) : device_id_(device_id), cuda_mem_limit_(cuda_mem_limit), arena_extend_strategy_(arena_extend_strategy), - cudnn_conv_algo_search_(cudnn_conv_algo_search) {} + cudnn_conv_algo_search_(cudnn_conv_algo_search), + do_copy_in_default_stream_(do_copy_in_default_stream) {} ~CUDAProviderFactory() override {} std::unique_ptr CreateProvider() override; @@ -32,6 +34,7 @@ struct CUDAProviderFactory : IExecutionProviderFactory { size_t cuda_mem_limit_; ArenaExtendStrategy arena_extend_strategy_; OrtCudnnConvAlgoSearch cudnn_conv_algo_search_; + bool do_copy_in_default_stream_; }; std::unique_ptr CUDAProviderFactory::CreateProvider() { @@ -40,14 +43,16 @@ std::unique_ptr CUDAProviderFactory::CreateProvider() { info.cuda_mem_limit = cuda_mem_limit_; info.arena_extend_strategy = arena_extend_strategy_; info.cudnn_conv_algo = cudnn_conv_algo_search_; + info.do_copy_in_default_stream = do_copy_in_default_stream_; return onnxruntime::make_unique(info); } std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE, size_t cuda_mem_limit = std::numeric_limits::max(), - ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo) { - return std::make_shared(device_id, cuda_mem_limit, arena_extend_strategy, cudnn_conv_algo_search); + ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, + bool do_copy_in_default_stream = true) { + return std::make_shared(device_id, cuda_mem_limit, arena_extend_strategy, cudnn_conv_algo_search, do_copy_in_default_stream); } } // namespace onnxruntime @@ -61,6 +66,7 @@ ORT_API_STATUS_IMPL(OrtApis::OrtSessionOptionsAppendExecutionProvider_CUDA, _In_ OrtSessionOptions* options, _In_ OrtCUDAProviderOptions* cuda_options) { options->provider_factories.push_back(onnxruntime::CreateExecutionProviderFactory_CUDA(static_cast(cuda_options->device_id), cuda_options->cudnn_conv_algo_search, cuda_options->cuda_mem_limit, - static_cast(cuda_options->arena_extend_strategy))); + static_cast(cuda_options->arena_extend_strategy), + cuda_options->do_copy_in_default_stream)); return nullptr; } diff --git a/onnxruntime/core/providers/cuda/gpu_data_transfer.cc b/onnxruntime/core/providers/cuda/gpu_data_transfer.cc index 5b7603277c..6618688087 100644 --- a/onnxruntime/core/providers/cuda/gpu_data_transfer.cc +++ b/onnxruntime/core/providers/cuda/gpu_data_transfer.cc @@ -4,22 +4,35 @@ #include "core/providers/cuda/gpu_data_transfer.h" #include "cuda_common.h" +// use default stream for copy for now, to avoid racing in BFC arena as in issue #4829 +// note this may cause some models to run slower if there are ops running on CPU +// so we leave it as optional, in case user need the previous behavior +// a full fix to BFC arena is being looked at, and once it's in, we can revert this change namespace onnxruntime { -GPUDataTransfer::GPUDataTransfer() { +GPUDataTransfer::GPUDataTransfer(bool do_copy_in_default_stream) { // create streams, default is nullptr streams_[kCudaStreamDefault] = nullptr; - CUDA_CALL_THROW(cudaStreamCreateWithFlags(&streams_[kCudaStreamCopyIn], cudaStreamNonBlocking)); - CUDA_CALL_THROW(cudaStreamCreateWithFlags(&streams_[kCudaStreamCopyOut], cudaStreamNonBlocking)); + if (do_copy_in_default_stream) { + streams_[kCudaStreamCopyIn] = nullptr; + streams_[kCudaStreamCopyOut] = nullptr; + } else { + CUDA_CALL_THROW(cudaStreamCreateWithFlags(&streams_[kCudaStreamCopyIn], cudaStreamNonBlocking)); + CUDA_CALL_THROW(cudaStreamCreateWithFlags(&streams_[kCudaStreamCopyOut], cudaStreamNonBlocking)); + } } GPUDataTransfer::~GPUDataTransfer() { - CUDA_CALL(cudaStreamDestroy(streams_[kCudaStreamCopyIn])); - CUDA_CALL(cudaStreamDestroy(streams_[kCudaStreamCopyOut])); + if (streams_[kCudaStreamCopyIn] != nullptr) { + CUDA_CALL(cudaStreamDestroy(streams_[kCudaStreamCopyIn])); + } + if (streams_[kCudaStreamCopyOut] != nullptr) { + CUDA_CALL(cudaStreamDestroy(streams_[kCudaStreamCopyOut])); + } } bool GPUDataTransfer::CanCopy(const OrtDevice& src_device, const OrtDevice& dst_device) const { - return src_device.Type() == OrtDevice::GPU || src_device.MemType() == OrtDevice::MemType::CUDA_PINNED - || dst_device.Type() == OrtDevice::GPU || dst_device.MemType() == OrtDevice::MemType::CUDA_PINNED; + return src_device.Type() == OrtDevice::GPU || src_device.MemType() == OrtDevice::MemType::CUDA_PINNED || + dst_device.Type() == OrtDevice::GPU || dst_device.MemType() == OrtDevice::MemType::CUDA_PINNED; } common::Status GPUDataTransfer::CopyTensor(const Tensor& src, Tensor& dst, int exec_queue_id) const { diff --git a/onnxruntime/core/providers/cuda/gpu_data_transfer.h b/onnxruntime/core/providers/cuda/gpu_data_transfer.h index 50ba91c441..055e2a90fd 100644 --- a/onnxruntime/core/providers/cuda/gpu_data_transfer.h +++ b/onnxruntime/core/providers/cuda/gpu_data_transfer.h @@ -17,7 +17,7 @@ enum CUDAStreamType : int { class GPUDataTransfer : public IDataTransfer { public: - GPUDataTransfer(); + GPUDataTransfer(bool do_copy_in_default_stream = true); ~GPUDataTransfer(); bool CanCopy(const OrtDevice& src_device, const OrtDevice& dst_device) const override; diff --git a/onnxruntime/python/onnxruntime_pybind_state.cc b/onnxruntime/python/onnxruntime_pybind_state.cc index 6c8de4b33e..6499389eee 100644 --- a/onnxruntime/python/onnxruntime_pybind_state.cc +++ b/onnxruntime/python/onnxruntime_pybind_state.cc @@ -145,6 +145,8 @@ OrtDevice::DeviceId cuda_device_id = 0; OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE; size_t cuda_mem_limit = std::numeric_limits::max(); onnxruntime::ArenaExtendStrategy arena_extend_strategy = onnxruntime::ArenaExtendStrategy::kNextPowerOfTwo; +bool do_copy_in_default_stream = true; + #endif #ifdef USE_TENSORRT #include "core/providers/tensorrt/tensorrt_provider_factory.h" @@ -189,7 +191,8 @@ std::shared_ptr CreateExecutionProviderFactory_CPU(in std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo_search, size_t cuda_mem_limit, - onnxruntime::ArenaExtendStrategy arena_extend_strategy); + onnxruntime::ArenaExtendStrategy arena_extend_strategy, + bool do_copy_in_default_stream); std::shared_ptr CreateExecutionProviderFactory_Tensorrt(int device_id); std::shared_ptr CreateExecutionProviderFactory_MIGraphX(int device_id); std::shared_ptr CreateExecutionProviderFactory_Dnnl(int use_arena); @@ -626,13 +629,15 @@ static void RegisterExecutionProviders(InferenceSession* sess, const std::vector sess, *onnxruntime::CreateExecutionProviderFactory_CUDA(cuda_provider_options.device_id, cuda_provider_options.cudnn_conv_algo, cuda_provider_options.cuda_mem_limit, - cuda_provider_options.arena_extend_strategy)); + cuda_provider_options.arena_extend_strategy, + cuda_provider_options.do_copy_in_default_stream)); } else { RegisterExecutionProvider( sess, *onnxruntime::CreateExecutionProviderFactory_CUDA(cuda_device_id, cudnn_conv_algo_search, cuda_mem_limit, - arena_extend_strategy)); + arena_extend_strategy, + do_copy_in_default_stream)); } #endif } else if (type == kDnnlExecutionProvider) { @@ -853,7 +858,7 @@ void addGlobalMethods(py::module& m, const Environment& env) { std::vector> factories = { onnxruntime::CreateExecutionProviderFactory_CPU(0), #ifdef USE_CUDA - onnxruntime::CreateExecutionProviderFactory_CUDA(cuda_device_id, cudnn_conv_algo_search, cuda_mem_limit, arena_extend_strategy), + onnxruntime::CreateExecutionProviderFactory_CUDA(cuda_device_id, cudnn_conv_algo_search, cuda_mem_limit, arena_extend_strategy, do_copy_in_default_stream), #endif #ifdef USE_DNNL onnxruntime::CreateExecutionProviderFactory_Dnnl(1), @@ -909,6 +914,7 @@ void addGlobalMethods(py::module& m, const Environment& env) { m.def("set_cudnn_conv_algo_search", [](const OrtCudnnConvAlgoSearch algo) { cudnn_conv_algo_search = algo; }); m.def("set_cuda_mem_limit", [](const int64_t limit) { cuda_mem_limit = static_cast(limit); }); m.def("set_arena_extend_strategy", [](const onnxruntime::ArenaExtendStrategy strategy) { arena_extend_strategy = strategy; }); + m.def("set_do_copy_in_default_stream", [](const bool use_single_stream) { do_copy_in_default_stream = use_single_stream; }); #endif } diff --git a/onnxruntime/test/onnx/main.cc b/onnxruntime/test/onnx/main.cc index 07a8436c13..43f253974d 100644 --- a/onnxruntime/test/onnx/main.cc +++ b/onnxruntime/test/onnx/main.cc @@ -336,6 +336,7 @@ int real_main(int argc, char* argv[], Ort::Env& env) { OrtCudnnConvAlgoSearch::EXHAUSTIVE, std::numeric_limits::max(), 0, + true }; Ort::ThrowOnError(sf.OrtSessionOptionsAppendExecutionProvider_CUDA(sf, &cuda_options)); #else diff --git a/onnxruntime/test/perftest/command_args_parser.cc b/onnxruntime/test/perftest/command_args_parser.cc index cc8daff820..50f833a9e2 100644 --- a/onnxruntime/test/perftest/command_args_parser.cc +++ b/onnxruntime/test/perftest/command_args_parser.cc @@ -49,12 +49,13 @@ namespace perftest { "\t\tPlease see onnxruntime_c_api.h (enum GraphOptimizationLevel) for the full list of all optimization levels. \n" "\t-u [optimized_model_path]: Specify the optimized model path for saving.\n" "\t-d [cudnn_conv_algorithm]: Specify CUDNN convolution algothrithms: 0(benchmark), 1(heuristic), 2(default). \n" + "\t-q: [CUDA only] use separate stream for copy. \n" "\t-h: help\n"); } /*static*/ bool CommandLineParser::ParseArguments(PerformanceTestConfig& test_config, int argc, ORTCHAR_T* argv[]) { int ch; - while ((ch = getopt(argc, argv, ORT_TSTR("b:m:e:r:t:p:x:y:c:d:o:u:AMPIvhs"))) != -1) { + while ((ch = getopt(argc, argv, ORT_TSTR("b:m:e:r:t:p:x:y:c:d:o:u:AMPIvhsq"))) != -1) { switch (ch) { case 'm': if (!CompareCString(optarg, ORT_TSTR("duration"))) { @@ -181,6 +182,9 @@ namespace perftest { case 'd': test_config.run_config.cudnn_conv_algo = static_cast(OrtStrtol(optarg, nullptr)); break; + case 'q': + test_config.run_config.do_cuda_copy_in_separate_stream = true; + break; case '?': case 'h': default: diff --git a/onnxruntime/test/perftest/ort_test_session.cc b/onnxruntime/test/perftest/ort_test_session.cc index e399bc359f..73ba1c30c4 100644 --- a/onnxruntime/test/perftest/ort_test_session.cc +++ b/onnxruntime/test/perftest/ort_test_session.cc @@ -48,9 +48,10 @@ OnnxRuntimeTestSession::OnnxRuntimeTestSession(Ort::Env& env, std::random_device #ifdef USE_CUDA OrtCUDAProviderOptions cuda_options{ 0, - OrtCudnnConvAlgoSearch::EXHAUSTIVE, + static_cast(performance_test_config.run_config.cudnn_conv_algo), std::numeric_limits::max(), 0, + !performance_test_config.run_config.do_cuda_copy_in_separate_stream }; Ort::ThrowOnError(session_options.OrtSessionOptionsAppendExecutionProvider_CUDA(session_options, &cuda_options)); #else diff --git a/onnxruntime/test/perftest/test_configuration.h b/onnxruntime/test/perftest/test_configuration.h index cd5b242962..2fb65dcf80 100644 --- a/onnxruntime/test/perftest/test_configuration.h +++ b/onnxruntime/test/perftest/test_configuration.h @@ -51,6 +51,7 @@ struct RunConfig { GraphOptimizationLevel optimization_level{ORT_ENABLE_ALL}; std::basic_string optimized_model_path; int cudnn_conv_algo{0}; + bool do_cuda_copy_in_separate_stream{false}; }; struct PerformanceTestConfig { diff --git a/onnxruntime/test/python/onnxruntime_test_python.py b/onnxruntime/test/python/onnxruntime_test_python.py index f1526a7e70..1d81da6b56 100644 --- a/onnxruntime/test/python/onnxruntime_test_python.py +++ b/onnxruntime/test/python/onnxruntime_test_python.py @@ -765,6 +765,24 @@ class TestInferenceSession(unittest.TestCase): # The constructed OrtValue should still be valid after being used in a session self.assertTrue(np.array_equal(ortvalue2.numpy(), numpy_arr_input)) + def testRunModelWithCudaCopyStream(self): + available_providers = onnxrt.get_available_providers() + + if (not 'CUDAExecutionProvider' in available_providers): + print("Skipping testRunModelWithCudaCopyStream when CUDA is not available") + else: + # adapted from issue #4829 for a race condition when copy is not on default stream + # note: + # 1. if there are intermittent failure in this test, something is wrong + # 2. it's easier to repro on slower GPU (like M60, Geforce 1070) + + # to repro #4829, uncomment the line below to run copy in a separate stream + #onnxrt.capi._pybind_state.set_do_copy_in_default_stream(False) + + session = onnxrt.InferenceSession(get_name("issue4829.onnx")) + shape = np.array([2,2], dtype=np.int64) + for iteration in range(100000): + result = session.run(output_names=['output'], input_feed={'shape': shape}) if __name__ == '__main__': unittest.main() diff --git a/onnxruntime/test/testdata/issue4829.onnx b/onnxruntime/test/testdata/issue4829.onnx new file mode 100755 index 0000000000000000000000000000000000000000..93090cbffbfa7f19e8af8ca5c12b37696ad72909 GIT binary patch literal 873 zcmZ{j!D`z;5Qa0dEv?6fsv8o_(vwIpN}YIROR{q+?xA2Pg_K@;S(LEFDaLlKRY>wQ zd8Ry3Mz$oJU5p= z?vq8P{jTfG*imskLleJXE1f56jUYJ$f_+cSTyK>0>|N$}u$esErM;tPuyTvxo<&e= z0S7Ipd{kKs4_lld1PCJzp=kZc7HZ@*=!WMAA%{qGelOGinbuEQgkc3r=Nd+$uL{fChmXxd??LPO*(uJyNhd{K;;nqy|0;w5NhN4@6 zRWhi>oa+P9M{o4QTK$0Z6N-=`qPkeLcVkgSVBi59bNK4Gyl=jbxvU0Q|wO&(Lr GZNCB{9mw1O literal 0 HcmV?d00001 diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index 5e61fa6a40..324c3dcd48 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -13,7 +13,8 @@ std::shared_ptr CreateExecutionProviderFactory_CPU(in std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo = OrtCudnnConvAlgoSearch::EXHAUSTIVE, size_t cuda_mem_limit = std::numeric_limits::max(), - ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo); + ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, + bool do_copy_in_default_stream = true); std::shared_ptr CreateExecutionProviderFactory_Dnnl(int use_arena); std::shared_ptr CreateExecutionProviderFactory_NGraph(const char* ng_backend_type); std::shared_ptr CreateExecutionProviderFactory_OpenVINO(const char* device_type, bool enable_vpu_fast_compile, const char* device_id); diff --git a/orttraining/orttraining/models/bert/main.cc b/orttraining/orttraining/models/bert/main.cc index a1c46e244c..a31e613178 100644 --- a/orttraining/orttraining/models/bert/main.cc +++ b/orttraining/orttraining/models/bert/main.cc @@ -22,7 +22,8 @@ namespace onnxruntime { std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE, size_t cuda_mem_limit = std::numeric_limits::max(), - onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo); + onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, + bool do_copy_in_default_stream = true); } using namespace onnxruntime; diff --git a/orttraining/orttraining/models/gpt2/main.cc b/orttraining/orttraining/models/gpt2/main.cc index c415bf9578..70610cf1d3 100644 --- a/orttraining/orttraining/models/gpt2/main.cc +++ b/orttraining/orttraining/models/gpt2/main.cc @@ -21,7 +21,8 @@ namespace onnxruntime { std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE, size_t cuda_mem_limit = std::numeric_limits::max(), - onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo); + onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, + bool do_copy_in_default_stream = true); } using namespace onnxruntime; diff --git a/orttraining/orttraining/models/mnist/main.cc b/orttraining/orttraining/models/mnist/main.cc index 52e02df82d..5e9c0502cb 100644 --- a/orttraining/orttraining/models/mnist/main.cc +++ b/orttraining/orttraining/models/mnist/main.cc @@ -21,7 +21,8 @@ namespace onnxruntime { std::shared_ptr CreateExecutionProviderFactory_CUDA(OrtDevice::DeviceId device_id, OrtCudnnConvAlgoSearch cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE, size_t cuda_mem_limit = std::numeric_limits::max(), - onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo); + onnxruntime::ArenaExtendStrategy arena_extend_strategy = ArenaExtendStrategy::kNextPowerOfTwo, + bool do_copy_in_default_stream = true); } using namespace onnxruntime;