modules/gpu/src/cuda/lbp.hpp

   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 bpied warranties, including, but not limited to, the bpied
  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 #ifndef __OPENCV_GPU_DEVICE_LBP_HPP_
  44 #define __OPENCV_GPU_DEVICE_LBP_HPP_
  45
  46 #include "internal_shared.hpp"
  47 #include <opencv2/gpu/device/emulation.hpp>
  48
  49 namespace cv { namespace gpu { namespace device {
  50
  51 namespace lbp {
  52
  53     struct Stage
  54     {
  55         int    first;
  56         int    ntrees;
  57         float  threshold;
  58     };
  59
  60     struct ClNode
  61     {
  62         int   left;
  63         int   right;
  64         int   featureIdx;
  65     };
  66
  67     struct InSameComponint
  68     {
  69     public:
  70         __device__ __forceinline__ InSameComponint(float _eps) : eps(_eps) {}
  71         __device__ __forceinline__ InSameComponint(const InSameComponint& other) : eps(other.eps) {}
  72
  73         __device__ __forceinline__ bool operator()(const int4& r1, const int4& r2) const
  74         {
  75             float delta = eps * (min(r1.z, r2.z) + min(r1.w, r2.w)) * 0.5f;
  76
  77             return abs(r1.x - r2.x) <= delta && abs(r1.y - r2.y) <= delta
  78                 && abs(r1.x + r1.z - r2.x - r2.z) <= delta && abs(r1.y + r1.w - r2.y - r2.w) <= delta;
  79         }
  80         float eps;
  81     };
  82
  83     template<typename Pr>
  84     __device__ __forceinline__ void partition(int4* vec, unsigned int n, int* labels, Pr predicate)
  85     {
  86         unsigned tid = threadIdx.x;
  87         labels[tid] = tid;
  88         __syncthreads();
  89         for (unsigned int id = 0; id < n; id++)
  90         {
  91             if (tid != id && predicate(vec[tid], vec[id]))
  92             {
  93                 int p = labels[tid];
  94                 int q = labels[id];
  95                 if (p < q)
  96                 {
  97                     Emulation::smem::atomicMin(labels + id, p);
  98                 }
  99                 else if (p > q)
 100                 {
 101                     Emulation::smem::atomicMin(labels + tid, q);
 102                 }
 103             }
 104         }
 105         __syncthreads();
 106     }
 107
 108 } // lbp
 109
 110 } } }// namespaces
 111
 112 #endif