From: Kirill Kornyakov <no@email>
Date: Mon, 6 Sep 2010 14:27:23 +0000 (+0000)
Subject: GPU module update: _GPU suffix removed, some namespaces renamed, minor refactorings.
X-Git-Tag: accepted/2.0/20130307.220821~4521
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=6960e1544d0104422092090698022aa544e8ea65;p=profile%2Fivi%2Fopencv.git

GPU module update: _GPU suffix removed, some namespaces renamed, minor refactorings.
---

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index f439a37..44ad9d2 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -344,10 +344,10 @@ namespace cv
         ////////////////////////////// Image processing //////////////////////////////
         // DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.
         // xymap.type() == xymap.type() == CV_32FC1
-        CV_EXPORTS void remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst);
+        CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap);
 
         // Does mean shift filtering on GPU.
-        CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, 
+        CV_EXPORTS void meanShiftFiltering(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.
@@ -362,18 +362,18 @@ namespace cv
         // The output is a 4-channel floating-point (CV_32FC4) matrix. 
         // Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.
         // Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.
-        CV_EXPORTS void reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q);
+        CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q);
         // Acync version
-        CV_EXPORTS void reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream);
+        CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream);
 
-        CV_EXPORTS void cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn = 0);
-        CV_EXPORTS void cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream);
+        CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0);
+        CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream);
 
         //////////////////////////////// StereoBM_GPU ////////////////////////////////
 
         class CV_EXPORTS StereoBM_GPU
         {
-        public:
+        public:                                    
             enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 };
 
             enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 };
diff --git a/modules/gpu/src/cuda/beliefpropagation.cu b/modules/gpu/src/cuda/beliefpropagation.cu
index ddf5bd3..063a9ad 100644
--- a/modules/gpu/src/cuda/beliefpropagation.cu
+++ b/modules/gpu/src/cuda/beliefpropagation.cu
@@ -45,7 +45,6 @@
 #include "safe_call.hpp"
 
 using namespace cv::gpu;
-using namespace cv::gpu::impl;
 
 #ifndef FLT_MAX
 #define FLT_MAX 3.402823466e+38F
diff --git a/modules/gpu/src/cuda/bilateral_filter.cu b/modules/gpu/src/cuda/bilateral_filter.cu
index 9eca3f4..132ca84 100644
--- a/modules/gpu/src/cuda/bilateral_filter.cu
+++ b/modules/gpu/src/cuda/bilateral_filter.cu
@@ -45,7 +45,6 @@
 #include "safe_call.hpp"
 
 using namespace cv::gpu;
-using namespace cv::gpu::impl;
 
 #ifndef FLT_MAX
 #define FLT_MAX 3.402823466e+30F
diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu
index b14a1b0..a53a4f8 100644
--- a/modules/gpu/src/cuda/color.cu
+++ b/modules/gpu/src/cuda/color.cu
@@ -44,7 +44,6 @@
 #include "saturate_cast.hpp"
 
 using namespace cv::gpu;
-using namespace cv::gpu::impl;
 
 #ifndef CV_DESCALE
 #define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
@@ -167,7 +166,7 @@ namespace imgproc
         }
 }
 
