Enable cuda graph in TensorRT EP #10423
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -435,6 +435,7 @@ TensorrtExecutionProvider::TensorrtExecutionProvider(const TensorrtExecutionProv | |
| engine_decryption_lib_path_ = info.engine_decryption_lib_path; | ||
| } | ||
| force_sequential_engine_build_ = info.force_sequential_engine_build; | ||
| cuda_graph_enable_ = info.cuda_graph_enable; | ||
| } else { | ||
| const std::string max_partition_iterations_env = onnxruntime::GetEnvironmentVar(tensorrt_env_vars::kMaxPartitionIterations); | ||
| if (!max_partition_iterations_env.empty()) { | ||
|
|
@@ -519,6 +520,11 @@ TensorrtExecutionProvider::TensorrtExecutionProvider(const TensorrtExecutionProv | |
| if (!force_sequential_engine_build_env.empty()) { | ||
| force_sequential_engine_build_ = (std::stoi(force_sequential_engine_build_env) == 0 ? false : true); | ||
| } | ||
|
|
||
| const std::string cuda_graph_enable_env = onnxruntime::GetEnvironmentVar(tensorrt_env_vars::kCUDAGraphEnable); | ||
| if (!cuda_graph_enable_env.empty()) { | ||
| cuda_graph_enable_ = (std::stoi(cuda_graph_enable_env) == 0 ? false : true); | ||
| } | ||
| } | ||
|
|
||
| // Validate setting | ||
|
|
@@ -579,7 +585,8 @@ TensorrtExecutionProvider::TensorrtExecutionProvider(const TensorrtExecutionProv | |
| << ", trt_cache_path: " << cache_path_ | ||
| << ", trt_engine_decryption_enable: " << engine_decryption_enable_ | ||
| << ", trt_engine_decryption_lib_path: " << engine_decryption_lib_path_ | ||
| << ", trt_force_sequential_engine_build: " << force_sequential_engine_build_ | ||
| << ", trt_cuda_graph_enable: " << cuda_graph_enable_; | ||
| } | ||
|
|
||
| TensorrtExecutionProvider::~TensorrtExecutionProvider() { | ||
|
|
@@ -1159,7 +1166,10 @@ std::unique_lock<OrtMutex> TensorrtExecutionProvider::GetEngineBuildLock() const | |
|
|
||
| common::Status TensorrtExecutionProvider::Compile(const std::vector<Node*>& fused_nodes, | ||
| std::vector<NodeComputeInfo>& node_compute_funcs) { | ||
| int fused_nodes_size = fused_nodes.size(); | ||
| cuda_graphs_.reserve(fused_nodes_size); | ||
| for (int node_idx = 0; node_idx < fused_nodes_size; node_idx++) { | ||
| const auto* fused_node = fused_nodes[node_idx]; | ||
|
There was a problem hiding this comment. In general, how does the TRT EP handle control flow nodes ? I fear we must explicitly not support using cuda graphs for models with control flow nodes as the graph captured for one input may not the same graph required for another input (because of the dynamic graph branching).
Contributor
There was a problem hiding this comment. agreed. we should explicitly exclude graphs with loops/conditionals.
Contributor
Author
There was a problem hiding this comment. We may highlight the constrains for dynamic shape cases in document, so that users can choose to enable cuda graph or not.
Contributor
There was a problem hiding this comment. it would be nice to enforce constraints (doesn't support dynamic shapes and dynamic graphs) in code rather than punting to user/documentation. Let's see if there's a reasonable balance that can be achieved here.
Contributor
Author
There was a problem hiding this comment. Here is a little bit tricky. Dynamic shape model is okay, but when input shapes of incoming data change, cuda graph needs to be recaptured, so the check has to be done in runtime. There is an API to update executable graph, but I haven't seen any APIs that can check existing cuda graph's profile, and we can't afford to update graph for every enqueue. |
||
| // Build map from input name to its index in input definitions | ||
| std::unordered_map<std::string, size_t> input_map; | ||
| const auto& input_defs = fused_node->InputDefs(); | ||
|
|
@@ -1407,6 +1417,7 @@ common::Status TensorrtExecutionProvider::Compile(const std::vector<Node*>& fuse | |
| contexts_.emplace(fused_node->Name(), std::move(trt_context)); | ||
| builders_.emplace(fused_node->Name(), std::move(trt_builder)); | ||
| networks_.emplace(fused_node->Name(), std::move(trt_network)); | ||
| cuda_graph_instances_.emplace(fused_node->Name(), cuda_graphs_[node_idx]); | ||
| input_info_[fused_node->Name()].push_back(input_indexes); | ||
| output_info_[fused_node->Name()].push_back(output_indexes); | ||
| output_info_[fused_node->Name()].push_back(output_types); | ||
|
|
@@ -1421,8 +1432,9 @@ common::Status TensorrtExecutionProvider::Compile(const std::vector<Node*>& fuse | |
| &engines_[context->node_name], &contexts_[context->node_name], &builders_[context->node_name], | ||
| &networks_[context->node_name], input_info_[context->node_name], output_info_[context->node_name], | ||
| input_shape_ranges_[context->node_name], &tensorrt_mu_, fp16_enable_, int8_enable_, int8_calibration_cache_available_, | ||
| dla_enable_, dla_core_, cuda_graph_enable_, nullptr, cuda_graph_instances_[context->node_name], &max_workspace_size_, | ||
| trt_node_name_with_precision, engine_cache_enable_, cache_path_, runtime_.get(), nullptr, allocator_, | ||
| dynamic_range_map, engine_decryption_enable_, engine_decryption_, engine_encryption_}; | ||
| *state = p.release(); | ||
| return 0; | ||
| }; | ||
|
|
@@ -1445,6 +1457,8 @@ common::Status TensorrtExecutionProvider::Compile(const std::vector<Node*>& fuse | |
| auto trt_engine = trt_state->engine->get(); | ||
| auto trt_context = trt_state->context->get(); | ||
| auto trt_profile = &(trt_state->trt_profile); | ||
| auto cuda_graph = &(trt_state->cuda_graph); | ||
|
|
||
| auto alloc = trt_state->scratch_allocator; | ||
| int num_inputs = static_cast<int>(input_indexes.size()); | ||
| int num_outputs = static_cast<int>(output_indexes.size()); | ||
|
|
@@ -2006,10 +2020,29 @@ common::Status TensorrtExecutionProvider::Compile(const std::vector<Node*>& fuse | |
| } | ||
| } | ||
|
|
||
| // Run TRT inference | ||
| if (trt_state->cuda_graph_enable) | ||
| { | ||
| if (*cuda_graph == nullptr) { | ||
|
Contributor
There was a problem hiding this comment. i think this entire path isn't thread safe as cuda graphs and associated api's don't seem to be thread safe. There was a problem hiding this comment. Yes, if cuda graphs are enabled, then Run() will no longer be thread-safe and calls to Run() needs to be serialized either by the caller or ORT itself should perform the graph replay within a critical section.
Contributor
Author
There was a problem hiding this comment. we have lock in the inference and the compute() is serialized already.
Contributor
There was a problem hiding this comment. ah yes tensorrt_mu_ already protects the compute_func body. |
||
| cudaGraph_t graph; | ||
| *cuda_graph = &(trt_state->cuda_graph_instance); | ||
| //warm up for cuda graph capturing | ||
| if (!trt_context->enqueueV2(&buffers[0], stream, nullptr)) { | ||
|
There was a problem hiding this comment. Does enqueueV2() synchronize with the GPU before returning ? If not, we may have to wait for the warm-up tasks queued on the stream to finish before the stream capture...
Contributor
There was a problem hiding this comment. why is this warm up even needed?
Contributor
There was a problem hiding this comment. seems like it's for handling a known issue with dynamic shapes? (please add a comment)
Contributor
Author
There was a problem hiding this comment. warm-up is needed to do initialization (flushing any old context) before graph capturing according to Nvidia. CUDA graph still has issue in some dynamic shape cases.
Contributor
There was a problem hiding this comment. in the note section of the api doc for enqueuev2 https://docs.nvidia.com/deeplearning/tensorrt/api/c_api/classnvinfer1_1_1_i_execution_context.html#a2f4429652736e8ef6e19f433400108c7
Contributor
Author
There was a problem hiding this comment. If shape changed, cuda graph needs to be recaptured, which is not desired because the capturing happens in inference. |
||
| return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "TensorRT EP execution context enqueue failed."); | ||
| } | ||
| cudaStreamBeginCapture(stream, cudaStreamCaptureModeRelaxed ); | ||
|
There was a problem hiding this comment. Should the CUDA calls have CUDA_CALL_THROW() to deal with CUDA call errors if any ? |
||
| if (!trt_context->enqueueV2(&buffers[0], stream, nullptr)) { | ||
| return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "TensorRT EP execution context enqueue failed."); | ||
| } | ||
| cudaStreamEndCapture(stream, &graph); | ||
|
hariharans29 marked this conversation as resolved.
Outdated
|
||
| cudaGraphInstantiate(*cuda_graph, graph, NULL, NULL, 0); | ||
| } | ||
|
Contributor
There was a problem hiding this comment. what happens to graph? do you need to destroy it? else we leak memory? |
||
| cudaGraphLaunch(**cuda_graph, stream); | ||
| } else { | ||
| if (!trt_context->enqueueV2(&buffers[0], stream, nullptr)) { | ||
| return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "TensorRT EP execution context enqueue failed."); | ||
| } | ||
| } | ||
|
|
||
| // Cast INT64 input to INT32 because TensorRT doesn't fully support INT64 | ||
| for (size_t i = 0, end = output_binding_names.size(); i < end; ++i) { | ||
|
|
||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Shouldn't the struct get versioned with the addition of a new option ?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I thought we agreed we shouldn't be updating provider structs in c api anymore (for the very reason that Hari brings up about versioning)
and instead only updating the opaque struct OrtTensorRTProviderOptionsV2
+Chi Lo (@chilo-ms) FYI
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yeah, I think it should become OrtTensorRTProviderOptionsV3 if V2 has shipped with the previous ORT release (This was my understanding)
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
we don't need to update the version of the opaque struct when adding fields right? since it's only accessed via api and not directly. if the newly added field can't be represented as a string, then we would need to add another api to access those.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
IIRC from the time we were discussing it, we need to version it for any mutation (addition or removal of fields). If we add support for a new field without versioning it, doesn't the UpdateTensorRTProviderOptions API behave differently in ORT 1.10 (where the V2 struct won't support the new field) and in ORT 1.11 (where the V2 struct will support the new field) ?
Uh oh!
There was an error while loading. Please reload this page.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
We had a consensus on using key-value strings for provider options that can be represented by string but we didn't explicitly say we support versioning. But we require error reporting for undocumented config keys.
stevenlix (@stevenlix), here is my "enable timing cache" PR, you can reference it to add new field to opaque struct OrtTensorRTProviderOptionsV2
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
No versioning of struct is required beyond V2. The availability or the unavailability of APIs to manipulate the V2 struct provides the versioning.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
"The availability or the unavailability of APIs to manipulate the V2 struct provides the versioning." - What APIs exists to manipulate the V2 struct will be the same in ORT 1.10 and ORT 1.11 won't they ? It is just that the
UpdateTensorRTProviderOptions()API will additionally support one more key (enable_cuda_graph) in 1.11 (which obviously won't be supported in the released 1.10).There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
yes, one uses CreateTensorRTProviderOptions(), UpdateTensorRTProviderOptions() to create and update the struct.
The second api deals with strings, and will recognize a new string key "enable_cuda_graph" in ort 1.11
api signatures don't change.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
talked to Chi Lo (@chilo-ms) offline. He is going to make a separate PR from the timing cache PR for OrtTensorRTProviderOptionsV2. After the PR merged, I will add cuda_graph option in OrtTensorRTProviderOptionsV2.