From 91ba48c5f99158a51883c4884567c68f450060ab Mon Sep 17 00:00:00 2001 From: Marina Kolpakova Date: Mon, 25 Jun 2012 16:39:40 +0000 Subject: [PATCH] added lbp related CUDA files --- modules/gpu/include/opencv2/gpu/gpu.hpp | 2 +- modules/gpu/src/cascadeclassifier.cpp | 41 +++++++++++++++++- modules/gpu/src/cuda/lbp.cu | 54 ++++++++++++++++++++++++ modules/gpu/src/opencv2/gpu/device/lbp.hpp | 68 ++++++++++++++++++++++++++++++ 4 files changed, 163 insertions(+), 2 deletions(-) create mode 100644 modules/gpu/src/cuda/lbp.cu create mode 100644 modules/gpu/src/opencv2/gpu/device/lbp.hpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 7ac3268..ea9dcca 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1435,7 +1435,7 @@ public: bool load(const std::string& filename); void release(); - int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, double scaleFactor = 1.2, int minNeighbors = 4, Size minSize = Size()); + int detectMultiScale(const GpuMat& image, GpuMat& scaledImageBuffer, GpuMat& objectsBuf, double scaleFactor = 1.2, int minNeighbors = 4/*, Size minSize = Size()*/); bool findLargestObject; bool visualizeInPlace; diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index b284818..5e15353 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -259,8 +259,47 @@ Size cv::gpu::CascadeClassifier_GPU_LBP::getClassifierSize() const return NxM; } -int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, double scaleFactor=1.2, int minNeighbors=4, Size minSize=Size()) +namespace cv { namespace gpu { namespace device { + namespace lbp + { + void CascadeClassify(DevMem2Db image, DevMem2Db objects, double scaleFactor = 1.2, int minNeighbors = 4, cudaStream_t stream = 0); + } +}}} + +int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, GpuMat& scaledImageBuffer, GpuMat& objects, double scaleFactor, int minNeighbors /*, Size minSize=Size()*/) +{ + CV_Assert( scaleFactor > 1 && image.depth() == CV_8U ); + CV_Assert(empty()); + + const int defaultObjSearchNum = 100; + + if( !objects.empty() && objects.depth() == CV_32S) + objects.reshape(4, 1); + else + objects.create(1 , defaultObjSearchNum, CV_32SC4); + + scaledImageBuffer.create(image.size(), image.type()); + + // TODO: specify max objects size + for( double factor = 1; ; factor *= scaleFactor ) + { + cv::Size windowSize(cvRound(NxM.width * factor), cvRound(NxM.height * factor)); + cv::Size scaledImageSize(cvRound( image.cols / factor ), cvRound( image.rows / factor )); + cv::Size processingRectSize( scaledImageSize.width - NxM.width + 1, scaledImageSize.height - NxM.height + 1 ); + + // nothing to do + if (processingRectSize.width <= 0 || processingRectSize.height <= 0 ) + break; + // TODO: min max object sizes cheching + cv::gpu::resize(image, scaledImageBuffer, scaledImageSize, 0, 0, INTER_NEAREST); + + int yStep = (factor > 2.) + 1; + int stripCount = 1, stripSize = processingRectSize.height; + + } + // TODO: reject levels + return 0; } diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu new file mode 100644 index 0000000..929077c --- /dev/null +++ b/modules/gpu/src/cuda/lbp.cu @@ -0,0 +1,54 @@ +/*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 bpied warranties, including, but not limited to, the bpied +// 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 + +namespace cv { namespace gpu { namespace device +{ + namespace lbp + { + void CascadeClassify(DevMem2Db image, DevMem2Db objects, double scaleFactor=1.2, int minNeighbors=4, cudaStream_t stream) + { + printf("CascadeClassify"); + } + } +}}} \ No newline at end of file diff --git a/modules/gpu/src/opencv2/gpu/device/lbp.hpp b/modules/gpu/src/opencv2/gpu/device/lbp.hpp new file mode 100644 index 0000000..3b104f6 --- /dev/null +++ b/modules/gpu/src/opencv2/gpu/device/lbp.hpp @@ -0,0 +1,68 @@ +/*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 bpied warranties, including, but not limited to, the bpied +// 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_GPU_DEVICE_LBP_HPP_ +#define __OPENCV_GPU_DEVICE_LBP_HPP_ + +// #define CALC_SUM_(p0, p1, p2, p3, offset) \ +// ((p0)[offset] - (p1)[offset] - (p2)[offset] + (p3)[offset]) + +// __device__ __forceinline__ int sum(p0, p1, p2, p3, offset) +// { + +// } + +namespace cv { namespace gpu { namespace device { + + struct Feature + { + __device__ __forceinline__ Feature(const Feature& other) {(void)other;} + __device__ __forceinline__ Feature() {} + __device__ __forceinline__ char operator() (volatile int* ptr, int offset) + { + return char(0); + } + + } +}// namespaces + +#endif \ No newline at end of file -- 2.7.4