-namespace cv { namespace gpu { namespace impl
+namespace cv { namespace gpu { namespace improc
 {
     template <typename T>
     void RGB2RGB_caller(const DevMem2D_<T>& src, int srccn, const DevMem2D_<T>& dst, int dstcn, int bidx, cudaStream_t stream)
@@ -377,7 +376,7 @@ namespace imgproc
     //};
 }
 
-namespace cv { namespace gpu { namespace impl
+namespace cv { namespace gpu { namespace improc
 {
     template <typename T>
     void Gray2RGB_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int dstcn, cudaStream_t stream)
@@ -627,7 +626,7 @@ namespace imgproc
     }
 }
 
-namespace cv { namespace gpu { namespace impl
+namespace cv { namespace gpu { namespace improc
 {
     void RGB2Gray_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream)
     {
diff --git a/modules/gpu/src/cuda/constantspacebp.cu b/modules/gpu/src/cuda/constantspacebp.cu
index 48f4359..7be1179 100644
--- a/modules/gpu/src/cuda/constantspacebp.cu
+++ b/modules/gpu/src/cuda/constantspacebp.cu
@@ -45,7 +45,6 @@
 #include "safe_call.hpp"
 
 using namespace cv::gpu;
-using namespace cv::gpu::impl;
 
 #ifndef FLT_MAX
 #define FLT_MAX 3.402823466e+30F
diff --git a/modules/gpu/src/cuda/cuda_shared.hpp b/modules/gpu/src/cuda/cuda_shared.hpp
index 73304f5..449fcb0 100644
--- a/modules/gpu/src/cuda/cuda_shared.hpp
+++ b/modules/gpu/src/cuda/cuda_shared.hpp
@@ -56,10 +56,10 @@ namespace cv
         typedef unsigned short ushort;
         typedef unsigned int uint;
 
-        namespace impl
-        {
-            static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }
+        static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }
 
+        namespace matrix_operations
+        {            
             extern "C" void copy_to_with_mask(const DevMem2D& src, DevMem2D dst, int depth, const DevMem2D& mask, int channels, const cudaStream_t & stream = 0);
 
             extern "C" void set_to_without_mask (DevMem2D dst, int depth, const double *scalar, int channels, const cudaStream_t & stream = 0);
diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 5fc28d5..1dd4616 100644
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -123,7 +123,7 @@ namespace imgproc
     }
 }
 
-namespace cv { namespace gpu { namespace impl 
+namespace cv { namespace gpu { namespace improc 
 {
     void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst)
     {
@@ -231,7 +231,7 @@ namespace imgproc
     }
 }
 
-namespace cv { namespace gpu { namespace impl 
+namespace cv { namespace gpu { namespace improc 
 {
     extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps)
     {                        
@@ -354,7 +354,7 @@ namespace imgproc
     }
 }
 
-namespace cv { namespace gpu { namespace impl 
+namespace cv { namespace gpu { namespace improc 
 {
     void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream)
     {
@@ -420,7 +420,7 @@ namespace imgproc
     }
 }
 
-namespace cv { namespace gpu { namespace impl 
+namespace cv { namespace gpu { namespace improc 
 {
     template <typename T>
     inline void reprojectImageTo3D_caller(const DevMem2D_<T>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream)
diff --git a/modules/gpu/src/cuda/matrix_operations.cu b/modules/gpu/src/cuda/matrix_operations.cu
index fe730cc..e16e553 100644
--- a/modules/gpu/src/cuda/matrix_operations.cu
+++ b/modules/gpu/src/cuda/matrix_operations.cu
@@ -47,7 +47,7 @@
 #include "saturate_cast.hpp"
 
 using namespace cv::gpu;
-using namespace cv::gpu::impl;
+using namespace cv::gpu::matrix_operations;
 
 
 namespace mat_operators
@@ -261,7 +261,7 @@ namespace cv
 {
     namespace gpu
     {
-        namespace impl
+        namespace matrix_operations
         {
 
             ///////////////////////////////////////////////////////////////////////////
diff --git a/modules/gpu/src/cuda/stereobm.cu b/modules/gpu/src/cuda/stereobm.cu
index 91f3294..73a70b8 100644
--- a/modules/gpu/src/cuda/stereobm.cu
+++ b/modules/gpu/src/cuda/stereobm.cu
@@ -316,7 +316,7 @@ __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t i
 }
 
 
-namespace cv { namespace gpu { namespace impl
+namespace cv { namespace gpu { namespace bm
 {
     template<int RADIUS> void kernel_caller(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int maxdisp, const cudaStream_t & stream)
     {
@@ -408,7 +408,7 @@ extern "C" __global__ void prefilter_kernel(unsigned char *output, size_t step,
 
 }
 
-namespace cv { namespace gpu  { namespace impl
+namespace cv { namespace gpu  { namespace bm
 {
     extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap)
     {
@@ -530,7 +530,7 @@ extern "C" __global__ void textureness_kernel(unsigned char *disp, size_t disp_s
 }
 }
 
-namespace cv { namespace gpu  { namespace impl
+namespace cv { namespace gpu  { namespace bm
 {
     extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp)
     {
diff --git a/modules/gpu/src/cudastream.cpp b/modules/gpu/src/cudastream.cpp
index df0a696..f933453 100644
--- a/modules/gpu/src/cudastream.cpp
+++ b/modules/gpu/src/cudastream.cpp
@@ -158,12 +158,12 @@ void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) { devcopy(src,
 
 void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val)
 {
-    impl::set_to_without_mask(src, src.depth(), val.val, src.channels(), impl->stream);
+    matrix_operations::set_to_without_mask(src, src.depth(), val.val, src.channels(), impl->stream);
 }
 
 void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val, const GpuMat& mask)
 {
-    impl::set_to_with_mask(src, src.depth(), val.val, mask, src.channels(), impl->stream);
+    matrix_operations::set_to_with_mask(src, src.depth(), val.val, mask, src.channels(), impl->stream);
 }
 
 void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype, double alpha, double beta)
@@ -188,7 +188,7 @@ void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype,
         psrc = &(temp = src);
 
     dst.create( src.size(), rtype );
-    impl::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta, impl->stream);
+    matrix_operations::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta, impl->stream);
 }
 
 
diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp
index e9e7d33..719f106 100644
--- a/modules/gpu/src/imgproc_gpu.cpp
+++ b/modules/gpu/src/imgproc_gpu.cpp
@@ -47,20 +47,20 @@ using namespace cv::gpu;
 
 #if !defined (HAVE_CUDA)
 
-void cv::gpu::remap(const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::meanShiftFiltering_GPU(const GpuMat&, GpuMat&, int, int, TermCriteria ) { throw_nogpu(); }
+void cv::gpu::remap( const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap ){ throw_nogpu(); }
+void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria ) { throw_nogpu(); }
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }
-void cv::gpu::reprojectImageTo3D_GPU(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }
-void cv::gpu::reprojectImageTo3D_GPU(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }
-void cv::gpu::cvtColor_GPU(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }
-void cv::gpu::cvtColor_GPU(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }
+void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }
+void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }
+void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }
+void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }
 
 #else /* !defined (HAVE_CUDA) */
 
 namespace cv { namespace gpu 
 { 
-    namespace impl 
+    namespace improc 
     {
         void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
         void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
@@ -90,10 +90,10 @@ namespace cv { namespace gpu
 ////////////////////////////////////////////////////////////////////////
 // remap
 
-void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst)
+void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap)
 {
     typedef void (*remap_gpu_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
-    static const remap_gpu_t callers[] = {impl::remap_gpu_1c, 0, impl::remap_gpu_3c};
+    static const remap_gpu_t callers[] = {improc::remap_gpu_1c, 0, improc::remap_gpu_3c};
 
     CV_Assert((src.type() == CV_8U || src.type() == CV_8UC3) && xmap.type() == CV_32F && ymap.type() == CV_32F);
 
@@ -111,7 +111,7 @@ void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, G
 ////////////////////////////////////////////////////////////////////////
 // meanShiftFiltering_GPU
 
-void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria)
+void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria)
 {       
     if( src.empty() )
         CV_Error( CV_StsBadArg, "The input image is empty" );
@@ -131,7 +131,7 @@ void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int
         eps = 1.f;
     eps = (float)std::max(criteria.epsilon, 0.0);        
 
-    impl::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);    
+    improc::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);    
 }
 
 ////////////////////////////////////////////////////////////////////////
@@ -147,7 +147,7 @@ namespace
             out = dst;
         out.create(src.size(), CV_8UC4);
 
-        impl::drawColorDisp_gpu((DevMem2D_<T>)src, out, ndisp, stream);
+        improc::drawColorDisp_gpu((DevMem2D_<T>)src, out, ndisp, stream);
 
         dst = out;
     }
@@ -180,7 +180,7 @@ namespace
     void reprojectImageTo3D_caller(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream)
     {        
         xyzw.create(disp.rows, disp.cols, CV_32FC4);
-        impl::reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream);
+        improc::reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream);
     }
     
     typedef void (*reprojectImageTo3D_caller_t)(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream);
@@ -188,14 +188,14 @@ namespace
     const reprojectImageTo3D_caller_t reprojectImageTo3D_callers[] = {reprojectImageTo3D_caller<unsigned char>, 0, 0, reprojectImageTo3D_caller<short>, 0, 0, 0, 0};
 }
 
-void cv::gpu::reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q)
+void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q)
 {
     CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);
     
     reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, 0);
 }
 
-void cv::gpu::reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream)
+void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream)
 {
     CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);
     
@@ -229,11 +229,11 @@ namespace
                 
                 out.create(sz, CV_MAKETYPE(depth, dcn));
                 if( depth == CV_8U )
-                    impl::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream);
+                    improc::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream);
                 else if( depth == CV_16U )
-                    impl::RGB2RGB_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, dcn, bidx, stream);
+                    improc::RGB2RGB_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, dcn, bidx, stream);
                 else
-                    impl::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream);
+                    improc::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream);
                 break;
                 
             //case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
@@ -270,11 +270,11 @@ namespace
                 bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
                 
                 if( depth == CV_8U )
-                    impl::RGB2Gray_gpu((DevMem2D)src, scn, (DevMem2D)out, bidx, stream);
+                    improc::RGB2Gray_gpu((DevMem2D)src, scn, (DevMem2D)out, bidx, stream);
                 else if( depth == CV_16U )
-                    impl::RGB2Gray_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, bidx, stream);
+                    improc::RGB2Gray_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, bidx, stream);
                 else
-                    impl::RGB2Gray_gpu((DevMem2Df)src, scn, (DevMem2Df)out, bidx, stream);
+                    improc::RGB2Gray_gpu((DevMem2Df)src, scn, (DevMem2Df)out, bidx, stream);
                 break;
             
             //case CV_BGR5652GRAY: case CV_BGR5552GRAY:
@@ -291,11 +291,11 @@ namespace
                 out.create(sz, CV_MAKETYPE(depth, dcn));
                 
                 if( depth == CV_8U )
-                    impl::Gray2RGB_gpu((DevMem2D)src, (DevMem2D)out, dcn, stream);
+                    improc::Gray2RGB_gpu((DevMem2D)src, (DevMem2D)out, dcn, stream);
                 else if( depth == CV_16U )
-                    impl::Gray2RGB_gpu((DevMem2D_<unsigned short>)src, (DevMem2D_<unsigned short>)out, dcn, stream);
+                    improc::Gray2RGB_gpu((DevMem2D_<unsigned short>)src, (DevMem2D_<unsigned short>)out, dcn, stream);
                 else
-                    impl::Gray2RGB_gpu((DevMem2Df)src, (DevMem2Df)out, dcn, stream);
+                    improc::Gray2RGB_gpu((DevMem2Df)src, (DevMem2Df)out, dcn, stream);
                 break;
                 
             //case CV_GRAY2BGR565: case CV_GRAY2BGR555:
@@ -516,12 +516,12 @@ namespace
     }
 }
 
-void cv::gpu::cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn)
+void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)
 {
     cvtColor_caller(src, dst, code, dcn, 0);
 }
 
-void cv::gpu::cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)
+void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)
 {
     cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));
 }
diff --git a/modules/gpu/src/matrix_operations.cpp b/modules/gpu/src/matrix_operations.cpp
index b075d80..97c1f10 100644
--- a/modules/gpu/src/matrix_operations.cpp
+++ b/modules/gpu/src/matrix_operations.cpp
@@ -120,7 +120,7 @@ void cv::gpu::GpuMat::copyTo( GpuMat& mat, const GpuMat& mask ) const
     else
     {
         mat.create(size(), type());
-        cv::gpu::impl::copy_to_with_mask(*this, mat, depth(), mask, channels());
+        cv::gpu::matrix_operations::copy_to_with_mask(*this, mat, depth(), mask, channels());
     }
 }
 
@@ -146,12 +146,12 @@ void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double be
         psrc = &(temp = *this);
 
     dst.create( size(), rtype );
-    impl::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta);
+    matrix_operations::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta);
 }
 
 GpuMat& GpuMat::operator = (const Scalar& s)
 {
-    impl::set_to_without_mask( *this, depth(), s.val, channels());
+    matrix_operations::set_to_without_mask( *this, depth(), s.val, channels());
     return *this;
 }
 
