← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/superres/src/cuda/btv_l1_gpu.cu

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/superres/src/cuda/btv_l1_gpu.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
#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
{
 
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
 
 
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>
 
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>
 
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
 
 
145
    __device__ __forceinline__ float diffSign(float a, float b)
 
146
    {
 
147
        return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
 
148
    }
 
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
    }
 
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
    {
 
169
        __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
{
 
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>
 
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
 
 
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>
 
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