diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 888f3a99a2..3468d1f968 100644 --- a/modules/gpu/CMakeLists.txt +++ b/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) @@ -215,4 +215,4 @@ if(BUILD_TESTS AND NOT ANDROID AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/test) if(WIN32) install(TARGETS ${the_test_target} RUNTIME DESTINATION bin COMPONENT main) endif() -endif() \ No newline at end of file +endif() diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 0a4c3487bb..518d95cb6c 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/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 //////////////////////////////// /*! diff --git a/modules/gpu/src/cuda/project_points.cu b/modules/gpu/src/cuda/project_points.cu new file mode 100644 index 0000000000..77144b5425 --- /dev/null +++ b/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_ src, const float* rot, + const float* transl, DevMem2D_ 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_)src, (DevMem2D_)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_ src, const float* rot, + const float* transl, const float* proj, DevMem2D_ 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_)src, (DevMem2D_)dst, ProjectOp()); + } + } // namespace project_points + +}} // namespace cv { namespace gpu diff --git a/modules/gpu/src/precomp.hpp b/modules/gpu/src/precomp.hpp index e86c1a9ea5..a1a3a01cd5 100644 --- a/modules/gpu/src/precomp.hpp +++ b/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) diff --git a/modules/gpu/src/project_points.cpp b/modules/gpu/src/project_points.cpp new file mode 100644 index 0000000000..61a28b1c9a --- /dev/null +++ b/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_ src, const float* rot, + const float* transl, DevMem2D_ 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(), tvec.ptr(), dst); +} + + +namespace cv { namespace gpu { namespace project_points { + void call(const DevMem2D_ src, const float* rot, + const float* transl, const float* proj, DevMem2D_ 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(), tvec.ptr(), + camera_mat.ptr(), dst); +} + +#endif diff --git a/modules/gpu/test/test_precomp.hpp b/modules/gpu/test/test_precomp.hpp index 89426c3bf0..9df4698b75 100644 --- a/modules/gpu/test/test_precomp.hpp +++ b/modules/gpu/test/test_precomp.hpp @@ -1,6 +1,7 @@ #ifndef __OPENCV_TEST_PRECOMP_HPP__ #define __OPENCV_TEST_PRECOMP_HPP__ +#include #include #include #include "opencv2/core/core.hpp" diff --git a/modules/gpu/test/test_project_points.cpp b/modules/gpu/test/test_project_points.cpp new file mode 100644 index 0000000000..95e16f74c0 --- /dev/null +++ b/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(0, 1) = 0.f; + camera_mat.at(1, 0) = 0.f; + camera_mat.at(2, 0) = 0.f; + camera_mat.at(2, 1) = 0.f; + + vector 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(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(0, i); + Point3f res_gold( + rot.at(0, 0) * p.x + rot.at(0, 1) * p.y + rot.at(0, 2) * p.z + tvec.at(0, 0), + rot.at(1, 0) * p.x + rot.at(1, 1) * p.y + rot.at(1, 2) * p.z + tvec.at(0, 1), + rot.at(2, 0) * p.x + rot.at(2, 1) * p.y + rot.at(2, 2) * p.z + tvec.at(0, 2)); + Point3f res_actual = h_dst.at(0, i); + Point3f err = res_actual - res_gold; + ASSERT_LT(err.dot(err) / res_gold.dot(res_gold), 1e-3f); + } +}