@@ -162,9 +162,9 @@ GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask)
     CV_DbgAssert(!this->empty());
 
     if (mask.empty())
-        impl::set_to_without_mask( *this, depth(), s.val, channels());
+        matrix_operations::set_to_without_mask( *this, depth(), s.val, channels());
     else
-        impl::set_to_with_mask( *this, depth(), s.val, mask, channels());
+        matrix_operations::set_to_with_mask( *this, depth(), s.val, mask, channels());
 
     return *this;
 }
diff --git a/modules/gpu/src/stereobm_gpu.cpp b/modules/gpu/src/stereobm_gpu.cpp
index 060a63b..1c5716f 100644
--- a/modules/gpu/src/stereobm_gpu.cpp
+++ b/modules/gpu/src/stereobm_gpu.cpp
@@ -58,7 +58,7 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat&, const GpuMat&, GpuMat&,
 
 namespace cv { namespace gpu
 {
-    namespace impl
+    namespace bm
     {
         //extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf);
         extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf, const cudaStream_t & stream);
@@ -115,17 +115,17 @@ static void stereo_bm_gpu_operator ( GpuMat& minSSD,  GpuMat& leBuf, GpuMat&  ri
         leBuf.create( left.size(),  left.type());
         riBuf.create(right.size(), right.type());
 
-        impl::prefilter_xsobel( left, leBuf);
-        impl::prefilter_xsobel(right, riBuf);
+        bm::prefilter_xsobel( left, leBuf);
+        bm::prefilter_xsobel(right, riBuf);
 
         le_for_bm = leBuf;
         ri_for_bm = riBuf;
     }
 
-    impl::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD, stream);
+    bm::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD, stream);
 
     if (avergeTexThreshold)
