added async version of postfilter_textureness and prefilter_xsobel, modified async...

diff --git a/modules/gpu/src/cuda/stereobm.cu b/modules/gpu/src/cuda/stereobm.cu
index 73a70b8..5685c54 100644 (file)
--- a/modules/gpu/src/cuda/stereobm.cu
+++ b/modules/gpu/src/cuda/stereobm.cu
@@ -410,7 +410,7 @@ extern "C" __global__ void prefilter_kernel(unsigned char *output, size_t step,
 \r
 namespace cv { namespace gpu  { namespace bm\r
 {\r
-    extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap)\r
+    extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap, const cudaStream_t & stream)\r
     {\r
         cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();\r
         cudaSafeCall( cudaBindTexture2D( 0, stereobm_gpu::texForSobel, input.ptr, desc, input.cols, input.rows, input.step ) );\r
@@ -421,10 +421,18 @@ namespace cv { namespace gpu  { namespace bm
         grid.x = divUp(input.cols, threads.x);\r
         grid.y = divUp(input.rows, threads.y);\r
 \r
-        stereobm_gpu::prefilter_kernel<<<grid, threads>>>(output.ptr, output.step, output.cols, output.rows, prefilterCap);\r
-        cudaSafeCall( cudaThreadSynchronize() );\r
+        if (stream == 0)
+        {
+                       stereobm_gpu::prefilter_kernel<<<grid, threads>>>(output.ptr, output.step, output.cols, output.rows, prefilterCap);\r
+                       cudaSafeCall( cudaThreadSynchronize() );\r
+        }
+        else
+        {
+            stereobm_gpu::prefilter_kernel<<<grid, threads, 0, stream>>>(output.ptr, output.step, output.cols, output.rows, prefilterCap);
+        }
 \r
         cudaSafeCall( cudaUnbindTexture (stereobm_gpu::texForSobel ) );\r
+\r
     }\r
 \r
 }}}\r
@@ -532,7 +540,7 @@ extern "C" __global__ void textureness_kernel(unsigned char *disp, size_t disp_s
 \r
 namespace cv { namespace gpu  { namespace bm\r
 {\r
-    extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp)\r
+    extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp, const cudaStream_t & stream)\r
     {\r
         avgTexturenessThreshold *= winsz * winsz;\r
 \r
@@ -551,8 +559,15 @@ namespace cv { namespace gpu  { namespace bm
 \r
         size_t smem_size = (threads.x + threads.x + (winsz/2) * 2 ) * sizeof(float);\r
 \r
-        stereobm_gpu::textureness_kernel<<<grid, threads, smem_size>>>(disp.ptr, disp.step, winsz, avgTexturenessThreshold, disp.cols, disp.rows);\r
-        cudaSafeCall( cudaThreadSynchronize() );\r
+               if (stream == 0)\r
+               {\r
+                       stereobm_gpu::textureness_kernel<<<grid, threads, smem_size>>>(disp.ptr, disp.step, winsz, avgTexturenessThreshold, disp.cols, disp.rows);\r
+                       cudaSafeCall( cudaThreadSynchronize() );\r
+               }\r
+               else\r
+               {\r
+                       stereobm_gpu::textureness_kernel<<<grid, threads, smem_size, stream>>>(disp.ptr, disp.step, winsz, avgTexturenessThreshold, disp.cols, disp.rows);              \r
+               }\r
 \r
         cudaSafeCall( cudaUnbindTexture (stereobm_gpu::texForTF) );\r
     }\r

diff --git a/modules/gpu/src/stereobm_gpu.cpp b/modules/gpu/src/stereobm_gpu.cpp
index 1c5716f..706e4ef 100644 (file)
--- a/modules/gpu/src/stereobm_gpu.cpp
+++ b/modules/gpu/src/stereobm_gpu.cpp
@@ -62,8 +62,8 @@ namespace cv { namespace gpu
     {\r
         //extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf);\r
         extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf, const cudaStream_t & stream);\r
-        extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap = 31);\r
-        extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avergeTexThreshold, const DevMem2D& disp);\r
+        extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap /*= 31*/, const cudaStream_t & stream);\r
+        extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp, const cudaStream_t & stream);\r
     }\r
 }}\r
 \r
