Home | History | Annotate | Download | only in cudalegacy
      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