Point Cloud Library (PCL)  1.9.1-dev
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()
142  {
143  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagTilted) != 0;
144  }
145 
146  __device__ __host__ NcvBool isLeftNodeLeaf()
147  {
148  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf) != 0;
149  }
150 
151  __device__ __host__ NcvBool isRightNodeLeaf()
152  {
153  return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf) != 0;
154  }
155 
156  __device__ __host__ Ncv32u getNumFeatures()
157  {
158  return (this->desc >> HaarFeatureDescriptor32_NumFeatures_Shift) & HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures;
159  }
160 
161  __device__ __host__ Ncv32u getFeaturesOffset()
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() // 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()
295  {
296  return (this->_ui2.y >> HaarStage64_Interpret_ShiftRootNodeOffset);
297  }
298 
299  __host__ __device__ Ncv32u getNumClassifierRootNodes()
300  {
301  return (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes);
302  }
303 };
304 
305 NCV_CT_ASSERT(sizeof(HaarFeature64) == 8);
306 NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4);
307 NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4);
308 NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16);
309 NCV_CT_ASSERT(sizeof(HaarStage64) == 8);
310 
311 /**
312  * \brief Classifier cascade descriptor
313  */
315 {
316  Ncv32u NumStages;
319  Ncv32u NumFeatures;
321  NcvBool bNeedsTiltedII;
322  NcvBool bHasStumpsOnly;
323 };
324 
325 //==============================================================================
326 //
327 // Functional interface
328 //
329 //==============================================================================
330 
331 enum
332 {
333  NCVPipeObjDet_Default = 0x000,
334  NCVPipeObjDet_UseFairImageScaling = 0x001,
335  NCVPipeObjDet_FindLargestObject = 0x002,
336  NCVPipeObjDet_VisualizeInPlace = 0x004,
337 };
338 
339 NCV_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 
364 NCV_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 
382 NCV_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 
397 NCV_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 
407 NCV_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 
416 NCV_EXPORTS NCVStatus ncvHaarGetClassifierSize(const std::string &filename, Ncv32u &numStages,
417  Ncv32u &numNodes, Ncv32u &numFeatures);
418 
419 NCV_EXPORTS NCVStatus ncvHaarLoadFromFile_host(const std::string &filename,
421  NCVVector<HaarStage64> &h_HaarStages,
423  NCVVector<HaarFeature64> &h_HaarFeatures);
424 
425 NCV_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_
__host__ __device__ Ncv32f getThreshold()
__host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val)
__host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc()
__device__ __host__ Ncv32f getWeight()
Ncv8u y
Definition: NCV.hpp:131
__host__ NCVStatus setWeight(Ncv32f weight)
NCVMatrixAlloc.
Definition: NCV.hpp:785
__host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u, Ncv32u)
__host__ NCVStatus setNumClassifierRootNodes(Ncv32u val)
__device__ __host__ NcvBool isLeftNodeLeaf()
__device__ __host__ Ncv32u getNextNodeOffset()
__host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl)
Ncv8u width
Definition: NCV.hpp:132
__host__ NCVStatus setThreshold(Ncv32f t)
Ncv8u x
Definition: NCV.hpp:130
__host__ NCVStatus create(Ncv32f leafValue)
Classifier cascade descriptor.
NCVVector (1D)
Definition: NCV.hpp:507
Ncv8u height
Definition: NCV.hpp:133
__device__ __host__ NcvBool isTilted()
__host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr)
__host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf, Ncv32u numFeatures, Ncv32u offsetFeatures)
__host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc()
INCVMemAllocator (Interface)
Definition: NCV.hpp:402
__device__ __host__ NcvBool isRightNodeLeaf()
__host__ NCVStatus create(Ncv32u offsetHaarClassifierNode)
__host__ __device__ Ncv32u getNumClassifierRootNodes()
__host__ __device__ Ncv32f getStageThreshold()
__device__ __host__ Ncv32u getFeaturesOffset()
NCVMatrix (2D)
Definition: NCV.hpp:679
__host__ __device__ Ncv32u getStartClassifierRootNodeOffset()
__host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f)
__device__ __host__ Ncv32u getNumFeatures()
__host__ __device__ HaarFeatureDescriptor32 getFeatureDesc()
__host__ NCVStatus setStageThreshold(Ncv32f t)
__device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)