added projectPoints and transformPoints into GPU module

modules/gpu/CMakeLists.txt

@@ -4,7 +4,7 @@ set(the_target "opencv_${name}")
 project(${the_target})
 
 
-set(DEPS "opencv_core" "opencv_imgproc" "opencv_objdetect" "opencv_features2d" "opencv_flann") #"opencv_features2d" "opencv_flann" "opencv_objdetect" - only headers needed 
+set(DEPS "opencv_core" "opencv_imgproc" "opencv_objdetect" "opencv_features2d" "opencv_flann" "opencv_calib3d") #"opencv_features2d" "opencv_flann" "opencv_objdetect" - only headers needed 
 set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} opencv_gpu)
 
 add_definitions(-DCVAPI_EXPORTS)

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -853,6 +853,14 @@ namespace cv
         CV_EXPORTS int countNonZero(const GpuMat& src, GpuMat& buf);
 
 
+        ///////////////////////////// Calibration 3D //////////////////////////////////
+
+        CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                                        GpuMat& dst);
+
+        CV_EXPORTS void projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                                      const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst);
+
         //////////////////////////////// Filter Engine ////////////////////////////////
 
         /*!

modules/gpu/src/cuda/project_points.cu

@@ -0,0 +1,115 @@
+/*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.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., 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 the copyright holders 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 "internal_shared.hpp"
+#include "opencv2/gpu/device/transform.hpp"
+
+namespace cv { namespace gpu
+{
+    namespace transform_points
+    {
+        __constant__ float3 crot0;
+        __constant__ float3 crot1;
+        __constant__ float3 crot2;
+        __constant__ float3 ctransl;
+
+        struct TransformOp
+        {
+            __device__ float3 operator()(float3 p) const
+            {
+                return make_float3(
+                        crot0.x * p.x + crot0.y * p.y + crot0.z * p.z + ctransl.x,
+                        crot1.x * p.x + crot1.y * p.y + crot1.z * p.z + ctransl.y,
+                        crot2.x * p.x + crot2.y * p.y + crot2.z * p.z + ctransl.z);
+            }
+        };
+
+        void call(const DevMem2D_<float> src, const float* rot,
+                  const float* transl, DevMem2D_<float> dst)
+        {
+            cudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3));
+            transform((const DevMem2D_<float3>)src, (DevMem2D_<float3>)dst, TransformOp());
+        }
+    } // namespace transform_points
+
+    namespace project_points
+    {
+        __constant__ float3 crot0;
+        __constant__ float3 crot1;
+        __constant__ float3 crot2;
+        __constant__ float3 ctransl;
+        __constant__ float3 cproj0;
+        __constant__ float3 cproj1;
+
+        struct ProjectOp
+        {
+            __device__ float2 operator()(float3 p) const
+            {
+                // Rotate and translate in 3D
+                float3 t = make_float3(
+                        crot0.x * p.x + crot0.y * p.y + crot0.z * p.z + ctransl.x,
+                        crot1.x * p.x + crot1.y * p.y + crot1.z * p.z + ctransl.y,
+                        crot2.x * p.x + crot2.y * p.y + crot2.z * p.z + ctransl.z);
+                // Project on 2D plane
+                return make_float2(
+                        (cproj0.x * t.x + cproj0.y * t.y) / t.z + cproj0.z,
+                        (cproj1.x * t.x + cproj1.y * t.y) / t.z + cproj1.z);
+            }
+        };
+
+        void call(const DevMem2D_<float> src, const float* rot,
+                  const float* transl, const float* proj, DevMem2D_<float> dst)
+        {
+            cudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(crot2, rot + 6, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3));
+            cudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3));
+            transform((const DevMem2D_<float3>)src, (DevMem2D_<float2>)dst, ProjectOp());
+        }
+    } // namespace project_points
+
+}} // namespace cv { namespace gpu

modules/gpu/src/precomp.hpp

@@ -61,6 +61,7 @@
 
 #include "opencv2/gpu/gpu.hpp"
 #include "opencv2/imgproc/imgproc.hpp"
+#include "opencv2/calib3d/calib3d.hpp"
 
 #if defined(HAVE_CUDA)
 

modules/gpu/src/project_points.cpp

@@ -0,0 +1,99 @@
+/*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.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., 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 the copyright holders 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 "precomp.hpp"
+
+#if !defined(HAVE_CUDA)
+
+void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&,
+                              GpuMat&) { throw_nogpu(); }
+
+void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&,
+                            const Mat&, const Mat&, GpuMat&) { throw_nogpu(); }
+
+#else
+
+namespace cv { namespace gpu { namespace transform_points {
+    void call(const DevMem2D_<float> src, const float* rot,
+              const float* transl, DevMem2D_<float> dst);
+}}}
+
+void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                              GpuMat& dst)
+{
+    CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
+    CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
+    CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
+
+    // Convert rotation vector into matrix
+    Mat rot;
+    Rodrigues(rvec, rot);
+
+    dst.create(src.size(), src.type());
+    transform_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), dst);
+}
+
+
+namespace cv { namespace gpu { namespace project_points {
+    void call(const DevMem2D_<float> src, const float* rot,
+              const float* transl, const float* proj, DevMem2D_<float> dst);
+}}}
+
+void cv::gpu::projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                            const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst)
+{
+    CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
+    CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
+    CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
+    CV_Assert(camera_mat.size() == Size(3, 3) && camera_mat.type() == CV_32F);
+    CV_Assert(dist_coef.empty()); // Undistortion isn't supported
+
+    // Convert rotation vector into matrix
+    Mat rot;
+    Rodrigues(rvec, rot);
+
+    dst.create(src.size(), CV_32FC2);
+    project_points::call(src, rot.ptr<float>(), tvec.ptr<float>(),
+                         camera_mat.ptr<float>(), dst);
+}
+
+#endif

modules/gpu/test/test_precomp.hpp

@@ -1,6 +1,7 @@
 #ifndef __OPENCV_TEST_PRECOMP_HPP__
 #define __OPENCV_TEST_PRECOMP_HPP__
 
+#include <iostream>
 #include <limits>
 #include <cvconfig.h>
 #include "opencv2/core/core.hpp"

modules/gpu/test/test_project_points.cpp

@@ -0,0 +1,107 @@
+/*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 "test_precomp.hpp"
+
+using namespace cv;
+using namespace cv::gpu;
+using namespace cvtest;
+
+TEST(projectPoints, accuracy)
+{
+    RNG& rng = TS::ptr()->get_rng();
+    Mat src = randomMat(rng, Size(1000, 1), CV_32FC3, 0, 10, false);
+    Mat rvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
+    Mat tvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
+    Mat camera_mat = randomMat(rng, Size(3, 3), CV_32F, 0, 1, false);
+    camera_mat.at<float>(0, 1) = 0.f;
+    camera_mat.at<float>(1, 0) = 0.f;
+    camera_mat.at<float>(2, 0) = 0.f;
+    camera_mat.at<float>(2, 1) = 0.f;
+
+    vector<Point2f> dst;
+    projectPoints(src, rvec, tvec, camera_mat, Mat(), dst);
+
+    GpuMat d_dst;
+    projectPoints(GpuMat(src), rvec, tvec, camera_mat, Mat(), d_dst);
+
+    ASSERT_EQ(dst.size(), (size_t)d_dst.cols);
+    ASSERT_EQ(1, d_dst.rows);
+    ASSERT_EQ(CV_32FC2, d_dst.type());
+
+    Mat h_dst(d_dst);
+    for (size_t i = 0; i < dst.size(); ++i)
+    {
+        Point2f res_gold = dst[i];
+        Point2f res_actual = h_dst.at<Point2f>(0, i);
+        Point2f err = res_actual - res_gold;
+        ASSERT_LT(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
+    }
+}
+
+
+TEST(transformPoints, accuracy)
+{
+    RNG& rng = TS::ptr()->get_rng();
+    Mat src = randomMat(rng, Size(1000, 1), CV_32FC3, 0, 10, false);
+    Mat rvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
+    Mat tvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
+
+    GpuMat d_dst;
+    transformPoints(GpuMat(src), rvec, tvec, d_dst);
+    ASSERT_TRUE(src.size() == d_dst.size());
+    ASSERT_EQ(src.type(), d_dst.type());
+
+    Mat h_dst(d_dst);
+    Mat rot;
+    Rodrigues(rvec, rot);
+    for (int i = 0; i < h_dst.cols; ++i)
+    {
+        Point3f p = src.at<Point3f>(0, i);
+        Point3f res_gold(
+                rot.at<float>(0, 0) * p.x + rot.at<float>(0, 1) * p.y + rot.at<float>(0, 2) * p.z + tvec.at<float>(0, 0),
+                rot.at<float>(1, 0) * p.x + rot.at<float>(1, 1) * p.y + rot.at<float>(1, 2) * p.z + tvec.at<float>(0, 1),
+                rot.at<float>(2, 0) * p.x + rot.at<float>(2, 1) * p.y + rot.at<float>(2, 2) * p.z + tvec.at<float>(0, 2));
+        Point3f res_actual = h_dst.at<Point3f>(0, i);
+        Point3f err = res_actual - res_gold;
+        ASSERT_LT(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
+    }
+}