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