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")
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)
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})
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})
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)
}
#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 */
#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
{
#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__)
#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__ */
--- /dev/null
+/*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")
#include "opencv2/gpu.hpp"
#include "interpolation.hpp"
-#include "main_test_nvidia.h"
#include "opencv2/core/gpu_private.hpp"
--- /dev/null
+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)
--- /dev/null
+/*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__ */
#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
#ifndef _ncvhaarobjectdetection_hpp_
#define _ncvhaarobjectdetection_hpp_
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
//==============================================================================
#include <memory>
#include <vector>
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
#if 0 //def _WIN32
#ifndef _npp_staging_hpp_
#define _npp_staging_hpp_
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
/**
//
//M*/
-#include <iostream>
-#include <vector>
-#include "NCV.hpp"
-
+#include "precomp.hpp"
//==============================================================================
//
debugOutputHandler = func;
}
-#if !defined CUDA_DISABLER
-
//==============================================================================
//
//===================================================================
-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;
}
-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");
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;
}
//===================================================================
-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", );
}
//
//===================================================================
+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,
{
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 */
--- /dev/null
+/*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);
+}
#ifndef _ncv_alg_hpp_
#define _ncv_alg_hpp_
-#include "NCV.hpp"
+#include "opencv2/gpunvidia/NCV.hpp"
template <class T>
//
////////////////////////////////////////////////////////////////////////////////
-#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;
return NCV_SUCCESS;
}
-
-#endif /* CUDA_DISABLER */
//
////////////////////////////////////////////////////////////////////////////////
-#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"
//==============================================================================
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
fclose(fp);
return NCV_SUCCESS;
}
-
-#endif /* CUDA_DISABLER */
#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;}
//
//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);};
template class NCVImagePyramid<float4>;
#endif //_WIN32
-
-#endif /* CUDA_DISABLER */
#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;
--- /dev/null
+/*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"
--- /dev/null
+/*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__ */
#include <vector>
#include "NCVTest.hpp"
-#include <main_test_nvidia.h>
-//enum OutputLevel
-//{
-// OutputLevelNone,
-// OutputLevelCompact,
-// OutputLevelFull
-//};
+#include "main_test_nvidia.h"
class NCVAutoTestLister
{
#include <fstream>
#include <cuda_runtime.h>
-#include "NPP_staging.hpp"
+
+#include "opencv2/gpunvidia.hpp"
struct NCVTestReport
#include <memory>
-#include "NCV.hpp"
-#include <opencv2/highgui.hpp>
+#include "opencv2/highgui.hpp"
+#include "opencv2/gpunvidia.hpp"
template <class T>
//
//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_)
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include "TestDrawRects.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
template <class T>
template class TestDrawRects<Ncv8u>;
template class TestDrawRects<Ncv32u>;
-
-#endif /* CUDA_DISABLER */
//
//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
{
}
}
-#include "TestHaarCascadeApplication.h"
-#include "NCVHaarObjectDetection.hpp"
-
-
TestHaarCascadeApplication::TestHaarCascadeApplication(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
std::string cascadeName_, Ncv32u width_, Ncv32u height_)
:
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include "TestHaarCascadeLoader.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
TestHaarCascadeLoader::TestHaarCascadeLoader(std::string testName_, std::string cascadeName_)
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include "TestHypothesesFilter.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
TestHypothesesFilter::TestHypothesesFilter(std::string testName_, NCVTestSourceProvider<Ncv32u> &src_,
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include "TestHypothesesGrow.h"
-#include "NCVHaarObjectDetection.hpp"
+#include "test_precomp.hpp"
TestHypothesesGrow::TestHypothesesGrow(std::string testName_, NCVTestSourceProvider<Ncv32u> &src_,
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-#include "TestIntegralImage.h"
+#include "test_precomp.hpp"
template <class T_in, class T_out>
template class TestIntegralImage<Ncv8u, Ncv32u>;
template class TestIntegralImage<Ncv32f, Ncv32f>;
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include "TestIntegralImageSquared.h"
+#include "test_precomp.hpp"
TestIntegralImageSquared::TestIntegralImageSquared(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestRectStdDev.h"
+#include "test_precomp.hpp"
TestRectStdDev::TestRectStdDev(std::string testName_, NCVTestSourceProvider<Ncv8u> &src_,
{
return true;
}
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestResize.h"
+#include "test_precomp.hpp"
template <class T>
template class TestResize<Ncv32u>;
template class TestResize<Ncv64u>;
-
-#endif /* CUDA_DISABLER */
//
//M*/
-#if !defined CUDA_DISABLER
-
-#include <math.h>
-
-#include "TestTranspose.h"
+#include "test_precomp.hpp"
template <class T>
template class TestTranspose<Ncv32u>;
template class TestTranspose<Ncv64u>;
-
-#endif /* CUDA_DISABLER */
//
//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 {
return testListerVisualize.invoke();
}
-
-#endif /* CUDA_DISABLER */
--- /dev/null
+/*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"
--- /dev/null
+/*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
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})
#include "opencv2/objdetect/objdetect_c.h"
#ifdef HAVE_CUDA
-#include "NCVHaarObjectDetection.hpp"
+#include "opencv2/gpunvidia.hpp"
#endif
using namespace std;
#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)
projects / platform / upstream / opencv.git / commitdiff