Merge pull request #11064 from tomoaki0705:fixCudaStreamAsync

modules/core/include/opencv2/core/private.cuda.hpp

@@ -108,6 +108,8 @@ static inline void throw_no_cuda() { CV_Error(cv::Error::GpuNotSupported, "The l
 
 #else // HAVE_CUDA
 
+#define nppSafeSetStream(oldStream, newStream) { if(oldStream != newStream) { cudaStreamSynchronize(oldStream); nppSetStream(newStream); } }
+
 static inline void throw_no_cuda() { CV_Error(cv::Error::StsNotImplemented, "The called functionality is disabled for current build or platform"); }
 
 namespace cv { namespace cuda
@@ -139,13 +141,13 @@ namespace cv { namespace cuda
         inline explicit NppStreamHandler(Stream& newStream)
         {
             oldStream = nppGetStream();
-            nppSetStream(StreamAccessor::getStream(newStream));
+            nppSafeSetStream(oldStream, StreamAccessor::getStream(newStream));
         }
 
         inline explicit NppStreamHandler(cudaStream_t newStream)
         {
             oldStream = nppGetStream();
-            nppSetStream(newStream);
+            nppSafeSetStream(oldStream, newStream);
         }
 
         inline ~NppStreamHandler()

modules/cudaarithm/src/reductions.cpp

@@ -137,11 +137,12 @@ void cv::cuda::meanStdDev(InputArray _src, OutputArray _dst, Stream& stream)
     if (!deviceSupports(FEATURE_SET_COMPUTE_13))
         CV_Error(cv::Error::StsNotImplemented, "Not sufficient compute capebility");
 
-    const GpuMat src = getInputMat(_src, stream);
+    GpuMat src = getInputMat(_src, stream);
 
     CV_Assert( src.type() == CV_8UC1 );
 
-    GpuMat dst = getOutputMat(_dst, 1, 2, CV_64FC1, stream);
+    _dst.create(1, 2, CV_64FC1);
+    GpuMat dst = _dst.getGpuMat();
 
     NppiSize sz;
     sz.width  = src.cols;
@@ -157,7 +158,8 @@ void cv::cuda::meanStdDev(InputArray _src, OutputArray _dst, Stream& stream)
     BufferPool pool(stream);
     GpuMat buf = pool.getBuffer(1, bufSize, CV_8UC1);
 
-    NppStreamHandler h(StreamAccessor::getStream(stream));
+    // detail: https://github.com/opencv/opencv/issues/11063
+    //NppStreamHandler h(StreamAccessor::getStream(stream));
 
     nppSafeCall( nppiMean_StdDev_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step), sz, buf.ptr<Npp8u>(), dst.ptr<Npp64f>(), dst.ptr<Npp64f>() + 1) );
 

modules/cudaimgproc/src/histogram.cpp

@@ -107,7 +107,7 @@ namespace hist
 
 void cv::cuda::equalizeHist(InputArray _src, OutputArray _dst, Stream& _stream)
 {
-    GpuMat src = _src.getGpuMat();
+    GpuMat src = getInputMat(_src, _stream);
 
     CV_Assert( src.type() == CV_8UC1 );
 

modules/cudaimgproc/test/test_histogram.cpp

@@ -194,6 +194,23 @@ PARAM_TEST_CASE(EqualizeHist, cv::cuda::DeviceInfo, cv::Size)
     }
 };
 
+CUDA_TEST_P(EqualizeHist, Async)
+{
+    cv::Mat src = randomMat(size, CV_8UC1);
+
+    cv::cuda::Stream stream;
+
+    cv::cuda::GpuMat dst;
+    cv::cuda::equalizeHist(loadMat(src), dst, stream);
+
+    stream.waitForCompletion();
+
+    cv::Mat dst_gold;
+    cv::equalizeHist(src, dst_gold);
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 3.0);
+}
+
 CUDA_TEST_P(EqualizeHist, Accuracy)
 {
     cv::Mat src = randomMat(size, CV_8UC1);