meanShiftFiltering_GPU output parameters changed to CV_8UC4. This is a start for...
authorKirill Kornyakov <no@email>
Fri, 3 Sep 2010 14:32:12 +0000 (14:32 +0000)
committerKirill Kornyakov <no@email>
Fri, 3 Sep 2010 14:32:12 +0000 (14:32 +0000)
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/src/cuda/imgproc.cu
modules/gpu/src/imgproc_gpu.cpp
tests/gpu/src/meanshift.cpp
tests/gpu/src/stereo_bp.cpp

index 0382ff2..f439a37 100644 (file)
@@ -347,7 +347,8 @@ namespace cv
         CV_EXPORTS void remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst);\r
 \r
         // Does mean shift filtering on GPU.\r
-        CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));\r
+        CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, \r
+            TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));\r
 \r
         // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.\r
         // Supported types of input disparity: CV_8U, CV_16S.\r
index 8c53ae2..5fc28d5 100644 (file)
@@ -163,7 +163,8 @@ namespace imgproc
 {\r
     texture<uchar4, 2> tex_meanshift;\r
 \r
-    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, int sp, int sr, int maxIter, float eps )\r
+    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, \r
+                                                 int sp, int sr, int maxIter, float eps )\r
     {\r
         int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
         int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
@@ -224,10 +225,8 @@ namespace imgproc
                     break;\r
             }\r
 \r
-            int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 3 * sizeof(uchar);\r
-            out[base+0] = c.x;\r
-            out[base+1] = c.y;\r
-            out[base+2] = c.z;\r
+            int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);\r
+            *(uchar4*)(out + base) = c;\r
         }\r
     }\r
 }\r
@@ -236,7 +235,7 @@ namespace cv { namespace gpu { namespace impl
 {\r
     extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps)\r
     {                        \r
-        dim3  grid(1, 1, 1);\r
+        dim3 grid(1, 1, 1);\r
         dim3 threads(32, 16, 1);\r
         grid.x = divUp(src.cols, threads.x);\r
         grid.y = divUp(src.rows, threads.y);\r
index 10d0225..e9e7d33 100644 (file)
@@ -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 )\r
         CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );\r
 \r
-    dst.create( src.size(), CV_8UC3 );\r
+    dst.create( src.size(), CV_8UC4 );\r
     \r
-    float eps;\r
     if( !(criteria.type & TermCriteria::MAX_ITER) )\r
         criteria.maxCount = 5;\r
     \r
     int maxIter = std::min(std::max(criteria.maxCount, 1), 100);\r
     \r
+    float eps;\r
     if( !(criteria.type & TermCriteria::EPS) )\r
         eps = 1.f;\r
-\r
     eps = (float)std::max(criteria.epsilon, 0.0);        \r
+\r
     impl::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);    \r
 }\r
 \r
index e3c2b82..84e24e3 100644 (file)
-/*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 <iostream>
-#include <string>
-
-#include <opencv2/opencv.hpp>
-#include <opencv2/gpu/gpu.hpp>
-
-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///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+//  By downloading, copying, installing or using the software you agree to this license.\r
+//  If you do not agree to this license, do not download, install,\r
+//  copy or use the software.\r
+//\r
+//\r
+//                        Intel License Agreement\r
+//                For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000, Intel Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+//   * Redistribution's of source code must retain the above copyright notice,\r
+//     this list of conditions and the following disclaimer.\r
+//\r
+//   * Redistribution's in binary form must reproduce the above copyright notice,\r
+//     this list of conditions and the following disclaimer in the documentation\r
+//     and/or other materials provided with the distribution.\r
+//\r
+//   * The name of Intel Corporation may not be used to endorse or promote products\r
+//     derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "gputest.hpp"\r
+#include <iostream>\r
+#include <string>\r
+\r
+#include <opencv2/opencv.hpp>\r
+#include <opencv2/gpu/gpu.hpp>\r
+\r
+class CV_GpuMeanShift : public CvTest\r
+{\r
+    public:\r
+        CV_GpuMeanShift();\r
+    protected:\r
+        void run(int);\r
+};\r
+\r
+CV_GpuMeanShift::CV_GpuMeanShift(): CvTest( "GPU-MeanShift", "MeanShift" ){}\r
+\r
+void CV_GpuMeanShift::run(int)\r
+{\r
+        int spatialRad = 30;\r
+        int colorRad = 30;\r
+\r
+        cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");\r
+        cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png");\r
+\r
+        if (img.empty() || img_template.empty())\r
+        {\r
+            ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);\r
+            return;\r
+        }\r
+\r
+        cv::Mat rgba;\r
+        cvtColor(img, rgba, CV_BGR2BGRA);\r
+\r
+        cv::gpu::GpuMat res;\r
+        cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );\r
+        if (res.type() != CV_8UC4)\r
+        {\r
+            ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+            return;\r
+        }\r
+\r
+        cv::Mat result;\r
+        res.download(result);\r
+\r
+        uchar maxDiff = 0;\r
+        for (int j = 0; j < result.rows; ++j)\r
+        {\r
+            const uchar* res_line = result.ptr<uchar>(j);\r
+            const uchar* ref_line = img_template.ptr<uchar>(j);\r
+\r
+            for (int i = 0; i < result.cols; ++i)\r
+            {\r
+                for (int k = 0; k < 3; ++k)\r
+                {\r
+                    const uchar& ch1 = res_line[result.channels()*i + k];\r
+                    const uchar& ch2 = ref_line[img_template.channels()*i + k];\r
+                    uchar diff = abs(ch1 - ch2);\r
+                    if (maxDiff < diff)\r
+                        maxDiff = diff;\r
+                }\r
+            }\r
+        }\r
+        if (maxDiff > 0) \r
+            ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff);\r
+        ts->set_failed_test_info((maxDiff == 0) ? CvTS::OK : CvTS::FAIL_GENERIC);\r
+}\r
+\r
+CV_GpuMeanShift CV_GpuMeanShift_test;
\ No newline at end of file
index 5bffb91..6aed6c3 100644 (file)
@@ -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);
 }