1 /*M///////////////////////////////////////////////////////////////////////////////////////
\r
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
\r
5 // By downloading, copying, installing or using the software you agree to this license.
\r
6 // If you do not agree to this license, do not download, install,
\r
7 // copy or use the software.
\r
10 // License Agreement
\r
11 // For Open Source Computer Vision Library
\r
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
\r
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
\r
15 // Third party copyrights are property of their respective owners.
\r
17 // Redistribution and use in source and binary forms, with or without modification,
\r
18 // are permitted provided that the following conditions are met:
\r
20 // * Redistribution's of source code must retain the above copyright notice,
\r
21 // this list of conditions and the following disclaimer.
\r
23 // * Redistribution's in binary form must reproduce the above copyright notice,
\r
24 // this list of conditions and the following disclaimer in the documentation
\r
25 // and/or other materials provided with the distribution.
\r
27 // * The name of the copyright holders may not be used to endorse or promote products
\r
28 // derived from this software without specific prior written permission.
\r
30 // This software is provided by the copyright holders and contributors "as is" and
\r
31 // any express or bpied warranties, including, but not limited to, the bpied
\r
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
\r
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
\r
34 // indirect, incidental, special, exemplary, or consequential damages
\r
35 // (including, but not limited to, procurement of substitute goods or services;
\r
36 // loss of use, data, or profits; or business interruption) however caused
\r
37 // and on any theory of liability, whether in contract, strict liability,
\r
38 // or tort (including negligence or otherwise) arising in any way out of
\r
39 // the use of this software, even if advised of the possibility of such damage.
\r
43 #include "internal_shared.hpp"
\r
44 #include "opencv2/gpu/device/limits.hpp"
\r
45 #include "opencv2/gpu/device/vec_distance.hpp"
\r
47 using namespace cv::gpu;
\r
48 using namespace cv::gpu::device;
\r
50 namespace cv { namespace gpu { namespace bf_knnmatch
\r
52 template <typename VecDiff, typename Dist, typename T, typename Mask>
\r
53 __device__ void distanceCalcLoop(const PtrStep_<T>& query, const DevMem2D_<T>& train, const Mask& m, int queryIdx,
\r
54 typename Dist::result_type& distMin1, typename Dist::result_type& distMin2, int& bestTrainIdx1, int& bestTrainIdx2,
\r
55 typename Dist::result_type* smem)
\r
57 const VecDiff vecDiff(query.ptr(queryIdx), train.cols, (typename Dist::value_type*)smem, threadIdx.y * blockDim.x + threadIdx.x, threadIdx.x);
\r
59 typename Dist::result_type* sdiffRow = smem + blockDim.x * threadIdx.y;
\r
61 distMin1 = numeric_limits<typename Dist::result_type>::max();
\r
62 distMin2 = numeric_limits<typename Dist::result_type>::max();
\r
67 for (int trainIdx = threadIdx.y; trainIdx < train.rows; trainIdx += blockDim.y)
\r
69 if (m(queryIdx, trainIdx))
\r
73 const T* trainRow = train.ptr(trainIdx);
\r
75 vecDiff.calc(trainRow, train.cols, dist, sdiffRow, threadIdx.x);
\r
77 const typename Dist::result_type val = dist;
\r
82 bestTrainIdx1 = trainIdx;
\r
84 else if (val < distMin2)
\r
87 bestTrainIdx2 = trainIdx;
\r
93 template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename VecDiff, typename Dist, typename T, typename Mask>
\r
94 __global__ void knnMatch2(const PtrStep_<T> query, const DevMem2D_<T> train, const Mask m, int2* trainIdx, float2* distance)
\r
96 typedef typename Dist::result_type result_type;
\r
97 typedef typename Dist::value_type value_type;
\r
99 __shared__ result_type smem[BLOCK_DIM_X * BLOCK_DIM_Y];
\r
101 const int queryIdx = blockIdx.x;
\r
103 result_type distMin1;
\r
104 result_type distMin2;
\r
109 distanceCalcLoop<VecDiff, Dist>(query, train, m, queryIdx, distMin1, distMin2, bestTrainIdx1, bestTrainIdx2, smem);
\r
112 volatile result_type* sdistMinRow = smem;
\r
113 volatile int* sbestTrainIdxRow = (int*)(sdistMinRow + 2 * BLOCK_DIM_Y);
\r
115 if (threadIdx.x == 0)
\r
117 sdistMinRow[threadIdx.y] = distMin1;
\r
118 sdistMinRow[threadIdx.y + BLOCK_DIM_Y] = distMin2;
\r
120 sbestTrainIdxRow[threadIdx.y] = bestTrainIdx1;
\r
121 sbestTrainIdxRow[threadIdx.y + BLOCK_DIM_Y] = bestTrainIdx2;
\r
125 if (threadIdx.x == 0 && threadIdx.y == 0)
\r
127 distMin1 = numeric_limits<result_type>::max();
\r
128 distMin2 = numeric_limits<result_type>::max();
\r
130 bestTrainIdx1 = -1;
\r
131 bestTrainIdx2 = -1;
\r
134 for (int i = 0; i < BLOCK_DIM_Y; ++i)
\r
136 result_type val = sdistMinRow[i];
\r
138 if (val < distMin1)
\r
141 bestTrainIdx1 = sbestTrainIdxRow[i];
\r
143 else if (val < distMin2)
\r
146 bestTrainIdx2 = sbestTrainIdxRow[i];
\r
151 for (int i = BLOCK_DIM_Y; i < 2 * BLOCK_DIM_Y; ++i)
\r
153 result_type val = sdistMinRow[i];
\r
155 if (val < distMin2)
\r
158 bestTrainIdx2 = sbestTrainIdxRow[i];
\r
162 trainIdx[queryIdx] = make_int2(bestTrainIdx1, bestTrainIdx2);
\r
163 distance[queryIdx] = make_float2(distMin1, distMin2);
\r
167 ///////////////////////////////////////////////////////////////////////////////
\r
168 // Knn 2 Match kernel caller
\r
170 template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>
\r
171 void knnMatch2Simple_caller(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask,
\r
172 const DevMem2D_<int2>& trainIdx, const DevMem2D_<float2>& distance,
\r
173 cudaStream_t stream)
\r
175 const dim3 grid(query.rows, 1, 1);
\r
176 const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);
\r
178 knnMatch2<BLOCK_DIM_X, BLOCK_DIM_Y, VecDiffGlobal<BLOCK_DIM_X, T>, Dist, T>
\r
179 <<<grid, threads, 0, stream>>>(query, train, mask, trainIdx, distance);
\r
180 cudaSafeCall( cudaGetLastError() );
\r
183 cudaSafeCall( cudaDeviceSynchronize() );
\r
186 template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int MAX_LEN, bool LEN_EQ_MAX_LEN, typename Dist, typename T, typename Mask>
\r
187 void knnMatch2Cached_caller(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask,
\r
188 const DevMem2D_<int2>& trainIdx, const DevMem2D_<float2>& distance,
\r
189 cudaStream_t stream)
\r
191 StaticAssert<BLOCK_DIM_X * BLOCK_DIM_Y >= MAX_LEN>::check(); // block size must be greter than descriptors length
\r
192 StaticAssert<MAX_LEN % BLOCK_DIM_X == 0>::check(); // max descriptors length must divide to blockDimX
\r
194 const dim3 grid(query.rows, 1, 1);
\r
195 const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);
\r
197 knnMatch2<BLOCK_DIM_X, BLOCK_DIM_Y, VecDiffCachedRegister<BLOCK_DIM_X, MAX_LEN, LEN_EQ_MAX_LEN, typename Dist::value_type>, Dist, T>
\r
198 <<<grid, threads, 0, stream>>>(query, train, mask, trainIdx.data, distance.data);
\r
199 cudaSafeCall( cudaGetLastError() );
\r
202 cudaSafeCall( cudaDeviceSynchronize() );
\r
205 ///////////////////////////////////////////////////////////////////////////////
\r
206 // Knn 2 Match Dispatcher
\r
208 template <typename Dist, typename T, typename Mask>
\r
209 void knnMatch2Dispatcher(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask,
\r
210 const DevMem2D& trainIdx, const DevMem2D& distance,
\r
211 int cc, cudaStream_t stream)
\r
213 if (query.cols < 64)
\r
215 knnMatch2Cached_caller<16, 16, 64, false, Dist>(
\r
216 query, train, mask,
\r
217 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
220 else if (query.cols == 64)
\r
222 knnMatch2Cached_caller<16, 16, 64, true, Dist>(
\r
223 query, train, mask,
\r
224 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
227 else if (query.cols < 128)
\r
229 knnMatch2Cached_caller<16, 16, 128, false, Dist>(
\r
230 query, train, mask,
\r
231 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
234 else if (query.cols == 128 && cc >= 12)
\r
236 knnMatch2Cached_caller<16, 16, 128, true, Dist>(
\r
237 query, train, mask,
\r
238 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
241 else if (query.cols < 256 && cc >= 12)
\r
243 knnMatch2Cached_caller<16, 16, 256, false, Dist>(
\r
244 query, train, mask,
\r
245 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
248 else if (query.cols == 256 && cc >= 12)
\r
250 knnMatch2Cached_caller<16, 16, 256, true, Dist>(
\r
251 query, train, mask,
\r
252 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
257 knnMatch2Simple_caller<16, 16, Dist>(
\r
258 query, train, mask,
\r
259 static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> >(distance),
\r
264 ///////////////////////////////////////////////////////////////////////////////
\r
265 // Calc distance kernel
\r
267 template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>
\r
268 __global__ void calcDistance(const PtrStep_<T> query, const DevMem2D_<T> train, const Mask mask, PtrStepf distance)
\r
270 __shared__ typename Dist::result_type sdiff[BLOCK_DIM_X * BLOCK_DIM_Y];
\r
272 typename Dist::result_type* sdiff_row = sdiff + BLOCK_DIM_X * threadIdx.y;
\r
274 const int queryIdx = blockIdx.x;
\r
275 const T* queryDescs = query.ptr(queryIdx);
\r
277 const int trainIdx = blockIdx.y * BLOCK_DIM_Y + threadIdx.y;
\r
279 if (trainIdx < train.rows)
\r
281 const T* trainDescs = train.ptr(trainIdx);
\r
283 typename Dist::result_type myDist = numeric_limits<typename Dist::result_type>::max();
\r
285 if (mask(queryIdx, trainIdx))
\r
289 calcVecDiffGlobal<BLOCK_DIM_X>(queryDescs, trainDescs, train.cols, dist, sdiff_row, threadIdx.x);
\r
294 if (threadIdx.x == 0)
\r
295 distance.ptr(queryIdx)[trainIdx] = myDist;
\r
299 ///////////////////////////////////////////////////////////////////////////////
\r
300 // Calc distance kernel caller
\r
302 template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>
\r
303 void calcDistance_caller(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask, const DevMem2Df& distance, cudaStream_t stream)
\r
305 const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);
\r
306 const dim3 grid(query.rows, divUp(train.rows, BLOCK_DIM_Y), 1);
\r
308 calcDistance<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, T><<<grid, threads, 0, stream>>>(query, train, mask, distance);
\r
309 cudaSafeCall( cudaGetLastError() );
\r
312 cudaSafeCall( cudaDeviceSynchronize() );
\r
315 template <typename Dist, typename T, typename Mask>
\r
316 void calcDistanceDispatcher(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask, const DevMem2D& allDist, cudaStream_t stream)
\r
318 calcDistance_caller<16, 16, Dist>(query, train, mask, static_cast<DevMem2Df>(allDist), stream);
\r
321 ///////////////////////////////////////////////////////////////////////////////
\r
322 // find knn match kernel
\r
324 template <int BLOCK_SIZE> __global__ void findBestMatch(DevMem2Df allDist_, int i, PtrStepi trainIdx_, PtrStepf distance_)
\r
326 const int SMEM_SIZE = BLOCK_SIZE > 64 ? BLOCK_SIZE : 64;
\r
327 __shared__ float sdist[SMEM_SIZE];
\r
328 __shared__ int strainIdx[SMEM_SIZE];
\r
330 const int queryIdx = blockIdx.x;
\r
332 float* allDist = allDist_.ptr(queryIdx);
\r
333 int* trainIdx = trainIdx_.ptr(queryIdx);
\r
334 float* distance = distance_.ptr(queryIdx);
\r
336 float dist = numeric_limits<float>::max();
\r
339 for (int i = threadIdx.x; i < allDist_.cols; i += BLOCK_SIZE)
\r
341 float reg = allDist[i];
\r
349 sdist[threadIdx.x] = dist;
\r
350 strainIdx[threadIdx.x] = bestIdx;
\r
353 reducePredVal<BLOCK_SIZE>(sdist, dist, strainIdx, bestIdx, threadIdx.x, less<volatile float>());
\r
355 if (threadIdx.x == 0)
\r
357 if (dist < numeric_limits<float>::max())
\r
359 allDist[bestIdx] = numeric_limits<float>::max();
\r
360 trainIdx[i] = bestIdx;
\r
361 distance[i] = dist;
\r
366 ///////////////////////////////////////////////////////////////////////////////
\r
367 // find knn match kernel caller
\r
369 template <int BLOCK_SIZE> void findKnnMatch_caller(int k, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream)
\r
371 const dim3 threads(BLOCK_SIZE, 1, 1);
\r
372 const dim3 grid(trainIdx.rows, 1, 1);
\r
374 for (int i = 0; i < k; ++i)
\r
376 findBestMatch<BLOCK_SIZE><<<grid, threads, 0, stream>>>(allDist, i, trainIdx, distance);
\r
377 cudaSafeCall( cudaGetLastError() );
\r
381 cudaSafeCall( cudaDeviceSynchronize() );
\r
384 void findKnnMatchDispatcher(int k, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, cudaStream_t stream)
\r
386 findKnnMatch_caller<256>(k, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), static_cast<DevMem2Df>(allDist), stream);
\r
389 ///////////////////////////////////////////////////////////////////////////////
\r
390 // knn match Dispatcher
\r
392 template <typename Dist, typename T>
\r
393 void knnMatchDispatcher(const DevMem2D_<T>& query, const DevMem2D_<T>& train, int k, const DevMem2D& mask,
\r
394 const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist,
\r
395 int cc, cudaStream_t stream)
\r
401 knnMatch2Dispatcher<Dist>(query, train, SingleMask(mask), trainIdx, distance, cc, stream);
\r
405 calcDistanceDispatcher<Dist>(query, train, SingleMask(mask), allDist, stream);
\r
411 knnMatch2Dispatcher<Dist>(query, train, WithOutMask(), trainIdx, distance, cc, stream);
\r
415 calcDistanceDispatcher<Dist>(query, train, WithOutMask(), allDist, stream);
\r
418 findKnnMatchDispatcher(k, trainIdx, distance, allDist, stream);
\r
421 ///////////////////////////////////////////////////////////////////////////////
\r
422 // knn match caller
\r
424 template <typename T> void knnMatchL1_gpu(const DevMem2D& query, const DevMem2D& train, int k, const DevMem2D& mask,
\r
425 const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist,
\r
426 int cc, cudaStream_t stream)
\r
428 knnMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);
\r
431 template void knnMatchL1_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
432 //template void knnMatchL1_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
433 template void knnMatchL1_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
434 template void knnMatchL1_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
435 template void knnMatchL1_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
436 template void knnMatchL1_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
438 template <typename T> void knnMatchL2_gpu(const DevMem2D& query, const DevMem2D& train, int k, const DevMem2D& mask,
\r
439 const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist,
\r
440 int cc, cudaStream_t stream)
\r
442 knnMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);
\r
445 //template void knnMatchL2_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
446 //template void knnMatchL2_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
447 //template void knnMatchL2_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
448 //template void knnMatchL2_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
449 //template void knnMatchL2_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
450 template void knnMatchL2_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
452 template <typename T> void knnMatchHamming_gpu(const DevMem2D& query, const DevMem2D& train, int k, const DevMem2D& mask,
\r
453 const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist,
\r
454 int cc, cudaStream_t stream)
\r
456 knnMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);
\r
459 template void knnMatchHamming_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
460 //template void knnMatchHamming_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
461 template void knnMatchHamming_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
462 //template void knnMatchHamming_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r
463 template void knnMatchHamming_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int k, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, int cc, cudaStream_t stream);
\r