added asynchronous versions of transform- and projectPoints into the GPU module, added docs

doc/gpu_calib3d.tex

@@ -421,3 +421,41 @@ void reprojectImageTo3D(const GpuMat\& disp, GpuMat\& xyzw, \par const Mat\& Q,
 \end{description}
 
 See also: \cvCppCross{reprojectImageTo3D}.
+
+
+\cvCppFunc{gpu::transformPoints}
+Rotates and translates points.
+
+\cvdefCpp{
+void transformPoints(const GpuMat\& src, const Mat\& rvec, \par const Mat\& tvec, GpuMat\& dst);\newline
+void transformPoints(const GpuMat\& src, const Mat\& rvec, \par const Mat\& tvec, GpuMat\& dst, const Stream\& stream);
+}
+
+\begin{description}
+\cvarg{src}{Source points. Single-row \texttt{CV\_32FC3} matrix.}
+\cvarg{rvec}{\texttt{CV\_32F} 3D rotation vector.}
+\cvarg{tvec}{\texttt{CV\_32F} 3D translation vector.}
+\cvarg{dst}{Transformed points. Single-row \texttt{CV\_32FC3} matrix.}
+\cvarg{stream}{Stream for the asynchronous version.}
+\end{description}
+
+
+\cvCppFunc{gpu::projectPoints}
+Projects points.
+
+\cvdefCpp{
+void projectPoints(const GpuMat\& src, const Mat\& rvec, \par const Mat\& tvec, const Mat\& camera\_mat, \par const Mat\& dist\_coef, GpuMat\& dst);\newline
+void projectPoints(const GpuMat\& src, const Mat\& rvec, \par const Mat\& tvec, const Mat\& camera\_mat, \par const Mat\& dist\_coef, GpuMat\& dst, const Stream\& stream);
+}
+
+\begin{description}
+\cvarg{src}{Source points. Single-row \texttt{CV\_32FC3} matrix.}
+\cvarg{rvec}{\texttt{CV\_32F} 3D rotation vector.}
+\cvarg{tvec}{\texttt{CV\_32F} 3D translation vector.}
+\cvarg{camera\_mat}{\texttt{CV\_32F} 3x3 camera matrix.}
+\cvarg{dist\_coef}{Distortion coefficients. This parameter isn't supported for now, must be empty matrix.}
+\cvarg{dst}{Projected points. Single-row \texttt{CV\_32FC2} matrix.}
+\cvarg{stream}{Stream for the asynchronous version.}
+\end{description}
+
+See also: \cvCppCross{projectPoints}.

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

@@ -858,9 +858,16 @@ namespace cv
         CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
                                         GpuMat& dst);
 
+        CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                                        GpuMat& dst, const Stream& stream);
+
         CV_EXPORTS void projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
                                       const Mat& camera_mat, const Mat& dist_coef, 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,
+                                      const Stream& stream);
+
         //////////////////////////////// Filter Engine ////////////////////////////////
 
         /*!

modules/gpu/src/cuda/project_points.cu

@@ -64,13 +64,14 @@ namespace cv { namespace gpu
         };
 
         void call(const DevMem2D_<float3> src, const float* rot,
-                  const float* transl, DevMem2D_<float3> dst)
+                  const float* transl, DevMem2D_<float3> dst,
+                  cudaStream_t stream)
         {
             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(src, dst, TransformOp());
+            transform(src, dst, TransformOp(), stream);
         }
     } // namespace transform_points
 
@@ -100,7 +101,8 @@ namespace cv { namespace gpu
         };
 
         void call(const DevMem2D_<float3> src, const float* rot,
-                  const float* transl, const float* proj, DevMem2D_<float2> dst)
+                  const float* transl, const float* proj, DevMem2D_<float2> dst,
+                  cudaStream_t stream)
         {
             cudaSafeCall(cudaMemcpyToSymbol(crot0, rot, sizeof(float) * 3));
             cudaSafeCall(cudaMemcpyToSymbol(crot1, rot + 3, sizeof(float) * 3));
@@ -108,7 +110,7 @@ namespace cv { namespace gpu
             cudaSafeCall(cudaMemcpyToSymbol(ctransl, transl, sizeof(float) * 3));
             cudaSafeCall(cudaMemcpyToSymbol(cproj0, proj, sizeof(float) * 3));
             cudaSafeCall(cudaMemcpyToSymbol(cproj1, proj + 3, sizeof(float) * 3));
-            transform(src, dst, ProjectOp());
+            transform(src, dst, ProjectOp(), stream);
         }
     } // namespace project_points
 

modules/gpu/src/project_points.cpp

@@ -47,52 +47,95 @@
 void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&,
                               GpuMat&) { throw_nogpu(); }
 
+void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&,
+                              GpuMat&, const Stream&) { throw_nogpu(); }
+
 void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&,
                             const Mat&, const Mat&, GpuMat&) { throw_nogpu(); }
 
+void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&,
+                            const Mat&, const Mat&, GpuMat&, const Stream&) { throw_nogpu(); }
+
 #else
 
+using namespace cv;
+using namespace cv::gpu;
+
 namespace cv { namespace gpu { namespace transform_points 
 {
-    void call(const DevMem2D_<float3> src, const float* rot, const float* transl, DevMem2D_<float3> dst);
+    void call(const DevMem2D_<float3> src, const float* rot, const float* transl,
+              DevMem2D_<float3> dst, cudaStream_t stream);
 }}}
 
-void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
-                              GpuMat& dst)
+namespace
 {
-    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);
+    void transformPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                               GpuMat& dst, cudaStream_t stream)
+    {
+        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);
 
-    // 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, stream);
+    }
+}
 
-    dst.create(src.size(), src.type());
-    transform_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), dst);
+void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                              GpuMat& dst)
+{
+    ::transformPointsCaller(src, rvec, tvec, dst, 0);
 }
 
+void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                              GpuMat& dst, const Stream& stream)
+{
+    ::transformPointsCaller(src, rvec, tvec, dst, StreamAccessor::getStream(stream));
+}
 
 namespace cv { namespace gpu { namespace project_points 
 {
-    void call(const DevMem2D_<float3> src, const float* rot, const float* transl, const float* proj, DevMem2D_<float2> dst);
+    void call(const DevMem2D_<float3> src, const float* rot, const float* transl,
+              const float* proj, DevMem2D_<float2> dst, cudaStream_t stream);
 }}}
 
+namespace
+{
+    void projectPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                             const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst,
+                             cudaStream_t stream)
+    {
+        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,stream);
+    }
+}
+
 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);
+    ::projectPointsCaller(src, rvec, tvec, camera_mat, dist_coef, dst, 0);
+}
+
+void cv::gpu::projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                            const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst,
+                            const Stream& stream)
+{
+    ::projectPointsCaller(src, rvec, tvec, camera_mat, dist_coef, dst, StreamAccessor::getStream(stream));
 }
 
 #endif