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 #include "opencv2/opencv_modules.hpp"
44 
45 #if defined(HAVE_OPENCV_CUDAARITHM) && defined(HAVE_OPENCV_CUDAWARPING) && defined(HAVE_OPENCV_CUDAFILTERS)
46 
47 #include "opencv2/core/cuda/common.hpp"
48 #include "opencv2/core/cuda/transform.hpp"
49 #include "opencv2/core/cuda/vec_traits.hpp"
50 #include "opencv2/core/cuda/vec_math.hpp"
51 
52 using namespace cv::cuda;
53 using namespace cv::cuda::device;
54 
55 namespace btv_l1_cudev
56 {
57     void buildMotionMaps(PtrStepSzf forwardMotionX, PtrStepSzf forwardMotionY,
58                          PtrStepSzf backwardMotionX, PtrStepSzf bacwardMotionY,
59                          PtrStepSzf forwardMapX, PtrStepSzf forwardMapY,
60                          PtrStepSzf backwardMapX, PtrStepSzf backwardMapY);
61 
62     template <int cn>
63     void upscale(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream);
64 
65     void diffSign(PtrStepSzf src1, PtrStepSzf src2, PtrStepSzf dst, cudaStream_t stream);
66 
67     void loadBtvWeights(const float* weights, size_t count);
68     template <int cn> void calcBtvRegularization(PtrStepSzb src, PtrStepSzb dst, int ksize);
69 }
70 
71 namespace btv_l1_cudev
72 {
buildMotionMapsKernel(const PtrStepSzf forwardMotionX,const PtrStepf forwardMotionY,PtrStepf backwardMotionX,PtrStepf backwardMotionY,PtrStepf forwardMapX,PtrStepf forwardMapY,PtrStepf backwardMapX,PtrStepf backwardMapY)73     __global__ void buildMotionMapsKernel(const PtrStepSzf forwardMotionX, const PtrStepf forwardMotionY,
74                                           PtrStepf backwardMotionX, PtrStepf backwardMotionY,
75                                           PtrStepf forwardMapX, PtrStepf forwardMapY,
76                                           PtrStepf backwardMapX, PtrStepf backwardMapY)
77     {
78         const int x = blockIdx.x * blockDim.x + threadIdx.x;
79         const int y = blockIdx.y * blockDim.y + threadIdx.y;
80 
81         if (x >= forwardMotionX.cols || y >= forwardMotionX.rows)
82             return;
83 
84         const float fx = forwardMotionX(y, x);
85         const float fy = forwardMotionY(y, x);
86 
87         const float bx = backwardMotionX(y, x);
88         const float by = backwardMotionY(y, x);
89 
90         forwardMapX(y, x) = x + bx;
91         forwardMapY(y, x) = y + by;
92 
93         backwardMapX(y, x) = x + fx;
94         backwardMapY(y, x) = y + fy;
95     }
96 
buildMotionMaps(PtrStepSzf forwardMotionX,PtrStepSzf forwardMotionY,PtrStepSzf backwardMotionX,PtrStepSzf bacwardMotionY,PtrStepSzf forwardMapX,PtrStepSzf forwardMapY,PtrStepSzf backwardMapX,PtrStepSzf backwardMapY)97     void buildMotionMaps(PtrStepSzf forwardMotionX, PtrStepSzf forwardMotionY,
98                          PtrStepSzf backwardMotionX, PtrStepSzf bacwardMotionY,
99                          PtrStepSzf forwardMapX, PtrStepSzf forwardMapY,
100                          PtrStepSzf backwardMapX, PtrStepSzf backwardMapY)
101     {
102         const dim3 block(32, 8);
103         const dim3 grid(divUp(forwardMapX.cols, block.x), divUp(forwardMapX.rows, block.y));
104 
105         buildMotionMapsKernel<<<grid, block>>>(forwardMotionX, forwardMotionY,
106                                                backwardMotionX, bacwardMotionY,
107                                                forwardMapX, forwardMapY,
108                                                backwardMapX, backwardMapY);
109         cudaSafeCall( cudaGetLastError() );
110 
111         cudaSafeCall( cudaDeviceSynchronize() );
112     }
113 
114     template <typename T>
upscaleKernel(const PtrStepSz<T> src,PtrStep<T> dst,const int scale)115     __global__ void upscaleKernel(const PtrStepSz<T> src, PtrStep<T> dst, const int scale)
116     {
117         const int x = blockIdx.x * blockDim.x + threadIdx.x;
118         const int y = blockIdx.y * blockDim.y + threadIdx.y;
119 
120         if (x >= src.cols || y >= src.rows)
121             return;
122 
123         dst(y * scale, x * scale) = src(y, x);
124     }
125 
126     template <int cn>
upscale(const PtrStepSzb src,PtrStepSzb dst,int scale,cudaStream_t stream)127     void upscale(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream)
128     {
129         typedef typename TypeVec<float, cn>::vec_type src_t;
130 
131         const dim3 block(32, 8);
132         const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
133 
134         upscaleKernel<src_t><<<grid, block, 0, stream>>>((PtrStepSz<src_t>) src, (PtrStepSz<src_t>) dst, scale);
135         cudaSafeCall( cudaGetLastError() );
136 
137         if (stream == 0)
138             cudaSafeCall( cudaDeviceSynchronize() );
139     }
140 
141     template void upscale<1>(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream);
142     template void upscale<3>(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream);
143     template void upscale<4>(const PtrStepSzb src, PtrStepSzb dst, int scale, cudaStream_t stream);
144 
diffSign(float a,float b)145     __device__ __forceinline__ float diffSign(float a, float b)
146     {
147         return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
148     }
diffSign(const float3 & a,const float3 & b)149     __device__ __forceinline__ float3 diffSign(const float3& a, const float3& b)
150     {
151         return make_float3(
152             a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f,
153             a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f,
154             a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f
155         );
156     }
diffSign(const float4 & a,const float4 & b)157     __device__ __forceinline__ float4 diffSign(const float4& a, const float4& b)
158     {
159         return make_float4(
160             a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f,
161             a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f,
162             a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f,
163             0.0f
164         );
165     }
166 
167     struct DiffSign : binary_function<float, float, float>
168     {
operator ()btv_l1_cudev::DiffSign169         __device__ __forceinline__ float operator ()(float a, float b) const
170         {
171             return diffSign(a, b);
172         }
173     };
174 }
175 
176 namespace cv { namespace cuda { namespace device
177 {
178     template <> struct TransformFunctorTraits<btv_l1_cudev::DiffSign> : DefaultTransformFunctorTraits<btv_l1_cudev::DiffSign>
179     {
180         enum { smart_block_dim_y = 8 };
181         enum { smart_shift = 4 };
182     };
183 }}}
184 
185 namespace btv_l1_cudev
186 {
diffSign(PtrStepSzf src1,PtrStepSzf src2,PtrStepSzf dst,cudaStream_t stream)187     void diffSign(PtrStepSzf src1, PtrStepSzf src2, PtrStepSzf dst, cudaStream_t stream)
188     {
189         transform(src1, src2, dst, DiffSign(), WithOutMask(), stream);
190     }
191 
192     __constant__ float c_btvRegWeights[16*16];
193 
194     template <typename T>
calcBtvRegularizationKernel(const PtrStepSz<T> src,PtrStep<T> dst,const int ksize)195     __global__ void calcBtvRegularizationKernel(const PtrStepSz<T> src, PtrStep<T> dst, const int ksize)
196     {
197         const int x = blockIdx.x * blockDim.x + threadIdx.x + ksize;
198         const int y = blockIdx.y * blockDim.y + threadIdx.y + ksize;
199 
200         if (y >= src.rows - ksize || x >= src.cols - ksize)
201             return;
202 
203         const T srcVal = src(y, x);
204 
205         T dstVal = VecTraits<T>::all(0);
206 
207         for (int m = 0, count = 0; m <= ksize; ++m)
208         {
209             for (int l = ksize; l + m >= 0; --l, ++count)
210                 dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src(y + m, x + l)) - diffSign(src(y - m, x - l), srcVal));
211         }
212 
213         dst(y, x) = dstVal;
214     }
215 
loadBtvWeights(const float * weights,size_t count)216     void loadBtvWeights(const float* weights, size_t count)
217     {
218         cudaSafeCall( cudaMemcpyToSymbol(c_btvRegWeights, weights, count * sizeof(float)) );
219     }
220 
221     template <int cn>
calcBtvRegularization(PtrStepSzb src,PtrStepSzb dst,int ksize)222     void calcBtvRegularization(PtrStepSzb src, PtrStepSzb dst, int ksize)
223     {
224         typedef typename TypeVec<float, cn>::vec_type src_t;
225 
226         const dim3 block(32, 8);
227         const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
228 
229         calcBtvRegularizationKernel<src_t><<<grid, block>>>((PtrStepSz<src_t>) src, (PtrStepSz<src_t>) dst, ksize);
230         cudaSafeCall( cudaGetLastError() );
231 
232         cudaSafeCall( cudaDeviceSynchronize() );
233     }
234 
235     template void calcBtvRegularization<1>(PtrStepSzb src, PtrStepSzb dst, int ksize);
236     template void calcBtvRegularization<3>(PtrStepSzb src, PtrStepSzb dst, int ksize);
237     template void calcBtvRegularization<4>(PtrStepSzb src, PtrStepSzb dst, int ksize);
238 }
239 
240 #endif
241