← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/cudaobjdetect/src/cascadeclassifier.cpp

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/cudaobjdetect/src/cascadeclassifier.cpp
 
1
/*M///////////////////////////////////////////////////////////////////////////////////////
 
2
//
 
3
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
 
4
//
 
5
//  By downloading, copying, installing or using the software you agree to this license.
 
6
//  If you do not agree to this license, do not download, install,
 
7
//  copy or use the software.
 
8
//
 
9
//
 
10
//                           License Agreement
 
11
//                For Open Source Computer Vision Library
 
12
//
 
13
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
 
14
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
 
15
// Third party copyrights are property of their respective owners.
 
16
//
 
17
// Redistribution and use in source and binary forms, with or without modification,
 
18
// are permitted provided that the following conditions are met:
 
19
//
 
20
//   * Redistribution's of source code must retain the above copyright notice,
 
21
//     this list of conditions and the following disclaimer.
 
22
//
 
23
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
24
//     this list of conditions and the following disclaimer in the documentation
 
25
//     and/or other materials provided with the distribution.
 
26
//
 
27
//   * The name of the copyright holders may not be used to endorse or promote products
 
28
//     derived from this software without specific prior written permission.
 
29
//
 
30
// This software is provided by the copyright holders and contributors "as is" and
 
31
// any express or implied warranties, including, but not limited to, the implied
 
32
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
33
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
34
// indirect, incidental, special, exemplary, or consequential damages
 
35
// (including, but not limited to, procurement of substitute goods or services;
 
36
// loss of use, data, or profits; or business interruption) however caused
 
37
// and on any theory of liability, whether in contract, strict liability,
 
38
// or tort (including negligence or otherwise) arising in any way out of
 
39
// the use of this software, even if advised of the possibility of such damage.
 
40
//
 
41
//M*/
 
42
 
 
43
#include "precomp.hpp"
 
44
#include "opencv2/objdetect/objdetect_c.h"
 
45
 
 
46
using namespace cv;
 
47
using namespace cv::cuda;
 
48
 
 
49
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
 
50
 
 
51
Ptr<cuda::CascadeClassifier> cv::cuda::CascadeClassifier::create(const String&) { throw_no_cuda(); return Ptr<cuda::CascadeClassifier>(); }
 
52
Ptr<cuda::CascadeClassifier> cv::cuda::CascadeClassifier::create(const FileStorage&) { throw_no_cuda(); return Ptr<cuda::CascadeClassifier>(); }
 
53
 
 
54
#else
 
55
 
 
56
//
 
57
// CascadeClassifierBase
 
58
//
 
59
 
 
60
namespace
 
61
{
 
62
    class CascadeClassifierBase : public cuda::CascadeClassifier
 
63
    {
 
64
    public:
 
65
        CascadeClassifierBase();
 
66
 
 
67
        virtual void setMaxObjectSize(Size maxObjectSize) { maxObjectSize_ = maxObjectSize; }
 
68
        virtual Size getMaxObjectSize() const { return maxObjectSize_; }
 
69
 
 
70
        virtual void setMinObjectSize(Size minSize) { minObjectSize_ = minSize; }
 
71
        virtual Size getMinObjectSize() const { return minObjectSize_; }
 
72
 
 
73
        virtual void setScaleFactor(double scaleFactor) { scaleFactor_ = scaleFactor; }
 
74
        virtual double getScaleFactor() const { return scaleFactor_; }
 
75
 
 
76
        virtual void setMinNeighbors(int minNeighbors) { minNeighbors_ = minNeighbors; }
 
77
        virtual int getMinNeighbors() const { return minNeighbors_; }
 
78
 
 
79
        virtual void setFindLargestObject(bool findLargestObject) { findLargestObject_ = findLargestObject; }
 
80
        virtual bool getFindLargestObject() { return findLargestObject_; }
 
81
 
 
82
        virtual void setMaxNumObjects(int maxNumObjects) { maxNumObjects_ = maxNumObjects; }
 
83
        virtual int getMaxNumObjects() const { return maxNumObjects_; }
 
84
 
 
85
    protected:
 
86
        Size maxObjectSize_;
 
87
        Size minObjectSize_;
 
88
        double scaleFactor_;
 
89
        int minNeighbors_;
 
90
        bool findLargestObject_;
 
91
        int maxNumObjects_;
 
92
    };
 
93
 
 
94
    CascadeClassifierBase::CascadeClassifierBase() :
 
95
        maxObjectSize_(),
 
96
        minObjectSize_(),
 
97
        scaleFactor_(1.2),
 
98
        minNeighbors_(4),
 
99
        findLargestObject_(false),
 
100
        maxNumObjects_(100)
 
101
    {
 
102
    }
 
103
}
 
104
 
 
105
//
 
106
// HaarCascade
 
107
//
 
108
 
 
109
#ifdef HAVE_OPENCV_CUDALEGACY
 
110
 
 
111
namespace
 
