From 39854ceda4040eba415f61d132b5a9124f5f9d56 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <vlad.vinogradov@itseez.com>
Date: Thu, 12 Nov 2015 13:38:29 +0300
Subject: [PATCH] cuda::StreamAccessor::wrapStream and
 cuda::EventAccessor::wrapEvent

to import existed CUDA stream or CUDA event to OpenCV
---
 modules/core/include/opencv2/core/cuda.hpp    |  1 +
 .../core/include/opencv2/core/cuda.inl.hpp    | 10 ++++
 .../opencv2/core/cuda_stream_accessor.hpp     |  7 ++-
 modules/core/src/cuda_stream.cpp              | 48 ++++++++++++++-----
 modules/cudaarithm/test/test_stream.cpp       | 22 +++++++++
 5 files changed, 72 insertions(+), 16 deletions(-)

diff --git a/modules/core/include/opencv2/core/cuda.hpp b/modules/core/include/opencv2/core/cuda.hpp
index a9c7a39a8f..64bc53ef51 100644
--- a/modules/core/include/opencv2/core/cuda.hpp
+++ b/modules/core/include/opencv2/core/cuda.hpp
@@ -528,6 +528,7 @@ public:
 
 private:
     Ptr<Impl> impl_;
+    Event(const Ptr<Impl>& impl);
 
     friend struct EventAccessor;
 };
diff --git a/modules/core/include/opencv2/core/cuda.inl.hpp b/modules/core/include/opencv2/core/cuda.inl.hpp
index d9ab2ae4f3..01dc6d7c0b 100644
--- a/modules/core/include/opencv2/core/cuda.inl.hpp
+++ b/modules/core/include/opencv2/core/cuda.inl.hpp
@@ -540,6 +540,16 @@ Stream::Stream(const Ptr<Impl>& impl)
 {
 }
 
+//===================================================================================
+// Event
+//===================================================================================
+
+inline
+Event::Event(const Ptr<Impl>& impl)
+    : impl_(impl)
+{
+}
+
 //===================================================================================
 // Initialization & Info
 //===================================================================================
diff --git a/modules/core/include/opencv2/core/cuda_stream_accessor.hpp b/modules/core/include/opencv2/core/cuda_stream_accessor.hpp
index dd6589bcb6..0f8ee9b2d4 100644
--- a/modules/core/include/opencv2/core/cuda_stream_accessor.hpp
+++ b/modules/core/include/opencv2/core/cuda_stream_accessor.hpp
@@ -52,7 +52,7 @@
  */
 
 #include <cuda_runtime.h>
