Point Cloud Library (PCL) 1.12.0
Loading...
Searching...
No Matches
NCVHaarObjectDetection.hpp
1/*
2 * Software License Agreement (BSD License)
3 *
4 * Point Cloud Library (PCL) - www.pointclouds.org
5 * Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.
6 * Third party copyrights are property of their respective owners.
7 *
8 * All rights reserved.
9 *
10 * Redistribution and use in source and binary forms, with or without
11 * modification, are permitted provided that the following conditions
12 * are met:
13 *
14 * * Redistributions of source code must retain the above copyright
15 * notice, this list of conditions and the following disclaimer.
16 * * Redistributions in binary form must reproduce the above
17 * copyright notice, this list of conditions and the following
18 * disclaimer in the documentation and/or other materials provided
19 * with the distribution.
20 * * Neither the name of Willow Garage, Inc. nor the names of its
21 * contributors may be used to endorse or promote products derived
22 * from this software without specific prior written permission.
23 *
24 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
25 * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
26 * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
27 * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
28 * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
29 * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
30 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
31 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
32 * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
33 * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
34 * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
35 * POSSIBILITY OF SUCH DAMAGE.
36 *
37 * $Id: $
38 * Ported to PCL by Koen Buys : Attention Work in progress!
39 */
40
41////////////////////////////////////////////////////////////////////////////////
42//
43// NVIDIA CUDA implementation of Viola-Jones Object Detection Framework
44//
45// The algorithm and code are explained in the upcoming GPU Computing Gems
46// chapter in detail:
47//
48// Anton Obukhov, "Haar Classifiers for Object Detection with CUDA"
49// PDF URL placeholder
50// email: aobukhov@nvidia.com, devsupport@nvidia.com
51//
52// Credits for help with the code to:
53// Alexey Mendelenko, Cyril Crassin, and Mikhail Smirnov.
54//
55////////////////////////////////////////////////////////////////////////////////
56
57#ifndef PCL_GPU_PEOPLE_NCVHAAROBJECTDETECTION_HPP_
58#define PCL_GPU_PEOPLE_NCVHAAROBJECTDETECTION_HPP_
59
60#include <string>
61#include "NCV.hpp"
62
63//==============================================================================
64//
65// Guaranteed size cross-platform classifier structures
66//
67//==============================================================================
68
70{
71 uint2 _ui2;
72
73#define HaarFeature64_CreateCheck_MaxRectField 0xFF
74
75 __host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u /*clsWidth*/, Ncv32u /*clsHeight*/)
76 {
77 ncvAssertReturn(rectWidth <= HaarFeature64_CreateCheck_MaxRectField && rectHeight <= HaarFeature64_CreateCheck_MaxRectField, NCV_HAAR_TOO_LARGE_FEATURES);
78 ((NcvRect8u*)&(this->_ui2.x))->x = (Ncv8u)rectX;
79 ((NcvRect8u*)&(this->_ui2.x))->y = (Ncv8u)rectY;
80 ((NcvRect8u*)&(this->_ui2.x))->width = (Ncv8u)rectWidth;
81 ((NcvRect8u*)&(this->_ui2.x))->height = (Ncv8u)rectHeight;
82 return NCV_SUCCESS;
83 }
84
85 __host__ NCVStatus setWeight(Ncv32f weight)
86 {
87 ((Ncv32f*)&(this->_ui2.y))[0] = weight;
88 return NCV_SUCCESS;
89 }
90
91 __device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)
92 {
93 NcvRect8u tmpRect = *(NcvRect8u*)(&this->_ui2.x);
94 *rectX = tmpRect.x;
95 *rectY = tmpRect.y;
96 *rectWidth = tmpRect.width;
97 *rectHeight = tmpRect.height;
98 }
99
100 __device__ __host__ Ncv32f getWeight()
101 {
102 return *(Ncv32f*)(&this->_ui2.y);
103 }
104};
105
107{
108 private:
109
110#define HaarFeatureDescriptor32_Interpret_MaskFlagTilted 0x80000000
111#define HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf 0x40000000
112#define HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf 0x20000000
113#define HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures 0x1F
114#define HaarFeatureDescriptor32_NumFeatures_Shift 24
115#define HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset 0x00FFFFFF
116
117 Ncv32u desc;
118
119 public:
120
121 __host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf,
122 Ncv32u numFeatures, Ncv32u offsetFeatures)
123 {
124 if (numFeatures > HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures)
125 {
126 return NCV_HAAR_TOO_MANY_FEATURES_IN_CLASSIFIER;
127 }
128 if (offsetFeatures > HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset)
129 {
130 return NCV_HAAR_TOO_MANY_FEATURES_IN_CASCADE;
131 }
132 this->desc = 0;
133 this->desc |= (bTilted ? HaarFeatureDescriptor32_Interpret_MaskFlagTilted : 0);
134 this->desc |= (bLeftLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf : 0);
135 this->desc |= (bRightLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf : 0);
136 this->desc |= (numFeatures << HaarFeatureDescriptor32_NumFeatures_Shift);
137 this->desc |= offsetFeatures;
138 return NCV_SUCCESS;
139 }
140
141 __device__ __host__ NcvBool isTilted() const
142 {
143 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagTilted) != 0;
144 }
145
146 __device__ __host__ NcvBool isLeftNodeLeaf() const
147 {
148 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf) != 0;
149 }
150
151 __device__ __host__ NcvBool isRightNodeLeaf() const
152 {
153 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf) != 0;
154 }
155
156 __device__ __host__ Ncv32u getNumFeatures() const
157 {
158 return (this->desc >> HaarFeatureDescriptor32_NumFeatures_Shift) & HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures;
159 }
160
161 __device__ __host__ Ncv32u getFeaturesOffset() const
162 {
163 return this->desc & HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset;
164 }
165};
166
168{
169 uint1 _ui1;
170
171 __host__ NCVStatus create(Ncv32f leafValue)
172 {
173 *(Ncv32f *)&this->_ui1 = leafValue;
174 return (NCV_SUCCESS);
175 }
176
177 __host__ NCVStatus create(Ncv32u offsetHaarClassifierNode)
178 {
179 this->_ui1.x = offsetHaarClassifierNode;
180 return (NCV_SUCCESS);
181 }
182
183 __host__ Ncv32f getLeafValueHost()
184 {
185 return (*(Ncv32f *)&this->_ui1.x);
186 }
187
188 __host__ bool isLeaf() const // TODO: check this hack don't know if is correct
189 {
190 return ( _ui1.x != 0);
191 }
192
193#ifdef __CUDACC__
194 __device__ Ncv32f getLeafValue(void)
195 {
196 return (__int_as_float(this->_ui1.x));
197 }
198#endif
199
200 __device__ __host__ Ncv32u getNextNodeOffset()
201 {
202 return (this->_ui1.x);
203 }
204};
205
207{
208 uint4 _ui4;
209
211 {
212 this->_ui4.x = *(Ncv32u *)&f;
213 return NCV_SUCCESS;
214 }
215
216 __host__ NCVStatus setThreshold(Ncv32f t)
217 {
218 this->_ui4.y = *(Ncv32u *)&t;
219 return NCV_SUCCESS;
220 }
221
223 {
224 this->_ui4.z = *(Ncv32u *)&nl;
225 return NCV_SUCCESS;
226 }
227
229 {
230 this->_ui4.w = *(Ncv32u *)&nr;
231 return NCV_SUCCESS;
232 }
233
235 {
236 return *(HaarFeatureDescriptor32 *)&this->_ui4.x;
237 }
238
239 __host__ __device__ Ncv32f getThreshold()
240 {
241 return *(Ncv32f*)&this->_ui4.y;
242 }
243
245 {
246 return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.z;
247 }
248
250 {
251 return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.w;
252 }
253};
254
256{
257#define HaarStage64_Interpret_MaskRootNodes 0x0000FFFF
258#define HaarStage64_Interpret_MaskRootNodeOffset 0xFFFF0000
259#define HaarStage64_Interpret_ShiftRootNodeOffset 16
260
261 uint2 _ui2;
262
263 __host__ NCVStatus setStageThreshold(Ncv32f t)
264 {
265 this->_ui2.x = *(Ncv32u *)&t;
266 return NCV_SUCCESS;
267 }
268
269 __host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val)
270 {
271 if (val > (HaarStage64_Interpret_MaskRootNodeOffset >> HaarStage64_Interpret_ShiftRootNodeOffset))
272 {
273 return NCV_HAAR_XML_LOADING_EXCEPTION;
274 }
275 this->_ui2.y = (val << HaarStage64_Interpret_ShiftRootNodeOffset) | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes);
276 return NCV_SUCCESS;
277 }
278
279 __host__ NCVStatus setNumClassifierRootNodes(Ncv32u val)
280 {
281 if (val > HaarStage64_Interpret_MaskRootNodes)
282 {
283 return NCV_HAAR_XML_LOADING_EXCEPTION;
284 }
285 this->_ui2.y = val | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodeOffset);
286 return NCV_SUCCESS;
287 }
288
289 __host__ __device__ Ncv32f getStageThreshold()
290 {
291 return *(Ncv32f*)&this->_ui2.x;
292 }
293
294 __host__ __device__ Ncv32u getStartClassifierRootNodeOffset() const
295 {
296 return (this->_ui2.y >> HaarStage64_Interpret_ShiftRootNodeOffset);
297 }
298
299 __host__ __device__ Ncv32u getNumClassifierRootNodes() const
300 {
301 return (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes);
302 }
303};
304
305NCV_CT_ASSERT(sizeof(HaarFeature64) == 8);
306NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4);
307NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4);
308NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16);
309NCV_CT_ASSERT(sizeof(HaarStage64) == 8);
310
311/**
312 * \brief Classifier cascade descriptor
313 */
315{
316 Ncv32u NumStages;
323};
324
325//==============================================================================
326//
327// Functional interface
328//
329//==============================================================================
330
331enum
332{
333 NCVPipeObjDet_Default = 0x000,
334 NCVPipeObjDet_UseFairImageScaling = 0x001,
335 NCVPipeObjDet_FindLargestObject = 0x002,
336 NCVPipeObjDet_VisualizeInPlace = 0x004,
337};
338
339NCV_EXPORTS NCVStatus ncvDetectObjectsMultiScale_device(NCVMatrix<Ncv8u> &d_srcImg,
340 NcvSize32u srcRoi,
341 NCVVector<NcvRect32u> &d_dstRects,
342 Ncv32u &dstNumRects,
343
345 NCVVector<HaarStage64> &h_HaarStages,
346 NCVVector<HaarStage64> &d_HaarStages,
348 NCVVector<HaarFeature64> &d_HaarFeatures,
349
350 NcvSize32u minObjSize,
351 Ncv32u minNeighbors, //default 4
352 Ncv32f scaleStep, //default 1.2f
353 Ncv32u pixelStep, //default 1
354 Ncv32u flags, //default NCVPipeObjDet_Default
355
356 INCVMemAllocator &gpuAllocator,
357 INCVMemAllocator &cpuAllocator,
358 cudaDeviceProp &devProp,
359 cudaStream_t cuStream);
360
361#define OBJDET_MASK_ELEMENT_INVALID_32U 0xFFFFFFFF
362#define HAAR_STDDEV_BORDER 1
363
364NCV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix<Ncv32u> &d_integralImage,
365 NCVMatrix<Ncv32f> &d_weights,
366 NCVMatrixAlloc<Ncv32u> &d_pixelMask,
367 Ncv32u &numDetections,
369 NCVVector<HaarStage64> &h_HaarStages,
370 NCVVector<HaarStage64> &d_HaarStages,
372 NCVVector<HaarFeature64> &d_HaarFeatures,
373 NcvBool bMaskElements,
374 NcvSize32u anchorsRoi,
375 Ncv32u pixelStep,
376 Ncv32f scaleArea,
377 INCVMemAllocator &gpuAllocator,
378 INCVMemAllocator &cpuAllocator,
379 cudaDeviceProp &devProp,
380 cudaStream_t cuStream);
381
382NCV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_host(NCVMatrix<Ncv32u> &h_integralImage,
383 NCVMatrix<Ncv32f> &h_weights,
384 NCVMatrixAlloc<Ncv32u> &h_pixelMask,
385 Ncv32u &numDetections,
387 NCVVector<HaarStage64> &h_HaarStages,
389 NCVVector<HaarFeature64> &h_HaarFeatures,
390 NcvBool bMaskElements,
391 NcvSize32u anchorsRoi,
392 Ncv32u pixelStep,
393 Ncv32f scaleArea);
394
395#define RECT_SIMILARITY_PROPORTION 0.2f
396
397NCV_EXPORTS NCVStatus ncvGrowDetectionsVector_device(NCVVector<Ncv32u> &pixelMask,
398 Ncv32u numPixelMaskDetections,
399 NCVVector<NcvRect32u> &hypotheses,
400 Ncv32u &totalDetections,
401 Ncv32u totalMaxDetections,
402 Ncv32u rectWidth,
403 Ncv32u rectHeight,
404 Ncv32f curScale,
405 cudaStream_t cuStream);
406
407NCV_EXPORTS NCVStatus ncvGrowDetectionsVector_host(NCVVector<Ncv32u> &pixelMask,
408 Ncv32u numPixelMaskDetections,
409 NCVVector<NcvRect32u> &hypotheses,
410 Ncv32u &totalDetections,
411 Ncv32u totalMaxDetections,
412 Ncv32u rectWidth,
413 Ncv32u rectHeight,
414 Ncv32f curScale);
415
416NCV_EXPORTS NCVStatus ncvHaarGetClassifierSize(const std::string &filename, Ncv32u &numStages,
417 Ncv32u &numNodes, Ncv32u &numFeatures);
418
419NCV_EXPORTS NCVStatus ncvHaarLoadFromFile_host(const std::string &filename,
421 NCVVector<HaarStage64> &h_HaarStages,
423 NCVVector<HaarFeature64> &h_HaarFeatures);
424
425NCV_EXPORTS NCVStatus ncvHaarStoreNVBIN_host(const std::string &filename,
427 NCVVector<HaarStage64> &h_HaarStages,
429 NCVVector<HaarFeature64> &h_HaarFeatures);
430
431#endif // PCL_GPU_PEOPLE_NCVHAAROBJECTDETECTION_HPP_
INCVMemAllocator (Interface)
Definition NCV.hpp:403
NCVMatrixAlloc.
Definition NCV.hpp:786
NCVMatrix (2D)
Definition NCV.hpp:680
NCVVector (1D)
Definition NCV.hpp:508
Classifier cascade descriptor.
__host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc()
__host__ __device__ HaarFeatureDescriptor32 getFeatureDesc()
__host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr)
__host__ NCVStatus setThreshold(Ncv32f t)
__host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f)
__host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc()
__host__ __device__ Ncv32f getThreshold()
__host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl)
__host__ NCVStatus create(Ncv32u offsetHaarClassifierNode)
__host__ NCVStatus create(Ncv32f leafValue)
__device__ __host__ Ncv32u getNextNodeOffset()
__device__ __host__ Ncv32f getWeight()
__host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u, Ncv32u)
__host__ NCVStatus setWeight(Ncv32f weight)
__device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)
__device__ __host__ Ncv32u getNumFeatures() const
__device__ __host__ NcvBool isRightNodeLeaf() const
__device__ __host__ NcvBool isTilted() const
__device__ __host__ Ncv32u getFeaturesOffset() const
__device__ __host__ NcvBool isLeftNodeLeaf() const
__host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf, Ncv32u numFeatures, Ncv32u offsetFeatures)
__host__ __device__ Ncv32u getNumClassifierRootNodes() const
__host__ NCVStatus setNumClassifierRootNodes(Ncv32u val)
__host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val)
__host__ NCVStatus setStageThreshold(Ncv32f t)
__host__ __device__ Ncv32f getStageThreshold()
__host__ __device__ Ncv32u getStartClassifierRootNodeOffset() const
Ncv8u y
Definition NCV.hpp:131
Ncv8u width
Definition NCV.hpp:132
Ncv8u height
Definition NCV.hpp:133
Ncv8u x
Definition NCV.hpp:130