112
{
 
113
    class HaarCascade_Impl : public CascadeClassifierBase
 
114
    {
 
115
    public:
 
116
        explicit HaarCascade_Impl(const String& filename);
 
117
 
 
118
        virtual Size getClassifierSize() const;
 
119
 
 
120
        virtual void detectMultiScale(InputArray image,
 
121
                                      OutputArray objects,
 
122
                                      Stream& stream);
 
123
 
 
124
        virtual void convert(OutputArray gpu_objects,
 
125
                             std::vector<Rect>& objects);
 
126
 
 
127
    private:
 
128
        NCVStatus load(const String& classifierFile);
 
129
        NCVStatus calculateMemReqsAndAllocate(const Size& frameSize);
 
130
        NCVStatus process(const GpuMat& src, GpuMat& objects, cv::Size ncvMinSize, /*out*/ unsigned int& numDetections);
 
131
 
 
132
        Size lastAllocatedFrameSize;
 
133
 
 
134
        Ptr<NCVMemStackAllocator> gpuAllocator;
 
135
        Ptr<NCVMemStackAllocator> cpuAllocator;
 
136
 
 
137
        cudaDeviceProp devProp;
 
138
        NCVStatus ncvStat;
 
139
 
 
140
        Ptr<NCVMemNativeAllocator> gpuCascadeAllocator;
 
141
        Ptr<NCVMemNativeAllocator> cpuCascadeAllocator;
 
142
 
 
143
        Ptr<NCVVectorAlloc<HaarStage64> >           h_haarStages;
 
144
        Ptr<NCVVectorAlloc<HaarClassifierNode128> > h_haarNodes;
 
145
        Ptr<NCVVectorAlloc<HaarFeature64> >         h_haarFeatures;
 
146
 
 
147
        HaarClassifierCascadeDescriptor haar;
 
148
 
 
149
        Ptr<NCVVectorAlloc<HaarStage64> >           d_haarStages;
 
150
        Ptr<NCVVectorAlloc<HaarClassifierNode128> > d_haarNodes;
 
151
        Ptr<NCVVectorAlloc<HaarFeature64> >         d_haarFeatures;
 
152
    };
 
153
 
 
154
    static void NCVDebugOutputHandler(const String &msg)
 
155
    {
 
156
        CV_Error(Error::GpuApiCallError, msg.c_str());
 
157
    }
 
158
 
 
159
    HaarCascade_Impl::HaarCascade_Impl(const String& filename) :
 
160
        lastAllocatedFrameSize(-1, -1)
 
161
    {
 
162
        ncvSetDebugOutputHandler(NCVDebugOutputHandler);
 
163
        ncvSafeCall( load(filename) );
 
164
    }
 
165
 
 
166
    Size HaarCascade_Impl::getClassifierSize() const
 
167
    {
 
168
        return Size(haar.ClassifierSize.width, haar.ClassifierSize.height);
 
169
    }
 
170
 
 
171
    void HaarCascade_Impl::detectMultiScale(InputArray _image,
 
172
                                            OutputArray _objects,
 
173
                                            Stream& stream)
 
174
    {
 
175
        const GpuMat image = _image.getGpuMat();
 
176
 
 
177
        CV_Assert( image.depth() == CV_8U);
 
178
        CV_Assert( scaleFactor_ > 1 );
 
179
        CV_Assert( !stream );
 
180
 
 
181
        Size ncvMinSize = getClassifierSize();
 
182
        if (ncvMinSize.width < minObjectSize_.width && ncvMinSize.height < minObjectSize_.height)
 
183
        {
 
184
            ncvMinSize.width = minObjectSize_.width;
 
185
            ncvMinSize.height = minObjectSize_.height;
 
186
        }
 
187
 
 
188
        BufferPool pool(stream);
 
189
        GpuMat objectsBuf = pool.getBuffer(1, maxNumObjects_, DataType<Rect>::type);
 
190
 
 
191
        unsigned int numDetections;
 
192
        ncvSafeCall( process(image, objectsBuf, ncvMinSize, numDetections) );
 
193
 
 
194
        if (numDetections > 0)
 
195
        {
 
196
            objectsBuf.colRange(0, numDetections).copyTo(_objects);
 
197
        }
 
198
        else
 
199
        {
 
200
            _objects.release();
 
201
        }
 
202
    }
 
203
 
 
204
    void HaarCascade_Impl::convert(OutputArray _gpu_objects, std::vector<Rect>& objects)
 
205
    {
 
206
        if (_gpu_objects.empty())
 
207
        {
 
208
            objects.clear();
 
209
            return;
 
210
        }
 
211
 
 
212
        Mat gpu_objects;
 
213
        if (_gpu_objects.kind() == _InputArray::CUDA_GPU_MAT)
 
214
        {
 
215
            _gpu_objects.getGpuMat().download(gpu_objects);
 
216
        }
 
217
        else
 
218
        {
 
219
            gpu_objects = _gpu_objects.getMat();
 
220
        }
 
221
 
 
222
        CV_Assert( gpu_objects.rows == 1 );
 
223
        CV_Assert( gpu_objects.type() == DataType<Rect>::type );
 
224
 
 
225
        Rect* ptr = gpu_objects.ptr<Rect>();
 
226
        objects.assign(ptr, ptr + gpu_objects.cols);
 
227
    }
 
228
 
 
229
    NCVStatus HaarCascade_Impl::load(const String& classifierFile)
 
230
    {
 
231
        int devId = cv::cuda::getDevice();
 
232
        ncvAssertCUDAReturn(cudaGetDeviceProperties(&devProp, devId), NCV_CUDA_ERROR);
 
233
 
 
234
        // Load the classifier from file (assuming its size is about 1 mb) using a simple allocator
 
235
        gpuCascadeAllocator = makePtr<NCVMemNativeAllocator>(NCVMemoryTypeDevice, static_cast<int>(devProp.textureAlignment));
 
236
        cpuCascadeAllocator = makePtr<NCVMemNativeAllocator>(NCVMemoryTypeHostPinned, static_cast<int>(devProp.textureAlignment));
 
237
 
 
238
        ncvAssertPrintReturn(gpuCascadeAllocator->isInitialized(), "Error creating cascade GPU allocator", NCV_CUDA_ERROR);
 
239
        ncvAssertPrintReturn(cpuCascadeAllocator->isInitialized(), "Error creating cascade CPU allocator", NCV_CUDA_ERROR);
 
240
 
 
241
        Ncv32u haarNumStages, haarNumNodes, haarNumFeatures;
 
242
        ncvStat = ncvHaarGetClassifierSize(classifierFile, haarNumStages, haarNumNodes, haarNumFeatures);
 
243
        ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "Error reading classifier size (check the file)", NCV_FILE_ERROR);
 
244
 
 
245
        h_haarStages.reset  (new NCVVectorAlloc<HaarStage64>(*cpuCascadeAllocator, haarNumStages));
 
246
        h_haarNodes.reset   (new NCVVectorAlloc<HaarClassifierNode128>(*cpuCascadeAllocator, haarNumNodes));
 
247
        h_haarFeatures.reset(new NCVVectorAlloc<HaarFeature64>(*cpuCascadeAllocator, haarNumFeatures));
 
248
 
 
249
        ncvAssertPrintReturn(h_haarStages->isMemAllocated(), "Error in cascade CPU allocator", NCV_CUDA_ERROR);
 
250
        ncvAssertPrintReturn(h_haarNodes->isMemAllocated(), "Error in cascade CPU allocator", NCV_CUDA_ERROR);
 
251
        ncvAssertPrintReturn(h_haarFeatures->isMemAllocated(), "Error in cascade CPU allocator", NCV_CUDA_ERROR);
 
252
 
 
253
        ncvStat = ncvHaarLoadFromFile_host(classifierFile, haar, *h_haarStages, *h_haarNodes, *h_haarFeatures);
 
254
        ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "Error loading classifier", NCV_FILE_ERROR);
 