-        impl::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity);
+        bm::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity);
 }
 
 
diff --git a/tests/gpu/src/meanshift.cpp b/tests/gpu/src/meanshift.cpp
index 072af98..e5443c7 100644
--- a/tests/gpu/src/meanshift.cpp
+++ b/tests/gpu/src/meanshift.cpp
@@ -76,7 +76,7 @@ void CV_GpuMeanShiftTest::run(int)
         cvtColor(img, rgba, CV_BGR2BGRA);
 
         cv::gpu::GpuMat res;
-        cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
+        cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
         if (res.type() != CV_8UC4)
         {
             ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
diff --git a/tests/gpu/src/stereo_bm.cpp b/tests/gpu/src/stereo_bm.cpp
index 45de01e..8631c57 100644
--- a/tests/gpu/src/stereo_bm.cpp
+++ b/tests/gpu/src/stereo_bm.cpp
@@ -51,6 +51,7 @@ class CV_GpuStereoBMTest : public CvTest
 {
 public:
 	CV_GpuStereoBMTest();
+
 protected:
 	void run(int);
 };
@@ -61,9 +62,9 @@ void CV_GpuStereoBMTest::run(int )
 {
 	cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0);
 	cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0);
-	cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
+	cv::Mat img_reference = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
 
-    if (img_l.empty() || img_r.empty() || img_template.empty())
+    if (img_l.empty() || img_r.empty() || img_reference.empty())
     {
         ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
         return;
@@ -71,14 +72,11 @@ void CV_GpuStereoBMTest::run(int )
 
 	cv::gpu::GpuMat disp;
 	cv::gpu::StereoBM_GPU bm(0, 128, 19);
-
 	bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
 
-	//cv::imwrite(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", disp);
-
-	disp.convertTo(disp, img_template.type());
+	disp.convertTo(disp, img_reference.type());
+	double norm = cv::norm(disp, img_reference, cv::NORM_INF);
 
-	double norm = cv::norm(disp, img_template, cv::NORM_INF);
 	if (norm >= 100) 
         ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
 	ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC);