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 BEGIN_OPENCV_DEVICE_NAMESPACE\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 \r
1159 } // namespace bf_knnmatch\r
1160 \r
1161 END_OPENCV_DEVICE_NAMESPACE\r