255
 
 
256
        d_haarStages.reset  (new NCVVectorAlloc<HaarStage64>(*gpuCascadeAllocator, haarNumStages));
 
257
        d_haarNodes.reset   (new NCVVectorAlloc<HaarClassifierNode128>(*gpuCascadeAllocator, haarNumNodes));
 
258
        d_haarFeatures.reset(new NCVVectorAlloc<HaarFeature64>(*gpuCascadeAllocator, haarNumFeatures));
 
259
 
 
260
        ncvAssertPrintReturn(d_haarStages->isMemAllocated(), "Error in cascade GPU allocator", NCV_CUDA_ERROR);
 
261
        ncvAssertPrintReturn(d_haarNodes->isMemAllocated(), "Error in cascade GPU allocator", NCV_CUDA_ERROR);
 
262
        ncvAssertPrintReturn(d_haarFeatures->isMemAllocated(), "Error in cascade GPU allocator", NCV_CUDA_ERROR);
 
263
 
 
264
        ncvStat = h_haarStages->copySolid(*d_haarStages, 0);
 
265
        ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "Error copying cascade to GPU", NCV_CUDA_ERROR);
 
266
        ncvStat = h_haarNodes->copySolid(*d_haarNodes, 0);
 
267
        ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "Error copying cascade to GPU", NCV_CUDA_ERROR);
 
268
        ncvStat = h_haarFeatures->copySolid(*d_haarFeatures, 0);
 
269
        ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "Error copying cascade to GPU", NCV_CUDA_ERROR);
 
270
 
 
271
        return NCV_SUCCESS;
 
272
    }
 
273
 
 
274
    NCVStatus HaarCascade_Impl::calculateMemReqsAndAllocate(const Size& frameSize)
 
