added lbp related CUDA files

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;

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;
 }
 

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 <opencv2/gpu/device/detail/lbp.hpp>
+
+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");
+        }
+    }
+}}}

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