\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
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
\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
\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
{\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
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
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
//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();
--- /dev/null
+/*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;