275
    {
 
276
        if (lastAllocatedFrameSize == frameSize)
 
277
        {
 
278
            return NCV_SUCCESS;
 
279
        }
 
280
 
 
281
        // Calculate memory requirements and create real allocators
 
282
        NCVMemStackAllocator gpuCounter(static_cast<int>(devProp.textureAlignment));
 
283
        NCVMemStackAllocator cpuCounter(static_cast<int>(devProp.textureAlignment));
 
284
 
 
285
        ncvAssertPrintReturn(gpuCounter.isInitialized(), "Error creating GPU memory counter", NCV_CUDA_ERROR);
 
286
        ncvAssertPrintReturn(cpuCounter.isInitialized(), "Error creating CPU memory counter", NCV_CUDA_ERROR);
 
287
 
 
288
        NCVMatrixAlloc<Ncv8u> d_src(gpuCounter, frameSize.width, frameSize.height);
 
289
        NCVMatrixAlloc<Ncv8u> h_src(cpuCounter, frameSize.width, frameSize.height);
 
290
 
 
291
        ncvAssertReturn(d_src.isMemAllocated(), NCV_ALLOCATOR_BAD_ALLOC);
 
292
        ncvAssertReturn(h_src.isMemAllocated(), NCV_ALLOCATOR_BAD_ALLOC);
 
293
 
 
294
        NCVVectorAlloc<NcvRect32u> d_rects(gpuCounter, 100);
 
295
        ncvAssertReturn(d_rects.isMemAllocated(), NCV_ALLOCATOR_BAD_ALLOC);
 
296
 
 
297
        NcvSize32u roi;
 
298
        roi.width = d_src.width();
 
299
        roi.height = d_src.height();
 
300
        Ncv32u numDetections;
 
301
        ncvStat = ncvDetectObjectsMultiScale_device(d_src, roi, d_rects, numDetections, haar, *h_haarStages,
 
302
            *d_haarStages, *d_haarNodes, *d_haarFeatures, haar.ClassifierSize, 4, 1.2f, 1, 0, gpuCounter, cpuCounter, devProp, 0);
 
303
 
 
304
        ncvAssertReturnNcvStat(ncvStat);
 
305
        ncvAssertCUDAReturn(cudaStreamSynchronize(0), NCV_CUDA_ERROR);
 
306
 
 
307
        gpuAllocator = makePtr<NCVMemStackAllocator>(NCVMemoryTypeDevice, gpuCounter.maxSize(), static_cast<int>(devProp.textureAlignment));
 
308
        cpuAllocator = makePtr<NCVMemStackAllocator>(NCVMemoryTypeHostPinned, cpuCounter.maxSize(), static_cast<int>(devProp.textureAlignment));
 
309
 
 
310
        ncvAssertPrintReturn(gpuAllocator->isInitialized(), "Error creating GPU memory allocator", NCV_CUDA_ERROR);
 
311
        ncvAssertPrintReturn(cpuAllocator->isInitialized(), "Error creating CPU memory allocator", NCV_CUDA_ERROR);
 
312
 
 
313
        lastAllocatedFrameSize = frameSize;
 
314
        return NCV_SUCCESS;
 
315
    }
 
316
 
 
317
    NCVStatus HaarCascade_Impl::process(const GpuMat& src, GpuMat& objects, cv::Size ncvMinSize, /*out*/ unsigned int& numDetections)
 
318
    {
 
319
        calculateMemReqsAndAllocate(src.size());
 
320
 
 
321
        NCVMemPtr src_beg;
 
322
        src_beg.ptr = (void*)src.ptr<Ncv8u>();
 
323
        src_beg.memtype = NCVMemoryTypeDevice;
 
324
 
 
325
        NCVMemSegment src_seg;
 
326
        src_seg.begin = src_beg;
 
327
        src_seg.size  = src.step * src.rows;
 
328
 
 
329
        NCVMatrixReuse<Ncv8u> d_src(src_seg, static_cast<int>(devProp.textureAlignment), src.cols, src.rows, static_cast<int>(src.step), true);
 
330
        ncvAssertReturn(d_src.isMemReused(), NCV_ALLOCATOR_BAD_REUSE);
 
331
 
 
332
        CV_Assert(objects.rows == 1);
 
333
 
 
334
        NCVMemPtr objects_beg;
 
335
        objects_beg.ptr = (void*)objects.ptr<NcvRect32u>();
 
336
        objects_beg.memtype = NCVMemoryTypeDevice;
 
337
 
 
338
        NCVMemSegment objects_seg;
 
339
        objects_seg.begin = objects_beg;
 
340
        objects_seg.size = objects.step * objects.rows;
 
341
        NCVVectorReuse<NcvRect32u> d_rects(objects_seg, objects.cols);
 
342
        ncvAssertReturn(d_rects.isMemReused(), NCV_ALLOCATOR_BAD_REUSE);
 
343
 
 
344
        NcvSize32u roi;
 
345
        roi.width = d_src.width();
 
346
        roi.height = d_src.height();
 
347
 
 
348
        NcvSize32u winMinSize(ncvMinSize.width, ncvMinSize.height);
 
349
 
 
350
        Ncv32u flags = 0;
 
351
        flags |= findLargestObject_ ? NCVPipeObjDet_FindLargestObject : 0;
 
352
 
 
353
        ncvStat = ncvDetectObjectsMultiScale_device(
 
354
            d_src, roi, d_rects, numDetections, haar, *h_haarStages,
 
355
            *d_haarStages, *d_haarNodes, *d_haarFeatures,
 
356
            winMinSize,
 
357
            minNeighbors_,
 
358
            scaleFactor_, 1,
 
359
            flags,
 
360
            *gpuAllocator, *cpuAllocator, devProp, 0);
 
361
        ncvAssertReturnNcvStat(ncvStat);
 
362
        ncvAssertCUDAReturn(cudaStreamSynchronize(0), NCV_CUDA_ERROR);
 
363
 
 
364
        return NCV_SUCCESS;
 
365
    }
 
366
}
 
367
 
 
368
#endif
 
369
 
 
370
//
 
371
// LbpCascade
 
372
//
 
373
 
 
374
namespace cv { namespace cuda { namespace device
 
375
{
 
376
    namespace lbp
 
377
    {
 
378
        void classifyPyramid(int frameW,
 
379
                             int frameH,
 
380
                             int windowW,
 
381
                             int windowH,
 
382
                             float initalScale,
 
383
                             float factor,
 
384
                             int total,
 
385
                             const PtrStepSzb& mstages,
 
386
                             const int nstages,
 
387
                             const PtrStepSzi& mnodes,
 
388
                             const PtrStepSzf& mleaves,
 
389
                             const PtrStepSzi& msubsets,
 
390
                             const PtrStepSzb& mfeatures,
 
391
                             const int subsetSize,
 
392
                             PtrStepSz<int4> objects,
 
393
                             unsigned int* classified,
 
394
                             PtrStepSzi integral);
 
395
 
 
396
        void connectedConmonents(PtrStepSz<int4> candidates,
 
397
                                 int ncandidates,
 
398
                                 PtrStepSz<int4> objects,
 
399
                                 int groupThreshold,
 
400
                                 float grouping_eps,
 
401
                                 unsigned int* nclasses);
 
402
    }
 
403
}}}
 
