vec_distance.hpp
Go to the documentation of this file.
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_GPU_VEC_DISTANCE_HPP__
44 #define __OPENCV_GPU_VEC_DISTANCE_HPP__
45 
46 #include "reduce.hpp"
47 #include "functional.hpp"
49 
50 namespace cv { namespace gpu { namespace device
51 {
52  template <typename T> struct L1Dist
53  {
54  typedef int value_type;
55  typedef int result_type;
56 
57  __device__ __forceinline__ L1Dist() : mySum(0) {}
58 
59  __device__ __forceinline__ void reduceIter(int val1, int val2)
60  {
61  mySum = __sad(val1, val2, mySum);
62  }
63 
64  template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(int* smem, int tid)
65  {
66  reduce<THREAD_DIM>(smem, mySum, tid, plus<int>());
67  }
68 
69  __device__ __forceinline__ operator int() const
70  {
71  return mySum;
72  }
73 
74  int mySum;
75  };
76  template <> struct L1Dist<float>
77  {
78  typedef float value_type;
79  typedef float result_type;
80 
81  __device__ __forceinline__ L1Dist() : mySum(0.0f) {}
82 
83  __device__ __forceinline__ void reduceIter(float val1, float val2)
84  {
85  mySum += ::fabs(val1 - val2);
86  }
87 
88  template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(float* smem, int tid)
89  {
90  reduce<THREAD_DIM>(smem, mySum, tid, plus<float>());
91  }
92 
93  __device__ __forceinline__ operator float() const
94  {
95  return mySum;
96  }
97 
98  float mySum;
99  };
100 
101  struct L2Dist
102  {
103  typedef float value_type;
104  typedef float result_type;
105 
106  __device__ __forceinline__ L2Dist() : mySum(0.0f) {}
107 
108  __device__ __forceinline__ void reduceIter(float val1, float val2)
109  {
110  float reg = val1 - val2;
111  mySum += reg * reg;
112  }
113 
114  template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(float* smem, int tid)
115  {
116  reduce<THREAD_DIM>(smem, mySum, tid, plus<float>());
117  }
118 
119  __device__ __forceinline__ operator float() const
120  {
121  return sqrtf(mySum);
122  }
123 
124  float mySum;
125  };
126 
127  struct HammingDist
128  {
129  typedef int value_type;
130  typedef int result_type;
131 
132  __device__ __forceinline__ HammingDist() : mySum(0) {}
133 
134  __device__ __forceinline__ void reduceIter(int val1, int val2)
135  {
136  mySum += __popc(val1 ^ val2);
137  }
138 
139  template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(int* smem, int tid)
140  {
141  reduce<THREAD_DIM>(smem, mySum, tid, plus<int>());
142  }
143 
144  __device__ __forceinline__ operator int() const
145  {
146  return mySum;
147  }
148 
149  int mySum;
150  };
151 
152  // calc distance between two vectors in global memory
153  template <int THREAD_DIM, typename Dist, typename T1, typename T2>
154  __device__ void calcVecDiffGlobal(const T1* vec1, const T2* vec2, int len, Dist& dist, typename Dist::result_type* smem, int tid)
155  {
156  for (int i = tid; i < len; i += THREAD_DIM)
157  {
158  T1 val1;
159  ForceGlob<T1>::Load(vec1, i, val1);
160 
161  T2 val2;
162  ForceGlob<T2>::Load(vec2, i, val2);
163 
164  dist.reduceIter(val1, val2);
165  }
166 
167  dist.reduceAll<THREAD_DIM>(smem, tid);
168  }
169 
170  // calc distance between two vectors, first vector is cached in register or shared memory, second vector is in global memory
171  template <int THREAD_DIM, int MAX_LEN, bool LEN_EQ_MAX_LEN, typename Dist, typename T1, typename T2>
172  __device__ __forceinline__ void calcVecDiffCached(const T1* vecCached, const T2* vecGlob, int len, Dist& dist, typename Dist::result_type* smem, int tid)
173  {
175 
176  dist.reduceAll<THREAD_DIM>(smem, tid);
177  }
178 
179  // calc distance between two vectors in global memory
180  template <int THREAD_DIM, typename T1> struct VecDiffGlobal
181  {
182  explicit __device__ __forceinline__ VecDiffGlobal(const T1* vec1_, int = 0, void* = 0, int = 0, int = 0)
183  {
184  vec1 = vec1_;
185  }
186 
187  template <typename T2, typename Dist>
188  __device__ __forceinline__ void calc(const T2* vec2, int len, Dist& dist, typename Dist::result_type* smem, int tid) const
189  {
190  calcVecDiffGlobal<THREAD_DIM>(vec1, vec2, len, dist, smem, tid);
191  }
192 
193  const T1* vec1;
194  };
195 
196  // calc distance between two vectors, first vector is cached in register memory, second vector is in global memory
197  template <int THREAD_DIM, int MAX_LEN, bool LEN_EQ_MAX_LEN, typename U> struct VecDiffCachedRegister
198  {
199  template <typename T1> __device__ __forceinline__ VecDiffCachedRegister(const T1* vec1, int len, U* smem, int glob_tid, int tid)
200  {
201  if (glob_tid < len)
202  smem[glob_tid] = vec1[glob_tid];
203  __syncthreads();
204 
205  U* vec1ValsPtr = vec1Vals;
206 
207  #pragma unroll
208  for (int i = tid; i < MAX_LEN; i += THREAD_DIM)
209  *vec1ValsPtr++ = smem[i];
210 
211  __syncthreads();
212  }
213 
214  template <typename T2, typename Dist>
215  __device__ __forceinline__ void calc(const T2* vec2, int len, Dist& dist, typename Dist::result_type* smem, int tid) const
216  {
217  calcVecDiffCached<THREAD_DIM, MAX_LEN, LEN_EQ_MAX_LEN>(vec1Vals, vec2, len, dist, smem, tid);
218  }
219 
220  U vec1Vals[MAX_LEN / THREAD_DIM];
221  };
222 }}} // namespace cv { namespace gpu { namespace device
223 
224 #endif // __OPENCV_GPU_VEC_DISTANCE_HPP__
__device__ __forceinline__ void reduceAll(float *smem, int tid)
Definition: vec_distance.hpp:114
__device__ __forceinline__ void reduceIter(int val1, int val2)
Definition: vec_distance.hpp:134
__device__ __forceinline__ VecDiffGlobal(const T1 *vec1_, int=0, void *=0, int=0, int=0)
Definition: vec_distance.hpp:182
const CvArr * U
Definition: core_c.h:733
__device__ __forceinline__ void reduceIter(float val1, float val2)
Definition: vec_distance.hpp:108
__device__ __forceinline__ L1Dist()
Definition: vec_distance.hpp:57
float mySum
Definition: vec_distance.hpp:124
__device__ __forceinline__ VecDiffCachedRegister(const T1 *vec1, int len, U *smem, int glob_tid, int tid)
Definition: vec_distance.hpp:199
__device__ __forceinline__ void reduceIter(float val1, float val2)
Definition: vec_distance.hpp:83
float mySum
Definition: vec_distance.hpp:98
__device__ __forceinline__ void reduceAll(int *smem, int tid)
Definition: vec_distance.hpp:139
Definition: functional.hpp:59
const T1 * vec1
Definition: vec_distance.hpp:193
const void int len
Definition: core_c.h:1568
__device__ __forceinline__ L2Dist()
Definition: vec_distance.hpp:106
float result_type
Definition: vec_distance.hpp:104
const CvArr * vec2
Definition: core_c.h:812
__device__ __forceinline__ void calc(const T2 *vec2, int len, Dist &dist, typename Dist::result_type *smem, int tid) const
Definition: vec_distance.hpp:188
int value_type
Definition: vec_distance.hpp:54
float value_type
Definition: vec_distance.hpp:78
__device__ __forceinline__ L1Dist()
Definition: vec_distance.hpp:81
Definition: vec_distance.hpp:101
__device__ __forceinline__ void calc(const T2 *vec2, int len, Dist &dist, typename Dist::result_type *smem, int tid) const
Definition: vec_distance.hpp:215
int mySum
Definition: vec_distance.hpp:74
int value_type
Definition: vec_distance.hpp:129
__device__ __forceinline__ void reduceAll(float *smem, int tid)
Definition: vec_distance.hpp:88
Definition: vec_distance.hpp:197
__device__ void calcVecDiffGlobal(const T1 *vec1, const T2 *vec2, int len, Dist &dist, typename Dist::result_type *smem, int tid)
Definition: vec_distance.hpp:154
Definition: vec_distance.hpp:180
int result_type
Definition: vec_distance.hpp:55
::max::max::max float
Definition: functional.hpp:326
Definition: vec_distance.hpp:127
GLenum GLsizei len
float result_type
Definition: vec_distance.hpp:79
::max::max int
Definition: functional.hpp:324
__device__ __forceinline__ void reduceIter(int val1, int val2)
Definition: vec_distance.hpp:59
GLclampf f
float value_type
Definition: vec_distance.hpp:103
int result_type
Definition: vec_distance.hpp:130
CvPoint3D64f double * dist
Definition: legacy.hpp:556
Definition: vec_distance.hpp:52
__device__ __forceinline__ void calcVecDiffCached(const T1 *vecCached, const T2 *vecGlob, int len, Dist &dist, typename Dist::result_type *smem, int tid)
Definition: vec_distance.hpp:172
U vec1Vals[MAX_LEN/THREAD_DIM]
Definition: vec_distance.hpp:220
__device__ __forceinline__ void reduceAll(int *smem, int tid)
Definition: vec_distance.hpp:64
__device__ __forceinline__ HammingDist()
Definition: vec_distance.hpp:132
__device__ static __forceinline__ void Load(const T *ptr, int offset, T &val)
Definition: datamov_utils.hpp:55
int mySum
Definition: vec_distance.hpp:149