diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt
index 1ed9806944..ecad9e4dd0 100644
--- a/modules/gpu/CMakeLists.txt
+++ b/modules/gpu/CMakeLists.txt
@@ -3,7 +3,8 @@ if(ANDROID OR IOS)
 endif()
 
 set(the_description "GPU-accelerated Computer Vision")
-ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy opencv_gpuarithm opencv_gpufilters)
+
+ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy opencv_gpuarithm opencv_gpufilters OPTIONAL opencv_gpunvidia)
 
 ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda")
 
@@ -18,12 +19,7 @@ source_group("Src\\Host"      FILES ${lib_srcs} ${lib_int_hdrs})
 source_group("Src\\Cuda"      FILES ${lib_cuda} ${lib_cuda_hdrs})
 
 if(HAVE_CUDA)
-  file(GLOB_RECURSE ncv_srcs "src/nvidia/*.cpp" "src/nvidia/*.h*")
-  file(GLOB_RECURSE ncv_cuda "src/nvidia/*.cu")
-  set(ncv_files ${ncv_srcs} ${ncv_cuda})
-
-  source_group("Src\\NVidia" FILES ${ncv_files})
-  ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging" ${CUDA_INCLUDE_DIRS})
+  ocv_include_directories(${CUDA_INCLUDE_DIRS})
   ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter /wd4211 /wd4201 /wd4100 /wd4505 /wd4408)
 
   if(MSVC)
@@ -43,12 +39,11 @@ else()
   set(lib_cuda "")
   set(cuda_objs "")
   set(cuda_link_libs "")
-  set(ncv_files "")
 endif()
 
 ocv_set_module_sources(
   HEADERS ${lib_hdrs}
-  SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_srcs} ${lib_cuda} ${ncv_files} ${cuda_objs}
+  SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_srcs} ${lib_cuda} ${cuda_objs}
   )
 
 ocv_create_module(${cuda_link_libs})
@@ -57,10 +52,6 @@ if(HAVE_CUDA)
   if(HAVE_CUFFT)
     CUDA_ADD_CUFFT_TO_TARGET(${the_module})
   endif()
-
-  install(FILES src/nvidia/NPP_staging/NPP_staging.hpp  src/nvidia/core/NCV.hpp
-    DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2/${name}
-    COMPONENT main)
 endif()
 
 ocv_add_precompiled_headers(${the_module})
@@ -71,15 +62,8 @@ ocv_add_precompiled_headers(${the_module})
 file(GLOB test_srcs "test/*.cpp")
 file(GLOB test_hdrs "test/*.hpp" "test/*.h")
 
-set(nvidia "")
-if(HAVE_CUDA)
-  file(GLOB nvidia "test/nvidia/*.cpp" "test/nvidia/*.hpp" "test/nvidia/*.h")
-  set(nvidia FILES "Src\\\\\\\\NVidia" ${nvidia}) # 8 ugly backslashes :'(
-endif()
-
 ocv_add_accuracy_tests(FILES "Include" ${test_hdrs}
-                       FILES "Src" ${test_srcs}
-                       ${nvidia})
+                       FILES "Src" ${test_srcs})
 ocv_add_perf_tests()
 
 if(HAVE_CUDA)
diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp
index c7514bb217..0b9f9aafc2 100644
--- a/modules/gpu/src/cascadeclassifier.cpp
+++ b/modules/gpu/src/cascadeclassifier.cpp
@@ -722,240 +722,3 @@ bool cv::gpu::CascadeClassifier_GPU::load(const String& filename)
 }
 
 #endif
-
-//////////////////////////////////////////////////////////////////////////////////////////////////////
-
-#if defined (HAVE_CUDA)
-
-struct RectConvert
-{
-    Rect operator()(const NcvRect32u& nr) const { return Rect(nr.x, nr.y, nr.width, nr.height); }
-    NcvRect32u operator()(const Rect& nr) const
-    {
-        NcvRect32u rect;
-        rect.x = nr.x;
-        rect.y = nr.y;
-        rect.width = nr.width;
-        rect.height = nr.height;
-        return rect;
-    }
-};
-
-void groupRectangles(std::vector<NcvRect32u> &hypotheses, int groupThreshold, double eps, std::vector<Ncv32u> *weights)
-{
-    std::vector<Rect> rects(hypotheses.size());
-    std::transform(hypotheses.begin(), hypotheses.end(), rects.begin(), RectConvert());
-
-    if (weights)
-    {
-        std::vector<int> weights_int;
-        weights_int.assign(weights->begin(), weights->end());
-        cv::groupRectangles(rects, weights_int, groupThreshold, eps);
-    }
-    else
-    {
-        cv::groupRectangles(rects, groupThreshold, eps);
-    }
-    std::transform(rects.begin(), rects.end(), hypotheses.begin(), RectConvert());
-    hypotheses.resize(rects.size());
-}
-
-NCVStatus loadFromXML(const String &filename,
-                      HaarClassifierCascadeDescriptor &haar,
-                      std::vector<HaarStage64> &haarStages,
-                      std::vector<HaarClassifierNode128> &haarClassifierNodes,
-                      std::vector<HaarFeature64> &haarFeatures)
-{
-    NCVStatus ncvStat;
-
-    haar.NumStages = 0;
-    haar.NumClassifierRootNodes = 0;
-    haar.NumClassifierTotalNodes = 0;
-    haar.NumFeatures = 0;
-    haar.ClassifierSize.width = 0;
-    haar.ClassifierSize.height = 0;
-    haar.bHasStumpsOnly = true;
-    haar.bNeedsTiltedII = false;
-    Ncv32u curMaxTreeDepth;
-
-    std::vector<char> xmlFileCont;
-
-    std::vector<HaarClassifierNode128> h_TmpClassifierNotRootNodes;
-    haarStages.resize(0);
-    haarClassifierNodes.resize(0);
-    haarFeatures.resize(0);
-
-    Ptr<CvHaarClassifierCascade> oldCascade = (CvHaarClassifierCascade*)cvLoad(filename.c_str(), 0, 0, 0);
-    if (oldCascade.empty())
-    {
-        return NCV_HAAR_XML_LOADING_EXCEPTION;
-    }
-
-    haar.ClassifierSize.width = oldCascade->orig_window_size.width;
-    haar.ClassifierSize.height = oldCascade->orig_window_size.height;
-
-    int stagesCound = oldCascade->count;
-    for(int s = 0; s < stagesCound; ++s) // by stages
-    {
-        HaarStage64 curStage;
-        curStage.setStartClassifierRootNodeOffset(static_cast<Ncv32u>(haarClassifierNodes.size()));
-
-        curStage.setStageThreshold(oldCascade->stage_classifier[s].threshold);
-
-        int treesCount = oldCascade->stage_classifier[s].count;
-        for(int t = 0; t < treesCount; ++t) // by trees
-        {
-            Ncv32u nodeId = 0;
-            CvHaarClassifier* tree = &oldCascade->stage_classifier[s].classifier[t];
-
-            int nodesCount = tree->count;
-            for(int n = 0; n < nodesCount; ++n)  //by features
-            {
-                CvHaarFeature* feature = &tree->haar_feature[n];
-
-                HaarClassifierNode128 curNode;
-                curNode.setThreshold(tree->threshold[n]);
-
-                NcvBool bIsLeftNodeLeaf = false;
-                NcvBool bIsRightNodeLeaf = false;
-
-                HaarClassifierNodeDescriptor32 nodeLeft;
-                if ( tree->left[n] <= 0 )
-                {
-                    Ncv32f leftVal = tree->alpha[-tree->left[n]];
-                    ncvStat = nodeLeft.create(leftVal);
-                    ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat);
-                    bIsLeftNodeLeaf = true;
-                }
-                else
-                {
-                    Ncv32u leftNodeOffset = tree->left[n];
-                    nodeLeft.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + leftNodeOffset - 1));
-                    haar.bHasStumpsOnly = false;
-                }
-                curNode.setLeftNodeDesc(nodeLeft);
-
-                HaarClassifierNodeDescriptor32 nodeRight;
-                if ( tree->right[n] <= 0 )
-                {
-                    Ncv32f rightVal = tree->alpha[-tree->right[n]];
-                    ncvStat = nodeRight.create(rightVal);
-                    ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat);
-                    bIsRightNodeLeaf = true;
-                }
-                else
-                {
-                    Ncv32u rightNodeOffset = tree->right[n];
-                    nodeRight.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + rightNodeOffset - 1));
-                    haar.bHasStumpsOnly = false;
-                }
-                curNode.setRightNodeDesc(nodeRight);
-
-                Ncv32u tiltedVal = feature->tilted;
-                haar.bNeedsTiltedII = (tiltedVal != 0);
-
-                Ncv32u featureId = 0;
-                for(int l = 0; l < CV_HAAR_FEATURE_MAX; ++l) //by rects
-                {
-                    Ncv32u rectX = feature->rect[l].r.x;
-                    Ncv32u rectY = feature->rect[l].r.y;
-                    Ncv32u rectWidth = feature->rect[l].r.width;
-                    Ncv32u rectHeight = feature->rect[l].r.height;
-
-                    Ncv32f rectWeight = feature->rect[l].weight;
-
-                    if (rectWeight == 0/* && rectX == 0 &&rectY == 0 && rectWidth == 0 && rectHeight == 0*/)
-                        break;
-
-                    HaarFeature64 curFeature;
-                    ncvStat = curFeature.setRect(rectX, rectY, rectWidth, rectHeight, haar.ClassifierSize.width, haar.ClassifierSize.height);
-                    curFeature.setWeight(rectWeight);
-                    ncvAssertReturn(NCV_SUCCESS == ncvStat, ncvStat);
-                    haarFeatures.push_back(curFeature);
-
-                    featureId++;
-                }
-
-                HaarFeatureDescriptor32 tmpFeatureDesc;
-                ncvStat = tmpFeatureDesc.create(haar.bNeedsTiltedII, bIsLeftNodeLeaf, bIsRightNodeLeaf,
-                    featureId, static_cast<Ncv32u>(haarFeatures.size()) - featureId);
-                ncvAssertReturn(NCV_SUCCESS == ncvStat, ncvStat);
-                curNode.setFeatureDesc(tmpFeatureDesc);
-
-                if (!nodeId)
-                {
-                    //root node
-                    haarClassifierNodes.push_back(curNode);
-                    curMaxTreeDepth = 1;
-                }
-                else
-                {
-                    //other node
-                    h_TmpClassifierNotRootNodes.push_back(curNode);
-                    curMaxTreeDepth++;
-                }
-
-                nodeId++;
-            }
-        }
-
-        curStage.setNumClassifierRootNodes(treesCount);
-        haarStages.push_back(curStage);
-    }
-
-    //fill in cascade stats
-    haar.NumStages = static_cast<Ncv32u>(haarStages.size());
-    haar.NumClassifierRootNodes = static_cast<Ncv32u>(haarClassifierNodes.size());
-    haar.NumClassifierTotalNodes = static_cast<Ncv32u>(haar.NumClassifierRootNodes + h_TmpClassifierNotRootNodes.size());
-    haar.NumFeatures = static_cast<Ncv32u>(haarFeatures.size());
-
-    //merge root and leaf nodes in one classifiers array
-    Ncv32u offsetRoot = static_cast<Ncv32u>(haarClassifierNodes.size());
-    for (Ncv32u i=0; i<haarClassifierNodes.size(); i++)
-    {
-        HaarFeatureDescriptor32 featureDesc = haarClassifierNodes[i].getFeatureDesc();
-
-        HaarClassifierNodeDescriptor32 nodeLeft = haarClassifierNodes[i].getLeftNodeDesc();
-        if (!featureDesc.isLeftNodeLeaf())
-        {
-            Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot;
-            nodeLeft.create(newOffset);
-        }
-        haarClassifierNodes[i].setLeftNodeDesc(nodeLeft);
-
-        HaarClassifierNodeDescriptor32 nodeRight = haarClassifierNodes[i].getRightNodeDesc();
-        if (!featureDesc.isRightNodeLeaf())
-        {
-            Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot;
-            nodeRight.create(newOffset);
-        }
-        haarClassifierNodes[i].setRightNodeDesc(nodeRight);
-    }
-
-    for (Ncv32u i=0; i<h_TmpClassifierNotRootNodes.size(); i++)
-    {
-        HaarFeatureDescriptor32 featureDesc = h_TmpClassifierNotRootNodes[i].getFeatureDesc();
-
-        HaarClassifierNodeDescriptor32 nodeLeft = h_TmpClassifierNotRootNodes[i].getLeftNodeDesc();
-        if (!featureDesc.isLeftNodeLeaf())
-        {
-            Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot;
-            nodeLeft.create(newOffset);
-        }
-        h_TmpClassifierNotRootNodes[i].setLeftNodeDesc(nodeLeft);
-
-        HaarClassifierNodeDescriptor32 nodeRight = h_TmpClassifierNotRootNodes[i].getRightNodeDesc();
-        if (!featureDesc.isRightNodeLeaf())
-        {
-            Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot;
-            nodeRight.create(newOffset);
-        }
-        h_TmpClassifierNotRootNodes[i].setRightNodeDesc(nodeRight);
-
-        haarClassifierNodes.push_back(h_TmpClassifierNotRootNodes[i]);
-    }
-
-    return NCV_SUCCESS;
-}
-
-#endif /* HAVE_CUDA */
diff --git a/modules/gpu/src/cuda/internal_shared.hpp b/modules/gpu/src/cuda/internal_shared.hpp
index 4362451f82..c8d4e5b7f6 100644
--- a/modules/gpu/src/cuda/internal_shared.hpp
+++ b/modules/gpu/src/cuda/internal_shared.hpp
@@ -45,10 +45,12 @@
 
 #include <cuda_runtime.h>
 #include <npp.h>