404
 
 
405
namespace
 
406
{
 
407
    cv::Size operator -(const cv::Size& a, const cv::Size& b)
 
408
    {
 
409
        return cv::Size(a.width - b.width, a.height - b.height);
 
410
    }
 
411
 
 
412
    cv::Size operator +(const cv::Size& a, const int& i)
 
413
    {
 
414
        return cv::Size(a.width + i, a.height + i);
 
415
    }
 
416
 
 
417
    cv::Size operator *(const cv::Size& a, const float& f)
 
418
    {
 
419
        return cv::Size(cvRound(a.width * f), cvRound(a.height * f));
 
420
    }
 
421
 
 
422
    cv::Size operator /(const cv::Size& a, const float& f)
 
423
    {
 
424
        return cv::Size(cvRound(a.width / f), cvRound(a.height / f));
 
425
    }
 
426
 
 
427
    bool operator <=(const cv::Size& a, const cv::Size& b)
 
428
    {
 
429
        return a.width <= b.width && a.height <= b.width;
 
430
    }
 
431
 
 
432
    struct PyrLavel
 
433
    {
 
434
        PyrLavel(int _order, float _scale, cv::Size frame, cv::Size window, cv::Size minObjectSize)
 
435
        {
 
436
            do
 
437
            {
 
438
                order = _order;
 
439
                scale = pow(_scale, order);
 
440
                sFrame = frame / scale;
 
441
                workArea = sFrame - window + 1;
 
442
                sWindow = window * scale;
 
443
                _order++;
 
444
            } while (sWindow <= minObjectSize);
 
445
        }
 
446
 
 
447
        bool isFeasible(cv::Size maxObj)
 
448
        {
 
449
            return workArea.width > 0 && workArea.height > 0 && sWindow <= maxObj;
 
450
        }
 
451
 
 
452
        PyrLavel next(float factor, cv::Size frame, cv::Size window, cv::Size minObjectSize)
 
453
        {
 
454
            return PyrLavel(order + 1, factor, frame, window, minObjectSize);
 
455
        }
 
456
 
 
457
        int order;
 
458
        float scale;
 
459
        cv::Size sFrame;
 
460
        cv::Size workArea;
 
461
        cv::Size sWindow;
 
462
    };
 
463
 
 
464
    class LbpCascade_Impl : public CascadeClassifierBase
 
465
    {
 
466
    public:
 
467
        explicit LbpCascade_Impl(const FileStorage& file);
 
468
 
 
469
        virtual Size getClassifierSize() const { return NxM; }
 
470
 
 
471
        virtual void detectMultiScale(InputArray image,
 
472
                                      OutputArray objects,
 
473
                                      Stream& stream);
 
474
 
 
475
        virtual void convert(OutputArray gpu_objects,
 
476
                             std::vector<Rect>& objects);
 
477
 
 
478
    private:
 
479
        bool load(const FileNode &root);
 
480
        void allocateBuffers(cv::Size frame);
 
481
 
 
482
    private:
 
483
        struct Stage
 
484
        {
 
485
            int    first;
 
486
            int    ntrees;
 
487
            float  threshold;
 
488
        };
 
489
 
 
490
        enum stage { BOOST = 0 };
 
491
        enum feature { LBP = 1, HAAR = 2 };
 
492
 
 
493
        static const stage stageType = BOOST;
 
494
        static const feature featureType = LBP;
 
495
 
 
496
        cv::Size NxM;
 
497
        bool isStumps;
 
498
        int ncategories;
 
499
        int subsetSize;
 
500
        int nodeStep;
 
501
 
 
502
        // gpu representation of classifier
 
503
        GpuMat stage_mat;
 
504
        GpuMat trees_mat;
 
505
        GpuMat nodes_mat;
 
506
        GpuMat leaves_mat;
 
507
        GpuMat subsets_mat;
 
508
        GpuMat features_mat;
 
509
 
 
510
        GpuMat integral;
 
511
        GpuMat integralBuffer;
 
512
        GpuMat resuzeBuffer;
 
513
 
 
514
        GpuMat candidates;
 
515
        static const int integralFactor = 4;
 
516
    };
 
517
 
 
518
    LbpCascade_Impl::LbpCascade_Impl(const FileStorage& file)
 
519
    {
 
520
        load(file.getFirstTopLevelNode());
 
521
    }
 
522
 
 
523
    void LbpCascade_Impl::detectMultiScale(InputArray _image,
 
524
                                           OutputArray _objects,
 
525
                                           Stream& stream)
 
526
    {
 
527
        const GpuMat image = _image.getGpuMat();
 
528
 
 
529
        CV_Assert( image.depth() == CV_8U);
 
530
        CV_Assert( scaleFactor_ > 1 );
 
531
        CV_Assert( !stream );
 
532
 
 
533
        const float grouping_eps = 0.2f;
 
534
 
 
535
        BufferPool pool(stream);
 
536
        GpuMat objects = pool.getBuffer(1, maxNumObjects_, DataType<Rect>::type);
 
537
 
 
538
        // used for debug
 
539
        // candidates.setTo(cv::Scalar::all(0));
 
540
        // objects.setTo(cv::Scalar::all(0));
 
541
 
 
542
        if (maxObjectSize_ == cv::Size())
 
543
            maxObjectSize_ = image.size();
 
544
 
 
545
        allocateBuffers(image.size());
 
546
 
 
547
        unsigned int classified = 0;
 
548
        GpuMat dclassified(1, 1, CV_32S);
 
549
        cudaSafeCall( cudaMemcpy(dclassified.ptr(), &classified, sizeof(int), cudaMemcpyHostToDevice) );
 
550
 
 
551
        PyrLavel level(0, scaleFactor_, image.size(), NxM, minObjectSize_);
 
552
 
 
553
        while (level.isFeasible(maxObjectSize_))
 
554
        {
 
555
            int acc = level.sFrame.width + 1;
 
556
            float iniScale = level.scale;
 
557
 
 
558
            cv::Size area = level.workArea;
 
559
            int step = 1 + (level.scale <= 2.f);
 
560
 
 
561
            int total = 0, prev  = 0;
 
562
 
 
563
            while (acc <= integralFactor * (image.cols + 1) && level.isFeasible(maxObjectSize_))
 
564
            {
 
565
                // create sutable matrix headers
 
566
                GpuMat src  = resuzeBuffer(cv::Rect(0, 0, level.sFrame.width, level.sFrame.height));
 
567
                GpuMat sint = integral(cv::Rect(prev, 0, level.sFrame.width + 1, level.sFrame.height + 1));
 
568
 
 
569
                // generate integral for scale
 
570
                cuda::resize(image, src, level.sFrame, 0, 0, cv::INTER_LINEAR);
 
571
                cuda::integral(src, sint);
 
572
 
 
573
                // calculate job
 
574
                int totalWidth = level.workArea.width / step;
 
575
                total += totalWidth * (level.workArea.height / step);
 
576
 
 
577
                // go to next pyramide level
 
578
                level = level.next(scaleFactor_, image.size(), NxM, minObjectSize_);
 
579
                area = level.workArea;
 
580
 
 
581
                step = (1 + (level.scale <= 2.f));
 
582
                prev = acc;
 
583
                acc += level.sFrame.width + 1;
 
584
            }
 
585
 
 
586
            device::lbp::classifyPyramid(image.cols, image.rows, NxM.width - 1, NxM.height - 1, iniScale, scaleFactor_, total, stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat,
 
587
                leaves_mat, subsets_mat, features_mat, subsetSize, candidates, dclassified.ptr<unsigned int>(), integral);
 
588
        }
 
589
 
 
590
        if (minNeighbors_ <= 0  || objects.empty())
 
591
            return;
 
592
 
 
593
        cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) );
 
594
        device::lbp::connectedConmonents(candidates, classified, objects, minNeighbors_, grouping_eps, dclassified.ptr<unsigned int>());
 
595
 
 
596
        cudaSafeCall( cudaMemcpy(&classified, dclassified.ptr(), sizeof(int), cudaMemcpyDeviceToHost) );
 
597
        cudaSafeCall( cudaDeviceSynchronize() );
 
598
 
 
599
        if (classified > 0)
 
600
        {
 
601
            objects.colRange(0, classified).copyTo(_objects);
 
602
        }
 
603
        else
 
604
        {
 
605
            _objects.release();
 
606
        }
 
607
    }
 
608
 
 
609
    void LbpCascade_Impl::convert(OutputArray _gpu_objects, std::vector<Rect>& objects)
 
610
    {
 
611
        if (_gpu_objects.empty())
 
612
        {
 
613
            objects.clear();
 
614
            return;
 
615
        }
 
616
 
 
617
        Mat gpu_objects;
 
618
        if (_gpu_objects.kind() == _InputArray::CUDA_GPU_MAT)
 
619
        {
 
620
            _gpu_objects.getGpuMat().download(gpu_objects);
 
621
        }
 
622
        else
 
623
        {
 
624
            gpu_objects = _gpu_objects.getMat();
 
625
        }
 
626
 
 
627
        CV_Assert( gpu_objects.rows == 1 );
 
628
        CV_Assert( gpu_objects.type() == DataType<Rect>::type );
 
629
 
 
630
        Rect* ptr = gpu_objects.ptr<Rect>();
 
631
        objects.assign(ptr, ptr + gpu_objects.cols);
 
632
    }
 
633
 
 
634
    bool LbpCascade_Impl::load(const FileNode &root)
 
635
    {
 
636
        const char *CUDA_CC_STAGE_TYPE       = "stageType";
 
637
        const char *CUDA_CC_FEATURE_TYPE     = "featureType";
 
638
        const char *CUDA_CC_BOOST            = "BOOST";
 
639
        const char *CUDA_CC_LBP              = "LBP";
 
640
        const char *CUDA_CC_MAX_CAT_COUNT    = "maxCatCount";
 
641
        const char *CUDA_CC_HEIGHT           = "height";
 
642
        const char *CUDA_CC_WIDTH            = "width";
 
643
        const char *CUDA_CC_STAGE_PARAMS     = "stageParams";
 
644
        const char *CUDA_CC_MAX_DEPTH        = "maxDepth";
 
645
        const char *CUDA_CC_FEATURE_PARAMS   = "featureParams";
 
646
        const char *CUDA_CC_STAGES           = "stages";
 
647
        const char *CUDA_CC_STAGE_THRESHOLD  = "stageThreshold";
 
648
        const float CUDA_THRESHOLD_EPS       = 1e-5f;
 
649
        const char *CUDA_CC_WEAK_CLASSIFIERS = "weakClassifiers";
 
650
        const char *CUDA_CC_INTERNAL_NODES   = "internalNodes";
 
651
        const char *CUDA_CC_LEAF_VALUES      = "leafValues";
 
652
        const char *CUDA_CC_FEATURES         = "features";
 
653
        const char *CUDA_CC_RECT             = "rect";
 
654
 
 
655
        String stageTypeStr = (String)root[CUDA_CC_STAGE_TYPE];
 
656
        CV_Assert(stageTypeStr == CUDA_CC_BOOST);
 
657
 
 
658
        String featureTypeStr = (String)root[CUDA_CC_FEATURE_TYPE];
 
659
        CV_Assert(featureTypeStr == CUDA_CC_LBP);
 
660
 
 
661
        NxM.width =  (int)root[CUDA_CC_WIDTH];
 
662
        NxM.height = (int)root[CUDA_CC_HEIGHT];
 
663
        CV_Assert( NxM.height > 0 && NxM.width > 0 );
 
664
 
 
665
        isStumps = ((int)(root[CUDA_CC_STAGE_PARAMS][CUDA_CC_MAX_DEPTH]) == 1) ? true : false;
 
666
        CV_Assert(isStumps);
 
667
 
 
668
        FileNode fn = root[CUDA_CC_FEATURE_PARAMS];
 
669
        if (fn.empty())
 
670
            return false;
 
671
 
 
672
        ncategories = fn[CUDA_CC_MAX_CAT_COUNT];
 
673
 
 
674
        subsetSize = (ncategories + 31) / 32;
 
675
        nodeStep = 3 + ( ncategories > 0 ? subsetSize : 1 );
 
676
 
 
677
        fn = root[CUDA_CC_STAGES];
 
678
        if (fn.empty())
 
679
            return false;
 
680
 
 
681
        std::vector<Stage> stages;
 
682
        stages.reserve(fn.size());
 
683
 
 
684
        std::vector<int> cl_trees;
 
685
        std::vector<int> cl_nodes;
 
686
        std::vector<float> cl_leaves;
 
687
        std::vector<int> subsets;
 
688
 
 
689
        FileNodeIterator it = fn.begin(), it_end = fn.end();
 
690
        for (size_t si = 0; it != it_end; si++, ++it )
 
691
        {
 
692
            FileNode fns = *it;
 
693
            Stage st;
 
694
            st.threshold = (float)fns[CUDA_CC_STAGE_THRESHOLD] - CUDA_THRESHOLD_EPS;
 
695
 
 
696
            fns = fns[CUDA_CC_WEAK_CLASSIFIERS];
 
697
            if (fns.empty())
 
698
                return false;
 
699
 
 
700
            st.ntrees = (int)fns.size();
 
701
            st.first = (int)cl_trees.size();
 
702
 
 
703
            stages.push_back(st);// (int, int, float)
 
704
 
 
705
            cl_trees.reserve(stages[si].first + stages[si].ntrees);
 
706
 
 
707
            // weak trees
 
708
            FileNodeIterator it1 = fns.begin(), it1_end = fns.end();
 
709
            for ( ; it1 != it1_end; ++it1 )
 
710
            {
 
711
                FileNode fnw = *it1;
 
712
 
 
713
                FileNode internalNodes = fnw[CUDA_CC_INTERNAL_NODES];
 
714
                FileNode leafValues = fnw[CUDA_CC_LEAF_VALUES];
 
715
                if ( internalNodes.empty() || leafValues.empty() )
 
716
                    return false;
 
717
 
 
718
                int nodeCount = (int)internalNodes.size()/nodeStep;
 
719
                cl_trees.push_back(nodeCount);
 
720
 
 
721
                cl_nodes.reserve((cl_nodes.size() + nodeCount) * 3);
 
722
                cl_leaves.reserve(cl_leaves.size() + leafValues.size());
 
723
 
 
724
                if( subsetSize > 0 )
 
725
                    subsets.reserve(subsets.size() + nodeCount * subsetSize);
 
726
 
 
727
                // nodes
 
728
                FileNodeIterator iIt = internalNodes.begin(), iEnd = internalNodes.end();
 
729
 
 
730
                for( ; iIt != iEnd; )
 
731
                {
 
732
                    cl_nodes.push_back((int)*(iIt++));
 
733
                    cl_nodes.push_back((int)*(iIt++));
 
734
                    cl_nodes.push_back((int)*(iIt++));
 
735
 
 
736
                    if( subsetSize > 0 )
 
737
                        for( int j = 0; j < subsetSize; j++, ++iIt )
 
738
                            subsets.push_back((int)*iIt);
 
739
                }
 
740
 
 
741
                // leaves
 
742
                iIt = leafValues.begin(), iEnd = leafValues.end();
 
743
                for( ; iIt != iEnd; ++iIt )
 
744
                    cl_leaves.push_back((float)*iIt);
 
745
            }
 
746
        }
 
747
 
 
748
        fn = root[CUDA_CC_FEATURES];
 
749
        if( fn.empty() )
 
750
            return false;
 
751
        std::vector<uchar> features;
 
752
        features.reserve(fn.size() * 4);
 
753
        FileNodeIterator f_it = fn.begin(), f_end = fn.end();
 
754
        for (; f_it != f_end; ++f_it)
 
755
        {
 
756
            FileNode rect = (*f_it)[CUDA_CC_RECT];
 
757
            FileNodeIterator r_it = rect.begin();
 
758
            features.push_back(saturate_cast<uchar>((int)*(r_it++)));
 
759
            features.push_back(saturate_cast<uchar>((int)*(r_it++)));
 
760
            features.push_back(saturate_cast<uchar>((int)*(r_it++)));
 
761
            features.push_back(saturate_cast<uchar>((int)*(r_it++)));
 
762
        }
 
763
 
 
764
        // copy data structures on gpu
 
765
        stage_mat.upload(cv::Mat(1, (int) (stages.size() * sizeof(Stage)), CV_8UC1, (uchar*)&(stages[0]) ));
 
766
        trees_mat.upload(cv::Mat(cl_trees).reshape(1,1));
 
767
        nodes_mat.upload(cv::Mat(cl_nodes).reshape(1,1));
 
768
        leaves_mat.upload(cv::Mat(cl_leaves).reshape(1,1));
 
769
        subsets_mat.upload(cv::Mat(subsets).reshape(1,1));
 
770
        features_mat.upload(cv::Mat(features).reshape(4,1));
 
771
 
 
772
        return true;
 
773
    }
 
774
 
 
775
    void LbpCascade_Impl::allocateBuffers(cv::Size frame)
 
776
    {
 
777
        if (frame == cv::Size())
 
778
            return;
 
779
 
 
780
        if (resuzeBuffer.empty() || frame.width > resuzeBuffer.cols || frame.height > resuzeBuffer.rows)
 
781
        {
 
782
            resuzeBuffer.create(frame, CV_8UC1);
 
783
 
 
784
            integral.create(frame.height + 1, integralFactor * (frame.width + 1), CV_32SC1);
 
785
 
 
786
        #ifdef HAVE_OPENCV_CUDALEGACY
 
787
            NcvSize32u roiSize;
 
788
            roiSize.width = frame.width;
 
789
            roiSize.height = frame.height;
 
790
 
 
791
            cudaDeviceProp prop;
 
792
            cudaSafeCall( cudaGetDeviceProperties(&prop, cv::cuda::getDevice()) );
 
793
 
 
794
            Ncv32u bufSize;
 
795
            ncvSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize, prop) );
 
796
            integralBuffer.create(1, bufSize, CV_8UC1);
 
797
        #endif
 
798
 
 
799
            candidates.create(1 , frame.width >> 1, CV_32SC4);
 
800
        }
 
801
    }
 
802
 
 
803
}
 
804
 
 
805
//
 
806
// create
 
807
//
 
808
 
 
809
Ptr<cuda::CascadeClassifier> cv::cuda::CascadeClassifier::create(const String& filename)
 
810
{
 
811
    String fext = filename.substr(filename.find_last_of(".") + 1);
 
812
    fext = fext.toLowerCase();
 
813
 
 
814
    if (fext == "nvbin")
 
815
    {
 
816
    #ifndef HAVE_OPENCV_CUDALEGACY
 
817
        CV_Error(Error::StsUnsupportedFormat, "OpenCV CUDA objdetect was built without HaarCascade");
 
818
        return Ptr<cuda::CascadeClassifier>();
 
819
    #else
 
820
        return makePtr<HaarCascade_Impl>(filename);
 
821
    #endif
 
822
    }
 
823
 
 
824
    FileStorage fs(filename, FileStorage::READ);
 
825
 
 
826
    if (!fs.isOpened())
 
827
    {
 
828
    #ifndef HAVE_OPENCV_CUDALEGACY
 
829
        CV_Error(Error::StsUnsupportedFormat, "OpenCV CUDA objdetect was built without HaarCascade");
 
830
        return Ptr<cuda::CascadeClassifier>();
 
831
    #else
 
832
        return makePtr<HaarCascade_Impl>(filename);
 
833
    #endif
 
834
    }
 
835
 
 
836
    const char *CUDA_CC_LBP = "LBP";
 
837
    String featureTypeStr = (String)fs.getFirstTopLevelNode()["featureType"];
 
838
    if (featureTypeStr == CUDA_CC_LBP)
 
839
    {
 
840
        return makePtr<LbpCascade_Impl>(fs);
 
841
    }
 
842
    else
 
843
    {
 
844
    #ifndef HAVE_OPENCV_CUDALEGACY
 
845
        CV_Error(Error::StsUnsupportedFormat, "OpenCV CUDA objdetect was built without HaarCascade");
 
846
        return Ptr<cuda::CascadeClassifier>();
 
847
    #else
 
848
        return makePtr<HaarCascade_Impl>(filename);
 
849
    #endif
 
850
    }
 
851
 
 
852
    CV_Error(Error::StsUnsupportedFormat, "Unsupported format for CUDA CascadeClassifier");
 
853
    return Ptr<cuda::CascadeClassifier>();
 
854
}
 
855
 
 
856
Ptr<cuda::CascadeClassifier> cv::cuda::CascadeClassifier::create(const FileStorage& file)
 
857
{
 
858
    return makePtr<LbpCascade_Impl>(file);
 
859
}
 
860
 
 
861
#endif