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 //////////////////////////////////////////////////////////////////////////////// 44 // 45 // NVIDIA CUDA implementation of Viola-Jones Object Detection Framework 46 // 47 // The algorithm and code are explained in the upcoming GPU Computing Gems 48 // chapter in detail: 49 // 50 // Anton Obukhov, "Haar Classifiers for Object Detection with CUDA" 51 // PDF URL placeholder 52 // email: aobukhov (at) nvidia.com, devsupport (at) nvidia.com 53 // 54 // Credits for help with the code to: 55 // Alexey Mendelenko, Cyril Crassin, and Mikhail Smirnov. 56 // 57 //////////////////////////////////////////////////////////////////////////////// 58 59 #ifndef _ncvhaarobjectdetection_hpp_ 60 #define _ncvhaarobjectdetection_hpp_ 61 62 #include "opencv2/cudalegacy/NCV.hpp" 63 64 //! @addtogroup cudalegacy 65 //! @{ 66 67 //============================================================================== 68 // 69 // Guaranteed size cross-platform classifier structures 70 // 71 //============================================================================== 72 #if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__ > 4 73 typedef Ncv32f __attribute__((__may_alias__)) Ncv32f_a; 74 #else 75 typedef Ncv32f Ncv32f_a; 76 #endif 77 78 struct HaarFeature64 79 { 80 uint2 _ui2; 81 82 #define HaarFeature64_CreateCheck_MaxRectField 0xFF 83 84 __host__ NCVStatus setRect(Ncv32u rectX, Ncv32u rectY, Ncv32u rectWidth, Ncv32u rectHeight, Ncv32u /*clsWidth*/, Ncv32u /*clsHeight*/) 85 { 86 ncvAssertReturn(rectWidth <= HaarFeature64_CreateCheck_MaxRectField && rectHeight <= HaarFeature64_CreateCheck_MaxRectField, NCV_HAAR_TOO_LARGE_FEATURES); 87 ((NcvRect8u*)&(this->_ui2.x))->x = (Ncv8u)rectX; 88 ((NcvRect8u*)&(this->_ui2.x))->y = (Ncv8u)rectY; 89 ((NcvRect8u*)&(this->_ui2.x))->width = (Ncv8u)rectWidth; 90 ((NcvRect8u*)&(this->_ui2.x))->height = (Ncv8u)rectHeight; 91 return NCV_SUCCESS; 92 } 93 94 __host__ NCVStatus setWeight(Ncv32f weight) 95 { 96 ((Ncv32f_a*)&(this->_ui2.y))[0] = weight; 97 return NCV_SUCCESS; 98 } 99 100 __device__ __host__ void getRect(Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight) 101 { 102 NcvRect8u tmpRect = *(NcvRect8u*)(&this->_ui2.x); 103 *rectX = tmpRect.x; 104 *rectY = tmpRect.y; 105 *rectWidth = tmpRect.width; 106 *rectHeight = tmpRect.height; 107 } 108 109 __device__ __host__ Ncv32f getWeight(void) 110 { 111 return *(Ncv32f_a*)(&this->_ui2.y); 112 } 113 }; 114 115 116 struct HaarFeatureDescriptor32 117 { 118 private: 119 120 #define HaarFeatureDescriptor32_Interpret_MaskFlagTilted 0x80000000 121 #define HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf 0x40000000 122 #define HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf 0x20000000 123 #define HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures 0x1F 124 #define HaarFeatureDescriptor32_NumFeatures_Shift 24 125 #define HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset 0x00FFFFFF 126 127 Ncv32u desc; 128 129 public: 130 131 __host__ NCVStatus create(NcvBool bTilted, NcvBool bLeftLeaf, NcvBool bRightLeaf, 132 Ncv32u numFeatures, Ncv32u offsetFeatures) 133 { 134 if (numFeatures > HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures) 135 { 136 return NCV_HAAR_TOO_MANY_FEATURES_IN_CLASSIFIER; 137 } 138 if (offsetFeatures > HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset) 139 { 140 return NCV_HAAR_TOO_MANY_FEATURES_IN_CASCADE; 141 } 142 this->desc = 0; 143 this->desc |= (bTilted ? HaarFeatureDescriptor32_Interpret_MaskFlagTilted : 0); 144 this->desc |= (bLeftLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf : 0); 145 this->desc |= (bRightLeaf ? HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf : 0); 146 this->desc |= (numFeatures << HaarFeatureDescriptor32_NumFeatures_Shift); 147 this->desc |= offsetFeatures; 148 return NCV_SUCCESS; 149 } 150 151 __device__ __host__ NcvBool isTilted(void) 152 { 153 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagTilted) != 0; 154 } 155 156 __device__ __host__ NcvBool isLeftNodeLeaf(void) 157 { 158 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagLeftNodeLeaf) != 0; 159 } 160 161 __device__ __host__ NcvBool isRightNodeLeaf(void) 162 { 163 return (this->desc & HaarFeatureDescriptor32_Interpret_MaskFlagRightNodeLeaf) != 0; 164 } 165 166 __device__ __host__ Ncv32u getNumFeatures(void) 167 { 168 return (this->desc >> HaarFeatureDescriptor32_NumFeatures_Shift) & HaarFeatureDescriptor32_CreateCheck_MaxNumFeatures; 169 } 170 171 __device__ __host__ Ncv32u getFeaturesOffset(void) 172 { 173 return this->desc & HaarFeatureDescriptor32_CreateCheck_MaxFeatureOffset; 174 } 175 }; 176 177 struct HaarClassifierNodeDescriptor32 178 { 179 uint1 _ui1; 180 181 __host__ NCVStatus create(Ncv32f leafValue) 182 { 183 *(Ncv32f_a *)&this->_ui1 = leafValue; 184 return NCV_SUCCESS; 185 } 186 187 __host__ NCVStatus create(Ncv32u offsetHaarClassifierNode) 188 { 189 this->_ui1.x = offsetHaarClassifierNode; 190 return NCV_SUCCESS; 191 } 192 193 __host__ Ncv32f getLeafValueHost(void) 194 { 195 return *(Ncv32f_a *)&this->_ui1.x; 196 } 197 198 #ifdef __CUDACC__ 199 __device__ Ncv32f getLeafValue(void) 200 { 201 return __int_as_float(this->_ui1.x); 202 } 203 #endif 204 205 __device__ __host__ Ncv32u getNextNodeOffset(void) 206 { 207 return this->_ui1.x; 208 } 209 }; 210 211 #if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__ > 4 212 typedef Ncv32u __attribute__((__may_alias__)) Ncv32u_a; 213 #else 214 typedef Ncv32u Ncv32u_a; 215 #endif 216 217 struct HaarClassifierNode128 218 { 219 uint4 _ui4; 220 221 __host__ NCVStatus setFeatureDesc(HaarFeatureDescriptor32 f) 222 { 223 this->_ui4.x = *(Ncv32u *)&f; 224 return NCV_SUCCESS; 225 } 226 227 __host__ NCVStatus setThreshold(Ncv32f t) 228 { 229 this->_ui4.y = *(Ncv32u_a *)&t; 230 return NCV_SUCCESS; 231 } 232 233 __host__ NCVStatus setLeftNodeDesc(HaarClassifierNodeDescriptor32 nl) 234 { 235 this->_ui4.z = *(Ncv32u_a *)&nl; 236 return NCV_SUCCESS; 237 } 238 239 __host__ NCVStatus setRightNodeDesc(HaarClassifierNodeDescriptor32 nr) 240 { 241 this->_ui4.w = *(Ncv32u_a *)&nr; 242 return NCV_SUCCESS; 243 } 244 245 __host__ __device__ HaarFeatureDescriptor32 getFeatureDesc(void) 246 { 247 return *(HaarFeatureDescriptor32 *)&this->_ui4.x; 248 } 249 250 __host__ __device__ Ncv32f getThreshold(void) 251 { 252 return *(Ncv32f_a*)&this->_ui4.y; 253 } 254 255 __host__ __device__ HaarClassifierNodeDescriptor32 getLeftNodeDesc(void) 256 { 257 return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.z; 258 } 259 260 __host__ __device__ HaarClassifierNodeDescriptor32 getRightNodeDesc(void) 261 { 262 return *(HaarClassifierNodeDescriptor32 *)&this->_ui4.w; 263 } 264 }; 265 266 267 struct HaarStage64 268 { 269 #define HaarStage64_Interpret_MaskRootNodes 0x0000FFFF 270 #define HaarStage64_Interpret_MaskRootNodeOffset 0xFFFF0000 271 #define HaarStage64_Interpret_ShiftRootNodeOffset 16 272 273 uint2 _ui2; 274 275 __host__ NCVStatus setStageThreshold(Ncv32f t) 276 { 277 this->_ui2.x = *(Ncv32u_a *)&t; 278 return NCV_SUCCESS; 279 } 280 281 __host__ NCVStatus setStartClassifierRootNodeOffset(Ncv32u val) 282 { 283 if (val > (HaarStage64_Interpret_MaskRootNodeOffset >> HaarStage64_Interpret_ShiftRootNodeOffset)) 284 { 285 return NCV_HAAR_XML_LOADING_EXCEPTION; 286 } 287 this->_ui2.y = (val << HaarStage64_Interpret_ShiftRootNodeOffset) | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes); 288 return NCV_SUCCESS; 289 } 290 291 __host__ NCVStatus setNumClassifierRootNodes(Ncv32u val) 292 { 293 if (val > HaarStage64_Interpret_MaskRootNodes) 294 { 295 return NCV_HAAR_XML_LOADING_EXCEPTION; 296 } 297 this->_ui2.y = val | (this->_ui2.y & HaarStage64_Interpret_MaskRootNodeOffset); 298 return NCV_SUCCESS; 299 } 300 301 __host__ __device__ Ncv32f getStageThreshold(void) 302 { 303 return *(Ncv32f_a*)&this->_ui2.x; 304 } 305 306 __host__ __device__ Ncv32u getStartClassifierRootNodeOffset(void) 307 { 308 return (this->_ui2.y >> HaarStage64_Interpret_ShiftRootNodeOffset); 309 } 310 311 __host__ __device__ Ncv32u getNumClassifierRootNodes(void) 312 { 313 return (this->_ui2.y & HaarStage64_Interpret_MaskRootNodes); 314 } 315 }; 316 317 318 NCV_CT_ASSERT(sizeof(HaarFeature64) == 8); 319 NCV_CT_ASSERT(sizeof(HaarFeatureDescriptor32) == 4); 320 NCV_CT_ASSERT(sizeof(HaarClassifierNodeDescriptor32) == 4); 321 NCV_CT_ASSERT(sizeof(HaarClassifierNode128) == 16); 322 NCV_CT_ASSERT(sizeof(HaarStage64) == 8); 323 324 325 //============================================================================== 326 // 327 // Classifier cascade descriptor 328 // 329 //============================================================================== 330 331 332 struct HaarClassifierCascadeDescriptor 333 { 334 Ncv32u NumStages; 335 Ncv32u NumClassifierRootNodes; 336 Ncv32u NumClassifierTotalNodes; 337 Ncv32u NumFeatures; 338 NcvSize32u ClassifierSize; 339 NcvBool bNeedsTiltedII; 340 NcvBool bHasStumpsOnly; 341 }; 342 343 344 //============================================================================== 345 // 346 // Functional interface 347 // 348 //============================================================================== 349 350 351 enum 352 { 353 NCVPipeObjDet_Default = 0x000, 354 NCVPipeObjDet_UseFairImageScaling = 0x001, 355 NCVPipeObjDet_FindLargestObject = 0x002, 356 NCVPipeObjDet_VisualizeInPlace = 0x004, 357 }; 358 359 360 CV_EXPORTS NCVStatus ncvDetectObjectsMultiScale_device(NCVMatrix<Ncv8u> &d_srcImg, 361 NcvSize32u srcRoi, 362 NCVVector<NcvRect32u> &d_dstRects, 363 Ncv32u &dstNumRects, 364 365 HaarClassifierCascadeDescriptor &haar, 366 NCVVector<HaarStage64> &h_HaarStages, 367 NCVVector<HaarStage64> &d_HaarStages, 368 NCVVector<HaarClassifierNode128> &d_HaarNodes, 369 NCVVector<HaarFeature64> &d_HaarFeatures, 370 371 NcvSize32u minObjSize, 372 Ncv32u minNeighbors, //default 4 373 Ncv32f scaleStep, //default 1.2f 374 Ncv32u pixelStep, //default 1 375 Ncv32u flags, //default NCVPipeObjDet_Default 376 377 INCVMemAllocator &gpuAllocator, 378 INCVMemAllocator &cpuAllocator, 379 cudaDeviceProp &devProp, 380 cudaStream_t cuStream); 381 382 383 #define OBJDET_MASK_ELEMENT_INVALID_32U 0xFFFFFFFF 384 #define HAAR_STDDEV_BORDER 1 385 386 387 CV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix<Ncv32u> &d_integralImage, 388 NCVMatrix<Ncv32f> &d_weights, 389 NCVMatrixAlloc<Ncv32u> &d_pixelMask, 390 Ncv32u &numDetections, 391 HaarClassifierCascadeDescriptor &haar, 392 NCVVector<HaarStage64> &h_HaarStages, 393 NCVVector<HaarStage64> &d_HaarStages, 394 NCVVector<HaarClassifierNode128> &d_HaarNodes, 395 NCVVector<HaarFeature64> &d_HaarFeatures, 396 NcvBool bMaskElements, 397 NcvSize32u anchorsRoi, 398 Ncv32u pixelStep, 399 Ncv32f scaleArea, 400 INCVMemAllocator &gpuAllocator, 401 INCVMemAllocator &cpuAllocator, 402 cudaDeviceProp &devProp, 403 cudaStream_t cuStream); 404 405 406 CV_EXPORTS NCVStatus ncvApplyHaarClassifierCascade_host(NCVMatrix<Ncv32u> &h_integralImage, 407 NCVMatrix<Ncv32f> &h_weights, 408 NCVMatrixAlloc<Ncv32u> &h_pixelMask, 409 Ncv32u &numDetections, 410 HaarClassifierCascadeDescriptor &haar, 411 NCVVector<HaarStage64> &h_HaarStages, 412 NCVVector<HaarClassifierNode128> &h_HaarNodes, 413 NCVVector<HaarFeature64> &h_HaarFeatures, 414 NcvBool bMaskElements, 415 NcvSize32u anchorsRoi, 416 Ncv32u pixelStep, 417 Ncv32f scaleArea); 418 419 420 #define RECT_SIMILARITY_PROPORTION 0.2f 421 422 423 CV_EXPORTS NCVStatus ncvGrowDetectionsVector_device(NCVVector<Ncv32u> &pixelMask, 424 Ncv32u numPixelMaskDetections, 425 NCVVector<NcvRect32u> &hypotheses, 426 Ncv32u &totalDetections, 427 Ncv32u totalMaxDetections, 428 Ncv32u rectWidth, 429 Ncv32u rectHeight, 430 Ncv32f curScale, 431 cudaStream_t cuStream); 432 433 434 CV_EXPORTS NCVStatus ncvGrowDetectionsVector_host(NCVVector<Ncv32u> &pixelMask, 435 Ncv32u numPixelMaskDetections, 436 NCVVector<NcvRect32u> &hypotheses, 437 Ncv32u &totalDetections, 438 Ncv32u totalMaxDetections, 439 Ncv32u rectWidth, 440 Ncv32u rectHeight, 441 Ncv32f curScale); 442 443 444 CV_EXPORTS NCVStatus ncvHaarGetClassifierSize(const cv::String &filename, Ncv32u &numStages, 445 Ncv32u &numNodes, Ncv32u &numFeatures); 446 447 448 CV_EXPORTS NCVStatus ncvHaarLoadFromFile_host(const cv::String &filename, 449 HaarClassifierCascadeDescriptor &haar, 450 NCVVector<HaarStage64> &h_HaarStages, 451 NCVVector<HaarClassifierNode128> &h_HaarNodes, 452 NCVVector<HaarFeature64> &h_HaarFeatures); 453 454 455 CV_EXPORTS NCVStatus ncvHaarStoreNVBIN_host(const cv::String &filename, 456 HaarClassifierCascadeDescriptor haar, 457 NCVVector<HaarStage64> &h_HaarStages, 458 NCVVector<HaarClassifierNode128> &h_HaarNodes, 459 NCVVector<HaarFeature64> &h_HaarFeatures); 460 461 //! @} 462 463 #endif // _ncvhaarobjectdetection_hpp_ 464