-#include "NPP_staging.hpp"
-#include "opencv2/gpu/devmem2d.hpp"
-#include "safe_call.hpp"
+
+#include "opencv2/core/cuda_devptrs.hpp"
 #include "opencv2/core/cuda/common.hpp"
+#include "opencv2/gpunvidia.hpp"
+
+#include "safe_call.hpp"
 
 namespace cv { namespace gpu
 {
diff --git a/modules/gpu/src/cuda/safe_call.hpp b/modules/gpu/src/cuda/safe_call.hpp
index 10b72e6230..fa62fb15aa 100644
--- a/modules/gpu/src/cuda/safe_call.hpp
+++ b/modules/gpu/src/cuda/safe_call.hpp
@@ -45,7 +45,7 @@
 
 #include <cuda_runtime_api.h>
 #include <cufft.h>
-#include "NCV.hpp"
+#include "opencv2/gpunvidia.hpp"
 
 #if defined(__GNUC__)
     #define ncvSafeCall(expr)  ___ncvSafeCall(expr, __FILE__, __LINE__, __func__)
diff --git a/modules/gpu/src/precomp.hpp b/modules/gpu/src/precomp.hpp
index 69ddeaed9e..0127bd28ea 100644
--- a/modules/gpu/src/precomp.hpp
+++ b/modules/gpu/src/precomp.hpp
@@ -79,10 +79,7 @@
     #include "internal_shared.hpp"
     #include "opencv2/core/stream_accessor.hpp"
 
-    #include "nvidia/core/NCV.hpp"
-    #include "nvidia/NPP_staging/NPP_staging.hpp"
-    #include "nvidia/NCVHaarObjectDetection.hpp"
-    #include "nvidia/NCVBroxOpticalFlow.hpp"
+    #include "opencv2/gpunvidia.hpp"
 #endif /* defined(HAVE_CUDA) */
 
 #endif /* __OPENCV_PRECOMP_H__ */
diff --git a/modules/gpu/test/test_main.cpp b/modules/gpu/test/test_main.cpp
new file mode 100644
index 0000000000..eea3d7c008
--- /dev/null
+++ b/modules/gpu/test/test_main.cpp
@@ -0,0 +1,45 @@
+/*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 "test_precomp.hpp"
+
+CV_GPU_TEST_MAIN("gpu")
diff --git a/modules/gpu/test/test_precomp.hpp b/modules/gpu/test/test_precomp.hpp
index f6d1392c76..08807d51e0 100644
--- a/modules/gpu/test/test_precomp.hpp
+++ b/modules/gpu/test/test_precomp.hpp
@@ -75,7 +75,6 @@
 #include "opencv2/gpu.hpp"
 
 #include "interpolation.hpp"
-#include "main_test_nvidia.h"
 
 #include "opencv2/core/gpu_private.hpp"
 
diff --git a/modules/gpunvidia/CMakeLists.txt b/modules/gpunvidia/CMakeLists.txt
new file mode 100644
index 0000000000..7c15424309
--- /dev/null
+++ b/modules/gpunvidia/CMakeLists.txt
@@ -0,0 +1,9 @@
+if(NOT HAVE_CUDA)
+  ocv_module_disable(gpunvidia)
+endif()
+
+set(the_description "GPU-accelerated Computer Vision (HAL module)")
+
+ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations)
+
+ocv_define_module(gpunvidia opencv_core opencv_objdetect)
diff --git a/modules/gpunvidia/include/opencv2/gpunvidia.hpp b/modules/gpunvidia/include/opencv2/gpunvidia.hpp
new file mode 100644
index 0000000000..c59dc64021
--- /dev/null
+++ b/modules/gpunvidia/include/opencv2/gpunvidia.hpp
@@ -0,0 +1,52 @@
+/*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*/
+
+#ifndef __OPENCV_GPUNVIDIA_HPP__
+#define __OPENCV_GPUNVIDIA_HPP__
+
+#include "opencv2/gpunvidia/NCV.hpp"
+#include "opencv2/gpunvidia/NPP_staging.hpp"
+#include "opencv2/gpunvidia/NCVPyramid.hpp"
+#include "opencv2/gpunvidia/NCVBroxOpticalFlow.hpp"
+#include "opencv2/gpunvidia/NCVHaarObjectDetection.hpp"
+
+#endif /* __OPENCV_GPUNVIDIA_HPP__ */
diff --git a/modules/gpu/src/nvidia/core/NCV.hpp b/modules/gpunvidia/include/opencv2/gpunvidia/NCV.hpp
similarity index 100%
rename from modules/gpu/src/nvidia/core/NCV.hpp
rename to modules/gpunvidia/include/opencv2/gpunvidia/NCV.hpp
diff --git a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.hpp b/modules/gpunvidia/include/opencv2/gpunvidia/NCVBroxOpticalFlow.hpp
similarity index 99%
rename from modules/gpu/src/nvidia/NCVBroxOpticalFlow.hpp
rename to modules/gpunvidia/include/opencv2/gpunvidia/NCVBroxOpticalFlow.hpp
index 139f7bea8c..3300e1006a 100644
--- a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.hpp
+++ b/modules/gpunvidia/include/opencv2/gpunvidia/NCVBroxOpticalFlow.hpp
@@ -60,7 +60,7 @@
 #ifndef _ncv_optical_flow_h_
 #define _ncv_optical_flow_h_
 
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 /// \brief Model and solver parameters
 struct NCVBroxOpticalFlowDescriptor
diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp b/modules/gpunvidia/include/opencv2/gpunvidia/NCVHaarObjectDetection.hpp
similarity index 99%
rename from modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
rename to modules/gpunvidia/include/opencv2/gpunvidia/NCVHaarObjectDetection.hpp
index bbda61fecd..c067a9135e 100644
--- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.hpp
+++ b/modules/gpunvidia/include/opencv2/gpunvidia/NCVHaarObjectDetection.hpp
@@ -59,7 +59,7 @@
 #ifndef _ncvhaarobjectdetection_hpp_
 #define _ncvhaarobjectdetection_hpp_
 
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 
 //==============================================================================
diff --git a/modules/gpu/src/nvidia/core/NCVPyramid.hpp b/modules/gpunvidia/include/opencv2/gpunvidia/NCVPyramid.hpp
similarity index 98%
rename from modules/gpu/src/nvidia/core/NCVPyramid.hpp
rename to modules/gpunvidia/include/opencv2/gpunvidia/NCVPyramid.hpp
index 183428afa6..c88dbc2712 100644
--- a/modules/gpu/src/nvidia/core/NCVPyramid.hpp
+++ b/modules/gpunvidia/include/opencv2/gpunvidia/NCVPyramid.hpp
@@ -45,7 +45,7 @@
 
 #include <memory>
 #include <vector>
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 #if 0 //def _WIN32
 
diff --git a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.hpp b/modules/gpunvidia/include/opencv2/gpunvidia/NPP_staging.hpp
similarity index 99%
rename from modules/gpu/src/nvidia/NPP_staging/NPP_staging.hpp
rename to modules/gpunvidia/include/opencv2/gpunvidia/NPP_staging.hpp
index 073448e17c..823be69432 100644
--- a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.hpp
+++ b/modules/gpunvidia/include/opencv2/gpunvidia/NPP_staging.hpp
@@ -43,7 +43,7 @@
 #ifndef _npp_staging_hpp_
 #define _npp_staging_hpp_
 
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 
 /**
diff --git a/modules/gpu/src/nvidia/core/NCV.cu b/modules/gpunvidia/src/NCV.cpp
similarity index 82%
rename from modules/gpu/src/nvidia/core/NCV.cu
rename to modules/gpunvidia/src/NCV.cpp
index 718b4faabd..979276c7c8 100644
--- a/modules/gpu/src/nvidia/core/NCV.cu
+++ b/modules/gpunvidia/src/NCV.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#include <iostream>
-#include <vector>
-#include "NCV.hpp"
-
+#include "precomp.hpp"
 
 //==============================================================================
 //
@@ -72,8 +69,6 @@ void ncvSetDebugOutputHandler(NCVDebugOutputHandler *func)
     debugOutputHandler = func;
 }
 
-#if !defined CUDA_DISABLER
-
 
 //==============================================================================
 //
@@ -251,15 +246,14 @@ NCVStatus memSegCopyHelper2D(void *dst, Ncv32u dstPitch, NCVMemoryType dstType,
 //===================================================================
 
 
-NCVMemStackAllocator::NCVMemStackAllocator(Ncv32u alignment_)
-    :
-    currentSize(0),
-    _maxSize(0),
+NCVMemStackAllocator::NCVMemStackAllocator(Ncv32u alignment_) :
+    _memType(NCVMemoryTypeNone),
+    _alignment(alignment_),
     allocBegin(NULL),
     begin(NULL),
     end(NULL),
-    _memType(NCVMemoryTypeNone),
-    _alignment(alignment_),
+    currentSize(0),
+    _maxSize(0),
     bReusesMemory(false)
 {
     NcvBool bProperAlignment = (alignment_ & (alignment_ - 1)) == 0;
@@ -267,13 +261,12 @@ NCVMemStackAllocator::NCVMemStackAllocator(Ncv32u alignment_)
 }
 
 
-NCVMemStackAllocator::NCVMemStackAllocator(NCVMemoryType memT, size_t capacity, Ncv32u alignment_, void *reusePtr)
-    :
-    currentSize(0),
-    _maxSize(0),
-    allocBegin(NULL),
+NCVMemStackAllocator::NCVMemStackAllocator(NCVMemoryType memT, size_t capacity, Ncv32u alignment_, void *reusePtr) :
     _memType(memT),
-    _alignment(alignment_)
+    _alignment(alignment_),
+    allocBegin(NULL),
+    currentSize(0),
+    _maxSize(0)
 {
     NcvBool bProperAlignment = (alignment_ & (alignment_ - 1)) == 0;
     ncvAssertPrintCheck(bProperAlignment, "NCVMemStackAllocator ctor:: _alignment not power of 2");
@@ -389,7 +382,7 @@ NCVStatus NCVMemStackAllocator::dealloc(NCVMemSegment &seg)
 
 NcvBool NCVMemStackAllocator::isInitialized(void) const
 {
-    return ((this->_alignment & (this->_alignment-1)) == 0) && isCounting() || this->allocBegin != NULL;
+    return (((this->_alignment & (this->_alignment-1)) == 0) && isCounting()) || this->allocBegin != NULL;
 }
 
 
@@ -424,12 +417,11 @@ size_t NCVMemStackAllocator::maxSize(void) const
 //===================================================================
 
 
-NCVMemNativeAllocator::NCVMemNativeAllocator(NCVMemoryType memT, Ncv32u alignment_)
-    :
-    currentSize(0),
-    _maxSize(0),
+NCVMemNativeAllocator::NCVMemNativeAllocator(NCVMemoryType memT, Ncv32u alignment_) :
     _memType(memT),
-    _alignment(alignment_)
+    _alignment(alignment_),
+    currentSize(0),
+    _maxSize(0)
 {
     ncvAssertPrintReturn(memT != NCVMemoryTypeNone, "NCVMemNativeAllocator ctor:: counting not permitted for this allocator type", );
 }
@@ -649,9 +641,46 @@ double ncvEndQueryTimerMs(NcvTimer t)
 //
 //===================================================================
 
+struct RectConvert
+{
+    cv::Rect operator()(const NcvRect32u& nr) const { return cv::Rect(nr.x, nr.y, nr.width, nr.height); }
+    NcvRect32u operator()(const cv::Rect& nr) const
+    {
+        NcvRect32u rect;
+        rect.x = nr.x;
+        rect.y = nr.y;
+        rect.width = nr.width;
+        rect.height = nr.height;
+        return rect;
+    }
+};
+
+static void groupRectangles(std::vector<NcvRect32u> &hypotheses, int groupThreshold, double eps, std::vector<Ncv32u> *weights)
+{
+    std::vector<cv::Rect> rects(hypotheses.size());
+    std::transform(hypotheses.begin(), hypotheses.end(), rects.begin(), RectConvert());
+
+    if (weights)
+    {
+        std::vector<int> weights_int;
+        weights_int.assign(weights->begin(), weights->end());
+        cv::groupRectangles(rects, weights_int, groupThreshold, eps);
+    }
+    else
+    {
+        cv::groupRectangles(rects, groupThreshold, eps);
+    }
+    std::transform(rects.begin(), rects.end(), hypotheses.begin(), RectConvert());
+    hypotheses.resize(rects.size());
+}
+
+
+//===================================================================
+//
+// Operations with rectangles
+//
+//===================================================================
 
-//from OpenCV
-void groupRectangles(std::vector<NcvRect32u> &hypotheses, int groupThreshold, double eps, std::vector<Ncv32u> *weights);
 
 
 NCVStatus ncvGroupRectangles_host(NCVVector<NcvRect32u> &hypotheses,
@@ -776,133 +805,3 @@ NCVStatus ncvDrawRects_32u_host(Ncv32u *h_dst,
 {
     return drawRectsWrapperHost(h_dst, dstStride, dstWidth, dstHeight, h_rects, numRects, color);
 }
-
-
-const Ncv32u NUMTHREADS_DRAWRECTS = 32;
-const Ncv32u NUMTHREADS_DRAWRECTS_LOG2 = 5;
-
-
-template <class T>
-__global__ void drawRects(T *d_dst,
-                          Ncv32u dstStride,
-                          Ncv32u dstWidth,
-                          Ncv32u dstHeight,
-                          NcvRect32u *d_rects,
-                          Ncv32u numRects,
-                          T color)
-{
-    Ncv32u blockId = blockIdx.y * 65535 + blockIdx.x;
-    if (blockId > numRects * 4)
-    {
-        return;
-    }
-
-    NcvRect32u curRect = d_rects[blockId >> 2];
-    NcvBool bVertical = blockId & 0x1;
-    NcvBool bTopLeft = blockId & 0x2;
-
-    Ncv32u pt0x, pt0y;
-    if (bVertical)
-    {
-        Ncv32u numChunks = (curRect.height + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2;
-
-        pt0x = bTopLeft ? curRect.x : curRect.x + curRect.width - 1;
-        pt0y = curRect.y;
-
-        if (pt0x < dstWidth)
-        {
-            for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++)
-            {
-                Ncv32u ptY = pt0y + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x;
-                if (ptY < pt0y + curRect.height && ptY < dstHeight)
-                {
-                    d_dst[ptY * dstStride + pt0x] = color;
-                }
-            }
-        }
-    }
-    else
-    {
-        Ncv32u numChunks = (curRect.width + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2;
-
-        pt0x = curRect.x;
-        pt0y = bTopLeft ? curRect.y : curRect.y + curRect.height - 1;
-
-        if (pt0y < dstHeight)
-        {
-            for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++)
-            {
-                Ncv32u ptX = pt0x + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x;
-                if (ptX < pt0x + curRect.width && ptX < dstWidth)
-                {
-                    d_dst[pt0y * dstStride + ptX] = color;
-                }
-            }
-        }
-    }
-}
-
-
-template <class T>
-static NCVStatus drawRectsWrapperDevice(T *d_dst,
-                                        Ncv32u dstStride,
-                                        Ncv32u dstWidth,
-                                        Ncv32u dstHeight,
-                                        NcvRect32u *d_rects,
-                                        Ncv32u numRects,
-                                        T color,
-                                        cudaStream_t cuStream)
-{
-    (void)cuStream;
-    ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR);
-    ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID);
-    ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP);
-    ncvAssertReturn(numRects <= dstWidth * dstHeight, NCV_DIMENSIONS_INVALID);
-
-    if (numRects == 0)
-    {
-        return NCV_SUCCESS;
-    }
-
-    dim3 grid(numRects * 4);
-    dim3 block(NUMTHREADS_DRAWRECTS);
-    if (grid.x > 65535)
-    {
-        grid.y = (grid.x + 65534) / 65535;
-        grid.x = 65535;
-    }
-
-    drawRects<T><<<grid, block>>>(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color);
-
-    ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR);
-
-    return NCV_SUCCESS;
-}
-
-
-NCVStatus ncvDrawRects_8u_device(Ncv8u *d_dst,
-                                 Ncv32u dstStride,
-                                 Ncv32u dstWidth,
-                                 Ncv32u dstHeight,
-                                 NcvRect32u *d_rects,
-                                 Ncv32u numRects,
-                                 Ncv8u color,
-                                 cudaStream_t cuStream)
-{
-    return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream);
-}
-
-
-NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst,
-                                  Ncv32u dstStride,
-                                  Ncv32u dstWidth,
-                                  Ncv32u dstHeight,
-                                  NcvRect32u *d_rects,
-                                  Ncv32u numRects,
-                                  Ncv32u color,
-                                  cudaStream_t cuStream)
-{
-    return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream);
-}
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpunvidia/src/cuda/NCV.cu b/modules/gpunvidia/src/cuda/NCV.cu
new file mode 100644
index 0000000000..0e5b50e9af
--- /dev/null
+++ b/modules/gpunvidia/src/cuda/NCV.cu
@@ -0,0 +1,180 @@
+/*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 <iostream>
+#include <vector>
+
+#include "opencv2/gpunvidia/NCV.hpp"
+
+//===================================================================
+//
+// Operations with rectangles
+//
+//===================================================================
+
+
+const Ncv32u NUMTHREADS_DRAWRECTS = 32;
+const Ncv32u NUMTHREADS_DRAWRECTS_LOG2 = 5;
+
+
+template <class T>
+__global__ void drawRects(T *d_dst,
+                          Ncv32u dstStride,
+                          Ncv32u dstWidth,
+                          Ncv32u dstHeight,
+                          NcvRect32u *d_rects,
+                          Ncv32u numRects,
+                          T color)
+{
+    Ncv32u blockId = blockIdx.y * 65535 + blockIdx.x;
+    if (blockId > numRects * 4)
+    {
+        return;
+    }
+
+    NcvRect32u curRect = d_rects[blockId >> 2];
+    NcvBool bVertical = blockId & 0x1;
+    NcvBool bTopLeft = blockId & 0x2;
+
+    Ncv32u pt0x, pt0y;
+    if (bVertical)
+    {
+        Ncv32u numChunks = (curRect.height + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2;
+
+        pt0x = bTopLeft ? curRect.x : curRect.x + curRect.width - 1;
+        pt0y = curRect.y;
+
+        if (pt0x < dstWidth)
+        {
+            for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++)
+            {
+                Ncv32u ptY = pt0y + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x;
+                if (ptY < pt0y + curRect.height && ptY < dstHeight)
+                {
+                    d_dst[ptY * dstStride + pt0x] = color;
+                }
+            }
+        }
+    }
+    else
+    {
+        Ncv32u numChunks = (curRect.width + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2;
+
+        pt0x = curRect.x;
+        pt0y = bTopLeft ? curRect.y : curRect.y + curRect.height - 1;
+
+        if (pt0y < dstHeight)
+        {
+            for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++)
+            {
+                Ncv32u ptX = pt0x + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x;
+                if (ptX < pt0x + curRect.width && ptX < dstWidth)
+                {
+                    d_dst[pt0y * dstStride + ptX] = color;
+                }
+            }
+        }
+    }
+}
+
+
+template <class T>
+static NCVStatus drawRectsWrapperDevice(T *d_dst,
+                                        Ncv32u dstStride,
+                                        Ncv32u dstWidth,
+                                        Ncv32u dstHeight,
+                                        NcvRect32u *d_rects,
+                                        Ncv32u numRects,
+                                        T color,
+                                        cudaStream_t cuStream)
+{
+    (void)cuStream;
+    ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR);
+    ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID);
+    ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP);
+    ncvAssertReturn(numRects <= dstWidth * dstHeight, NCV_DIMENSIONS_INVALID);
+
+    if (numRects == 0)
+    {
+        return NCV_SUCCESS;
+    }
+
+    dim3 grid(numRects * 4);
+    dim3 block(NUMTHREADS_DRAWRECTS);
+    if (grid.x > 65535)
+    {
+        grid.y = (grid.x + 65534) / 65535;
+        grid.x = 65535;
+    }
+
+    drawRects<T><<<grid, block>>>(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color);
+
+    ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR);
+
+    return NCV_SUCCESS;
+}
+
+
+NCVStatus ncvDrawRects_8u_device(Ncv8u *d_dst,
+                                 Ncv32u dstStride,
+                                 Ncv32u dstWidth,
+                                 Ncv32u dstHeight,
+                                 NcvRect32u *d_rects,
+                                 Ncv32u numRects,
+                                 Ncv8u color,
+                                 cudaStream_t cuStream)
+{
+    return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream);
+}
+
+
+NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst,
+                                  Ncv32u dstStride,
+                                  Ncv32u dstWidth,
+                                  Ncv32u dstHeight,
+                                  NcvRect32u *d_rects,
+                                  Ncv32u numRects,
+                                  Ncv32u color,
+                                  cudaStream_t cuStream)
+{
+    return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream);
+}
diff --git a/modules/gpu/src/nvidia/core/NCVAlg.hpp b/modules/gpunvidia/src/cuda/NCVAlg.hpp
similarity index 99%
rename from modules/gpu/src/nvidia/core/NCVAlg.hpp
rename to modules/gpunvidia/src/cuda/NCVAlg.hpp
index 3a0a282204..ad14d749fd 100644
--- a/modules/gpu/src/nvidia/core/NCVAlg.hpp
+++ b/modules/gpunvidia/src/cuda/NCVAlg.hpp
@@ -43,7 +43,7 @@
 #ifndef _ncv_alg_hpp_
 #define _ncv_alg_hpp_
 
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 
 template <class T>
diff --git a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu b/modules/gpunvidia/src/cuda/NCVBroxOpticalFlow.cu
similarity index 99%
rename from modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu
rename to modules/gpunvidia/src/cuda/NCVBroxOpticalFlow.cu
index f865f957f6..4faba6331f 100644
--- a/modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu
+++ b/modules/gpunvidia/src/cuda/NCVBroxOpticalFlow.cu
@@ -57,16 +57,15 @@
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-#if !defined CUDA_DISABLER
-
 #include <iostream>
 #include <vector>
 #include <memory>
 
-#include "NPP_staging/NPP_staging.hpp"
-#include "NCVBroxOpticalFlow.hpp"
 #include "opencv2/core/cuda/utility.hpp"
 
+#include "opencv2/gpunvidia/NPP_staging.hpp"
+#include "opencv2/gpunvidia/NCVBroxOpticalFlow.hpp"
+
 
 typedef NCVVectorAlloc<Ncv32f> FloatVector;
 
@@ -1163,5 +1162,3 @@ NCVStatus NCVBroxOpticalFlow(const NCVBroxOpticalFlowDescriptor desc,
 
     return NCV_SUCCESS;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/src/nvidia/core/NCVColorConversion.hpp b/modules/gpunvidia/src/cuda/NCVColorConversion.hpp
similarity index 100%
rename from modules/gpu/src/nvidia/core/NCVColorConversion.hpp
rename to modules/gpunvidia/src/cuda/NCVColorConversion.hpp
diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu b/modules/gpunvidia/src/cuda/NCVHaarObjectDetection.cu
similarity index 92%
rename from modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
rename to modules/gpunvidia/src/cuda/NCVHaarObjectDetection.cu
index f7a913bc69..9ab0194a4e 100644
--- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
+++ b/modules/gpunvidia/src/cuda/NCVHaarObjectDetection.cu
@@ -56,18 +56,19 @@
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-#if !defined CUDA_DISABLER
-
 #include <algorithm>
 #include <cstdio>
 
-#include "NCV.hpp"
-#include "NCVAlg.hpp"
-#include "NPP_staging/NPP_staging.hpp"
-#include "NCVRuntimeTemplates.hpp"
-#include "NCVHaarObjectDetection.hpp"
 #include "opencv2/core/cuda/warp.hpp"
 #include "opencv2/core/cuda/warp_shuffle.hpp"
+#include "opencv2/objdetect.hpp"
+
+#include "opencv2/gpunvidia/NCV.hpp"
+#include "opencv2/gpunvidia/NPP_staging.hpp"
+#include "opencv2/gpunvidia/NCVHaarObjectDetection.hpp"
+
+#include "NCVRuntimeTemplates.hpp"
+#include "NCVAlg.hpp"
 
 
 //==============================================================================
@@ -2099,12 +2100,201 @@ NCVStatus ncvGrowDetectionsVector_host(NCVVector<Ncv32u> &pixelMask,
     return ncvStat;
 }
 
-
-NCVStatus loadFromXML(const cv::String &filename,
+static NCVStatus loadFromXML(const cv::String &filename,
                       HaarClassifierCascadeDescriptor &haar,
                       std::vector<HaarStage64> &haarStages,
                       std::vector<HaarClassifierNode128> &haarClassifierNodes,
-                      std::vector<HaarFeature64> &haarFeatures);
+                      std::vector<HaarFeature64> &haarFeatures)
+{
+    NCVStatus ncvStat;
+
+    haar.NumStages = 0;
+    haar.NumClassifierRootNodes = 0;
+    haar.NumClassifierTotalNodes = 0;
+    haar.NumFeatures = 0;
+    haar.ClassifierSize.width = 0;
+    haar.ClassifierSize.height = 0;
+    haar.bHasStumpsOnly = true;
+    haar.bNeedsTiltedII = false;
+    Ncv32u curMaxTreeDepth;
+
+    std::vector<HaarClassifierNode128> h_TmpClassifierNotRootNodes;
+    haarStages.resize(0);
+    haarClassifierNodes.resize(0);
+    haarFeatures.resize(0);
+
+    cv::Ptr<CvHaarClassifierCascade> oldCascade = (CvHaarClassifierCascade*)cvLoad(filename.c_str(), 0, 0, 0);
+    if (oldCascade.empty())
+    {
+        return NCV_HAAR_XML_LOADING_EXCEPTION;
+    }
+
+    haar.ClassifierSize.width = oldCascade->orig_window_size.width;
+    haar.ClassifierSize.height = oldCascade->orig_window_size.height;
+
+    int stagesCound = oldCascade->count;
+    for(int s = 0; s < stagesCound; ++s) // by stages
+    {
+        HaarStage64 curStage;
+        curStage.setStartClassifierRootNodeOffset(static_cast<Ncv32u>(haarClassifierNodes.size()));
+
+        curStage.setStageThreshold(oldCascade->stage_classifier[s].threshold);
+
+        int treesCount = oldCascade->stage_classifier[s].count;
+        for(int t = 0; t < treesCount; ++t) // by trees
+        {
+            Ncv32u nodeId = 0;
+            CvHaarClassifier* tree = &oldCascade->stage_classifier[s].classifier[t];
+
+            int nodesCount = tree->count;
+            for(int n = 0; n < nodesCount; ++n)  //by features
+            {
+                CvHaarFeature* feature = &tree->haar_feature[n];
+
+                HaarClassifierNode128 curNode;
+                curNode.setThreshold(tree->threshold[n]);
+
+                NcvBool bIsLeftNodeLeaf = false;
+                NcvBool bIsRightNodeLeaf = false;
+
+                HaarClassifierNodeDescriptor32 nodeLeft;
+                if ( tree->left[n] <= 0 )
+                {
+                    Ncv32f leftVal = tree->alpha[-tree->left[n]];
+                    ncvStat = nodeLeft.create(leftVal);
+                    ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat);
+                    bIsLeftNodeLeaf = true;
+                }
+                else
+                {
+                    Ncv32u leftNodeOffset = tree->left[n];
+                    nodeLeft.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + leftNodeOffset - 1));
+                    haar.bHasStumpsOnly = false;
+                }
+                curNode.setLeftNodeDesc(nodeLeft);
+
+                HaarClassifierNodeDescriptor32 nodeRight;
+                if ( tree->right[n] <= 0 )
+                {
+                    Ncv32f rightVal = tree->alpha[-tree->right[n]];
+                    ncvStat = nodeRight.create(rightVal);
+                    ncvAssertReturn(ncvStat == NCV_SUCCESS, ncvStat);
+                    bIsRightNodeLeaf = true;
+                }
+                else
+                {
+                    Ncv32u rightNodeOffset = tree->right[n];
+                    nodeRight.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + rightNodeOffset - 1));
+                    haar.bHasStumpsOnly = false;
+                }
+                curNode.setRightNodeDesc(nodeRight);
+
+                Ncv32u tiltedVal = feature->tilted;
+                haar.bNeedsTiltedII = (tiltedVal != 0);
+
+                Ncv32u featureId = 0;
+                for(int l = 0; l < CV_HAAR_FEATURE_MAX; ++l) //by rects
+                {
+                    Ncv32u rectX = feature->rect[l].r.x;
+                    Ncv32u rectY = feature->rect[l].r.y;
+                    Ncv32u rectWidth = feature->rect[l].r.width;
+                    Ncv32u rectHeight = feature->rect[l].r.height;
+
+                    Ncv32f rectWeight = feature->rect[l].weight;
+
+                    if (rectWeight == 0/* && rectX == 0 &&rectY == 0 && rectWidth == 0 && rectHeight == 0*/)
+                        break;
+
+                    HaarFeature64 curFeature;
+                    ncvStat = curFeature.setRect(rectX, rectY, rectWidth, rectHeight, haar.ClassifierSize.width, haar.ClassifierSize.height);
+                    curFeature.setWeight(rectWeight);
+                    ncvAssertReturn(NCV_SUCCESS == ncvStat, ncvStat);
+                    haarFeatures.push_back(curFeature);
+
+                    featureId++;
+                }
+
+                HaarFeatureDescriptor32 tmpFeatureDesc;
+                ncvStat = tmpFeatureDesc.create(haar.bNeedsTiltedII, bIsLeftNodeLeaf, bIsRightNodeLeaf,
+                    featureId, static_cast<Ncv32u>(haarFeatures.size()) - featureId);
+                ncvAssertReturn(NCV_SUCCESS == ncvStat, ncvStat);
+                curNode.setFeatureDesc(tmpFeatureDesc);
+
+                if (!nodeId)
+                {
+                    //root node
+                    haarClassifierNodes.push_back(curNode);
+                    curMaxTreeDepth = 1;
+                }
+                else
+                {
+                    //other node
+                    h_TmpClassifierNotRootNodes.push_back(curNode);
+                    curMaxTreeDepth++;
+                }
+
+                nodeId++;
+            }
+        }
+
+        curStage.setNumClassifierRootNodes(treesCount);
+        haarStages.push_back(curStage);
+    }
+
+    //fill in cascade stats
+    haar.NumStages = static_cast<Ncv32u>(haarStages.size());
+    haar.NumClassifierRootNodes = static_cast<Ncv32u>(haarClassifierNodes.size());
+    haar.NumClassifierTotalNodes = static_cast<Ncv32u>(haar.NumClassifierRootNodes + h_TmpClassifierNotRootNodes.size());
+    haar.NumFeatures = static_cast<Ncv32u>(haarFeatures.size());
+
+    //merge root and leaf nodes in one classifiers array
+    Ncv32u offsetRoot = static_cast<Ncv32u>(haarClassifierNodes.size());
+    for (Ncv32u i=0; i<haarClassifierNodes.size(); i++)
+    {
+        HaarFeatureDescriptor32 featureDesc = haarClassifierNodes[i].getFeatureDesc();
+
+        HaarClassifierNodeDescriptor32 nodeLeft = haarClassifierNodes[i].getLeftNodeDesc();
+        if (!featureDesc.isLeftNodeLeaf())
+        {
+            Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot;
+            nodeLeft.create(newOffset);
+        }
+        haarClassifierNodes[i].setLeftNodeDesc(nodeLeft);
+
+        HaarClassifierNodeDescriptor32 nodeRight = haarClassifierNodes[i].getRightNodeDesc();
+        if (!featureDesc.isRightNodeLeaf())
+        {
+            Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot;
+            nodeRight.create(newOffset);
+        }
+        haarClassifierNodes[i].setRightNodeDesc(nodeRight);
+    }
+
+    for (Ncv32u i=0; i<h_TmpClassifierNotRootNodes.size(); i++)
+    {
+        HaarFeatureDescriptor32 featureDesc = h_TmpClassifierNotRootNodes[i].getFeatureDesc();
+
+        HaarClassifierNodeDescriptor32 nodeLeft = h_TmpClassifierNotRootNodes[i].getLeftNodeDesc();
+        if (!featureDesc.isLeftNodeLeaf())
+        {
+            Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot;
+            nodeLeft.create(newOffset);
+        }
+        h_TmpClassifierNotRootNodes[i].setLeftNodeDesc(nodeLeft);
+
+        HaarClassifierNodeDescriptor32 nodeRight = h_TmpClassifierNotRootNodes[i].getRightNodeDesc();
+        if (!featureDesc.isRightNodeLeaf())
+        {
+            Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot;
+            nodeRight.create(newOffset);
+        }
+        h_TmpClassifierNotRootNodes[i].setRightNodeDesc(nodeRight);
+
+        haarClassifierNodes.push_back(h_TmpClassifierNotRootNodes[i]);
+    }
+
+    return NCV_SUCCESS;
+}
 
 
 #define NVBIN_HAAR_SIZERESERVED     16
@@ -2334,5 +2524,3 @@ NCVStatus ncvHaarStoreNVBIN_host(const cv::String &filename,
     fclose(fp);
     return NCV_SUCCESS;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/src/nvidia/core/NCVPixelOperations.hpp b/modules/gpunvidia/src/cuda/NCVPixelOperations.hpp
similarity index 99%
rename from modules/gpu/src/nvidia/core/NCVPixelOperations.hpp
rename to modules/gpunvidia/src/cuda/NCVPixelOperations.hpp
index ec2f16ebb7..2acdfb6829 100644
--- a/modules/gpu/src/nvidia/core/NCVPixelOperations.hpp
+++ b/modules/gpunvidia/src/cuda/NCVPixelOperations.hpp
@@ -45,7 +45,7 @@
 
 #include <limits.h>
 #include <float.h>
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
 
 template<typename TBase> inline __host__ __device__ TBase _pixMaxVal();
 template<> static inline __host__ __device__ Ncv8u  _pixMaxVal<Ncv8u>()  {return UCHAR_MAX;}
diff --git a/modules/gpu/src/nvidia/core/NCVPyramid.cu b/modules/gpunvidia/src/cuda/NCVPyramid.cu
similarity index 99%
rename from modules/gpu/src/nvidia/core/NCVPyramid.cu
rename to modules/gpunvidia/src/cuda/NCVPyramid.cu
index 380916ccee..6b76c644b6 100644
--- a/modules/gpu/src/nvidia/core/NCVPyramid.cu
+++ b/modules/gpunvidia/src/cuda/NCVPyramid.cu
@@ -40,15 +40,16 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <cuda_runtime.h>
 #include <stdio.h>
-#include "NCV.hpp"
+#include <cuda_runtime.h>
+
+#include "opencv2/core/cuda/common.hpp"
+
+#include "opencv2/gpunvidia/NCV.hpp"
+#include "opencv2/gpunvidia/NCVPyramid.hpp"
+
 #include "NCVAlg.hpp"
-#include "NCVPyramid.hpp"
 #include "NCVPixelOperations.hpp"
-#include "opencv2/core/cuda/common.hpp"
 
 template<typename T, Ncv32u CN> struct __average4_CN {static __host__ __device__ T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11);};
 
@@ -602,5 +603,3 @@ template class NCVImagePyramid<float3>;
 template class NCVImagePyramid<float4>;
 
 #endif //_WIN32
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp b/modules/gpunvidia/src/cuda/NCVRuntimeTemplates.hpp
similarity index 100%
rename from modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
rename to modules/gpunvidia/src/cuda/NCVRuntimeTemplates.hpp
diff --git a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu b/modules/gpunvidia/src/cuda/NPP_staging.cu
similarity index 99%
rename from modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu
rename to modules/gpunvidia/src/cuda/NPP_staging.cu
index 6d0f9c5c8d..31f7adc1d9 100644
--- a/modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu
+++ b/modules/gpunvidia/src/cuda/NPP_staging.cu
@@ -44,10 +44,12 @@
 
 #include <vector>
 #include <cuda_runtime.h>
-#include "NPP_staging.hpp"
+
 #include "opencv2/core/cuda/warp.hpp"
 #include "opencv2/core/cuda/warp_shuffle.hpp"
 
+#include "opencv2/gpunvidia/NPP_staging.hpp"
+
 
 texture<Ncv8u,  1, cudaReadModeElementType> tex8u;
 texture<Ncv32u, 1, cudaReadModeElementType> tex32u;
diff --git a/modules/gpunvidia/src/precomp.cpp b/modules/gpunvidia/src/precomp.cpp
new file mode 100644
index 0000000000..3c01a2596d
--- /dev/null
+++ b/modules/gpunvidia/src/precomp.cpp
@@ -0,0 +1,43 @@
+/*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"
diff --git a/modules/gpunvidia/src/precomp.hpp b/modules/gpunvidia/src/precomp.hpp
new file mode 100644
index 0000000000..613670009e
--- /dev/null
+++ b/modules/gpunvidia/src/precomp.hpp
@@ -0,0 +1,56 @@
+/*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*/
+
+#ifndef __OPENCV_PRECOMP_H__
+#define __OPENCV_PRECOMP_H__
+
+#include <limits>
+#include <iostream>
+#include <algorithm>
+
+#include "opencv2/gpunvidia.hpp"
+#include "opencv2/core/utility.hpp"
+#include "opencv2/objdetect.hpp"
+
+#include "opencv2/core/gpu_private.hpp"
+
+#endif /* __OPENCV_PRECOMP_H__ */
diff --git a/modules/gpu/test/nvidia/NCVAutoTestLister.hpp b/modules/gpunvidia/test/NCVAutoTestLister.hpp
similarity index 97%
rename from modules/gpu/test/nvidia/NCVAutoTestLister.hpp
rename to modules/gpunvidia/test/NCVAutoTestLister.hpp
index 6ac5bc0cf6..8730eeea7e 100644
--- a/modules/gpu/test/nvidia/NCVAutoTestLister.hpp
+++ b/modules/gpunvidia/test/NCVAutoTestLister.hpp
@@ -46,13 +46,7 @@
 #include <vector>
 
 #include "NCVTest.hpp"
-#include <main_test_nvidia.h>
-//enum OutputLevel
-//{
-//    OutputLevelNone,
-//    OutputLevelCompact,
-//    OutputLevelFull
-//};
+#include "main_test_nvidia.h"
 
 class NCVAutoTestLister
 {
diff --git a/modules/gpu/test/nvidia/NCVTest.hpp b/modules/gpunvidia/test/NCVTest.hpp
similarity index 99%
rename from modules/gpu/test/nvidia/NCVTest.hpp
rename to modules/gpunvidia/test/NCVTest.hpp
index 22958e565d..d08044db01 100644
--- a/modules/gpu/test/nvidia/NCVTest.hpp
+++ b/modules/gpunvidia/test/NCVTest.hpp
@@ -55,7 +55,8 @@
 #include <fstream>
 
 #include <cuda_runtime.h>
-#include "NPP_staging.hpp"
+
+#include "opencv2/gpunvidia.hpp"
 
 
 struct NCVTestReport
diff --git a/modules/gpu/test/nvidia/NCVTestSourceProvider.hpp b/modules/gpunvidia/test/NCVTestSourceProvider.hpp
similarity index 99%
rename from modules/gpu/test/nvidia/NCVTestSourceProvider.hpp
rename to modules/gpunvidia/test/NCVTestSourceProvider.hpp
index df245bb827..38b9d814c3 100644
--- a/modules/gpu/test/nvidia/NCVTestSourceProvider.hpp
+++ b/modules/gpunvidia/test/NCVTestSourceProvider.hpp
@@ -45,8 +45,8 @@
 
 #include <memory>
 
-#include "NCV.hpp"
-#include <opencv2/highgui.hpp>
+#include "opencv2/highgui.hpp"
+#include "opencv2/gpunvidia.hpp"
 
 
 template <class T>
diff --git a/modules/gpu/test/nvidia/TestCompact.cpp b/modules/gpunvidia/test/TestCompact.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestCompact.cpp
rename to modules/gpunvidia/test/TestCompact.cpp
index 9154101079..70640f37db 100644
--- a/modules/gpu/test/nvidia/TestCompact.cpp
+++ b/modules/gpunvidia/test/TestCompact.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestCompact.h"
-
+#include "test_precomp.hpp"
 
 TestCompact::TestCompact(std::string testName_, NCVTestSourceProvider<Ncv32u> &src_,
                                              Ncv32u length_, Ncv32u badElem_, Ncv32u badElemPercentage_)
@@ -160,5 +157,3 @@ bool TestCompact::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestCompact.h b/modules/gpunvidia/test/TestCompact.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestCompact.h
rename to modules/gpunvidia/test/TestCompact.h
diff --git a/modules/gpu/test/nvidia/TestDrawRects.cpp b/modules/gpunvidia/test/TestDrawRects.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestDrawRects.cpp
rename to modules/gpunvidia/test/TestDrawRects.cpp
index 3da458694a..40d8e21b49 100644
--- a/modules/gpu/test/nvidia/TestDrawRects.cpp
+++ b/modules/gpunvidia/test/TestDrawRects.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestDrawRects.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
 
 
 template <class T>
@@ -195,5 +192,3 @@ bool TestDrawRects<T>::deinit()
 
 template class TestDrawRects<Ncv8u>;
 template class TestDrawRects<Ncv32u>;
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestDrawRects.h b/modules/gpunvidia/test/TestDrawRects.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestDrawRects.h
rename to modules/gpunvidia/test/TestDrawRects.h
diff --git a/modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp b/modules/gpunvidia/test/TestHaarCascadeApplication.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp
rename to modules/gpunvidia/test/TestHaarCascadeApplication.cpp
index 01a9637b8d..121f31e434 100644
--- a/modules/gpu/test/nvidia/TestHaarCascadeApplication.cpp
+++ b/modules/gpunvidia/test/TestHaarCascadeApplication.cpp
@@ -40,13 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <float.h>
-
-#if defined(__GNUC__) && !defined(__APPLE__) && !defined(__arm__)
-    #include <fpu_control.h>
-#endif
+#include "test_precomp.hpp"
 
 namespace
 {
@@ -88,10 +82,6 @@ namespace
     }
 }
 