-#include "opencv2/core/cvdef.h"
+#include "opencv2/core/cuda.hpp"
 
 namespace cv
 {
@@ -62,14 +62,12 @@ namespace cv
 //! @addtogroup cudacore_struct
 //! @{
 
-        class Stream;
-        class Event;
-
         /** @brief Class that enables getting cudaStream_t from cuda::Stream
          */
         struct StreamAccessor
         {
             CV_EXPORTS static cudaStream_t getStream(const Stream& stream);
+            CV_EXPORTS static Stream wrapStream(cudaStream_t stream);
         };
 
         /** @brief Class that enables getting cudaEvent_t from cuda::Event
@@ -77,6 +75,7 @@ namespace cv
         struct EventAccessor
         {
             CV_EXPORTS static cudaEvent_t getEvent(const Event& event);
+            CV_EXPORTS static Event wrapEvent(cudaEvent_t event);
         };
 
 //! @}
diff --git a/modules/core/src/cuda_stream.cpp b/modules/core/src/cuda_stream.cpp
index d3b5545e94..1ea8df37b9 100644
--- a/modules/core/src/cuda_stream.cpp
+++ b/modules/core/src/cuda_stream.cpp
@@ -280,32 +280,37 @@ class cv::cuda::Stream::Impl
 {
 public:
     cudaStream_t stream;
-    Ptr<StackAllocator> stackAllocator_;
+    bool ownStream;
+
+    Ptr<StackAllocator> stackAllocator;
 
     Impl();
-    Impl(cudaStream_t stream);
+    explicit Impl(cudaStream_t stream);
 
     ~Impl();
 };
 
-cv::cuda::Stream::Impl::Impl() : stream(0)
+cv::cuda::Stream::Impl::Impl() : stream(0), ownStream(false)
 {
     cudaSafeCall( cudaStreamCreate(&stream) );
+    ownStream = true;
 
-    stackAllocator_ = makePtr<StackAllocator>(stream);
+    stackAllocator = makePtr<StackAllocator>(stream);
 }
 
-cv::cuda::Stream::Impl::Impl(cudaStream_t stream_) : stream(stream_)
+cv::cuda::Stream::Impl::Impl(cudaStream_t stream_) : stream(stream_), ownStream(false)
 {
-    stackAllocator_ = makePtr<StackAllocator>(stream);
+    stackAllocator = makePtr<StackAllocator>(stream);
 }
 
 cv::cuda::Stream::Impl::~Impl()
 {
-    stackAllocator_.release();
+    stackAllocator.release();
 
-    if (stream)
+    if (stream && ownStream)
+    {
         cudaStreamDestroy(stream);
+    }
 }
 
 #endif
@@ -516,6 +521,11 @@ cudaStream_t cv::cuda::StreamAccessor::getStream(const Stream& stream)
     return stream.impl_->stream;
 }
 
+Stream cv::cuda::StreamAccessor::wrapStream(cudaStream_t stream)
+{
+    return Stream(makePtr<Stream::Impl>(stream));
+}
+
 #endif
 
 /////////////////////////////////////////////////////////////
@@ -660,7 +670,7 @@ void cv::cuda::setBufferPoolConfig(int deviceId, size_t stackSize, int stackCoun
 
 #ifdef HAVE_CUDA
 
-cv::cuda::BufferPool::BufferPool(Stream& stream) : allocator_(stream.impl_->stackAllocator_.get())
+cv::cuda::BufferPool::BufferPool(Stream& stream) : allocator_(stream.impl_->stackAllocator.get())
 {
 }
 
@@ -693,20 +703,29 @@ class cv::cuda::Event::Impl
 {
 public:
     cudaEvent_t event;
+    bool ownEvent;
 
-    Impl(unsigned int flags);
+    explicit Impl(unsigned int flags);
+    explicit Impl(cudaEvent_t event);
     ~Impl();
 };
 
-cv::cuda::Event::Impl::Impl(unsigned int flags) : event(0)
+cv::cuda::Event::Impl::Impl(unsigned int flags) : event(0), ownEvent(false)
 {
     cudaSafeCall( cudaEventCreateWithFlags(&event, flags) );
+    ownEvent = true;
+}
+
+cv::cuda::Event::Impl::Impl(cudaEvent_t e) : event(e), ownEvent(false)
+{
 }
 
 cv::cuda::Event::Impl::~Impl()
 {
-    if (event)
+    if (event && ownEvent)
+    {
         cudaEventDestroy(event);
+    }
 }
 
 cudaEvent_t cv::cuda::EventAccessor::getEvent(const Event& event)
@@ -714,6 +733,11 @@ cudaEvent_t cv::cuda::EventAccessor::getEvent(const Event& event)
     return event.impl_->event;
 }
 
+Event cv::cuda::EventAccessor::wrapEvent(cudaEvent_t event)
+{
+    return Event(makePtr<Event::Impl>(event));
+}
+
 #endif
 
 cv::cuda::Event::Event(CreateFlags flags)
diff --git a/modules/cudaarithm/test/test_stream.cpp b/modules/cudaarithm/test/test_stream.cpp
index c9a5e694f2..785b10e748 100644
--- a/modules/cudaarithm/test/test_stream.cpp
+++ b/modules/cudaarithm/test/test_stream.cpp
@@ -47,6 +47,7 @@
 #include <cuda_runtime.h>
 
 #include "opencv2/core/cuda.hpp"
+#include "opencv2/core/cuda_stream_accessor.hpp"
 #include "opencv2/ts/cuda_test.hpp"
 
 using namespace cvtest;
@@ -129,6 +130,27 @@ CUDA_TEST_P(Async, Convert)
     stream.waitForCompletion();
 }
 
+CUDA_TEST_P(Async, WrapStream)
+{
+    cudaStream_t cuda_stream = NULL;
+    ASSERT_EQ(cudaSuccess, cudaStreamCreate(&cuda_stream));
+
+    {
+        cv::cuda::Stream stream = cv::cuda::StreamAccessor::wrapStream(cuda_stream);
+
+        d_src.upload(src, stream);
+        d_src.convertTo(d_dst, CV_32S, stream);
+        d_dst.download(dst, stream);
+
+        Async* test = this;
+        stream.enqueueHostCallback(checkConvert, test);
+
+        stream.waitForCompletion();
+    }
+
+    ASSERT_EQ(cudaSuccess, cudaStreamDestroy(cuda_stream));
+}
+
 CUDA_TEST_P(Async, HostMemAllocator)
 {
     cv::cuda::Stream stream;