Point Cloud Library (PCL)  1.9.1
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(void)
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(void)
142  {
143  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagTilted) != 0;
144  }
145 
146  __device__ __host__ NcvBool isLeftNodeLeaf(void)
147  {
148  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf) != 0;
149  }
150 
151  __device__ __host__ NcvBool isRightNodeLeaf(void)
152  {
153  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf) != 0;
154  }
155 
156  __device__ __host__ Ncv32u getNumFeatures(void)
157  {
158  return (this->desc >> HaarFeatureDescriptor32_NumFeatures_Shift) & HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures;
159  }
160 
161  __device__ __host__ Ncv32u getFeaturesOffset(void)
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(void)
184  {
185  return (*(Ncv32f *)&this->_ui1.x);
186  }
187 
188  __host__ bool isLeaf() // TODO: check this hack don't know if is correct
189  {
190  if( _ui1.x == 0)
191  return (false);
192  else
193  return (true);
194  }
195 
196 #ifdef __CUDACC__
197  __device__ Ncv32f getLeafValue(void)
198  {
199  return (__int_as_float(this->_ui1.x));
200  }
201 #endif
202 
203  __device__ __host__ Ncv32u getNextNodeOffset(void)
204  {
205  return (this->_ui1.x);
206  }
207 };
208 
210 {
211  uint4 _ui4;
212 
214  {
215  this->_ui4.x = *(Ncv32u *)&f;
216  return NCV_SUCCESS;
217  }
218 
219  __host__ NCVStatus setThreshold(Ncv32f t)
220  {
221  this->_ui4.y = *(Ncv32u *)&t;
222  return NCV_SUCCESS;
223  }
224 
226  {
227  this->_ui4.z = *(Ncv32u *)&nl;
228  return NCV_SUCCESS;
229  }
230 
232  {
233  this->_ui4.w = *(Ncv32u *)&nr;
234  return NCV_SUCCESS;
235  }
236 
237  __host__ __device__ HaarFeatureDescriptor32 getFeatureDesc(void)
238  {
239  return *(HaarFeatureDescriptor32 *)&this->_ui4.x;
240  }
241 
242  __host__ __device__ Ncv32f getThreshold(void)
243  {
244  return *(Ncv32f*)&this->_ui4.y;
245  }
246 
248  {
249  return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.z;
250  }
251 
253  {
254  return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.w;
255  }
256 };
257 
259 {
260 #define HaarStage64_Interpret_MaskRootNodes 0x0000FFFF
261 #define HaarStage64_Interpret_MaskRootNodeOffset 0xFFFF0000
262 #define HaarStage64_Interpret_ShiftRootNodeOffset 16
263 
264  uint2 _ui2;
265 
266  __host__ NCVStatus setStageThreshold(Ncv32f t)
267  {
268  this->_ui2.x = *(Ncv32u *)&t;
269  return NCV_SUCCESS;
270  }
271 
272  __host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val)
273  {
274  if (val > (HaarStage64_Interpret_MaskRootNodeOffset >> HaarStage64_Interpret_ShiftRootNodeOffset))
275  {
276  return NCV_HAAR_XML_LOADING_EXCEPTION;
277  }
278  this->_ui2.y = (val << HaarStage64_Interpret_ShiftRootNodeOffset) | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes);
279  return NCV_SUCCESS;
280  }
281 
282  __host__ NCVStatus setNumClassifierRootNodes(Ncv32u val)
283  {
284  if (val > HaarStage64_Interpret_MaskRootNodes)
285  {
286  return NCV_HAAR_XML_LOADING_EXCEPTION;
287  }
288  this->_ui2.y = val | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodeOffset);
289  return NCV_SUCCESS;
290  }
291 
292  __host__ __device__ Ncv32f getStageThreshold(void)
293  {
294  return *(Ncv32f*)&this->_ui2.x;
295  }
296 
297  __host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void)
298  {
299  return (this->_ui2.y >> HaarStage64_Interpret_ShiftRootNodeOffset);
300  }
301 
302  __host__ __device__ Ncv32u getNumClassifierRootNodes(void)
303  {
304  return (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes);
305  }
306 };
307 
308 NCV_CT_ASSERT(sizeof(HaarFeature64) == 8);
309 NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4);
310 NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4);
311 NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16);
312 NCV_CT_ASSERT(sizeof(HaarStage64) == 8);
313 
314 /**
315  * \brief Classifier cascade descriptor
316  */
318 {
319  Ncv32u NumStages;
322  Ncv32u NumFeatures;
324  NcvBool bNeedsTiltedII;
325  NcvBool bHasStumpsOnly;
326 };
327 
328 //==============================================================================
329 //
330 // Functional interface
331 //
332 //==============================================================================
333 
334 enum
335 {
336  NCVPipeObjDet_Default = 0x000,
337  NCVPipeObjDet_UseFairImageScaling = 0x001,
338  NCVPipeObjDet_FindLargestObject = 0x002,
339  NCVPipeObjDet_VisualizeInPlace = 0x004,
340 };
341 
342 NCV_EXPORTS NCVStatus ncvDetectObjectsMultiScale_device(NCVMatrix<Ncv8u> &d_srcImg,
343  NcvSize32u srcRoi,
344  NCVVector<NcvRect32u> &d_dstRects,
345  Ncv32u &dstNumRects,
346 
348  NCVVector<HaarStage64> &h_HaarStages,
349  NCVVector<HaarStage64> &d_HaarStages,
351  NCVVector<HaarFeature64> &d_HaarFeatures,
352 
353  NcvSize32u minObjSize,
354  Ncv32u minNeighbors, //default 4
355  Ncv32f scaleStep, //default 1.2f
356  Ncv32u pixelStep, //default 1
357  Ncv32u flags, //default NCVPipeObjDet_Default
358 
359  INCVMemAllocator &gpuAllocator,
360  INCVMemAllocator &cpuAllocator,
361  cudaDeviceProp &devProp,
362  cudaStream_t cuStream);
363 
364 #define OBJDET_MASK_ELEMENT_INVALID_32U 0xFFFFFFFF
365 #define HAAR_STDDEV_BORDER 1
366 
367 NCV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix<Ncv32u> &d_integralImage,
368  NCVMatrix<Ncv32f> &d_weights,
369  NCVMatrixAlloc<Ncv32u> &d_pixelMask,
370  Ncv32u &numDetections,
372  NCVVector<HaarStage64> &h_HaarStages,
373  NCVVector<HaarStage64> &d_HaarStages,
375  NCVVector<HaarFeature64> &d_HaarFeatures,
376  NcvBool bMaskElements,
377  NcvSize32u anchorsRoi,
378  Ncv32u pixelStep,
379  Ncv32f scaleArea,
380  INCVMemAllocator &gpuAllocator,
381  INCVMemAllocator &cpuAllocator,
382  cudaDeviceProp &devProp,
383  cudaStream_t cuStream);
384 
385 NCV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_host(NCVMatrix<Ncv32u> &h_integralImage,
386  NCVMatrix<Ncv32f> &h_weights,
387  NCVMatrixAlloc<Ncv32u> &h_pixelMask,
388  Ncv32u &numDetections,
390  NCVVector<HaarStage64> &h_HaarStages,
392  NCVVector<HaarFeature64> &h_HaarFeatures,
393  NcvBool bMaskElements,
394  NcvSize32u anchorsRoi,
395  Ncv32u pixelStep,
396  Ncv32f scaleArea);
397 
398 #define RECT_SIMILARITY_PROPORTION 0.2f
399 
400 NCV_EXPORTS NCVStatus ncvGrowDetectionsVector_device(NCVVector<Ncv32u> &pixelMask,
401  Ncv32u numPixelMaskDetections,
402  NCVVector<NcvRect32u> &hypotheses,
403  Ncv32u &totalDetections,
404  Ncv32u totalMaxDetections,
405  Ncv32u rectWidth,
406  Ncv32u rectHeight,
407  Ncv32f curScale,
408  cudaStream_t cuStream);
409 
410 NCV_EXPORTS NCVStatus ncvGrowDetectionsVector_host(NCVVector<Ncv32u> &pixelMask,
411  Ncv32u numPixelMaskDetections,
412  NCVVector<NcvRect32u> &hypotheses,
413  Ncv32u &totalDetections,
414  Ncv32u totalMaxDetections,
415  Ncv32u rectWidth,
416  Ncv32u rectHeight,
417  Ncv32f curScale);
418 
419 NCV_EXPORTS NCVStatus ncvHaarGetClassifierSize(const std::string &filename, Ncv32u &numStages,
420  Ncv32u &numNodes, Ncv32u &numFeatures);
421 
422 NCV_EXPORTS NCVStatus ncvHaarLoadFromFile_host(const std::string &filename,
424  NCVVector<HaarStage64> &h_HaarStages,
426  NCVVector<HaarFeature64> &h_HaarFeatures);
427 
428 NCV_EXPORTS NCVStatus ncvHaarStoreNVBIN_host(const std::string &filename,
430  NCVVector<HaarStage64> &h_HaarStages,
432  NCVVector<HaarFeature64> &h_HaarFeatures);
433 
434 #endif // PCL_GPU_PEOPLE_NCVHAAROBJECTDETECTION_HPP_
__host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val)
Ncv8u y
Definition: NCV.hpp:131
__host__ NCVStatus setWeight(Ncv32f weight)
NCVMatrixAlloc.
Definition: NCV.hpp:801
__host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u, Ncv32u)
__host__ NCVStatus setNumClassifierRootNodes(Ncv32u val)
__device__ __host__ NcvBool isTilted(void)
__host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc(void)
__device__ __host__ Ncv32u getFeaturesOffset(void)
__host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl)
Ncv8u width
Definition: NCV.hpp:132
__device__ __host__ NcvBool isLeftNodeLeaf(void)
__host__ __device__ Ncv32f getThreshold(void)
__host__ NCVStatus setThreshold(Ncv32f t)
Ncv8u x
Definition: NCV.hpp:130
__host__ NCVStatus create(Ncv32f leafValue)
__device__ __host__ NcvBool isRightNodeLeaf(void)
Classifier cascade descriptor.
NCVVector (1D)
Definition: NCV.hpp:523
__host__ __device__ HaarFeatureDescriptor32 getFeatureDesc(void)
__host__ __device__ Ncv32f getStageThreshold(void)
Ncv8u height
Definition: NCV.hpp:133
__host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr)
__host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf, Ncv32u numFeatures, Ncv32u offsetFeatures)
INCVMemAllocator (Interface)
Definition: NCV.hpp:418
__device__ __host__ Ncv32u getNextNodeOffset(void)
__host__ NCVStatus create(Ncv32u offsetHaarClassifierNode)
__host__ __device__ Ncv32u getNumClassifierRootNodes(void)
__device__ __host__ Ncv32u getNumFeatures(void)
NCVMatrix (2D)
Definition: NCV.hpp:695
__host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f)
__host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc(void)
__host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void)
__host__ NCVStatus setStageThreshold(Ncv32f t)
__device__ __host__ Ncv32f getWeight(void)
__device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)