modules/gpu/src/cuda/bf_knnmatch.cu

   1 /*M///////////////////////////////////////////////////////////////////////////////////////\r
   2 //\r
   3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
   4 //\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
   8 //\r
   9 //\r
  10 //                           License Agreement\r
  11 //                For Open Source Computer Vision Library\r
  12 //\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
  16 //\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
  19 //\r
  20 //   * Redistribution's of source code must retain the above copyright notice,\r
  21 //     this list of conditions and the following disclaimer.\r
  22 //\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
  26 //\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
  29 //\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
  40 //\r
  41 //M*/\r
  42 \r
  43 #include "internal_shared.hpp"\r
  44 #include "opencv2/gpu/device/limits.hpp"\r
  45 #include "opencv2/gpu/device/vec_distance.hpp"\r
  46 \r
  47 using namespace cv::gpu;\r
  48 using namespace cv::gpu::device;\r
  49 \r
  50 namespace cv { namespace gpu { namespace bfmatcher\r
  51 {\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
  56     {\r
  57         const VecDiff vecDiff(query.ptr(queryIdx), train.cols, (typename Dist::value_type*)smem, threadIdx.y * blockDim.x + threadIdx.x, threadIdx.x);\r
  58         \r
  59         typename Dist::result_type* sdiffRow = smem + blockDim.x * threadIdx.y;\r
  60         \r
  61         distMin1 = numeric_limits<typename Dist::result_type>::max();\r
  62         distMin2 = numeric_limits<typename Dist::result_type>::max();\r
  63 \r
  64         bestTrainIdx1 = -1;\r
  65         bestTrainIdx2 = -1;\r
  66 \r
  67         for (int trainIdx = threadIdx.y; trainIdx < train.rows; trainIdx += blockDim.y)\r
  68         {\r
  69             if (m(queryIdx, trainIdx))\r
  70             {\r
  71                 Dist dist;\r
  72 \r
  73                 const T* trainRow = train.ptr(trainIdx);\r
  74                 \r
  75                 vecDiff.calc(trainRow, train.cols, dist, sdiffRow, threadIdx.x);\r
  76 \r
  77                 const typename Dist::result_type val = dist;\r
  78 \r
  79                 if (val < distMin1)\r
  80                 {\r
  81                     distMin1 = val;\r
  82                     bestTrainIdx1 = trainIdx;\r
  83                 }\r
  84                 else if (val < distMin2)\r
  85                 {\r
  86                     distMin2 = val;\r
  87                     bestTrainIdx2 = trainIdx;\r
  88                 }\r
  89             }\r
  90         }\r
  91     }\r
  92 \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
  95     {\r
  96         typedef typename Dist::result_type result_type;\r
  97         typedef typename Dist::value_type value_type;\r
  98 \r
  99         __shared__ result_type smem[BLOCK_DIM_X * BLOCK_DIM_Y];\r
 100 \r
 101         const int queryIdx = blockIdx.x;\r
 102 \r
 103         result_type distMin1;\r
 104         result_type distMin2;\r
 105 \r
 106         int bestTrainIdx1;\r
 107         int bestTrainIdx2;\r
 108 \r
 109         distanceCalcLoop<VecDiff, Dist>(query, train, m, queryIdx, distMin1, distMin2, bestTrainIdx1, bestTrainIdx2, smem);\r
 110         __syncthreads();\r
 111 \r
 112         volatile result_type* sdistMinRow = smem;\r
 113         volatile int* sbestTrainIdxRow = (int*)(sdistMinRow + 2 * BLOCK_DIM_Y);\r
 114 \r
 115         if (threadIdx.x == 0)\r
 116         {\r
 117             sdistMinRow[threadIdx.y] = distMin1;\r
 118             sdistMinRow[threadIdx.y + BLOCK_DIM_Y] = distMin2;\r
 119 \r
 120             sbestTrainIdxRow[threadIdx.y] = bestTrainIdx1;            \r
 121             sbestTrainIdxRow[threadIdx.y + BLOCK_DIM_Y] = bestTrainIdx2;\r
 122         }\r
 123         __syncthreads();\r
 124 \r
 125         if (threadIdx.x == 0 && threadIdx.y == 0)\r
 126         {\r
 127             distMin1 = numeric_limits<result_type>::max();\r
 128             distMin2 = numeric_limits<result_type>::max();\r
 129 \r
 130             bestTrainIdx1 = -1;\r
 131             bestTrainIdx2 = -1;\r
 132 \r
 133             #pragma unroll\r
 134             for (int i = 0; i < BLOCK_DIM_Y; ++i)\r
 135             {\r
 136                 result_type val = sdistMinRow[i];\r
 137 \r
 138                 if (val < distMin1)\r
 139                 {\r
 140                     distMin1 = val;\r
 141                     bestTrainIdx1 = sbestTrainIdxRow[i];\r
 142                 }\r
 143                 else if (val < distMin2)\r
 144                 {\r
 145                     distMin2 = val;\r
 146                     bestTrainIdx2 = sbestTrainIdxRow[i];\r
 147                 }\r
 148             }\r
 149 \r
 150             #pragma unroll\r
 151             for (int i = BLOCK_DIM_Y; i < 2 * BLOCK_DIM_Y; ++i)\r
 152             {\r
 153                 result_type val = sdistMinRow[i];\r
 154 \r
 155                 if (val < distMin2)\r
 156                 {\r
 157                     distMin2 = val;\r
 158                     bestTrainIdx2 = sbestTrainIdxRow[i];\r
 159                 }\r
 160             }\r
 161 \r
 162             trainIdx[queryIdx] = make_int2(bestTrainIdx1, bestTrainIdx2);\r
 163             distance[queryIdx] = make_float2(distMin1, distMin2);\r
 164         }\r
 165     }\r
 166 \r
 167     ///////////////////////////////////////////////////////////////////////////////\r
 168     // Knn 2 Match kernel caller\r
 169 \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
 174     {\r
 175         const dim3 grid(query.rows, 1, 1);\r
 176         const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
 177 \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
 181 \r
 182         if (stream == 0)\r
 183             cudaSafeCall( cudaDeviceSynchronize() );\r
 184     }\r
 185 \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
 190     {\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
 193 \r
 194         const dim3 grid(query.rows, 1, 1);\r
 195         const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
 196 \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
 200 \r
 201         if (stream == 0)\r
 202             cudaSafeCall( cudaDeviceSynchronize() );\r
 203     }\r
 204 \r
 205     ///////////////////////////////////////////////////////////////////////////////\r
 206     // Knn 2 Match Dispatcher\r
 207     \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
 212     {\r
 213         if (query.cols < 64)\r
 214         {\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
 218                 stream);\r
 219         }\r
 220         else if (query.cols == 64)\r
 221         {\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
 225                 stream);\r
 226         }\r
 227         else if (query.cols < 128)\r
 228         {\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
 232                 stream);\r
 233         }\r
 234         else if (query.cols == 128 && cc >= 12)\r
 235         {\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
 239                 stream);\r
 240         }\r
 241         else if (query.cols < 256 && cc >= 12)\r
 242         {\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
 246                 stream);\r
 247         }\r
 248         else if (query.cols == 256 && cc >= 12)\r
 249         {\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
 253                 stream);\r
 254         }\r
 255         else\r
 256         {\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
 260                 stream);\r
 261         }\r
 262     }\r
 263     \r
 264     ///////////////////////////////////////////////////////////////////////////////\r
 265     // Calc distance kernel\r
 266 \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
 269     {\r
 270         __shared__ typename Dist::result_type sdiff[BLOCK_DIM_X * BLOCK_DIM_Y];\r
 271 \r
 272         typename Dist::result_type* sdiff_row = sdiff + BLOCK_DIM_X * threadIdx.y;\r
 273         \r
 274         const int queryIdx = blockIdx.x;\r
 275         const T* queryDescs = query.ptr(queryIdx);\r
 276 \r
 277         const int trainIdx = blockIdx.y * BLOCK_DIM_Y + threadIdx.y;\r
 278 \r
 279         if (trainIdx < train.rows)\r
 280         {\r
 281             const T* trainDescs = train.ptr(trainIdx);\r
 282 \r
 283             typename Dist::result_type myDist = numeric_limits<typename Dist::result_type>::max();\r
 284 \r
 285             if (mask(queryIdx, trainIdx))\r
 286             {\r
 287                 Dist dist;\r
 288 \r
 289                 calcVecDiffGlobal<BLOCK_DIM_X>(queryDescs, trainDescs, train.cols, dist, sdiff_row, threadIdx.x);\r
 290 \r
 291                 myDist = dist;\r
 292             }\r
 293             \r
 294             if (threadIdx.x == 0)\r
 295                 distance.ptr(queryIdx)[trainIdx] = myDist;\r
 296         }\r
 297     }\r
 298 \r
 299     ///////////////////////////////////////////////////////////////////////////////\r
 300     // Calc distance kernel caller\r
 301 \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
 304     {\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
 307 \r
 308         calcDistance<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, T><<<grid, threads, 0, stream>>>(query, train, mask, distance);\r
 309         cudaSafeCall( cudaGetLastError() );\r
 310 \r
 311         if (stream == 0)\r
 312             cudaSafeCall( cudaDeviceSynchronize() );\r
 313     }\r
 314 \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
 317     {\r
 318         calcDistance_caller<16, 16, Dist>(query, train, mask, static_cast<DevMem2Df>(allDist), stream);\r
 319     }\r
 320 \r
 321     ///////////////////////////////////////////////////////////////////////////////\r
 322     // find knn match kernel\r
 323 \r
 324     template <int BLOCK_SIZE> __global__ void findBestMatch(DevMem2Df allDist_, int i, PtrStepi trainIdx_, PtrStepf distance_)\r
 325     {\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
 329 \r
 330         const int queryIdx = blockIdx.x;\r
 331 \r
 332         float* allDist = allDist_.ptr(queryIdx);\r
 333         int* trainIdx = trainIdx_.ptr(queryIdx);\r
 334         float* distance = distance_.ptr(queryIdx);\r
 335 \r
 336         float dist = numeric_limits<float>::max();\r
 337         int bestIdx = -1;\r
 338         \r
 339         for (int i = threadIdx.x; i < allDist_.cols; i += BLOCK_SIZE)\r
 340         {\r
 341             float reg = allDist[i];\r
 342             if (reg < dist)\r
 343             {\r
 344                 dist = reg;\r
 345                 bestIdx = i;\r
 346             }\r
 347         }\r
 348 \r
 349         sdist[threadIdx.x] = dist;\r
 350         strainIdx[threadIdx.x] = bestIdx;\r
 351         __syncthreads();\r
 352 \r
 353         reducePredVal<BLOCK_SIZE>(sdist, dist, strainIdx, bestIdx, threadIdx.x, less<volatile float>());\r
 354 \r
 355         if (threadIdx.x == 0)\r
 356         {\r
 357             if (dist < numeric_limits<float>::max())\r
 358             {\r
 359                 allDist[bestIdx] = numeric_limits<float>::max();\r
 360                 trainIdx[i] = bestIdx;\r
 361                 distance[i] = dist;\r
 362             }\r
 363         }\r
 364     }\r
 365     \r
 366     ///////////////////////////////////////////////////////////////////////////////\r
 367     // find knn match kernel caller\r
 368 \r
 369     template <int BLOCK_SIZE> void findKnnMatch_caller(int k, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream)\r
 370     {\r
 371         const dim3 threads(BLOCK_SIZE, 1, 1);\r
 372         const dim3 grid(trainIdx.rows, 1, 1);\r
 373 \r
 374         for (int i = 0; i < k; ++i)\r
 375         {\r
 376             findBestMatch<BLOCK_SIZE><<<grid, threads, 0, stream>>>(allDist, i, trainIdx, distance);\r
 377             cudaSafeCall( cudaGetLastError() );\r
 378         }\r
 379 \r
 380         if (stream == 0)\r
 381             cudaSafeCall( cudaDeviceSynchronize() );\r
 382     }\r
 383 \r
 384     void findKnnMatchDispatcher(int k, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& allDist, cudaStream_t stream)\r
 385     {\r
 386         findKnnMatch_caller<256>(k, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), static_cast<DevMem2Df>(allDist), stream);\r
 387     }\r
 388     \r
 389     ///////////////////////////////////////////////////////////////////////////////\r
 390     // knn match Dispatcher\r
 391 \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
 396     {\r
 397         if (mask.data)\r
 398         {\r
 399             if (k == 2)\r
 400             {\r
 401                 knnMatch2Dispatcher<Dist>(query, train, SingleMask(mask), trainIdx, distance, cc, stream);\r
 402                 return;\r
 403             }\r
 404 \r
 405             calcDistanceDispatcher<Dist>(query, train, SingleMask(mask), allDist, stream);\r
 406         }\r
 407         else\r
 408         {\r
 409             if (k == 2)\r
 410             {\r
 411                 knnMatch2Dispatcher<Dist>(query, train, WithOutMask(), trainIdx, distance, cc, stream);\r
 412                 return;\r
 413             }\r
 414 \r
 415             calcDistanceDispatcher<Dist>(query, train, WithOutMask(), allDist, stream);\r
 416         }\r
 417 \r
 418         findKnnMatchDispatcher(k, trainIdx, distance, allDist, stream);\r
 419     }\r
 420     \r
 421     ///////////////////////////////////////////////////////////////////////////////\r
 422     // knn match caller\r
 423 \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
 427     {\r
 428         knnMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);\r
 429     }\r
 430 \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
 437 \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
 441     {\r
 442         knnMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);\r
 443     }\r
 444 \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
 451 \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
 455     {\r
 456         knnMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), k, mask, trainIdx, distance, allDist, cc, stream);\r
 457     }\r
 458 \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
 464 }}}\r