modules/gpu/src/cuda/imgproc.cu

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