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