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 #ifndef __OPENCV_CUDA_DEVICE_LBP_HPP_ 44 #define __OPENCV_CUDA_DEVICE_LBP_HPP_ 45 46 #include "opencv2/core/cuda/common.hpp" 47 #include "opencv2/core/cuda/emulation.hpp" 48 49 namespace cv { namespace cuda { 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: InSameComponintcv::cuda::device::lbp::InSameComponint70 __device__ __forceinline__ InSameComponint(float _eps) : eps(_eps) {} InSameComponintcv::cuda::device::lbp::InSameComponint71 __device__ __forceinline__ InSameComponint(const InSameComponint& other) : eps(other.eps) {} 72 operator ()cv::cuda::device::lbp::InSameComponint73 __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> partition(int4 * vec,unsigned int n,int * labels,Pr predicate)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 113