-#include "TestHaarCascadeApplication.h"
-#include "NCVHaarObjectDetection.hpp"
-
-
 TestHaarCascadeApplication::TestHaarCascadeApplication(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
                                                        std::string cascadeName_, Ncv32u width_, Ncv32u height_)
     :
@@ -343,5 +333,3 @@ bool TestHaarCascadeApplication::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestHaarCascadeApplication.h b/modules/gpunvidia/test/TestHaarCascadeApplication.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestHaarCascadeApplication.h
rename to modules/gpunvidia/test/TestHaarCascadeApplication.h
diff --git a/modules/gpu/test/nvidia/TestHaarCascadeLoader.cpp b/modules/gpunvidia/test/TestHaarCascadeLoader.cpp
similarity index 97%
rename from modules/gpu/test/nvidia/TestHaarCascadeLoader.cpp
rename to modules/gpunvidia/test/TestHaarCascadeLoader.cpp
index 42552295d2..b1e840a54c 100644
--- a/modules/gpu/test/nvidia/TestHaarCascadeLoader.cpp
+++ b/modules/gpunvidia/test/TestHaarCascadeLoader.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestHaarCascadeLoader.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
 
 
 TestHaarCascadeLoader::TestHaarCascadeLoader(std::string testName_, std::string cascadeName_)
