From ec7e93748110b6816ef75445639c029eb1584d6e Mon Sep 17 00:00:00 2001 From: Kirill Kornyakov Date: Fri, 3 Sep 2010 14:32:12 +0000 Subject: [PATCH] meanShiftFiltering_GPU output parameters changed to CV_8UC4. This is a start for moving from 3 channel to C4 images within GPU module. --- modules/gpu/include/opencv2/gpu/gpu.hpp | 3 +- modules/gpu/src/cuda/imgproc.cu | 11 +- modules/gpu/src/imgproc_gpu.cpp | 6 +- tests/gpu/src/meanshift.cpp | 192 ++++++++++++++++++-------------- tests/gpu/src/stereo_bp.cpp | 28 +++-- 5 files changed, 137 insertions(+), 103 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 0382ff2..f439a37 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -347,7 +347,8 @@ namespace cv CV_EXPORTS void remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst); // Does mean shift filtering on GPU. - CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); + CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, + TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. // Supported types of input disparity: CV_8U, CV_16S. diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 8c53ae2..5fc28d5 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -163,7 +163,8 @@ namespace imgproc { texture tex_meanshift; - extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, int sp, int sr, int maxIter, float eps ) + extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, + int sp, int sr, int maxIter, float eps ) { int x0 = blockIdx.x * blockDim.x + threadIdx.x; int y0 = blockIdx.y * blockDim.y + threadIdx.y; @@ -224,10 +225,8 @@ namespace imgproc break; } - int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 3 * sizeof(uchar); - out[base+0] = c.x; - out[base+1] = c.y; - out[base+2] = c.z; + int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar); + *(uchar4*)(out + base) = c; } } } @@ -236,7 +235,7 @@ namespace cv { namespace gpu { namespace impl { extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps) { - dim3 grid(1, 1, 1); + dim3 grid(1, 1, 1); dim3 threads(32, 16, 1); grid.x = divUp(src.cols, threads.x); grid.y = divUp(src.rows, threads.y); diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 10d0225..e9e7d33 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -119,18 +119,18 @@ void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int if( src.depth() != CV_8U || src.channels() != 4 ) CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" ); - dst.create( src.size(), CV_8UC3 ); + dst.create( src.size(), CV_8UC4 ); - float eps; if( !(criteria.type & TermCriteria::MAX_ITER) ) criteria.maxCount = 5; int maxIter = std::min(std::max(criteria.maxCount, 1), 100); + float eps; if( !(criteria.type & TermCriteria::EPS) ) eps = 1.f; - eps = (float)std::max(criteria.epsilon, 0.0); + impl::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps); } diff --git a/tests/gpu/src/meanshift.cpp b/tests/gpu/src/meanshift.cpp index e3c2b82..84e24e3 100644 --- a/tests/gpu/src/meanshift.cpp +++ b/tests/gpu/src/meanshift.cpp @@ -1,82 +1,110 @@ -/*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_GpuMeanShift : public CvTest -{ - public: - CV_GpuMeanShift(); - protected: - void run(int); -}; - -CV_GpuMeanShift::CV_GpuMeanShift(): CvTest( "GPU-MeanShift", "MeanShift" ){} - -void CV_GpuMeanShift::run(int ) -{ - int spatialRad = 30; - int colorRad = 30; - - cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png"); - cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png"); - - cv::Mat rgba; - cvtColor(img, rgba, CV_BGR2BGRA); - - cv::gpu::GpuMat res; - - cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad ); - - res.convertTo(res, img_template.type()); - - double norm = cv::norm(res, img_template, cv::NORM_INF); - if (norm >= 0.5) std::cout << "MeanShift norm = " << norm << std::endl; - ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); -} - - -CV_GpuMeanShift CV_GpuMeanShift_test; +/*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_GpuMeanShift : public CvTest +{ + public: + CV_GpuMeanShift(); + protected: + void run(int); +}; + +CV_GpuMeanShift::CV_GpuMeanShift(): CvTest( "GPU-MeanShift", "MeanShift" ){} + +void CV_GpuMeanShift::run(int) +{ + int spatialRad = 30; + int colorRad = 30; + + cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png"); + cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png"); + + if (img.empty() || img_template.empty()) + { + ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); + return; + } + + cv::Mat rgba; + cvtColor(img, rgba, CV_BGR2BGRA); + + cv::gpu::GpuMat res; + cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad ); + if (res.type() != CV_8UC4) + { + ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT); + return; + } + + cv::Mat result; + res.download(result); + + uchar maxDiff = 0; + for (int j = 0; j < result.rows; ++j) + { + const uchar* res_line = result.ptr(j); + const uchar* ref_line = img_template.ptr(j); + + for (int i = 0; i < result.cols; ++i) + { + for (int k = 0; k < 3; ++k) + { + const uchar& ch1 = res_line[result.channels()*i + k]; + const uchar& ch2 = ref_line[img_template.channels()*i + k]; + uchar diff = abs(ch1 - ch2); + if (maxDiff < diff) + maxDiff = diff; + } + } + } + if (maxDiff > 0) + ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff); + ts->set_failed_test_info((maxDiff == 0) ? CvTS::OK : CvTS::FAIL_GENERIC); +} + +CV_GpuMeanShift CV_GpuMeanShift_test; \ No newline at end of file diff --git a/tests/gpu/src/stereo_bp.cpp b/tests/gpu/src/stereo_bp.cpp index 5bffb91..6aed6c3 100644 --- a/tests/gpu/src/stereo_bp.cpp +++ b/tests/gpu/src/stereo_bp.cpp @@ -58,22 +58,28 @@ CV_GpuStereoBP::CV_GpuStereoBP(): CvTest( "GPU-StereoBP", "StereoBP" ){} void CV_GpuStereoBP::run(int ) { - cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png"); - cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-R.png"); - cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", 0); + cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png"); + cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-R.png"); + cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", 0); - cv::gpu::GpuMat disp; - cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S); + if (img_l.empty() || img_r.empty() || img_template.empty()) + { + ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); + return; + } - bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); + cv::gpu::GpuMat disp; + cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S); - //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp); + bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); - disp.convertTo(disp, img_template.type()); + //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp); - double norm = cv::norm(disp, img_template, cv::NORM_INF); - if (norm >= 0.5) std::cout << "StereoBP norm = " << norm << std::endl; - ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); + disp.convertTo(disp, img_template.type()); + + double norm = cv::norm(disp, img_template, cv::NORM_INF); + if (norm >= 0.5) std::cout << "StereoBP norm = " << norm << std::endl; + ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); } -- 2.7.4