modules/gpu/src/cuda/imgproc.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 implied warranties, including, but not limited to, the implied\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/vec_traits.hpp"\r
  45 #include "opencv2/gpu/device/vec_math.hpp"\r
  46 #include "opencv2/gpu/device/saturate_cast.hpp"\r
  47 #include "opencv2/gpu/device/border_interpolate.hpp"\r
  48 \r
  49 namespace cv { namespace gpu { namespace device\r
  50 {\r
  51     namespace imgproc\r
  52     {\r
  53         /////////////////////////////////// MeanShiftfiltering ///////////////////////////////////////////////\r
  54 \r
  55         texture<uchar4, 2> tex_meanshift;\r
  56 \r
  57         __device__ short2 do_mean_shift(int x0, int y0, unsigned char* out,\r
  58                                         size_t out_step, int cols, int rows,\r
  59                                         int sp, int sr, int maxIter, float eps)\r
  60         {\r
  61             int isr2 = sr*sr;\r
  62             uchar4 c = tex2D(tex_meanshift, x0, y0 );\r
  63 \r
  64             // iterate meanshift procedure\r
  65             for( int iter = 0; iter < maxIter; iter++ )\r
  66             {\r
  67                 int count = 0;\r
  68                 int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0;\r
  69                 float icount;\r
  70 \r
  71                 //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp)\r
  72                 int minx = x0-sp;\r
  73                 int miny = y0-sp;\r
  74                 int maxx = x0+sp;\r
  75                 int maxy = y0+sp;\r
  76 \r
  77                 for( int y = miny; y <= maxy; y++)\r
  78                 {\r
  79                     int rowCount = 0;\r
  80                     for( int x = minx; x <= maxx; x++ )\r
  81                     {\r
  82                         uchar4 t = tex2D( tex_meanshift, x, y );\r
  83 \r
  84                         int norm2 = (t.x - c.x) * (t.x - c.x) + (t.y - c.y) * (t.y - c.y) + (t.z - c.z) * (t.z - c.z);\r
  85                         if( norm2 <= isr2 )\r
  86                         {\r
  87                             s0 += t.x; s1 += t.y; s2 += t.z;\r
  88                             sx += x; rowCount++;\r
  89                         }\r
  90                     }\r
  91                     count += rowCount;\r
  92                     sy += y*rowCount;\r
  93                 }\r
  94 \r
  95                 if( count == 0 )\r
  96                     break;\r
  97 \r
  98                 icount = 1.f/count;\r
  99                 int x1 = __float2int_rz(sx*icount);\r
 100                 int y1 = __float2int_rz(sy*icount);\r
 101                 s0 = __float2int_rz(s0*icount);\r
 102                 s1 = __float2int_rz(s1*icount);\r
 103                 s2 = __float2int_rz(s2*icount);\r
 104 \r
 105                 int norm2 = (s0 - c.x) * (s0 - c.x) + (s1 - c.y) * (s1 - c.y) + (s2 - c.z) * (s2 - c.z);\r
 106 \r
 107                 bool stopFlag = (x0 == x1 && y0 == y1) || (::abs(x1-x0) + ::abs(y1-y0) + norm2 <= eps);\r
 108 \r
 109                 x0 = x1; y0 = y1;\r
 110                 c.x = s0; c.y = s1; c.z = s2;\r
 111 \r
 112                 if( stopFlag )\r
 113                     break;\r
 114             }\r
 115 \r
 116             int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);\r
 117             *(uchar4*)(out + base) = c;\r
 118 \r
 119             return make_short2((short)x0, (short)y0);\r
 120         }\r
 121 \r
 122         __global__ void meanshift_kernel(unsigned char* out, size_t out_step, int cols, int rows, int sp, int sr, int maxIter, float eps )\r
 123         {\r
 124             int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
 125             int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
 126 \r
 127             if( x0 < cols && y0 < rows )\r
 128                 do_mean_shift(x0, y0, out, out_step, cols, rows, sp, sr, maxIter, eps);\r
 129         }\r
 130 \r
 131         __global__ void meanshiftproc_kernel(unsigned char* outr, size_t outrstep,\r
 132                                              unsigned char* outsp, size_t outspstep,\r
 133                                              int cols, int rows,\r
 134                                              int sp, int sr, int maxIter, float eps)\r
 135         {\r
 136             int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
 137             int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
 138 \r
 139             if( x0 < cols && y0 < rows )\r
 140             {\r
 141                 int basesp = (blockIdx.y * blockDim.y + threadIdx.y) * outspstep + (blockIdx.x * blockDim.x + threadIdx.x) * 2 * sizeof(short);\r
 142                 *(short2*)(outsp + basesp) = do_mean_shift(x0, y0, outr, outrstep, cols, rows, sp, sr, maxIter, eps);\r
 143             }\r
 144         }\r
 145 \r
 146         void meanShiftFiltering_gpu(const DevMem2Db& src, DevMem2Db dst, int sp, int sr, int maxIter, float eps, cudaStream_t stream)\r
 147         {\r
 148             dim3 grid(1, 1, 1);\r
 149             dim3 threads(32, 8, 1);\r
 150             grid.x = divUp(src.cols, threads.x);\r
 151             grid.y = divUp(src.rows, threads.y);\r
 152 \r
 153             cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();\r
 154             cudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) );\r
 155 \r
 156             meanshift_kernel<<< grid, threads, 0, stream >>>( dst.data, dst.step, dst.cols, dst.rows, sp, sr, maxIter, eps );\r
 157             cudaSafeCall( cudaGetLastError() );\r
 158 \r
 159             if (stream == 0)\r
 160                 cudaSafeCall( cudaDeviceSynchronize() );\r
 161 \r
 162             //cudaSafeCall( cudaUnbindTexture( tex_meanshift ) );\r
 163         }\r
 164 \r
 165         void meanShiftProc_gpu(const DevMem2Db& src, DevMem2Db dstr, DevMem2Db dstsp, int sp, int sr, int maxIter, float eps, cudaStream_t stream)\r
 166         {\r
 167             dim3 grid(1, 1, 1);\r
 168             dim3 threads(32, 8, 1);\r
 169             grid.x = divUp(src.cols, threads.x);\r
 170             grid.y = divUp(src.rows, threads.y);\r
 171 \r
 172             cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();\r
 173             cudaSafeCall( cudaBindTexture2D( 0, tex_meanshift, src.data, desc, src.cols, src.rows, src.step ) );\r
 174 \r
 175             meanshiftproc_kernel<<< grid, threads, 0, stream >>>( dstr.data, dstr.step, dstsp.data, dstsp.step, dstr.cols, dstr.rows, sp, sr, maxIter, eps );\r
 176             cudaSafeCall( cudaGetLastError() );\r
 177 \r
 178             if (stream == 0)\r
 179                 cudaSafeCall( cudaDeviceSynchronize() );\r
 180 \r
 181             //cudaSafeCall( cudaUnbindTexture( tex_meanshift ) );\r
 182         }\r
 183 \r
 184         /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////\r
 185 \r
 186         template <typename T>\r
 187         __device__ unsigned int cvtPixel(T d, int ndisp, float S = 1, float V = 1)\r
 188         {\r
 189             unsigned int H = ((ndisp-d) * 240)/ndisp;\r
 190 \r
 191             unsigned int hi = (H/60) % 6;\r
 192             float f = H/60.f - H/60;\r
 193             float p = V * (1 - S);\r
 194             float q = V * (1 - f * S);\r
 195             float t = V * (1 - (1 - f) * S);\r
 196 \r
 197             float3 res;\r
 198 \r
 199             if (hi == 0) //R = V,       G = t,  B = p\r
 200             {\r
 201                 res.x = p;\r
 202                 res.y = t;\r
 203                 res.z = V;\r
 204             }\r
 205 \r
 206             if (hi == 1) // R = q,      G = V,  B = p\r
 207             {\r
 208                 res.x = p;\r
 209                 res.y = V;\r
 210                 res.z = q;\r
 211             }\r
 212 \r
 213             if (hi == 2) // R = p,      G = V,  B = t\r
 214             {\r
 215                 res.x = t;\r
 216                 res.y = V;\r
 217                 res.z = p;\r
 218             }\r
 219 \r
 220             if (hi == 3) // R = p,      G = q,  B = V\r
 221             {\r
 222                 res.x = V;\r
 223                 res.y = q;\r
 224                 res.z = p;\r
 225             }\r
 226 \r
 227             if (hi == 4) // R = t,      G = p,  B = V\r
 228             {\r
 229                 res.x = V;\r
 230                 res.y = p;\r
 231                 res.z = t;\r
 232             }\r
 233 \r
 234             if (hi == 5) // R = V,      G = p,  B = q\r
 235             {\r
 236                 res.x = q;\r
 237                 res.y = p;\r
 238                 res.z = V;\r
 239             }\r
 240             const unsigned int b = (unsigned int)(::max(0.f, ::min(res.x, 1.f)) * 255.f);\r
 241             const unsigned int g = (unsigned int)(::max(0.f, ::min(res.y, 1.f)) * 255.f);\r
 242             const unsigned int r = (unsigned int)(::max(0.f, ::min(res.z, 1.f)) * 255.f);\r
 243             const unsigned int a = 255U;\r
 244 \r
 245             return (a << 24) + (r << 16) + (g << 8) + b;\r
 246         }\r
 247 \r
 248         __global__ void drawColorDisp(uchar* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)\r
 249         {\r
 250             const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;\r
 251             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 252 \r
 253             if(x < width && y < height)\r
 254             {\r
 255                 uchar4 d4 = *(uchar4*)(disp + y * disp_step + x);\r
 256 \r
 257                 uint4 res;\r
 258                 res.x = cvtPixel(d4.x, ndisp);\r
 259                 res.y = cvtPixel(d4.y, ndisp);\r
 260                 res.z = cvtPixel(d4.z, ndisp);\r
 261                 res.w = cvtPixel(d4.w, ndisp);\r
 262 \r
 263                 uint4* line = (uint4*)(out_image + y * out_step);\r
 264                 line[x >> 2] = res;\r
 265             }\r
 266         }\r
 267 \r
 268         __global__ void drawColorDisp(short* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)\r
 269         {\r
 270             const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;\r
 271             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 272 \r
 273             if(x < width && y < height)\r
 274             {\r
 275                 short2 d2 = *(short2*)(disp + y * disp_step + x);\r
 276 \r
 277                 uint2 res;\r
 278                 res.x = cvtPixel(d2.x, ndisp);\r
 279                 res.y = cvtPixel(d2.y, ndisp);\r
 280 \r
 281                 uint2* line = (uint2*)(out_image + y * out_step);\r
 282                 line[x >> 1] = res;\r
 283             }\r
 284         }\r
 285 \r
 286 \r
 287         void drawColorDisp_gpu(const DevMem2Db& src, const DevMem2Db& dst, int ndisp, const cudaStream_t& stream)\r
 288         {\r
 289             dim3 threads(16, 16, 1);\r
 290             dim3 grid(1, 1, 1);\r
 291             grid.x = divUp(src.cols, threads.x << 2);\r
 292             grid.y = divUp(src.rows, threads.y);\r
 293 \r
 294             drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp);\r
 295             cudaSafeCall( cudaGetLastError() );\r
 296 \r
 297             if (stream == 0)\r
 298                 cudaSafeCall( cudaDeviceSynchronize() );\r
 299         }\r
 300 \r
 301         void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2Db& dst, int ndisp, const cudaStream_t& stream)\r
 302         {\r
 303             dim3 threads(32, 8, 1);\r
 304             dim3 grid(1, 1, 1);\r
 305             grid.x = divUp(src.cols, threads.x << 1);\r
 306             grid.y = divUp(src.rows, threads.y);\r
 307 \r
 308             drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp);\r
 309             cudaSafeCall( cudaGetLastError() );\r
 310 \r
 311             if (stream == 0)\r
 312                 cudaSafeCall( cudaDeviceSynchronize() );\r
 313         }\r
 314 \r
 315         /////////////////////////////////// reprojectImageTo3D ///////////////////////////////////////////////\r
 316 \r
 317         __constant__ float cq[16];\r
 318 \r
 319         template <typename T, typename D>\r
 320         __global__ void reprojectImageTo3D(const DevMem2D_<T> disp, PtrStep<D> xyz)\r
 321         {\r
 322             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 323             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 324 \r
 325             if (y >= disp.rows || x >= disp.cols)\r
 326                 return;\r
 327 \r
 328             const float qx = x * cq[ 0] + y * cq[ 1] + cq[ 3];\r
 329             const float qy = x * cq[ 4] + y * cq[ 5] + cq[ 7];\r
 330             const float qz = x * cq[ 8] + y * cq[ 9] + cq[11];\r
 331             const float qw = x * cq[12] + y * cq[13] + cq[15];\r
 332 \r
 333             const T d = disp(y, x);\r
 334 \r
 335             const float iW = 1.f / (qw + cq[14] * d);\r
 336 \r
 337             D v = VecTraits<D>::all(1.0f);\r
 338             v.x = (qx + cq[2] * d) * iW;\r
 339             v.y = (qy + cq[6] * d) * iW;\r
 340             v.z = (qz + cq[10] * d) * iW;\r
 341 \r
 342             xyz(y, x) = v;\r
 343         }\r
 344 \r
 345         template <typename T, typename D>\r
 346         void reprojectImageTo3D_gpu(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream)\r
 347         {\r
 348             dim3 block(32, 8);\r
 349             dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));\r
 350 \r
 351             cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );\r
 352 \r
 353             reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((DevMem2D_<T>)disp, (DevMem2D_<D>)xyz);\r
 354             cudaSafeCall( cudaGetLastError() );\r
 355 \r
 356             if (stream == 0)\r
 357                 cudaSafeCall( cudaDeviceSynchronize() );\r
 358         }\r
 359 \r
 360         template void reprojectImageTo3D_gpu<uchar, float3>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);\r
 361         template void reprojectImageTo3D_gpu<uchar, float4>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);\r
 362         template void reprojectImageTo3D_gpu<short, float3>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);\r
 363         template void reprojectImageTo3D_gpu<short, float4>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);\r
 364 \r
 365         /////////////////////////////////////////// Corner Harris /////////////////////////////////////////////////\r
 366 \r
 367         texture<float, cudaTextureType2D, cudaReadModeElementType> harrisDxTex(0, cudaFilterModePoint, cudaAddressModeClamp);\r
 368         texture<float, cudaTextureType2D, cudaReadModeElementType> harrisDyTex(0, cudaFilterModePoint, cudaAddressModeClamp);\r
 369 \r
 370         __global__ void cornerHarris_kernel(const int block_size, const float k, DevMem2Df dst)\r
 371         {\r
 372             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 373             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 374 \r
 375             if (x < dst.cols && y < dst.rows)\r
 376             {\r
 377                 float a = 0.f;\r
 378                 float b = 0.f;\r
 379                 float c = 0.f;\r
 380 \r
 381                 const int ibegin = y - (block_size / 2);\r
 382                 const int jbegin = x - (block_size / 2);\r
 383                 const int iend = ibegin + block_size;\r
 384                 const int jend = jbegin + block_size;\r
 385 \r
 386                 for (int i = ibegin; i < iend; ++i)\r
 387                 {\r
 388                     for (int j = jbegin; j < jend; ++j)\r
 389                     {\r
 390                         float dx = tex2D(harrisDxTex, j, i);\r
 391                         float dy = tex2D(harrisDyTex, j, i);\r
 392 \r
 393                         a += dx * dx;\r
 394                         b += dx * dy;\r
 395                         c += dy * dy;\r
 396                     }\r
 397                 }\r
 398 \r
 399                 dst(y, x) = a * c - b * b - k * (a + c) * (a + c);\r
 400             }\r
 401         }\r
 402 \r
 403         template <typename BR, typename BC>\r
 404         __global__ void cornerHarris_kernel(const int block_size, const float k, DevMem2Df dst, const BR border_row, const BC border_col)\r
 405         {\r
 406             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 407             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 408 \r
 409             if (x < dst.cols && y < dst.rows)\r
 410             {\r
 411                 float a = 0.f;\r
 412                 float b = 0.f;\r
 413                 float c = 0.f;\r
 414 \r
 415                 const int ibegin = y - (block_size / 2);\r
 416                 const int jbegin = x - (block_size / 2);\r
 417                 const int iend = ibegin + block_size;\r
 418                 const int jend = jbegin + block_size;\r
 419 \r
 420                 for (int i = ibegin; i < iend; ++i)\r
 421                 {\r
 422                     const int y = border_col.idx_row(i);\r
 423 \r
 424                     for (int j = jbegin; j < jend; ++j)\r
 425                     {\r
 426                         const int x = border_row.idx_col(j);\r
 427 \r
 428                         float dx = tex2D(harrisDxTex, x, y);\r
 429                         float dy = tex2D(harrisDyTex, x, y);\r
 430 \r
 431                         a += dx * dx;\r
 432                         b += dx * dy;\r
 433                         c += dy * dy;\r
 434                     }\r
 435                 }\r
 436 \r
 437                 dst(y, x) = a * c - b * b - k * (a + c) * (a + c);\r
 438             }\r
 439         }\r
 440 \r
 441         void cornerHarris_gpu(int block_size, float k, DevMem2Df Dx, DevMem2Df Dy, DevMem2Df dst, int border_type, cudaStream_t stream)\r
 442         {\r
 443             dim3 block(32, 8);\r
 444             dim3 grid(divUp(Dx.cols, block.x), divUp(Dx.rows, block.y));\r
 445 \r
 446             bindTexture(&harrisDxTex, Dx);\r
 447             bindTexture(&harrisDyTex, Dy);\r
 448 \r
 449             switch (border_type)\r
 450             {\r
 451             case BORDER_REFLECT101_GPU:\r
 452                 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));\r
 453                 break;\r
 454 \r
 455             case BORDER_REFLECT_GPU:\r
 456                 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));\r
 457                 break;\r
 458 \r
 459             case BORDER_REPLICATE_GPU:\r
 460                 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst);\r
 461                 break;\r
 462             }\r
 463 \r
 464             cudaSafeCall( cudaGetLastError() );\r
 465 \r
 466             if (stream == 0)\r
 467                 cudaSafeCall( cudaDeviceSynchronize() );\r
 468         }\r
 469 \r
 470         /////////////////////////////////////////// Corner Min Eigen Val /////////////////////////////////////////////////\r
 471 \r
 472         texture<float, cudaTextureType2D, cudaReadModeElementType> minEigenValDxTex(0, cudaFilterModePoint, cudaAddressModeClamp);\r
 473         texture<float, cudaTextureType2D, cudaReadModeElementType> minEigenValDyTex(0, cudaFilterModePoint, cudaAddressModeClamp);\r
 474 \r
 475         __global__ void cornerMinEigenVal_kernel(const int block_size, DevMem2Df dst)\r
 476         {\r
 477             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 478             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 479 \r
 480             if (x < dst.cols && y < dst.rows)\r
 481             {\r
 482                 float a = 0.f;\r
 483                 float b = 0.f;\r
 484                 float c = 0.f;\r
 485 \r
 486                 const int ibegin = y - (block_size / 2);\r
 487                 const int jbegin = x - (block_size / 2);\r
 488                 const int iend = ibegin + block_size;\r
 489                 const int jend = jbegin + block_size;\r
 490 \r
 491                 for (int i = ibegin; i < iend; ++i)\r
 492                 {\r
 493                     for (int j = jbegin; j < jend; ++j)\r
 494                     {\r
 495                         float dx = tex2D(minEigenValDxTex, j, i);\r
 496                         float dy = tex2D(minEigenValDyTex, j, i);\r
 497 \r
 498                         a += dx * dx;\r
 499                         b += dx * dy;\r
 500                         c += dy * dy;\r
 501                     }\r
 502                 }\r
 503 \r
 504                 a *= 0.5f;\r
 505                 c *= 0.5f;\r
 506 \r
 507                 dst(y, x) = (a + c) - sqrtf((a - c) * (a - c) + b * b);\r
 508             }\r
 509         }\r
 510 \r
 511 \r
 512         template <typename BR, typename BC>\r
 513         __global__ void cornerMinEigenVal_kernel(const int block_size, DevMem2Df dst, const BR border_row, const BC border_col)\r
 514         {\r
 515             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 516             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 517 \r
 518             if (x < dst.cols && y < dst.rows)\r
 519             {\r
 520                 float a = 0.f;\r
 521                 float b = 0.f;\r
 522                 float c = 0.f;\r
 523 \r
 524                 const int ibegin = y - (block_size / 2);\r
 525                 const int jbegin = x - (block_size / 2);\r
 526                 const int iend = ibegin + block_size;\r
 527                 const int jend = jbegin + block_size;\r
 528 \r
 529                 for (int i = ibegin; i < iend; ++i)\r
 530                 {\r
 531                     int y = border_col.idx_row(i);\r
 532 \r
 533                     for (int j = jbegin; j < jend; ++j)\r
 534                     {\r
 535                         int x = border_row.idx_col(j);\r
 536 \r
 537                         float dx = tex2D(minEigenValDxTex, x, y);\r
 538                         float dy = tex2D(minEigenValDyTex, x, y);\r
 539 \r
 540                         a += dx * dx;\r
 541                         b += dx * dy;\r
 542                         c += dy * dy;\r
 543                     }\r
 544                 }\r
 545 \r
 546                 a *= 0.5f;\r
 547                 c *= 0.5f;\r
 548 \r
 549                 dst(y, x) = (a + c) - sqrtf((a - c) * (a - c) + b * b);\r
 550             }\r
 551         }\r
 552 \r
 553         void cornerMinEigenVal_gpu(int block_size, DevMem2Df Dx, DevMem2Df Dy, DevMem2Df dst, int border_type, cudaStream_t stream)\r
 554         {\r
 555             dim3 block(32, 8);\r
 556             dim3 grid(divUp(Dx.cols, block.x), divUp(Dx.rows, block.y));\r
 557 \r
 558             bindTexture(&minEigenValDxTex, Dx);\r
 559             bindTexture(&minEigenValDyTex, Dy);\r
 560 \r
 561             switch (border_type)\r
 562             {\r
 563             case BORDER_REFLECT101_GPU:\r
 564                 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));\r
 565                 break;\r
 566 \r
 567             case BORDER_REFLECT_GPU:\r
 568                 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));\r
 569                 break;\r
 570 \r
 571             case BORDER_REPLICATE_GPU:\r
 572                 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst);\r
 573                 break;\r
 574             }\r
 575 \r
 576             cudaSafeCall( cudaGetLastError() );\r
 577 \r
 578             if (stream == 0)\r
 579                 cudaSafeCall(cudaDeviceSynchronize());\r
 580         }\r
 581 \r
 582         ////////////////////////////// Column Sum //////////////////////////////////////\r
 583 \r
 584         __global__ void column_sumKernel_32F(int cols, int rows, const PtrStepb src, const PtrStepb dst)\r
 585         {\r
 586             int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 587 \r
 588             if (x < cols)\r
 589             {\r
 590                 const unsigned char* src_data = src.data + x * sizeof(float);\r
 591                 unsigned char* dst_data = dst.data + x * sizeof(float);\r
 592 \r
 593                 float sum = 0.f;\r
 594                 for (int y = 0; y < rows; ++y)\r
 595                 {\r
 596                     sum += *(const float*)src_data;\r
 597                     *(float*)dst_data = sum;\r
 598                     src_data += src.step;\r
 599                     dst_data += dst.step;\r
 600                 }\r
 601             }\r
 602         }\r
 603 \r
 604 \r
 605         void columnSum_32F(const DevMem2Db src, const DevMem2Db dst)\r
 606         {\r
 607             dim3 threads(256);\r
 608             dim3 grid(divUp(src.cols, threads.x));\r
 609 \r
 610             column_sumKernel_32F<<<grid, threads>>>(src.cols, src.rows, src, dst);\r
 611             cudaSafeCall( cudaGetLastError() );\r
 612 \r
 613             cudaSafeCall( cudaDeviceSynchronize() );\r
 614         }\r
 615 \r
 616 \r
 617         //////////////////////////////////////////////////////////////////////////\r
 618         // mulSpectrums\r
 619 \r
 620         __global__ void mulSpectrumsKernel(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, DevMem2D_<cufftComplex> c)\r
 621         {\r
 622             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 623             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 624 \r
 625             if (x < c.cols && y < c.rows)\r
 626             {\r
 627                 c.ptr(y)[x] = cuCmulf(a.ptr(y)[x], b.ptr(y)[x]);\r
 628             }\r
 629         }\r
 630 \r
 631 \r
 632         void mulSpectrums(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, DevMem2D_<cufftComplex> c, cudaStream_t stream)\r
 633         {\r
 634             dim3 threads(256);\r
 635             dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y));\r
 636 \r
 637             mulSpectrumsKernel<<<grid, threads, 0, stream>>>(a, b, c);\r
 638             cudaSafeCall( cudaGetLastError() );\r
 639 \r
 640             if (stream == 0)\r
 641                 cudaSafeCall( cudaDeviceSynchronize() );\r
 642         }\r
 643 \r
 644 \r
 645         //////////////////////////////////////////////////////////////////////////\r
 646         // mulSpectrums_CONJ\r
 647 \r
 648         __global__ void mulSpectrumsKernel_CONJ(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, DevMem2D_<cufftComplex> c)\r
 649         {\r
 650             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 651             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 652 \r
 653             if (x < c.cols && y < c.rows)\r
 654             {\r
 655                 c.ptr(y)[x] = cuCmulf(a.ptr(y)[x], cuConjf(b.ptr(y)[x]));\r
 656             }\r
 657         }\r
 658 \r
 659 \r
 660         void mulSpectrums_CONJ(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, DevMem2D_<cufftComplex> c, cudaStream_t stream)\r
 661         {\r
 662             dim3 threads(256);\r
 663             dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y));\r
 664 \r
 665             mulSpectrumsKernel_CONJ<<<grid, threads, 0, stream>>>(a, b, c);\r
 666             cudaSafeCall( cudaGetLastError() );\r
 667 \r
 668             if (stream == 0)\r
 669                 cudaSafeCall( cudaDeviceSynchronize() );\r
 670         }\r
 671 \r
 672 \r
 673         //////////////////////////////////////////////////////////////////////////\r
 674         // mulAndScaleSpectrums\r
 675 \r
 676         __global__ void mulAndScaleSpectrumsKernel(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, float scale, DevMem2D_<cufftComplex> c)\r
 677         {\r
 678             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 679             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 680 \r
 681             if (x < c.cols && y < c.rows)\r
 682             {\r
 683                 cufftComplex v = cuCmulf(a.ptr(y)[x], b.ptr(y)[x]);\r
 684                 c.ptr(y)[x] = make_cuFloatComplex(cuCrealf(v) * scale, cuCimagf(v) * scale);\r
 685             }\r
 686         }\r
 687 \r
 688 \r
 689         void mulAndScaleSpectrums(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, float scale, DevMem2D_<cufftComplex> c, cudaStream_t stream)\r
 690         {\r
 691             dim3 threads(256);\r
 692             dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y));\r
 693 \r
 694             mulAndScaleSpectrumsKernel<<<grid, threads, 0, stream>>>(a, b, scale, c);\r
 695             cudaSafeCall( cudaGetLastError() );\r
 696 \r
 697             if (stream)\r
 698                 cudaSafeCall( cudaDeviceSynchronize() );\r
 699         }\r
 700 \r
 701 \r
 702         //////////////////////////////////////////////////////////////////////////\r
 703         // mulAndScaleSpectrums_CONJ\r
 704 \r
 705         __global__ void mulAndScaleSpectrumsKernel_CONJ(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, float scale, DevMem2D_<cufftComplex> c)\r
 706         {\r
 707             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 708             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 709 \r
 710             if (x < c.cols && y < c.rows)\r
 711             {\r
 712                 cufftComplex v = cuCmulf(a.ptr(y)[x], cuConjf(b.ptr(y)[x]));\r
 713                 c.ptr(y)[x] = make_cuFloatComplex(cuCrealf(v) * scale, cuCimagf(v) * scale);\r
 714             }\r
 715         }\r
 716 \r
 717 \r
 718         void mulAndScaleSpectrums_CONJ(const PtrStep<cufftComplex> a, const PtrStep<cufftComplex> b, float scale, DevMem2D_<cufftComplex> c, cudaStream_t stream)\r
 719         {\r
 720             dim3 threads(256);\r
 721             dim3 grid(divUp(c.cols, threads.x), divUp(c.rows, threads.y));\r
 722 \r
 723             mulAndScaleSpectrumsKernel_CONJ<<<grid, threads, 0, stream>>>(a, b, scale, c);\r
 724             cudaSafeCall( cudaGetLastError() );\r
 725 \r
 726             if (stream == 0)\r
 727                 cudaSafeCall( cudaDeviceSynchronize() );\r
 728         }\r
 729 \r
 730         //////////////////////////////////////////////////////////////////////////\r
 731         // buildWarpMaps\r
 732 \r
 733         // TODO use intrinsics like __sinf and so on\r
 734 \r
 735         namespace build_warp_maps\r
 736         {\r
 737 \r
 738             __constant__ float ck_rinv[9];\r
 739             __constant__ float cr_kinv[9];\r
 740             __constant__ float ct[3];\r
 741             __constant__ float cscale;\r
 742         }\r
 743 \r
 744 \r
 745         class PlaneMapper\r
 746         {\r
 747         public:\r
 748             static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)\r
 749             {\r
 750                 using namespace build_warp_maps;\r
 751 \r
 752                 float x_ = u / cscale - ct[0];\r
 753                 float y_ = v / cscale - ct[1];\r
 754 \r
 755                 float z;\r
 756                 x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]);\r
 757                 y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]);\r
 758                 z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]);\r
 759 \r
 760                 x /= z;\r
 761                 y /= z;\r
 762             }\r
 763         };\r
 764 \r
 765 \r
 766         class CylindricalMapper\r
 767         {\r
 768         public:\r
 769             static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)\r
 770             {\r
 771                 using namespace build_warp_maps;\r
 772 \r
 773                 u /= cscale;\r
 774                 float x_ = ::sinf(u);\r
 775                 float y_ = v / cscale;\r
 776                 float z_ = ::cosf(u);\r
 777 \r
 778                 float z;\r
 779                 x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_;\r
 780                 y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_;\r
 781                 z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_;\r
 782 \r
 783                 if (z > 0) { x /= z; y /= z; }\r
 784                 else x = y = -1;\r
 785             }\r
 786         };\r
 787 \r
 788 \r
 789         class SphericalMapper\r
 790         {\r
 791         public:\r
 792             static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)\r
 793             {\r
 794                 using namespace build_warp_maps;\r
 795 \r
 796                 v /= cscale;\r
 797                 u /= cscale;\r
 798 \r
 799                 float sinv = ::sinf(v);\r
 800                 float x_ = sinv * ::sinf(u);\r
 801                 float y_ = -::cosf(v);\r
 802                 float z_ = sinv * ::cosf(u);\r
 803 \r
 804                 float z;\r
 805                 x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_;\r
 806                 y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_;\r
 807                 z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_;\r
 808 \r
 809                 if (z > 0) { x /= z; y /= z; }\r
 810                 else x = y = -1;\r
 811             }\r
 812         };\r
 813 \r
 814 \r
 815         template <typename Mapper>\r
 816         __global__ void buildWarpMapsKernel(int tl_u, int tl_v, int cols, int rows,\r
 817                                             PtrStepf map_x, PtrStepf map_y)\r
 818         {\r
 819             int du = blockIdx.x * blockDim.x + threadIdx.x;\r
 820             int dv = blockIdx.y * blockDim.y + threadIdx.y;\r
 821             if (du < cols && dv < rows)\r
 822             {\r
 823                 float u = tl_u + du;\r
 824                 float v = tl_v + dv;\r
 825                 float x, y;\r
 826                 Mapper::mapBackward(u, v, x, y);\r
 827                 map_x.ptr(dv)[du] = x;\r
 828                 map_y.ptr(dv)[du] = y;\r
 829             }\r
 830         }\r
 831 \r
 832 \r
 833         void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 834                                 const float k_rinv[9], const float r_kinv[9], const float t[3],\r
 835                                 float scale, cudaStream_t stream)\r
 836         {\r
 837             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
 838             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
 839             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float)));\r
 840             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));\r
 841 \r
 842             int cols = map_x.cols;\r
 843             int rows = map_x.rows;\r
 844 \r
 845             dim3 threads(32, 8);\r
 846             dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
 847 \r
 848             buildWarpMapsKernel<PlaneMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);\r
 849             cudaSafeCall(cudaGetLastError());\r
 850             if (stream == 0)\r
 851                 cudaSafeCall(cudaDeviceSynchronize());\r
 852         }\r
 853 \r
 854 \r
 855         void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 856                                       const float k_rinv[9], const float r_kinv[9], float scale,\r
 857                                       cudaStream_t stream)\r
 858         {\r
 859             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
 860             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
 861             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));\r
 862 \r
 863             int cols = map_x.cols;\r
 864             int rows = map_x.rows;\r
 865 \r
 866             dim3 threads(32, 8);\r
 867             dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
 868 \r
 869             buildWarpMapsKernel<CylindricalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);\r
 870             cudaSafeCall(cudaGetLastError());\r
 871             if (stream == 0)\r
 872                 cudaSafeCall(cudaDeviceSynchronize());\r
 873         }\r
 874 \r
 875 \r
 876         void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
 877                                     const float k_rinv[9], const float r_kinv[9], float scale,\r
 878                                     cudaStream_t stream)\r
 879         {\r
 880             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
 881             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
 882             cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));\r
 883 \r
 884             int cols = map_x.cols;\r
 885             int rows = map_x.rows;\r
 886 \r
 887             dim3 threads(32, 8);\r
 888             dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
 889 \r
 890             buildWarpMapsKernel<SphericalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);\r
 891             cudaSafeCall(cudaGetLastError());\r
 892             if (stream == 0)\r
 893                 cudaSafeCall(cudaDeviceSynchronize());\r
 894         }\r
 895 \r
 896         //////////////////////////////////////////////////////////////////////////\r
 897         // filter2D\r
 898 \r
 899         #define FILTER2D_MAX_KERNEL_SIZE 16\r
 900 \r
 901         __constant__ float c_filter2DKernel[FILTER2D_MAX_KERNEL_SIZE * FILTER2D_MAX_KERNEL_SIZE];\r
 902 \r
 903         template <class SrcT, typename D>\r
 904         __global__ void filter2D(const SrcT src, DevMem2D_<D> dst, const int kWidth, const int kHeight, const int anchorX, const int anchorY)\r
 905         {\r
 906             typedef typename TypeVec<float, VecTraits<D>::cn>::vec_type sum_t;\r
 907 \r
 908             const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
 909             const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 910 \r
 911             if (x >= dst.cols || y >= dst.rows)\r
 912                 return;\r
 913 \r
 914             sum_t res = VecTraits<sum_t>::all(0);\r
 915             int kInd = 0;\r
 916 \r
 917             for (int i = 0; i < kHeight; ++i)\r
 918             {\r
 919                 for (int j = 0; j < kWidth; ++j)\r
 920                     res = res + src(y - anchorY + i, x - anchorX + j) * c_filter2DKernel[kInd++];\r
 921             }\r
 922 \r
 923             dst(y, x) = saturate_cast<D>(res);\r
 924         }\r
 925 \r
 926         template <typename T, typename D, template <typename> class Brd> struct Filter2DCaller;\r
 927 \r
 928         #define IMPLEMENT_FILTER2D_TEX_READER(type) \\r
 929             texture< type , cudaTextureType2D, cudaReadModeElementType> tex_filter2D_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \\r
 930             struct tex_filter2D_ ## type ## _reader \\r
 931             { \\r
 932                 typedef type elem_type; \\r
 933                 typedef int index_type; \\r
 934                 const int xoff; \\r
 935                 const int yoff; \\r
 936                 tex_filter2D_ ## type ## _reader (int xoff_, int yoff_) : xoff(xoff_), yoff(yoff_) {} \\r
 937                 __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \\r
 938                 { \\r
 939                     return tex2D(tex_filter2D_ ## type , x + xoff, y + yoff); \\r
 940                 } \\r
 941             }; \\r
 942             template <typename D, template <typename> class Brd> struct Filter2DCaller< type , D, Brd> \\r
 943             { \\r
 944                 static void call(const DevMem2D_< type > srcWhole, int xoff, int yoff, DevMem2D_<D> dst, \\r
 945                     int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream) \\r
 946                 { \\r
 947                     typedef typename TypeVec<float, VecTraits< type >::cn>::vec_type work_type; \\r
 948                     dim3 block(16, 16); \\r
 949                     dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \\r
 950                     bindTexture(&tex_filter2D_ ## type , srcWhole); \\r
 951                     tex_filter2D_ ## type ##_reader texSrc(xoff, yoff); \\r
 952                     Brd<work_type> brd(dst.rows, dst.cols, VecTraits<work_type>::make(borderValue)); \\r
 953                     BorderReader< tex_filter2D_ ## type ##_reader, Brd<work_type> > brdSrc(texSrc, brd); \\r
 954                     filter2D<<<grid, block, 0, stream>>>(brdSrc, dst, kWidth, kHeight, anchorX, anchorY); \\r
 955                     cudaSafeCall( cudaGetLastError() ); \\r
 956                     if (stream == 0) \\r
 957                         cudaSafeCall( cudaDeviceSynchronize() ); \\r
 958                 } \\r
 959             };\r
 960 \r
 961         IMPLEMENT_FILTER2D_TEX_READER(uchar);\r
 962         IMPLEMENT_FILTER2D_TEX_READER(uchar4);\r
 963 \r
 964         IMPLEMENT_FILTER2D_TEX_READER(ushort);\r
 965         IMPLEMENT_FILTER2D_TEX_READER(ushort4);\r
 966 \r
 967         IMPLEMENT_FILTER2D_TEX_READER(float);\r
 968         IMPLEMENT_FILTER2D_TEX_READER(float4);\r
 969 \r
 970         #undef IMPLEMENT_FILTER2D_TEX_READER\r
 971 \r
 972         template <typename T, typename D>\r
 973         void filter2D_gpu(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, \r
 974                           int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, \r
 975                           int borderMode, const float* borderValue, cudaStream_t stream)\r
 976         {\r
 977             typedef void (*func_t)(const DevMem2D_<T> srcWhole, int xoff, int yoff, DevMem2D_<D> dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream);\r
 978             static const func_t funcs[] = \r
 979             {\r
 980                 Filter2DCaller<T, D, BrdReflect101>::call,\r
 981                 Filter2DCaller<T, D, BrdReplicate>::call,\r
 982                 Filter2DCaller<T, D, BrdConstant>::call,\r
 983                 Filter2DCaller<T, D, BrdReflect>::call,\r
 984                 Filter2DCaller<T, D, BrdWrap>::call\r
 985             };\r
 986 \r
 987             cudaSafeCall(cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
 988 \r
 989             funcs[borderMode](static_cast< DevMem2D_<T> >(srcWhole), ofsX, ofsY, static_cast< DevMem2D_<D> >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream);\r
 990         }\r
 991 \r
 992         template void filter2D_gpu<uchar, uchar>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 993         template void filter2D_gpu<uchar4, uchar4>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 994         template void filter2D_gpu<ushort, ushort>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 995         template void filter2D_gpu<ushort4, ushort4>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 996         template void filter2D_gpu<float, float>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 997         template void filter2D_gpu<float4, float4>(DevMem2Db srcWhole, int ofsX, int ofsY, DevMem2Db dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);\r
 998     } // namespace imgproc\r
 999 }}} // namespace cv { namespace gpu { namespace device {\r