OpenCV 2.4.8 components for OpenCVgrabber.
[mmanager-3rdparty.git] / OpenCV2.4.8 / build / include / opencv2 / gpu / device / vec_distance.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 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"
48 #include "detail/vec_distance_detail.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     {
174         vec_distance_detail::VecDiffCachedCalculator<THREAD_DIM, MAX_LEN, LEN_EQ_MAX_LEN>::calc(vecCached, vecGlob, len, dist, tid);
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__