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 #if !defined CUDA_DISABLER\r
  44 \r
  45 #include "internal_shared.hpp"\r
  46 #include "opencv2/gpu/device/limits.hpp"\r
  47 #include "opencv2/gpu/device/vec_distance.hpp"\r
  48 #include "opencv2/gpu/device/datamov_utils.hpp"\r
  49 \r
  50 namespace cv { namespace gpu { namespace device\r
  51 {\r
  52     namespace bf_knnmatch\r
  53     {\r
  54         ///////////////////////////////////////////////////////////////////////////////\r
  55         // Reduction\r
  56 \r
  57         template <int BLOCK_SIZE>\r
  58         __device__ void findBestMatch(float& bestDistance1, float& bestDistance2,\r
  59                                       int& bestTrainIdx1, int& bestTrainIdx2,\r
  60                                       float* s_distance, int* s_trainIdx)\r
  61         {\r
  62             float myBestDistance1 = numeric_limits<float>::max();\r
  63             float myBestDistance2 = numeric_limits<float>::max();\r
  64             int myBestTrainIdx1 = -1;\r
  65             int myBestTrainIdx2 = -1;\r
  66 \r
  67             s_distance += threadIdx.y * BLOCK_SIZE;\r
  68             s_trainIdx += threadIdx.y * BLOCK_SIZE;\r
  69 \r
  70             s_distance[threadIdx.x] = bestDistance1;\r
  71             s_trainIdx[threadIdx.x] = bestTrainIdx1;\r
  72 \r
  73             __syncthreads();\r
  74 \r
  75             if (threadIdx.x == 0)\r
  76             {\r
  77                 #pragma unroll\r
  78                 for (int i = 0; i < BLOCK_SIZE; ++i)\r
  79                 {\r
  80                     float val = s_distance[i];\r
  81 \r
  82                     if (val < myBestDistance1)\r
  83                     {\r
  84                         myBestDistance2 = myBestDistance1;\r
  85                         myBestTrainIdx2 = myBestTrainIdx1;\r
  86 \r
  87                         myBestDistance1 = val;\r
  88                         myBestTrainIdx1 = s_trainIdx[i];\r
  89                     }\r
  90                     else if (val < myBestDistance2)\r
  91                     {\r
  92                         myBestDistance2 = val;\r
  93                         myBestTrainIdx2 = s_trainIdx[i];\r
  94                     }\r
  95                 }\r
  96             }\r
  97 \r
  98             __syncthreads();\r
  99 \r
 100             s_distance[threadIdx.x] = bestDistance2;\r
 101             s_trainIdx[threadIdx.x] = bestTrainIdx2;\r
 102 \r
 103             __syncthreads();\r
 104 \r
 105             if (threadIdx.x == 0)\r
 106             {\r
 107                 #pragma unroll\r
 108                 for (int i = 0; i < BLOCK_SIZE; ++i)\r
 109                 {\r
 110                     float val = s_distance[i];\r
 111 \r
 112                     if (val < myBestDistance2)\r
 113                     {\r
 114                         myBestDistance2 = val;\r
 115                         myBestTrainIdx2 = s_trainIdx[i];\r
 116                     }\r
 117                 }\r
 118             }\r
 119 \r
 120             bestDistance1 = myBestDistance1;\r
 121             bestDistance2 = myBestDistance2;\r
 122 \r
 123             bestTrainIdx1 = myBestTrainIdx1;\r
 124             bestTrainIdx2 = myBestTrainIdx2;\r
 125         }\r
 126 \r
 127         template <int BLOCK_SIZE>\r
 128         __device__ void findBestMatch(float& bestDistance1, float& bestDistance2,\r
 129                                        int& bestTrainIdx1, int& bestTrainIdx2,\r
 130                                        int& bestImgIdx1, int& bestImgIdx2,\r
 131                                        float* s_distance, int* s_trainIdx, int* s_imgIdx)\r
 132         {\r
 133             float myBestDistance1 = numeric_limits<float>::max();\r
 134             float myBestDistance2 = numeric_limits<float>::max();\r
 135             int myBestTrainIdx1 = -1;\r
 136             int myBestTrainIdx2 = -1;\r
 137             int myBestImgIdx1 = -1;\r
 138             int myBestImgIdx2 = -1;\r
 139 \r
 140             s_distance += threadIdx.y * BLOCK_SIZE;\r
 141             s_trainIdx += threadIdx.y * BLOCK_SIZE;\r
 142             s_imgIdx   += threadIdx.y * BLOCK_SIZE;\r
 143 \r
 144             s_distance[threadIdx.x] = bestDistance1;\r
 145             s_trainIdx[threadIdx.x] = bestTrainIdx1;\r
 146             s_imgIdx[threadIdx.x]   = bestImgIdx1;\r
 147 \r
 148             __syncthreads();\r
 149 \r
 150             if (threadIdx.x == 0)\r
 151             {\r
 152                 #pragma unroll\r
 153                 for (int i = 0; i < BLOCK_SIZE; ++i)\r
 154                 {\r
 155                     float val = s_distance[i];\r
 156 \r
 157                     if (val < myBestDistance1)\r
 158                     {\r
 159                         myBestDistance2 = myBestDistance1;\r
 160                         myBestTrainIdx2 = myBestTrainIdx1;\r
 161                         myBestImgIdx2   = myBestImgIdx1;\r
 162 \r
 163                         myBestDistance1 = val;\r
 164                         myBestTrainIdx1 = s_trainIdx[i];\r
 165                         myBestImgIdx1   = s_imgIdx[i];\r
 166                     }\r
 167                     else if (val < myBestDistance2)\r
 168                     {\r
 169                         myBestDistance2 = val;\r
 170                         myBestTrainIdx2 = s_trainIdx[i];\r
 171                         myBestImgIdx2   = s_imgIdx[i];\r
 172                     }\r
 173                 }\r
 174             }\r
 175 \r
 176             __syncthreads();\r
 177 \r
 178             s_distance[threadIdx.x] = bestDistance2;\r
 179             s_trainIdx[threadIdx.x] = bestTrainIdx2;\r
 180             s_imgIdx[threadIdx.x]   = bestImgIdx2;\r
 181 \r
 182             __syncthreads();\r
 183 \r
 184             if (threadIdx.x == 0)\r
 185             {\r
 186                 #pragma unroll\r
 187                 for (int i = 0; i < BLOCK_SIZE; ++i)\r
 188                 {\r
 189                     float val = s_distance[i];\r
 190 \r
 191                     if (val < myBestDistance2)\r
 192                     {\r
 193                         myBestDistance2 = val;\r
 194                         myBestTrainIdx2 = s_trainIdx[i];\r
 195                         myBestImgIdx2   = s_imgIdx[i];\r
 196                     }\r
 197                 }\r
 198             }\r
 199 \r
 200             bestDistance1 = myBestDistance1;\r
 201             bestDistance2 = myBestDistance2;\r
 202 \r
 203             bestTrainIdx1 = myBestTrainIdx1;\r
 204             bestTrainIdx2 = myBestTrainIdx2;\r
 205 \r
 206             bestImgIdx1 = myBestImgIdx1;\r
 207             bestImgIdx2 = myBestImgIdx2;\r
 208         }\r
 209 \r
 210         ///////////////////////////////////////////////////////////////////////////////\r
 211         // Match Unrolled Cached\r
 212 \r
 213         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename T, typename U>\r
 214         __device__ void loadQueryToSmem(int queryIdx, const PtrStepSz<T>& query, U* s_query)\r
 215         {\r
 216             #pragma unroll\r
 217             for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; ++i)\r
 218             {\r
 219                 const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 220                 s_query[threadIdx.y * MAX_DESC_LEN + loadX] = loadX < query.cols ? query.ptr(::min(queryIdx, query.rows - 1))[loadX] : 0;\r
 221             }\r
 222         }\r
 223 \r
 224         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 225         __device__ void loopUnrolledCached(int queryIdx, const PtrStepSz<T>& query, int imgIdx, const PtrStepSz<T>& train, const Mask& mask,\r
 226                                            typename Dist::value_type* s_query, typename Dist::value_type* s_train,\r
 227                                            float& bestDistance1, float& bestDistance2,\r
 228                                            int& bestTrainIdx1, int& bestTrainIdx2,\r
 229                                            int& bestImgIdx1, int& bestImgIdx2)\r
 230         {\r
 231             for (int t = 0, endt = (train.rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; ++t)\r
 232             {\r
 233                 Dist dist;\r
 234 \r
 235                 #pragma unroll\r
 236                 for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; ++i)\r
 237                 {\r
 238                     const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 239 \r
 240                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = 0;\r
 241 \r
 242                     if (loadX < train.cols)\r
 243                     {\r
 244                         T val;\r
 245 \r
 246                         ForceGlob<T>::Load(train.ptr(::min(t * BLOCK_SIZE + threadIdx.y, train.rows - 1)), loadX, val);\r
 247                         s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = val;\r
 248                     }\r
 249 \r
 250                     __syncthreads();\r
 251 \r
 252                     #pragma unroll\r
 253                     for (int j = 0; j < BLOCK_SIZE; ++j)\r
 254                         dist.reduceIter(s_query[threadIdx.y * MAX_DESC_LEN + i * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + threadIdx.x]);\r
 255 \r
 256                     __syncthreads();\r
 257                 }\r
 258 \r
 259                 typename Dist::result_type distVal = dist;\r
 260 \r
 261                 const int trainIdx = t * BLOCK_SIZE + threadIdx.x;\r
 262 \r
 263                 if (queryIdx < query.rows && trainIdx < train.rows && mask(queryIdx, trainIdx))\r
 264                 {\r
 265                     if (distVal < bestDistance1)\r
 266                     {\r
 267                         bestImgIdx2   = bestImgIdx1;\r
 268                         bestDistance2 = bestDistance1;\r
 269                         bestTrainIdx2 = bestTrainIdx1;\r
 270 \r
 271                         bestImgIdx1   = imgIdx;\r
 272                         bestDistance1 = distVal;\r
 273                         bestTrainIdx1 = trainIdx;\r
 274                     }\r
 275                     else if (distVal < bestDistance2)\r
 276                     {\r
 277                         bestImgIdx2   = imgIdx;\r
 278                         bestDistance2 = distVal;\r
 279                         bestTrainIdx2 = trainIdx;\r
 280                     }\r
 281                 }\r
 282             }\r
 283         }\r
 284 \r
 285         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 286         __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)\r
 287         {\r
 288             extern __shared__ int smem[];\r
 289 \r
 290             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 291 \r
 292             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 293             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * MAX_DESC_LEN);\r
 294 \r
 295             loadQueryToSmem<BLOCK_SIZE, MAX_DESC_LEN>(queryIdx, query, s_query);\r
 296 \r
 297             float myBestDistance1 = numeric_limits<float>::max();\r
 298             float myBestDistance2 = numeric_limits<float>::max();\r
 299             int myBestTrainIdx1 = -1;\r
 300             int myBestTrainIdx2 = -1;\r
 301 \r
 302             loopUnrolledCached<BLOCK_SIZE, MAX_DESC_LEN, Dist>(queryIdx, query, 0, train, mask, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestTrainIdx1, myBestTrainIdx2);\r
 303 \r
 304             __syncthreads();\r
 305 \r
 306             float* s_distance = (float*)(smem);\r
 307             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 308 \r
 309             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, s_distance, s_trainIdx);\r
 310 \r
 311             if (queryIdx < query.rows && threadIdx.x == 0)\r
 312             {\r
 313                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 314                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 315             }\r
 316         }\r
 317 \r
 318         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 319         void matchUnrolledCached(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask,\r
 320                                  const PtrStepSz<int2>& trainIdx, const PtrStepSz<float2>& distance,\r
 321                                  cudaStream_t stream)\r
 322         {\r
 323             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 324             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 325 \r
 326             const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= BLOCK_SIZE ? MAX_DESC_LEN : BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 327 \r
 328             matchUnrolledCached<BLOCK_SIZE, MAX_DESC_LEN, Dist><<<grid, block, smemSize, stream>>>(query, train, mask, trainIdx.data, distance.data);\r
 329             cudaSafeCall( cudaGetLastError() );\r
 330 \r
 331             if (stream == 0)\r
 332                 cudaSafeCall( cudaDeviceSynchronize() );\r
 333         }\r
 334 \r
 335         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 336         __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)\r
 337         {\r
 338             extern __shared__ int smem[];\r
 339 \r
 340             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 341 \r
 342             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 343             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * MAX_DESC_LEN);\r
 344 \r
 345             loadQueryToSmem<BLOCK_SIZE, MAX_DESC_LEN>(queryIdx, query, s_query);\r
 346 \r
 347             float myBestDistance1 = numeric_limits<float>::max();\r
 348             float myBestDistance2 = numeric_limits<float>::max();\r
 349             int myBestTrainIdx1 = -1;\r
 350             int myBestTrainIdx2 = -1;\r
 351             int myBestImgIdx1 = -1;\r
 352             int myBestImgIdx2 = -1;\r
 353 \r
 354             Mask m = mask;\r
 355 \r
 356             for (int imgIdx = 0; imgIdx < n; ++imgIdx)\r
 357             {\r
 358                 const PtrStepSz<T> train = trains[imgIdx];\r
 359                 m.next();\r
 360                 loopUnrolledCached<BLOCK_SIZE, MAX_DESC_LEN, Dist>(queryIdx, query, imgIdx, train, m, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2);\r
 361             }\r
 362 \r
 363             __syncthreads();\r
 364 \r
 365             float* s_distance = (float*)(smem);\r
 366             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 367             int* s_imgIdx = (int*)(smem + 2 * BLOCK_SIZE * BLOCK_SIZE);\r
 368 \r
 369             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2, s_distance, s_trainIdx, s_imgIdx);\r
 370 \r
 371             if (queryIdx < query.rows && threadIdx.x == 0)\r
 372             {\r
 373                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 374                 bestImgIdx[queryIdx] = make_int2(myBestImgIdx1, myBestImgIdx2);\r
 375                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 376             }\r
 377         }\r
 378 \r
 379         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 380         void matchUnrolledCached(const PtrStepSz<T>& query, const PtrStepSz<T>* trains, int n, const Mask& mask,\r
 381                                  const PtrStepSz<int2>& trainIdx, const PtrStepSz<int2>& imgIdx, const PtrStepSz<float2>& distance,\r
 382                                  cudaStream_t stream)\r
 383         {\r
 384             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 385             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 386 \r
 387             const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= 2 * BLOCK_SIZE ? MAX_DESC_LEN : 2 * BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 388 \r
 389             matchUnrolledCached<BLOCK_SIZE, MAX_DESC_LEN, Dist><<<grid, block, smemSize, stream>>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data);\r
 390             cudaSafeCall( cudaGetLastError() );\r
 391 \r
 392             if (stream == 0)\r
 393                 cudaSafeCall( cudaDeviceSynchronize() );\r
 394         }\r
 395 \r
 396         ///////////////////////////////////////////////////////////////////////////////\r
 397         // Match Unrolled\r
 398 \r
 399         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 400         __device__ void loopUnrolled(int queryIdx, const PtrStepSz<T>& query, int imgIdx, const PtrStepSz<T>& train, const Mask& mask,\r
 401                                      typename Dist::value_type* s_query, typename Dist::value_type* s_train,\r
 402                                      float& bestDistance1, float& bestDistance2,\r
 403                                      int& bestTrainIdx1, int& bestTrainIdx2,\r
 404                                      int& bestImgIdx1, int& bestImgIdx2)\r
 405         {\r
 406             for (int t = 0, endt = (train.rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; ++t)\r
 407             {\r
 408                 Dist dist;\r
 409 \r
 410                 #pragma unroll\r
 411                 for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; ++i)\r
 412                 {\r
 413                     const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 414 \r
 415                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = 0;\r
 416                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = 0;\r
 417 \r
 418                     if (loadX < query.cols)\r
 419                     {\r
 420                         T val;\r
 421 \r
 422                         ForceGlob<T>::Load(query.ptr(::min(queryIdx, query.rows - 1)), loadX, val);\r
 423                         s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = val;\r
 424 \r
 425                         ForceGlob<T>::Load(train.ptr(::min(t * BLOCK_SIZE + threadIdx.y, train.rows - 1)), loadX, val);\r
 426                         s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = val;\r
 427                     }\r
 428 \r
 429                     __syncthreads();\r
 430 \r
 431                     #pragma unroll\r
 432                     for (int j = 0; j < BLOCK_SIZE; ++j)\r
 433                         dist.reduceIter(s_query[threadIdx.y * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + threadIdx.x]);\r
 434 \r
 435                     __syncthreads();\r
 436                 }\r
 437 \r
 438                 typename Dist::result_type distVal = dist;\r
 439 \r
 440                 const int trainIdx = t * BLOCK_SIZE + threadIdx.x;\r
 441 \r
 442                 if (queryIdx < query.rows && trainIdx < train.rows && mask(queryIdx, trainIdx))\r
 443                 {\r
 444                     if (distVal < bestDistance1)\r
 445                     {\r
 446                         bestImgIdx2   = bestImgIdx1;\r
 447                         bestDistance2 = bestDistance1;\r
 448                         bestTrainIdx2 = bestTrainIdx1;\r
 449 \r
 450                         bestImgIdx1   = imgIdx;\r
 451                         bestDistance1 = distVal;\r
 452                         bestTrainIdx1 = trainIdx;\r
 453                     }\r
 454                     else if (distVal < bestDistance2)\r
 455                     {\r
 456                         bestImgIdx2   = imgIdx;\r
 457                         bestDistance2 = distVal;\r
 458                         bestTrainIdx2 = trainIdx;\r
 459                     }\r
 460                 }\r
 461             }\r
 462         }\r
 463 \r
 464         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 465         __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)\r
 466         {\r
 467             extern __shared__ int smem[];\r
 468 \r
 469             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 470 \r
 471             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 472             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 473 \r
 474             float myBestDistance1 = numeric_limits<float>::max();\r
 475             float myBestDistance2 = numeric_limits<float>::max();\r
 476             int myBestTrainIdx1 = -1;\r
 477             int myBestTrainIdx2 = -1;\r
 478 \r
 479             loopUnrolled<BLOCK_SIZE, MAX_DESC_LEN, Dist>(queryIdx, query, 0, train, mask, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestTrainIdx1, myBestTrainIdx2);\r
 480 \r
 481             __syncthreads();\r
 482 \r
 483             float* s_distance = (float*)(smem);\r
 484             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 485 \r
 486             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, s_distance, s_trainIdx);\r
 487 \r
 488             if (queryIdx < query.rows && threadIdx.x == 0)\r
 489             {\r
 490                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 491                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 492             }\r
 493         }\r
 494 \r
 495         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 496         void matchUnrolled(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask,\r
 497                            const PtrStepSz<int2>& trainIdx, const PtrStepSz<float2>& distance,\r
 498                            cudaStream_t stream)\r
 499         {\r
 500             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 501             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 502 \r
 503             const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 504 \r
 505             matchUnrolled<BLOCK_SIZE, MAX_DESC_LEN, Dist><<<grid, block, smemSize, stream>>>(query, train, mask, trainIdx.data, distance.data);\r
 506             cudaSafeCall( cudaGetLastError() );\r
 507 \r
 508             if (stream == 0)\r
 509                 cudaSafeCall( cudaDeviceSynchronize() );\r
 510         }\r
 511 \r
 512         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 513         __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)\r
 514         {\r
 515             extern __shared__ int smem[];\r
 516 \r
 517             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 518 \r
 519             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 520             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 521 \r
 522             float myBestDistance1 = numeric_limits<float>::max();\r
 523             float myBestDistance2 = numeric_limits<float>::max();\r
 524             int myBestTrainIdx1 = -1;\r
 525             int myBestTrainIdx2 = -1;\r
 526             int myBestImgIdx1 = -1;\r
 527             int myBestImgIdx2 = -1;\r
 528 \r
 529             Mask m = mask;\r
 530 \r
 531             for (int imgIdx = 0; imgIdx < n; ++imgIdx)\r
 532             {\r
 533                 const PtrStepSz<T> train = trains[imgIdx];\r
 534                 m.next();\r
 535                 loopUnrolled<BLOCK_SIZE, MAX_DESC_LEN, Dist>(queryIdx, query, imgIdx, train, m, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2);\r
 536             }\r
 537 \r
 538             __syncthreads();\r
 539 \r
 540             float* s_distance = (float*)(smem);\r
 541             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 542             int* s_imgIdx = (int*)(smem + 2 * BLOCK_SIZE * BLOCK_SIZE);\r
 543 \r
 544             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2, s_distance, s_trainIdx, s_imgIdx);\r
 545 \r
 546             if (queryIdx < query.rows && threadIdx.x == 0)\r
 547             {\r
 548                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 549                 bestImgIdx[queryIdx] = make_int2(myBestImgIdx1, myBestImgIdx2);\r
 550                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 551             }\r
 552         }\r
 553 \r
 554         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 555         void matchUnrolled(const PtrStepSz<T>& query, const PtrStepSz<T>* trains, int n, const Mask& mask,\r
 556                            const PtrStepSz<int2>& trainIdx, const PtrStepSz<int2>& imgIdx, const PtrStepSz<float2>& distance,\r
 557                            cudaStream_t stream)\r
 558         {\r
 559             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 560             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 561 \r
 562             const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 563 \r
 564             matchUnrolled<BLOCK_SIZE, MAX_DESC_LEN, Dist><<<grid, block, smemSize, stream>>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data);\r
 565             cudaSafeCall( cudaGetLastError() );\r
 566 \r
 567             if (stream == 0)\r
 568                 cudaSafeCall( cudaDeviceSynchronize() );\r
 569         }\r
 570 \r
 571         ///////////////////////////////////////////////////////////////////////////////\r
 572         // Match\r
 573 \r
 574         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 575         __device__ void loop(int queryIdx, const PtrStepSz<T>& query, int imgIdx, const PtrStepSz<T>& train, const Mask& mask,\r
 576                              typename Dist::value_type* s_query, typename Dist::value_type* s_train,\r
 577                              float& bestDistance1, float& bestDistance2,\r
 578                              int& bestTrainIdx1, int& bestTrainIdx2,\r
 579                              int& bestImgIdx1, int& bestImgIdx2)\r
 580         {\r
 581             for (int t = 0, endt = (train.rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; ++t)\r
 582             {\r
 583                 Dist dist;\r
 584 \r
 585                 for (int i = 0, endi = (query.cols + BLOCK_SIZE - 1) / BLOCK_SIZE; i < endi; ++i)\r
 586                 {\r
 587                     const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 588 \r
 589                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = 0;\r
 590                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = 0;\r
 591 \r
 592                     if (loadX < query.cols)\r
 593                     {\r
 594                         T val;\r
 595 \r
 596                         ForceGlob<T>::Load(query.ptr(::min(queryIdx, query.rows - 1)), loadX, val);\r
 597                         s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = val;\r
 598 \r
 599                         ForceGlob<T>::Load(train.ptr(::min(t * BLOCK_SIZE + threadIdx.y, train.rows - 1)), loadX, val);\r
 600                         s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = val;\r
 601                     }\r
 602 \r
 603                     __syncthreads();\r
 604 \r
 605                     #pragma unroll\r
 606                     for (int j = 0; j < BLOCK_SIZE; ++j)\r
 607                         dist.reduceIter(s_query[threadIdx.y * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + threadIdx.x]);\r
 608 \r
 609                     __syncthreads();\r
 610                 }\r
 611 \r
 612                 typename Dist::result_type distVal = dist;\r
 613 \r
 614                 const int trainIdx = t * BLOCK_SIZE + threadIdx.x;\r
 615 \r
 616                 if (queryIdx < query.rows && trainIdx < train.rows && mask(queryIdx, trainIdx))\r
 617                 {\r
 618                     if (distVal < bestDistance1)\r
 619                     {\r
 620                         bestImgIdx2   = bestImgIdx1;\r
 621                         bestDistance2 = bestDistance1;\r
 622                         bestTrainIdx2 = bestTrainIdx1;\r
 623 \r
 624                         bestImgIdx1   = imgIdx;\r
 625                         bestDistance1 = distVal;\r
 626                         bestTrainIdx1 = trainIdx;\r
 627                     }\r
 628                     else if (distVal < bestDistance2)\r
 629                     {\r
 630                         bestImgIdx2   = imgIdx;\r
 631                         bestDistance2 = distVal;\r
 632                         bestTrainIdx2 = trainIdx;\r
 633                     }\r
 634                 }\r
 635             }\r
 636         }\r
 637 \r
 638         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 639         __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)\r
 640         {\r
 641             extern __shared__ int smem[];\r
 642 \r
 643             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 644 \r
 645             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 646             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 647 \r
 648             float myBestDistance1 = numeric_limits<float>::max();\r
 649             float myBestDistance2 = numeric_limits<float>::max();\r
 650             int myBestTrainIdx1 = -1;\r
 651             int myBestTrainIdx2 = -1;\r
 652 \r
 653             loop<BLOCK_SIZE, Dist>(queryIdx, query, 0, train, mask, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestTrainIdx1, myBestTrainIdx2);\r
 654 \r
 655             __syncthreads();\r
 656 \r
 657             float* s_distance = (float*)(smem);\r
 658             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 659 \r
 660             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, s_distance, s_trainIdx);\r
 661 \r
 662             if (queryIdx < query.rows && threadIdx.x == 0)\r
 663             {\r
 664                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 665                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 666             }\r
 667         }\r
 668 \r
 669         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 670         void match(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask,\r
 671                    const PtrStepSz<int2>& trainIdx, const PtrStepSz<float2>& distance,\r
 672                    cudaStream_t stream)\r
 673         {\r
 674             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 675             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 676 \r
 677             const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 678 \r
 679             match<BLOCK_SIZE, Dist><<<grid, block, smemSize, stream>>>(query, train, mask, trainIdx.data, distance.data);\r
 680             cudaSafeCall( cudaGetLastError() );\r
 681 \r
 682             if (stream == 0)\r
 683                 cudaSafeCall( cudaDeviceSynchronize() );\r
 684         }\r
 685 \r
 686         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 687         __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)\r
 688         {\r
 689             extern __shared__ int smem[];\r
 690 \r
 691             const int queryIdx = blockIdx.x * BLOCK_SIZE + threadIdx.y;\r
 692 \r
 693             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 694             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 695 \r
 696             float myBestDistance1 = numeric_limits<float>::max();\r
 697             float myBestDistance2 = numeric_limits<float>::max();\r
 698             int myBestTrainIdx1 = -1;\r
 699             int myBestTrainIdx2 = -1;\r
 700             int myBestImgIdx1 = -1;\r
 701             int myBestImgIdx2 = -1;\r
 702 \r
 703             Mask m = mask;\r
 704 \r
 705             for (int imgIdx = 0; imgIdx < n; ++imgIdx)\r
 706             {\r
 707                 const PtrStepSz<T> train = trains[imgIdx];\r
 708                 m.next();\r
 709                 loop<BLOCK_SIZE, Dist>(queryIdx, query, imgIdx, train, m, s_query, s_train, myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2);\r
 710             }\r
 711 \r
 712             __syncthreads();\r
 713 \r
 714             float* s_distance = (float*)(smem);\r
 715             int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 716             int* s_imgIdx = (int*)(smem + 2 * BLOCK_SIZE * BLOCK_SIZE);\r
 717 \r
 718             findBestMatch<BLOCK_SIZE>(myBestDistance1, myBestDistance2, myBestTrainIdx1, myBestTrainIdx2, myBestImgIdx1, myBestImgIdx2, s_distance, s_trainIdx, s_imgIdx);\r
 719 \r
 720             if (queryIdx < query.rows && threadIdx.x == 0)\r
 721             {\r
 722                 bestTrainIdx[queryIdx] = make_int2(myBestTrainIdx1, myBestTrainIdx2);\r
 723                 bestImgIdx[queryIdx] = make_int2(myBestImgIdx1, myBestImgIdx2);\r
 724                 bestDistance[queryIdx] = make_float2(myBestDistance1, myBestDistance2);\r
 725             }\r
 726         }\r
 727 \r
 728         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 729         void match(const PtrStepSz<T>& query, const PtrStepSz<T>* trains, int n, const Mask& mask,\r
 730                    const PtrStepSz<int2>& trainIdx, const PtrStepSz<int2>& imgIdx, const PtrStepSz<float2>& distance,\r
 731                    cudaStream_t stream)\r
 732         {\r
 733             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 734             const dim3 grid(divUp(query.rows, BLOCK_SIZE));\r
 735 \r
 736             const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 737 \r
 738             match<BLOCK_SIZE, Dist><<<grid, block, smemSize, stream>>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data);\r
 739             cudaSafeCall( cudaGetLastError() );\r
 740 \r
 741             if (stream == 0)\r
 742                 cudaSafeCall( cudaDeviceSynchronize() );\r
 743         }\r
 744 \r
 745         ///////////////////////////////////////////////////////////////////////////////\r
 746         // knnMatch 2 dispatcher\r
 747 \r
 748         template <typename Dist, typename T, typename Mask>\r
 749         void match2Dispatcher(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask,\r
 750                               const PtrStepSzb& trainIdx, const PtrStepSzb& distance,\r
 751                               int cc, cudaStream_t stream)\r
 752         {\r
 753             (void)cc;\r
 754             if (query.cols <= 64)\r
 755             {\r
 756                 matchUnrolledCached<16, 64, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 757             }\r
 758             else if (query.cols <= 128)\r
 759             {\r
 760                 matchUnrolledCached<16, 128, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 761             }\r
 762             /*else if (query.cols <= 256)\r
 763             {\r
 764                 matchUnrolled<16, 256, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 765             }\r
 766             else if (query.cols <= 512)\r
 767             {\r
 768                 matchUnrolled<16, 512, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 769             }\r
 770             else if (query.cols <= 1024)\r
 771             {\r
 772                 matchUnrolled<16, 1024, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 773             }*/\r
 774             else\r
 775             {\r
 776                 match<16, Dist>(query, train, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 777             }\r
 778         }\r
 779 \r
 780         template <typename Dist, typename T, typename Mask>\r
 781         void match2Dispatcher(const PtrStepSz<T>& query, const PtrStepSz<T>* trains, int n, const Mask& mask,\r
 782                               const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance,\r
 783                               int cc, cudaStream_t stream)\r
 784         {\r
 785             (void)cc;\r
 786             if (query.cols <= 64)\r
 787             {\r
 788                 matchUnrolledCached<16, 64, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 789             }\r
 790             else if (query.cols <= 128)\r
 791             {\r
 792                 matchUnrolledCached<16, 128, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 793             }\r
 794             /*else if (query.cols <= 256)\r
 795             {\r
 796                 matchUnrolled<16, 256, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 797             }\r
 798             else if (query.cols <= 512)\r
 799             {\r
 800                 matchUnrolled<16, 512, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 801             }\r
 802             else if (query.cols <= 1024)\r
 803             {\r
 804                 matchUnrolled<16, 1024, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 805             }*/\r
 806             else\r
 807             {\r
 808                 match<16, Dist>(query, trains, n, mask, static_cast< PtrStepSz<int2> >(trainIdx), static_cast< PtrStepSz<int2> >(imgIdx), static_cast< PtrStepSz<float2> > (distance), stream);\r
 809             }\r
 810         }\r
 811 \r
 812         ///////////////////////////////////////////////////////////////////////////////\r
 813         // Calc distance kernel\r
 814 \r
 815         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 816         __global__ void calcDistanceUnrolled(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, PtrStepf allDist)\r
 817         {\r
 818             extern __shared__ int smem[];\r
 819 \r
 820             const int queryIdx = blockIdx.y * BLOCK_SIZE + threadIdx.y;\r
 821             const int trainIdx = blockIdx.x * BLOCK_SIZE + threadIdx.x;\r
 822 \r
 823             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 824             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 825 \r
 826             Dist dist;\r
 827 \r
 828             #pragma unroll\r
 829             for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; ++i)\r
 830             {\r
 831                 const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 832 \r
 833                 if (loadX < query.cols)\r
 834                 {\r
 835                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = query.ptr(::min(queryIdx, query.rows - 1))[loadX];\r
 836                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = train.ptr(::min(blockIdx.x * BLOCK_SIZE + threadIdx.y, train.rows - 1))[loadX];\r
 837                 }\r
 838                 else\r
 839                 {\r
 840                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = 0;\r
 841                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = 0;\r
 842                 }\r
 843 \r
 844                 __syncthreads();\r
 845 \r
 846                 #pragma unroll\r
 847                 for (int j = 0; j < BLOCK_SIZE; ++j)\r
 848                     dist.reduceIter(s_query[threadIdx.y * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + threadIdx.x]);\r
 849 \r
 850                 __syncthreads();\r
 851             }\r
 852 \r
 853             if (queryIdx < query.rows && trainIdx < train.rows)\r
 854             {\r
 855                 float distVal = numeric_limits<float>::max();\r
 856 \r
 857                 if (mask(queryIdx, trainIdx))\r
 858                     distVal = (typename Dist::result_type)dist;\r
 859 \r
 860                 allDist.ptr(queryIdx)[trainIdx] = distVal;\r
 861             }\r
 862         }\r
 863 \r
 864         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>\r
 865         void calcDistanceUnrolled(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask, const PtrStepSzf& allDist, cudaStream_t stream)\r
 866         {\r
 867             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 868             const dim3 grid(divUp(train.rows, BLOCK_SIZE), divUp(query.rows, BLOCK_SIZE));\r
 869 \r
 870             const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 871 \r
 872             calcDistanceUnrolled<BLOCK_SIZE, MAX_DESC_LEN, Dist><<<grid, block, smemSize, stream>>>(query, train, mask, allDist);\r
 873             cudaSafeCall( cudaGetLastError() );\r
 874 \r
 875             if (stream == 0)\r
 876                 cudaSafeCall( cudaDeviceSynchronize() );\r
 877         }\r
 878 \r
 879         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 880         __global__ void calcDistance(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, PtrStepf allDist)\r
 881         {\r
 882             extern __shared__ int smem[];\r
 883 \r
 884             const int queryIdx = blockIdx.y * BLOCK_SIZE + threadIdx.y;\r
 885             const int trainIdx = blockIdx.x * BLOCK_SIZE + threadIdx.x;\r
 886 \r
 887             typename Dist::value_type* s_query = (typename Dist::value_type*)(smem);\r
 888             typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * BLOCK_SIZE);\r
 889 \r
 890             Dist dist;\r
 891 \r
 892             for (int i = 0, endi = (query.cols + BLOCK_SIZE - 1) / BLOCK_SIZE; i < endi; ++i)\r
 893             {\r
 894                 const int loadX = threadIdx.x + i * BLOCK_SIZE;\r
 895 \r
 896                 if (loadX < query.cols)\r
 897                 {\r
 898                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = query.ptr(::min(queryIdx, query.rows - 1))[loadX];\r
 899                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = train.ptr(::min(blockIdx.x * BLOCK_SIZE + threadIdx.y, train.rows - 1))[loadX];\r
 900                 }\r
 901                 else\r
 902                 {\r
 903                     s_query[threadIdx.y * BLOCK_SIZE + threadIdx.x] = 0;\r
 904                     s_train[threadIdx.x * BLOCK_SIZE + threadIdx.y] = 0;\r
 905                 }\r
 906 \r
 907                 __syncthreads();\r
 908 \r
 909                 #pragma unroll\r
 910                 for (int j = 0; j < BLOCK_SIZE; ++j)\r
 911                     dist.reduceIter(s_query[threadIdx.y * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + threadIdx.x]);\r
 912 \r
 913                 __syncthreads();\r
 914             }\r
 915 \r
 916             if (queryIdx < query.rows && trainIdx < train.rows)\r
 917             {\r
 918                 float distVal = numeric_limits<float>::max();\r
 919 \r
 920                 if (mask(queryIdx, trainIdx))\r
 921                     distVal = (typename Dist::result_type)dist;\r
 922 \r
 923                 allDist.ptr(queryIdx)[trainIdx] = distVal;\r
 924             }\r
 925         }\r
 926 \r
 927         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>\r
 928         void calcDistance(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask, const PtrStepSzf& allDist, cudaStream_t stream)\r
 929         {\r
 930             const dim3 block(BLOCK_SIZE, BLOCK_SIZE);\r
 931             const dim3 grid(divUp(train.rows, BLOCK_SIZE), divUp(query.rows, BLOCK_SIZE));\r
 932 \r
 933             const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);\r
 934 \r
 935             calcDistance<BLOCK_SIZE, Dist><<<grid, block, smemSize, stream>>>(query, train, mask, allDist);\r
 936             cudaSafeCall( cudaGetLastError() );\r
 937 \r
 938             if (stream == 0)\r
 939                 cudaSafeCall( cudaDeviceSynchronize() );\r
 940         }\r
 941 \r
 942         ///////////////////////////////////////////////////////////////////////////////\r
 943         // Calc Distance dispatcher\r
 944 \r
 945         template <typename Dist, typename T, typename Mask>\r
 946         void calcDistanceDispatcher(const PtrStepSz<T>& query, const PtrStepSz<T>& train, const Mask& mask,\r
 947                                     const PtrStepSzf& allDist,\r
 948                                     int cc, cudaStream_t stream)\r
 949         {\r
 950             (void)cc;\r
 951             if (query.cols <= 64)\r
 952             {\r
 953                 calcDistanceUnrolled<16, 64, Dist>(query, train, mask, allDist, stream);\r
 954             }\r
 955             else if (query.cols <= 128)\r
 956             {\r
 957                 calcDistanceUnrolled<16, 128, Dist>(query, train, mask, allDist, stream);\r
 958             }\r
 959             /*else if (query.cols <= 256)\r
 960             {\r
 961                 calcDistanceUnrolled<16, 256, Dist>(query, train, mask, allDist, stream);\r
 962             }\r
 963             else if (query.cols <= 512)\r
 964             {\r
 965                 calcDistanceUnrolled<16, 512, Dist>(query, train, mask, allDist, stream);\r
 966             }\r
 967             else if (query.cols <= 1024)\r
 968             {\r
 969                 calcDistanceUnrolled<16, 1024, Dist>(query, train, mask, allDist, stream);\r
 970             }*/\r
 971             else\r
 972             {\r
 973                 calcDistance<16, Dist>(query, train, mask, allDist, stream);\r
 974             }\r
 975         }\r
 976 \r
 977         ///////////////////////////////////////////////////////////////////////////////\r
 978         // find knn match kernel\r
 979 \r
 980         template <int BLOCK_SIZE>\r
 981         __global__ void findBestMatch(PtrStepSzf allDist, int i, PtrStepi trainIdx, PtrStepf distance)\r
 982         {\r
 983             const int SMEM_SIZE = BLOCK_SIZE > 64 ? BLOCK_SIZE : 64;\r
 984             __shared__ float s_dist[SMEM_SIZE];\r
 985             __shared__ int s_trainIdx[SMEM_SIZE];\r
 986 \r
 987             const int queryIdx = blockIdx.x;\r
 988 \r
 989             float* allDistRow = allDist.ptr(queryIdx);\r
 990 \r
 991             float dist = numeric_limits<float>::max();\r
 992             int bestIdx = -1;\r
 993 \r
 994             for (int i = threadIdx.x; i < allDist.cols; i += BLOCK_SIZE)\r
 995             {\r
 996                 float reg = allDistRow[i];\r
 997                 if (reg < dist)\r
 998                 {\r
 999                     dist = reg;\r
1000                     bestIdx = i;\r
1001                 }\r
1002             }\r
1003 \r
1004             s_dist[threadIdx.x] = dist;\r
1005             s_trainIdx[threadIdx.x] = bestIdx;\r
1006             __syncthreads();\r
1007 \r
1008             reducePredVal<BLOCK_SIZE>(s_dist, dist, s_trainIdx, bestIdx, threadIdx.x, less<volatile float>());\r
1009 \r
1010             if (threadIdx.x == 0)\r
1011             {\r
1012                 if (dist < numeric_limits<float>::max())\r
1013                 {\r
1014                     allDistRow[bestIdx] = numeric_limits<float>::max();\r
1015                     trainIdx.ptr(queryIdx)[i] = bestIdx;\r
1016                     distance.ptr(queryIdx)[i] = dist;\r
1017                 }\r
1018             }\r
1019         }\r
1020 \r
1021         template <int BLOCK_SIZE>\r
1022         void findKnnMatch(int k, const PtrStepSzi& trainIdx, const PtrStepSzf& distance, const PtrStepSzf& allDist, cudaStream_t stream)\r
1023         {\r
1024             const dim3 block(BLOCK_SIZE, 1, 1);\r
1025             const dim3 grid(trainIdx.rows, 1, 1);\r
1026 \r
1027             for (int i = 0; i < k; ++i)\r
1028             {\r
1029                 findBestMatch<BLOCK_SIZE><<<grid, block, 0, stream>>>(allDist, i, trainIdx, distance);\r
1030                 cudaSafeCall( cudaGetLastError() );\r
1031             }\r
1032 \r
1033             if (stream == 0)\r
1034                 cudaSafeCall( cudaDeviceSynchronize() );\r
1035         }\r
1036 \r
1037         void findKnnMatchDispatcher(int k, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream)\r
1038         {\r
1039             findKnnMatch<256>(k, static_cast<PtrStepSzi>(trainIdx), static_cast<PtrStepSzf>(distance), allDist, stream);\r
1040         }\r
1041 \r
1042         ///////////////////////////////////////////////////////////////////////////////\r
1043         // knn match Dispatcher\r
1044 \r
1045         template <typename Dist, typename T, typename Mask>\r
1046         void matchDispatcher(const PtrStepSz<T>& query, const PtrStepSz<T>& train, int k, const Mask& mask,\r
1047             const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist,\r
1048             int cc, cudaStream_t stream)\r
1049         {\r
1050             if (k == 2)\r
1051             {\r
1052                 match2Dispatcher<Dist>(query, train, mask, trainIdx, distance, cc, stream);\r
1053             }\r
1054             else\r
1055             {\r
1056                 calcDistanceDispatcher<Dist>(query, train, mask, allDist, cc, stream);\r
1057                 findKnnMatchDispatcher(k, trainIdx, distance, allDist, cc, stream);\r
1058             }\r
1059         }\r
1060 \r
1061         ///////////////////////////////////////////////////////////////////////////////\r
1062         // knn match caller\r
1063 \r
1064         template <typename T> void matchL1_gpu(const PtrStepSzb& query, const PtrStepSzb& train, int k, const PtrStepSzb& mask,\r
1065             const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist,\r
1066             int cc, cudaStream_t stream)\r
1067         {\r
1068             if (mask.data)\r
1069                 matchDispatcher< L1Dist<T> >(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, SingleMask(mask), trainIdx, distance, allDist, cc, stream);\r
1070             else\r
1071                 matchDispatcher< L1Dist<T> >(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, WithOutMask(), trainIdx, distance, allDist, cc, stream);\r
1072         }\r
1073 \r
1074         template void matchL1_gpu<uchar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1075         //template void matchL1_gpu<schar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1076         template void matchL1_gpu<ushort>(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1077         template void matchL1_gpu<short >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1078         template void matchL1_gpu<int   >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1079         template void matchL1_gpu<float >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1080 \r
1081         template <typename T> void matchL2_gpu(const PtrStepSzb& query, const PtrStepSzb& train, int k, const PtrStepSzb& mask,\r
1082             const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist,\r
1083             int cc, cudaStream_t stream)\r
1084         {\r
1085             if (mask.data)\r
1086                 matchDispatcher<L2Dist>(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, SingleMask(mask), trainIdx, distance, allDist, cc, stream);\r
1087             else\r
1088                 matchDispatcher<L2Dist>(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, WithOutMask(), trainIdx, distance, allDist, cc, stream);\r
1089         }\r
1090 \r
1091         //template void matchL2_gpu<uchar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1092         //template void matchL2_gpu<schar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1093         //template void matchL2_gpu<ushort>(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1094         //template void matchL2_gpu<short >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1095         //template void matchL2_gpu<int   >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1096         template void matchL2_gpu<float >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1097 \r
1098         template <typename T> void matchHamming_gpu(const PtrStepSzb& query, const PtrStepSzb& train, int k, const PtrStepSzb& mask,\r
1099             const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist,\r
1100             int cc, cudaStream_t stream)\r
1101         {\r
1102             if (mask.data)\r
1103                 matchDispatcher<HammingDist>(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, SingleMask(mask), trainIdx, distance, allDist, cc, stream);\r
1104             else\r
1105                 matchDispatcher<HammingDist>(static_cast< PtrStepSz<T> >(query), static_cast< PtrStepSz<T> >(train), k, WithOutMask(), trainIdx, distance, allDist, cc, stream);\r
1106         }\r
1107 \r
1108         template void matchHamming_gpu<uchar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1109         //template void matchHamming_gpu<schar >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1110         template void matchHamming_gpu<ushort>(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1111         //template void matchHamming_gpu<short >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1112         template void matchHamming_gpu<int   >(const PtrStepSzb& queryDescs, const PtrStepSzb& trainDescs, int k, const PtrStepSzb& mask, const PtrStepSzb& trainIdx, const PtrStepSzb& distance, const PtrStepSzf& allDist, int cc, cudaStream_t stream);\r
1113 \r
1114         template <typename T> void match2L1_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks,\r
1115             const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance,\r
1116             int cc, cudaStream_t stream)\r
1117         {\r
1118             if (masks.data)\r
1119                 match2Dispatcher< L1Dist<T> >(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, MaskCollection(masks.data), trainIdx, imgIdx, distance, cc, stream);\r
1120             else\r
1121                 match2Dispatcher< L1Dist<T> >(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, WithOutMask(), trainIdx, imgIdx, distance, cc, stream);\r
1122         }\r
1123 \r
1124         template void match2L1_gpu<uchar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1125         //template void match2L1_gpu<schar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1126         template void match2L1_gpu<ushort>(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1127         template void match2L1_gpu<short >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1128         template void match2L1_gpu<int   >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1129         template void match2L1_gpu<float >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1130 \r
1131         template <typename T> void match2L2_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks,\r
1132             const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance,\r
1133             int cc, cudaStream_t stream)\r
1134         {\r
1135             if (masks.data)\r
1136                 match2Dispatcher<L2Dist>(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, MaskCollection(masks.data), trainIdx, imgIdx, distance, cc, stream);\r
1137             else\r
1138                 match2Dispatcher<L2Dist>(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, WithOutMask(), trainIdx, imgIdx, distance, cc, stream);\r
1139         }\r
1140 \r
1141         //template void match2L2_gpu<uchar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1142         //template void match2L2_gpu<schar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1143         //template void match2L2_gpu<ushort>(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1144         //template void match2L2_gpu<short >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1145         //template void match2L2_gpu<int   >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzi& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1146         template void match2L2_gpu<float >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1147 \r
1148         template <typename T> void match2Hamming_gpu(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks,\r
1149             const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance,\r
1150             int cc, cudaStream_t stream)\r
1151         {\r
1152             if (masks.data)\r
1153                 match2Dispatcher<HammingDist>(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, MaskCollection(masks.data), trainIdx, imgIdx, distance, cc, stream);\r
1154             else\r
1155                 match2Dispatcher<HammingDist>(static_cast< PtrStepSz<T> >(query), (const PtrStepSz<T>*)trains.ptr(), trains.cols, WithOutMask(), trainIdx, imgIdx, distance, cc, stream);\r
1156         }\r
1157 \r
1158         template void match2Hamming_gpu<uchar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1159         //template void match2Hamming_gpu<schar >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1160         template void match2Hamming_gpu<ushort>(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1161         //template void match2Hamming_gpu<short >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1162         template void match2Hamming_gpu<int   >(const PtrStepSzb& query, const PtrStepSzb& trains, const PtrStepSz<PtrStepb>& masks, const PtrStepSzb& trainIdx, const PtrStepSzb& imgIdx, const PtrStepSzb& distance, int cc, cudaStream_t stream);\r
1163     } // namespace bf_knnmatch\r
1164 }}} // namespace cv { namespace gpu { namespace device {\r
1165 \r
1166 \r
1167 #endif /* CUDA_DISABLER */