@@ -115,8 +115,8 @@ static void stereo_bm_gpu_operator ( GpuMat& minSSD,  GpuMat& leBuf, GpuMat&  ri
         leBuf.create( left.size(),  left.type());\r
         riBuf.create(right.size(), right.type());\r
 \r
-        bm::prefilter_xsobel( left, leBuf);\r
-        bm::prefilter_xsobel(right, riBuf);\r
+               bm::prefilter_xsobel( left, leBuf, 31, stream);\r
+        bm::prefilter_xsobel(right, riBuf, 31, stream);\r
 \r
         le_for_bm = leBuf;\r
         ri_for_bm = riBuf;\r
@@ -125,7 +125,7 @@ static void stereo_bm_gpu_operator ( GpuMat& minSSD,  GpuMat& leBuf, GpuMat&  ri
     bm::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD, stream);\r
 \r
     if (avergeTexThreshold)\r
-        bm::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity);\r
+        bm::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity, stream);\r
 }\r
 \r
 \r

diff --git a/tests/gpu/src/operator_async_call.cpp b/tests/gpu/src/operator_async_call.cpp
index 6033e67..1ecddc6 100644 (file)
--- a/tests/gpu/src/operator_async_call.cpp
+++ b/tests/gpu/src/operator_async_call.cpp
@@ -99,9 +99,9 @@ bool CV_GpuMatAsyncCallTest::compare_matrix(cv::Mat & cpumat)
     //int64 time = getTickCount();
 
     Stream stream;
-    stream.enqueueCopy(gmat0, gmat1);
-    stream.enqueueCopy(gmat0, gmat2);
-    stream.enqueueCopy(gmat0, gmat3);
+       stream.enqueueMemSet(gmat0, cv::Scalar::all(1), gmat1);
+       stream.enqueueMemSet(gmat0, cv::Scalar::all(1), gmat2);
+       stream.enqueueMemSet(gmat0, cv::Scalar::all(1), gmat3);
     stream.waitForCompletion();
 
     //int64 time1 = getTickCount();

diff --git a/tests/gpu/src/stereo_bm_async.cpp b/tests/gpu/src/stereo_bm_async.cpp
new file mode 100644 (file)
index 0000000..fdbe951
--- /dev/null
+++ b/tests/gpu/src/stereo_bm_async.cpp
@@ -0,0 +1,97 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                        Intel License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of Intel Corporation may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "gputest.hpp"
+#include "highgui.h"
+#include "cv.h"
+
+
+using namespace cv;
+using namespace std;
+
+class CV_GpuMatAsyncCallStereoBMTest : public CvTest
+{
+    public:
+        CV_GpuMatAsyncCallStereoBMTest();
+        ~CV_GpuMatAsyncCallStereoBMTest();
+
+    protected:
+        void run(int);
+};
+
+CV_GpuMatAsyncCallStereoBMTest::CV_GpuMatAsyncCallStereoBMTest(): CvTest( "GPU-MatAsyncCallStereoBM", "asyncStereoBM" )
+{}
+
+CV_GpuMatAsyncCallStereoBMTest::~CV_GpuMatAsyncCallStereoBMTest() {}
+
+void CV_GpuMatAsyncCallStereoBMTest::run( int /* start_from */)
+{
+       cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0);
+       cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0);
+       cv::Mat img_reference = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
+
+    if (img_l.empty() || img_r.empty() || img_reference.empty())
+    {
+        ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
+        return;
+    }
+
+       cv::gpu::GpuMat disp;
+       cv::gpu::StereoBM_GPU bm(0, 128, 19);
+
+       cv::gpu::Stream stream;
+
+       for (size_t i = 0; i < 50; i++)
+       {
+               bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream);
+       }
+
+       stream.waitForCompletion();
+       disp.convertTo(disp, img_reference.type());
+       double norm = cv::norm(disp, img_reference, cv::NORM_INF);
+
+       if (norm >= 100) 
+        ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
+       ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC);
+
+}
+
+CV_GpuMatAsyncCallStereoBMTest CV_GpuMatAsyncCallStereoBMTest_test;