From 73578524348463b636203e00693ecabf13df04f2 Mon Sep 17 00:00:00 2001 From: Andrey Morozov Date: Mon, 30 Aug 2010 12:30:08 +0000 Subject: [PATCH] added test stereo_bm and added "volatile" in constantspacebp --- modules/gpu/src/cuda/constantspacebp.cu | 28 +++++---- tests/gpu/src/stereo_bm.cpp | 80 +++++++++++++++++++++++++ 2 files changed, 96 insertions(+), 12 deletions(-) create mode 100644 tests/gpu/src/stereo_bm.cpp diff --git a/modules/gpu/src/cuda/constantspacebp.cu b/modules/gpu/src/cuda/constantspacebp.cu index a404e0eaed..48f435934f 100644 --- a/modules/gpu/src/cuda/constantspacebp.cu +++ b/modules/gpu/src/cuda/constantspacebp.cu @@ -322,12 +322,14 @@ namespace csbp_krnls if (winsz >= 256) { if (tid < 128) { dline[tid] += dline[tid + 128]; } __syncthreads(); } if (winsz >= 128) { if (tid < 64) { dline[tid] += dline[tid + 64]; } __syncthreads(); } - if (winsz >= 64) if (tid < 32) dline[tid] += dline[tid + 32]; - if (winsz >= 32) if (tid < 16) dline[tid] += dline[tid + 16]; - if (winsz >= 16) if (tid < 8) dline[tid] += dline[tid + 8]; - if (winsz >= 8) if (tid < 4) dline[tid] += dline[tid + 4]; - if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2]; - if (winsz >= 2) if (tid < 1) dline[tid] += dline[tid + 1]; + volatile float* vdline = smem + winsz * threadIdx.z; + + if (winsz >= 64) if (tid < 32) vdline[tid] += vdline[tid + 32]; + if (winsz >= 32) if (tid < 16) vdline[tid] += vdline[tid + 16]; + if (winsz >= 16) if (tid < 8) vdline[tid] += vdline[tid + 8]; + if (winsz >= 8) if (tid < 4) vdline[tid] += vdline[tid + 4]; + if (winsz >= 4) if (tid < 2) vdline[tid] += vdline[tid + 2]; + if (winsz >= 2) if (tid < 1) vdline[tid] += vdline[tid + 1]; T* data_cost = (T*)ctemp + y_out * cmsg_step1 + x_out; @@ -524,12 +526,14 @@ namespace csbp_krnls if (winsz >= 256) { if (tid < 128) { dline[tid] += dline[tid + 128]; } __syncthreads(); } if (winsz >= 128) { if (tid < 64) { dline[tid] += dline[tid + 64]; } __syncthreads(); } - if (winsz >= 64) if (tid < 32) dline[tid] += dline[tid + 32]; - if (winsz >= 32) if (tid < 16) dline[tid] += dline[tid + 16]; - if (winsz >= 16) if (tid < 8) dline[tid] += dline[tid + 8]; - if (winsz >= 8) if (tid < 4) dline[tid] += dline[tid + 4]; - if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2]; - if (winsz >= 2) if (tid < 1) dline[tid] += dline[tid + 1]; + volatile float* vdline = smem + winsz * threadIdx.z; + + if (winsz >= 64) if (tid < 32) vdline[tid] += vdline[tid + 32]; + if (winsz >= 32) if (tid < 16) vdline[tid] += vdline[tid + 16]; + if (winsz >= 16) if (tid < 8) vdline[tid] += vdline[tid + 8]; + if (winsz >= 8) if (tid < 4) vdline[tid] += vdline[tid + 4]; + if (winsz >= 4) if (tid < 2) vdline[tid] += vdline[tid + 2]; + if (winsz >= 2) if (tid < 1) vdline[tid] += vdline[tid + 1]; if (tid == 0) data_cost[cdisp_step1 * d] = saturate_cast(dline[0]); diff --git a/tests/gpu/src/stereo_bm.cpp b/tests/gpu/src/stereo_bm.cpp new file mode 100644 index 0000000000..f2eb4b3a4b --- /dev/null +++ b/tests/gpu/src/stereo_bm.cpp @@ -0,0 +1,80 @@ +/*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 +#include + +#include +#include + +class CV_GpuStereoBM : public CvTest +{ +public: + CV_GpuStereoBM(); +protected: + void run(int); +}; + +CV_GpuStereoBM::CV_GpuStereoBM(): CvTest( "GPU-StereoBM", "StereoBM" ){} + +void CV_GpuStereoBM::run(int ) +{ + 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_template = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0); + + cv::gpu::GpuMat disp; + cv::gpu::StereoBM_GPU bm; + + bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); + + cv::imwrite(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", disp); + + disp.convertTo(disp, img_template.type()); + + double norm = cv::norm(disp, img_template, cv::NORM_INF); + ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); +} + + +CV_GpuStereoBM CV_GpuStereoBM_test; +