modules/gpu/src/cuda/global_motion.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 // Copyright (C) 1993-2011, NVIDIA Corporation, all rights reserved.
  16 // Third party copyrights are property of their respective owners.
  17 //
  18 // Redistribution and use in source and binary forms, with or without modification,
  19 // are permitted provided that the following conditions are met:
  20 //
  21 //   * Redistribution's of source code must retain the above copyright notice,
  22 //     this list of conditions and the following disclaimer.
  23 //
  24 //   * Redistribution's in binary form must reproduce the above copyright notice,
  25 //     this list of conditions and the following disclaimer in the documentation
  26 //     and/or other materials provided with the distribution.
  27 //
  28 //   * The name of the copyright holders may not be used to endorse or promote products
  29 //     derived from this software without specific prior written permission.
  30 //
  31 // This software is provided by the copyright holders and contributors "as is" and
  32 // any express or bpied warranties, including, but not limited to, the bpied
  33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  34 // In no event shall the Intel Corporation or contributors be liable for any direct,
  35 // indirect, incidental, special, exemplary, or consequential damages
  36 // (including, but not limited to, procurement of substitute goods or services;
  37 // loss of use, data, or profits; or business interruption) however caused
  38 // and on any theory of liability, whether in contract, strict liability,
  39 // or tort (including negligence or otherwise) arising in any way out of
  40 // the use of this software, even if advised of the possibility of such damage.
  41 //
  42 //M*/
  43
  44 #if !defined CUDA_DISABLER
  45
  46 #include <thrust/device_ptr.h>
  47 #include <thrust/remove.h>
  48 #include <thrust/functional.h>
  49 #include "opencv2/gpu/device/common.hpp"
  50
  51 namespace cv { namespace gpu { namespace device { namespace globmotion {
  52
  53 __constant__ float cml[9];
  54 __constant__ float cmr[9];
  55
  56 int compactPoints(int N, float *points0, float *points1, const uchar *mask)
  57 {
  58     thrust::device_ptr<float2> dpoints0((float2*)points0);
  59     thrust::device_ptr<float2> dpoints1((float2*)points1);
  60     thrust::device_ptr<const uchar> dmask(mask);
  61
  62     return (int)(thrust::remove_if(thrust::make_zip_iterator(thrust::make_tuple(dpoints0, dpoints1)),
  63                              thrust::make_zip_iterator(thrust::make_tuple(dpoints0 + N, dpoints1 + N)),
  64                              dmask, thrust::not1(thrust::identity<uchar>()))
  65            - thrust::make_zip_iterator(make_tuple(dpoints0, dpoints1)));
  66 }
  67
  68
  69 __global__ void calcWobbleSuppressionMapsKernel(
  70         const int left, const int idx, const int right, const int width, const int height,
  71         PtrStepf mapx, PtrStepf mapy)
  72 {
  73     const int x = blockDim.x * blockIdx.x + threadIdx.x;
  74     const int y = blockDim.y * blockIdx.y + threadIdx.y;
  75
  76     if (x < width && y < height)
  77     {
  78         float xl = cml[0]*x + cml[1]*y + cml[2];
  79         float yl = cml[3]*x + cml[4]*y + cml[5];
  80         float izl = 1.f / (cml[6]*x + cml[7]*y + cml[8]);
  81         xl *= izl;
  82         yl *= izl;
  83
  84         float xr = cmr[0]*x + cmr[1]*y + cmr[2];
  85         float yr = cmr[3]*x + cmr[4]*y + cmr[5];
  86         float izr = 1.f / (cmr[6]*x + cmr[7]*y + cmr[8]);
  87         xr *= izr;
  88         yr *= izr;
  89
  90         float wl = idx - left;
  91         float wr = right - idx;
  92         mapx(y,x) = (wr * xl + wl * xr) / (wl + wr);
  93         mapy(y,x) = (wr * yl + wl * yr) / (wl + wr);
  94     }
  95 }
  96
  97
  98 void calcWobbleSuppressionMaps(
  99         int left, int idx, int right, int width, int height,
 100         const float *ml, const float *mr, PtrStepSzf mapx, PtrStepSzf mapy)
 101 {
 102     cudaSafeCall(cudaMemcpyToSymbol(cml, ml, 9*sizeof(float)));
 103     cudaSafeCall(cudaMemcpyToSymbol(cmr, mr, 9*sizeof(float)));
 104
 105     dim3 threads(32, 8);
 106     dim3 grid(divUp(width, threads.x), divUp(height, threads.y));
 107
 108     calcWobbleSuppressionMapsKernel<<<grid, threads>>>(
 109             left, idx, right, width, height, mapx, mapy);
 110
 111     cudaSafeCall(cudaGetLastError());
 112     cudaSafeCall(cudaDeviceSynchronize());
 113 }
 114
 115 }}}}
 116
 117
 118 #endif /* CUDA_DISABLER */