Add CUDA plugin EP Sync support for IOBinding - #27919
Conversation
There was a problem hiding this comment.
Pull request overview
Wires the CUDA plugin execution provider into ONNX Runtime’s sync surface so IOBinding can coordinate CUDA device work safely, and adds Python tests plus design documentation updates to cover the new EP-level callbacks.
Changes:
- Register and implement
OrtEp::CreateSyncStreamForDeviceandOrtEp::SyncinCudaEp(per-session stream creation + device synchronization). - Add Python
IOBindingtests (Add/MatMul) that bind CUDA inputs/outputs and exercise the plugin EP sync path. - Update the CUDA plugin EP design doc to clarify EP-level vs factory-level sync responsibilities.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 5 comments.
| File | Description |
|---|---|
| onnxruntime/core/providers/cuda/plugin/cuda_ep.cc | Adds EP-level sync-stream creation and EP-level sync implementation. |
| onnxruntime/core/providers/cuda/plugin/cuda_ep.h | Declares new EP-level callback entry points. |
| onnxruntime/test/python/transformers/test_cuda_plugin_ep.py | Adds CUDA device-id helper and new IOBinding-based tests. |
| docs/cuda_plugin_ep/cuda_plugin_ep_design.md | Documents the new sync ownership model and test coverage. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Dmitri Smirnov (yuslepukhin)
left a comment
Dmitri Smirnov (yuslepukhin)
left a comment
There was a problem hiding this comment.
See OrtStatus comments.
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Need copilot comment addressed
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 7 out of 7 changed files in this pull request and generated no new comments.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
…ghts_are_all_positive_ logic PR #27552 removed the weights_are_all_positive_ flag which determined whether to use the > 0.5 threshold with 1-score complement (write_additional_scores 0/1) or the > 0 threshold with -score complement (write_additional_scores 2/3). For RandomForestClassifier models where all leaf weights are non-negative (representing probabilities), removing this distinction caused: - Wrong probabilities (negative values instead of 1-p complement) - Wrong labels (when score is between 0 and 0.5) The LOGISTIC post_transform case (which issue #27533 was about) handles both paths identically (applying sigmoid to score and -score), so restoring this flag does not break the fix for #27533. Fixes #27919
Description
This change wires the CUDA plugin EP into ORT's sync surface (see #27538) so
IOBindingcan safely coordinate device work when inputs and outputs are bound on CUDA. It also clarifies the split between EP-level and factory-level sync-stream creation in the design doc and adds Python coverage to validate the new path with simple CUDA-bound models.Summary of Changes
CUDA plugin EP implementation
onnxruntime/core/providers/cuda/plugin/cuda_ep.ccOrtEp::CreateSyncStreamForDeviceandOrtEp::SyncinCudaEp, adds per-session CUDA sync-stream creation, and implements a conservative device-wide sync viacudaSetDevice+cudaDeviceSynchronize().onnxruntime/core/providers/cuda/plugin/cuda_ep.hCreateSyncStreamForDeviceImplandSyncImplentry points onCudaEp.Tests
onnxruntime/test/python/transformers/test_cuda_plugin_ep.pyIOBinding-based Add and MatMul tests that bind CUDA inputs/outputs and exercise the plugin EP sync path.Documentation
docs/cuda_plugin_ep/cuda_plugin_ep_design.mdCudaEpowns the preferredOrtEp::CreateSyncStreamForDeviceandOrtEp::Syncimplementations, whileCudaEpFactory::CreateSyncStreamForDeviceremains a fallback path; also records the newIOBindingtest coverage.Testing
ORT_CUDA_PLUGIN_PATHto the rebuilt CUDA plugin library underbuild/cuda/Releaseand runpython -m pytest onnxruntime/test/python/transformers/test_cuda_plugin_ep.py.IOBindingAdd and MatMul tests pass with CUDA-boundOrtValueinputs and outputs.IOBindingexecution paths.Motivation and Context
IOBindingrelies on provider synchronization to ensure asynchronous device copies are complete before dependent kernel execution continues. The CUDA plugin EP already supported sync-stream creation at the factory layer, but the staged changes connect the per-sessionOrtEpcallbacks that ORT prefers when coordinating bound CUDA execution. The documentation updates make that ownership model explicit so future plugin work does not conflate the fallback factory hook with the primary EP hook.Checklist