@@ -154,5 +151,3 @@ bool TestHaarCascadeLoader::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestHaarCascadeLoader.h b/modules/gpunvidia/test/TestHaarCascadeLoader.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestHaarCascadeLoader.h
rename to modules/gpunvidia/test/TestHaarCascadeLoader.h
diff --git a/modules/gpu/test/nvidia/TestHypothesesFilter.cpp b/modules/gpunvidia/test/TestHypothesesFilter.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestHypothesesFilter.cpp
rename to modules/gpunvidia/test/TestHypothesesFilter.cpp
index f71a5f72d0..39d6556616 100644
--- a/modules/gpu/test/nvidia/TestHypothesesFilter.cpp
+++ b/modules/gpunvidia/test/TestHypothesesFilter.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestHypothesesFilter.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
 
 
 TestHypothesesFilter::TestHypothesesFilter(std::string testName_, NCVTestSourceProvider<Ncv32u> &src_,
@@ -207,5 +204,3 @@ bool TestHypothesesFilter::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestHypothesesFilter.h b/modules/gpunvidia/test/TestHypothesesFilter.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestHypothesesFilter.h
rename to modules/gpunvidia/test/TestHypothesesFilter.h
diff --git a/modules/gpu/test/nvidia/TestHypothesesGrow.cpp b/modules/gpunvidia/test/TestHypothesesGrow.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestHypothesesGrow.cpp
rename to modules/gpunvidia/test/TestHypothesesGrow.cpp
index 5cb81ddf49..e7fe4d939d 100644
--- a/modules/gpu/test/nvidia/TestHypothesesGrow.cpp
+++ b/modules/gpunvidia/test/TestHypothesesGrow.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestHypothesesGrow.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
 
 
 TestHypothesesGrow::TestHypothesesGrow(std::string testName_, NCVTestSourceProvider<Ncv32u> &src_,
@@ -165,5 +162,3 @@ bool TestHypothesesGrow::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestHypothesesGrow.h b/modules/gpunvidia/test/TestHypothesesGrow.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestHypothesesGrow.h
rename to modules/gpunvidia/test/TestHypothesesGrow.h
diff --git a/modules/gpu/test/nvidia/TestIntegralImage.cpp b/modules/gpunvidia/test/TestIntegralImage.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestIntegralImage.cpp
rename to modules/gpunvidia/test/TestIntegralImage.cpp
index a0820a8216..c04edff7c9 100644
--- a/modules/gpu/test/nvidia/TestIntegralImage.cpp
+++ b/modules/gpunvidia/test/TestIntegralImage.cpp
@@ -40,10 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-#include "TestIntegralImage.h"
+#include "test_precomp.hpp"
 
 
 template <class T_in, class T_out>
@@ -216,5 +213,3 @@ bool TestIntegralImage<T_in, T_out>::deinit()
 
 template class TestIntegralImage<Ncv8u, Ncv32u>;
 template class TestIntegralImage<Ncv32f, Ncv32f>;
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestIntegralImage.h b/modules/gpunvidia/test/TestIntegralImage.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestIntegralImage.h
rename to modules/gpunvidia/test/TestIntegralImage.h
diff --git a/modules/gpu/test/nvidia/TestIntegralImageSquared.cpp b/modules/gpunvidia/test/TestIntegralImageSquared.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestIntegralImageSquared.cpp
rename to modules/gpunvidia/test/TestIntegralImageSquared.cpp
index 4f4e45d406..5481fa2e3a 100644
--- a/modules/gpu/test/nvidia/TestIntegralImageSquared.cpp
+++ b/modules/gpunvidia/test/TestIntegralImageSquared.cpp
@@ -40,9 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include "TestIntegralImageSquared.h"
+#include "test_precomp.hpp"
 
 
 TestIntegralImageSquared::TestIntegralImageSquared(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
@@ -148,5 +146,3 @@ bool TestIntegralImageSquared::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestIntegralImageSquared.h b/modules/gpunvidia/test/TestIntegralImageSquared.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestIntegralImageSquared.h
rename to modules/gpunvidia/test/TestIntegralImageSquared.h
diff --git a/modules/gpu/test/nvidia/TestRectStdDev.cpp b/modules/gpunvidia/test/TestRectStdDev.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestRectStdDev.cpp
rename to modules/gpunvidia/test/TestRectStdDev.cpp
index c019b0ee3c..86bb9ed23b 100644
--- a/modules/gpu/test/nvidia/TestRectStdDev.cpp
+++ b/modules/gpunvidia/test/TestRectStdDev.cpp
@@ -40,11 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestRectStdDev.h"
+#include "test_precomp.hpp"
 
 
 TestRectStdDev::TestRectStdDev(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
@@ -211,5 +207,3 @@ bool TestRectStdDev::deinit()
 {
     return true;
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestRectStdDev.h b/modules/gpunvidia/test/TestRectStdDev.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestRectStdDev.h
rename to modules/gpunvidia/test/TestRectStdDev.h
diff --git a/modules/gpu/test/nvidia/TestResize.cpp b/modules/gpunvidia/test/TestResize.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestResize.cpp
rename to modules/gpunvidia/test/TestResize.cpp
index 83443c88b8..d2080f06de 100644
--- a/modules/gpu/test/nvidia/TestResize.cpp
+++ b/modules/gpunvidia/test/TestResize.cpp
@@ -40,11 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestResize.h"
+#include "test_precomp.hpp"
 
 
 template <class T>
@@ -192,5 +188,3 @@ bool TestResize<T>::deinit()
 
 template class TestResize<Ncv32u>;
 template class TestResize<Ncv64u>;
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestResize.h b/modules/gpunvidia/test/TestResize.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestResize.h
rename to modules/gpunvidia/test/TestResize.h
diff --git a/modules/gpu/test/nvidia/TestTranspose.cpp b/modules/gpunvidia/test/TestTranspose.cpp
similarity index 98%
rename from modules/gpu/test/nvidia/TestTranspose.cpp
rename to modules/gpunvidia/test/TestTranspose.cpp
index 5f71da4f8f..3322a0758b 100644
--- a/modules/gpu/test/nvidia/TestTranspose.cpp
+++ b/modules/gpunvidia/test/TestTranspose.cpp
@@ -40,11 +40,7 @@
 //
 //M*/
 
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestTranspose.h"
+#include "test_precomp.hpp"
 
 
 template <class T>
@@ -179,5 +175,3 @@ bool TestTranspose<T>::deinit()
 
 template class TestTranspose<Ncv32u>;
 template class TestTranspose<Ncv64u>;
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/nvidia/TestTranspose.h b/modules/gpunvidia/test/TestTranspose.h
similarity index 100%
rename from modules/gpu/test/nvidia/TestTranspose.h
rename to modules/gpunvidia/test/TestTranspose.h
diff --git a/modules/gpu/test/nvidia/main_nvidia.cpp b/modules/gpunvidia/test/main_nvidia.cpp
similarity index 96%
rename from modules/gpu/test/nvidia/main_nvidia.cpp
rename to modules/gpunvidia/test/main_nvidia.cpp
index 2e36fbec94..1179b5b96f 100644
--- a/modules/gpu/test/nvidia/main_nvidia.cpp
+++ b/modules/gpunvidia/test/main_nvidia.cpp
@@ -40,35 +40,12 @@
 //
 //M*/
 
+#include "test_precomp.hpp"
+
 #if defined _MSC_VER && _MSC_VER >= 1200
 # pragma warning (disable : 4408 4201 4100)
 #endif
 
-#if !defined CUDA_DISABLER
-
-#include <cstdio>
-
-#include "NCV.hpp"
-#include "NCVHaarObjectDetection.hpp"
-
-#include "TestIntegralImage.h"
-#include "TestIntegralImageSquared.h"
-#include "TestRectStdDev.h"
-#include "TestResize.h"
-#include "TestCompact.h"
-#include "TestTranspose.h"
-
-#include "TestDrawRects.h"
-#include "TestHypothesesGrow.h"
-#include "TestHypothesesFilter.h"
-#include "TestHaarCascadeLoader.h"
-#include "TestHaarCascadeApplication.h"
-
-#include "NCVAutoTestLister.hpp"
-#include "NCVTestSourceProvider.hpp"
-
-#include "main_test_nvidia.h"
-
 static std::string path;
 
 namespace {
@@ -480,5 +457,3 @@ bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel out
 
     return testListerVisualize.invoke();
 }
-
-#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/main_test_nvidia.h b/modules/gpunvidia/test/main_test_nvidia.h
similarity index 100%
rename from modules/gpu/test/main_test_nvidia.h
rename to modules/gpunvidia/test/main_test_nvidia.h
diff --git a/modules/gpu/test/main.cpp b/modules/gpunvidia/test/test_main.cpp
similarity index 100%
rename from modules/gpu/test/main.cpp
rename to modules/gpunvidia/test/test_main.cpp
diff --git a/modules/gpu/test/test_nvidia.cpp b/modules/gpunvidia/test/test_nvidia.cpp
similarity index 100%
rename from modules/gpu/test/test_nvidia.cpp
rename to modules/gpunvidia/test/test_nvidia.cpp
diff --git a/modules/gpunvidia/test/test_precomp.cpp b/modules/gpunvidia/test/test_precomp.cpp
new file mode 100644
index 0000000000..0fb6521809
--- /dev/null
+++ b/modules/gpunvidia/test/test_precomp.cpp
@@ -0,0 +1,43 @@
+/*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 "test_precomp.hpp"
diff --git a/modules/gpunvidia/test/test_precomp.hpp b/modules/gpunvidia/test/test_precomp.hpp
new file mode 100644
index 0000000000..46acfc2eca
--- /dev/null
+++ b/modules/gpunvidia/test/test_precomp.hpp
@@ -0,0 +1,95 @@
+/*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*/
+
+#ifdef __GNUC__
+#  pragma GCC diagnostic ignored "-Wmissing-declarations"
+#  if defined __clang__ || defined __APPLE__
+#    pragma GCC diagnostic ignored "-Wmissing-prototypes"
+#    pragma GCC diagnostic ignored "-Wextra"
+#  endif
+#endif
+
+#ifndef __OPENCV_TEST_PRECOMP_HPP__
+#define __OPENCV_TEST_PRECOMP_HPP__
+
+#if defined(__GNUC__) && !defined(__APPLE__) && !defined(__arm__)
+    #include <fpu_control.h>
+#endif
+
+#include <cfloat>
+#include <cstdio>
+#include <cmath>
+#include <vector>
+#include <string>
+#include <map>
+#include <memory>
+#include <algorithm>
+#include <fstream>
+
+#include "opencv2/ts.hpp"
+#include "opencv2/ts/gpu_test.hpp"
+
+#include "opencv2/core/gpumat.hpp"
+#include "opencv2/gpunvidia.hpp"
+#include "opencv2/highgui.hpp"
+
+#include "opencv2/core/gpu_private.hpp"
+
+#include "NCVTest.hpp"
+#include "NCVAutoTestLister.hpp"
+#include "NCVTestSourceProvider.hpp"
+
+#include "TestIntegralImage.h"
+#include "TestIntegralImageSquared.h"
+#include "TestRectStdDev.h"
+#include "TestResize.h"
+#include "TestCompact.h"
+#include "TestTranspose.h"
+#include "TestDrawRects.h"
+#include "TestHypothesesGrow.h"
+#include "TestHypothesesFilter.h"
+#include "TestHaarCascadeLoader.h"
+#include "TestHaarCascadeApplication.h"
+
+#include "main_test_nvidia.h"
+
+#endif
diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt
index 760bc26e6a..3b05553666 100644
--- a/samples/gpu/CMakeLists.txt
+++ b/samples/gpu/CMakeLists.txt
@@ -2,7 +2,7 @@ SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc ope
                                      opencv_ml opencv_video opencv_objdetect opencv_features2d
                                      opencv_calib3d opencv_legacy opencv_contrib opencv_gpu
                                      opencv_nonfree opencv_softcascade opencv_superres
-                                     opencv_gpucodec opencv_gpuarithm opencv_gpufilters)
+                                     opencv_gpucodec opencv_gpuarithm opencv_gpufilters opencv_gpunvidia)
 
 ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})
 
diff --git a/samples/gpu/cascadeclassifier_nvidia_api.cpp b/samples/gpu/cascadeclassifier_nvidia_api.cpp
index 694825424a..a4bc6a9736 100644
--- a/samples/gpu/cascadeclassifier_nvidia_api.cpp
+++ b/samples/gpu/cascadeclassifier_nvidia_api.cpp
@@ -11,7 +11,7 @@
 #include "opencv2/objdetect/objdetect_c.h"
 
 #ifdef HAVE_CUDA
-#include "NCVHaarObjectDetection.hpp"
+#include "opencv2/gpunvidia.hpp"
 #endif
 
 using namespace std;
diff --git a/samples/gpu/opticalflow_nvidia_api.cpp b/samples/gpu/opticalflow_nvidia_api.cpp
index 10663ce45a..e4fc93cd5d 100644
--- a/samples/gpu/opticalflow_nvidia_api.cpp
+++ b/samples/gpu/opticalflow_nvidia_api.cpp
@@ -16,8 +16,7 @@
 #include "opencv2/highgui/highgui_c.h"
 
 #ifdef HAVE_CUDA
-#include "NPP_staging/NPP_staging.hpp"
-#include "NCVBroxOpticalFlow.hpp"
+#include "opencv2/gpunvidia.hpp"
 #endif
 
 #if !defined(HAVE_CUDA)