← 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/cudev/include/opencv2/cudev/grid/detail/reduce.hpp

  • 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/cudev/include/opencv2/cudev/grid/detail/reduce.hpp
 
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) 2013, OpenCV Foundation, 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 implied warranties, including, but not limited to, the implied
 
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
#pragma once
 
45
 
 
46
#ifndef __OPENCV_CUDEV_GRID_REDUCE_DETAIL_HPP__
 
47
#define __OPENCV_CUDEV_GRID_REDUCE_DETAIL_HPP__
 
48
 
 
49
#include "../../common.hpp"
 
50
#include "../../util/tuple.hpp"
 
51
#include "../../util/saturate_cast.hpp"
 
52
#include "../../util/atomic.hpp"
 
53
#include "../../util/vec_traits.hpp"
 
54
#include "../../util/type_traits.hpp"
 
55
#include "../../util/limits.hpp"
 
56
#include "../../block/reduce.hpp"
 
57
#include "../../functional/functional.hpp"
 
58
#include "../../ptr2d/traits.hpp"
 
59
 
 
60
namespace cv { namespace cudev {
 
61
 
 
62
namespace grid_reduce_detail
 
63
{
 
64
    // Unroll
 
65
 
 
66
    template <int cn> struct Unroll;
 
67
 
 
68
    template <> struct Unroll<1>
 
69
    {
 
70
        template <int BLOCK_SIZE, typename R>
 
71
        __device__ __forceinline__ static volatile R* smem(R* ptr)
 
72
        {
 
73
            return ptr;
 
74
        }
 
75
 
 
76
        template <typename R>
 
77
        __device__ __forceinline__ static R& res(R& val)
 
78
        {
 
79
            return val;
 
80
        }
 
81
 
 
82
        template <class Op>
 
83
        __device__ __forceinline__ static const Op& op(const Op& aop)
 
84
        {
 
85
            return aop;
 
86
        }
 
87
    };
 
88
 
 
89
    template <> struct Unroll<2>
 
90
    {
 
91
        template <int BLOCK_SIZE, typename R>
 
92
        __device__ __forceinline__ static tuple<volatile R*, volatile R*> smem(R* ptr)
 
93
        {
 
94
            return smem_tuple(ptr, ptr + BLOCK_SIZE);
 
95
        }
 
96
 
 
97
        template <typename R>
 
98
        __device__ __forceinline__ static tuple<typename VecTraits<R>::elem_type&, typename VecTraits<R>::elem_type&> res(R& val)
 
99
        {
 
100
            return tie(val.x, val.y);
 
101
        }
 
102
 
 
103
        template <class Op>
 
104
        __device__ __forceinline__ static tuple<Op, Op> op(const Op& aop)
 
105
        {
 
106
            return make_tuple(aop, aop);
 
107
        }
 
108
    };
 
109
 
 
110
    template <> struct Unroll<3>
 
111
    {
 
112
        template <int BLOCK_SIZE, typename R>
 
113
        __device__ __forceinline__ static tuple<volatile R*, volatile R*, volatile R*> smem(R* ptr)
 
114
        {
 
115
            return smem_tuple(ptr, ptr + BLOCK_SIZE, ptr + 2 * BLOCK_SIZE);
 
116
        }
 
117
 
 
118
        template <typename R>
 
119
        __device__ __forceinline__ static tuple<typename VecTraits<R>::elem_type&,
 
120
                                                typename VecTraits<R>::elem_type&,
 
121
                                                typename VecTraits<R>::elem_type&> res(R& val)
 
122
        {
 
123
            return tie(val.x, val.y, val.z);
 
124
        }
 
125
 
 
126
        template <class Op>
 
127
        __device__ __forceinline__ static tuple<Op, Op, Op> op(const Op& aop)
 
128
        {
 
129
            return make_tuple(aop, aop, aop);
 
130
        }
 
131
    };
 
132
 
 
133
    template <> struct Unroll<4>
 
134
    {
 
135
        template <int BLOCK_SIZE, typename R>
 
136
        __device__ __forceinline__ static tuple<volatile R*, volatile R*, volatile R*, volatile R*> smem(R* ptr)
 
137
        {
 
138
            return smem_tuple(ptr, ptr + BLOCK_SIZE, ptr + 2 * BLOCK_SIZE, ptr + 3 * BLOCK_SIZE);
 
139
        }
 
140
 
 
141
        template <typename R>
 
142
        __device__ __forceinline__ static tuple<typename VecTraits<R>::elem_type&,
 
143
                                                typename VecTraits<R>::elem_type&,
 
144
                                                typename VecTraits<R>::elem_type&,
 
145
                                                typename VecTraits<R>::elem_type&> res(R& val)
 
146
        {
 
147
            return tie(val.x, val.y, val.z, val.w);
 
148
        }
 
149
 
 
150
        template <class Op>
 
151
        __device__ __forceinline__ static tuple<Op, Op, Op, Op> op(const Op& aop)
 
152
        {
 
153
            return make_tuple(aop, aop, aop, aop);
 
154
        }
 
155
    };
 
156
 
 
157
    // AtomicUnroll
 
158
 
 
159
    template <typename R, int cn> struct AtomicUnroll;
 
160
 
 
161
    template <typename R> struct AtomicUnroll<R, 1>
 
162
    {
 
163
        __device__ __forceinline__ static void add(R* ptr, R val)
 
164
        {
 
165
            atomicAdd(ptr, val);
 
166
        }
 
167
 
 
168
        __device__ __forceinline__ static void min(R* ptr, R val)
 
169
        {
 
170
            atomicMin(ptr, val);
 
171
        }
 
172
 
 
173
        __device__ __forceinline__ static void max(R* ptr, R val)
 
174
        {
 
175
            atomicMax(ptr, val);
 
176
        }
 
177
    };
 
178
 
 
179
    template <typename R> struct AtomicUnroll<R, 2>
 
180
    {
 
181
        typedef typename MakeVec<R, 2>::type val_type;
 
182
 
 
183
        __device__ __forceinline__ static void add(R* ptr, val_type val)
 
184
        {
 
185
            atomicAdd(ptr, val.x);
 
186
            atomicAdd(ptr + 1, val.y);
 
187
        }
 
188
 
 
189
        __device__ __forceinline__ static void min(R* ptr, val_type val)
 
190
        {
 
191
            atomicMin(ptr, val.x);
 
192
            atomicMin(ptr + 1, val.y);
 
193
        }
 
194
 
 
195
        __device__ __forceinline__ static void max(R* ptr, val_type val)
 
196
        {
 
197
            atomicMax(ptr, val.x);
 
198
            atomicMax(ptr + 1, val.y);
 
199
        }
 
200
    };
 
201
 
 
202
    template <typename R> struct AtomicUnroll<R, 3>
 
203
    {
 
204
        typedef typename MakeVec<R, 3>::type val_type;
 
205
 
 
206
        __device__ __forceinline__ static void add(R* ptr, val_type val)
 
207
        {
 
208
            atomicAdd(ptr, val.x);
 
209
            atomicAdd(ptr + 1, val.y);
 
210
            atomicAdd(ptr + 2, val.z);
 
211
        }
 
212
 
 
213
        __device__ __forceinline__ static void min(R* ptr, val_type val)
 
214
        {
 
215
            atomicMin(ptr, val.x);
 
216
            atomicMin(ptr + 1, val.y);
 
217
            atomicMin(ptr + 2, val.z);
 
218
        }
 
219
 
 
220
        __device__ __forceinline__ static void max(R* ptr, val_type val)
 
221
        {
 
222
            atomicMax(ptr, val.x);
 
223
            atomicMax(ptr + 1, val.y);
 
224
            atomicMax(ptr + 2, val.z);
 
225
        }
 
226
    };
 
227
 
 
228
    template <typename R> struct AtomicUnroll<R, 4>
 
229
    {
 
230
        typedef typename MakeVec<R, 4>::type val_type;
 
231
 
 
232
        __device__ __forceinline__ static void add(R* ptr, val_type val)
 
233
        {
 
234
            atomicAdd(ptr, val.x);
 
235
            atomicAdd(ptr + 1, val.y);
 
236
            atomicAdd(ptr + 2, val.z);
 
237
            atomicAdd(ptr + 3, val.w);
 
238
        }
 
239
 
 
240
        __device__ __forceinline__ static void min(R* ptr, val_type val)
 
241
        {
 
242
            atomicMin(ptr, val.x);
 
243
            atomicMin(ptr + 1, val.y);
 
244
            atomicMin(ptr + 2, val.z);
 
245
            atomicMin(ptr + 3, val.w);
 
246
        }
 
247
 
 
248
        __device__ __forceinline__ static void max(R* ptr, val_type val)
 
249
        {
 
250
            atomicMax(ptr, val.x);
 
251
            atomicMax(ptr + 1, val.y);
 
252
            atomicMax(ptr + 2, val.z);
 
253
            atomicMax(ptr + 3, val.w);
 
254
        }
 
255
    };
 
256
 
 
257
    // SumReductor
 
258
 
 
259
    template <typename src_type, typename work_type> struct SumReductor
 
260
    {
 
261
        typedef typename VecTraits<work_type>::elem_type work_elem_type;
 
262
        enum { cn = VecTraits<src_type>::cn };
 
263
 
 
264
        work_type sum;
 
265
 
 
266
        __device__ __forceinline__ SumReductor()
 
267
        {
 
268
            sum = VecTraits<work_type>::all(0);
 
269
        }
 
270
 
 
271
        __device__ __forceinline__ void reduceVal(typename TypeTraits<src_type>::parameter_type srcVal)
 
272
        {
 
273
            sum = sum + saturate_cast<work_type>(srcVal);
 
274
        }
 
275
 
 
276
        template <int BLOCK_SIZE>
 
277
        __device__ void reduceGrid(work_elem_type* result, int tid)
 
278
        {
 
279
            __shared__ work_elem_type smem[BLOCK_SIZE * cn];
 
280
 
 
281
            blockReduce<BLOCK_SIZE>(Unroll<cn>::template smem<BLOCK_SIZE>(smem), Unroll<cn>::res(sum), tid, Unroll<cn>::op(plus<work_elem_type>()));
 
282
 
 
283
            if (tid == 0)
 
284
                AtomicUnroll<work_elem_type, cn>::add(result, sum);
 
285
        }
 
286
    };
 
287
 
 
288
    // MinMaxReductor
 
289
 
 
290
    template <typename T> struct minop : minimum<T>
 
291
    {
 
292
        __device__ __forceinline__ static T initial()
 
293
        {
 
294
            return numeric_limits<T>::max();
 
295
        }
 
296
 
 
297
        __device__ __forceinline__ static void atomic(T* result, T myval)
 
298
        {
 
299
            atomicMin(result, myval);
 
300
        }
 
301
    };
 
302
 
 
303
    template <typename T> struct maxop : maximum<T>
 
304
    {
 
305
        __device__ __forceinline__ static T initial()
 
306
        {
 
307
            return -numeric_limits<T>::max();
 
308
        }
 
309
 
 
310
        __device__ __forceinline__ static void atomic(T* result, T myval)
 
311
        {
 
312
            atomicMax(result, myval);
 
313
        }
 
314
    };
 
315
 
 
316
    struct both
 
317
    {
 
318
    };
 
319
 
 
320
    template <class Op, typename src_type, typename work_type> struct MinMaxReductor
 
321
    {
 
322
        work_type myval;
 
323
 
 
324
        __device__ __forceinline__ MinMaxReductor()
 
325
        {
 
326
            myval = Op::initial();
 
327
        }
 
328
 
 
329
        __device__ __forceinline__ void reduceVal(typename TypeTraits<src_type>::parameter_type srcVal)
 
330
        {
 
331
            Op op;
 
332
 
 
333
            myval = op(myval, srcVal);
 
334
        }
 
335
 
 
336
        template <int BLOCK_SIZE>
 
337
        __device__ void reduceGrid(work_type* result, int tid)
 
338
        {
 
339
            __shared__ work_type smem[BLOCK_SIZE];
 
340
 
 
341
            Op op;
 
342
 
 
343
            blockReduce<BLOCK_SIZE>(smem, myval, tid, op);
 
344
 
 
345
            if (tid == 0)
 
346
                Op::atomic(result, myval);
 
347
        }
 
348
    };
 
349
 
 
350
    template <typename src_type, typename work_type> struct MinMaxReductor<both, src_type, work_type>
 
351
    {
 
352
        work_type mymin;
 
353
        work_type mymax;
 
354
 
 
355
        __device__ __forceinline__ MinMaxReductor()
 
356
        {
 
357
            mymin = numeric_limits<work_type>::max();
 
358
            mymax = -numeric_limits<work_type>::max();
 
359
        }
 
360
 
 
361
        __device__ __forceinline__ void reduceVal(typename TypeTraits<src_type>::parameter_type srcVal)
 
362
        {
 
363
            minimum<work_type> minOp;
 
364
            maximum<work_type> maxOp;
 
365
 
 
366
            mymin = minOp(mymin, srcVal);
 
367
            mymax = maxOp(mymax, srcVal);
 
368
        }
 
369
 
 
370
        template <int BLOCK_SIZE>
 
371
        __device__ void reduceGrid(work_type* result, int tid)
 
372
        {
 
373
            __shared__ work_type sminval[BLOCK_SIZE];
 
374
            __shared__ work_type smaxval[BLOCK_SIZE];
 
375
 
 
376
            minimum<work_type> minOp;
 
377
            maximum<work_type> maxOp;
 
378
 
 
379
            blockReduce<BLOCK_SIZE>(smem_tuple(sminval, smaxval), tie(mymin, mymax), tid, make_tuple(minOp, maxOp));
 
380
 
 
381
            if (tid == 0)
 
382
            {
 
383
                atomicMin(result, mymin);
 
384
                atomicMax(result + 1, mymax);
 
385
            }
 
386
        }
 
387
    };
 
388
 
 
389
    // glob_reduce
 
390
 
 
391
    template <class Reductor, int BLOCK_SIZE, int PATCH_X, int PATCH_Y, class SrcPtr, typename ResType, class MaskPtr>
 
392
    __global__ void reduce(const SrcPtr src, ResType* result, const MaskPtr mask, const int rows, const int cols)
 
393
    {
 
394
        const int x0 = blockIdx.x * blockDim.x * PATCH_X + threadIdx.x;
 
395
        const int y0 = blockIdx.y * blockDim.y * PATCH_Y + threadIdx.y;
 
396
 
 
397
        Reductor reductor;
 
398
 
 
399
        for (int i = 0, y = y0; i < PATCH_Y && y < rows; ++i, y += blockDim.y)
 
400
        {
 
401
            for (int j = 0, x = x0; j < PATCH_X && x < cols; ++j, x += blockDim.x)
 
402
            {
 
403
                if (mask(y, x))
 
404
                {
 
405
                    reductor.reduceVal(src(y, x));
 
406
                }
 
407
            }
 
408
        }
 
409
 
 
410
        const int tid = threadIdx.y * blockDim.x + threadIdx.x;
 
411
 
 
412
        reductor.template reduceGrid<BLOCK_SIZE>(result, tid);
 
413
    }
 
414
 
 
415
    template <class Reductor, class Policy, class SrcPtr, typename ResType, class MaskPtr>
 
416
    __host__ void reduce(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
 
417
    {
 
418
        const dim3 block(Policy::block_size_x, Policy::block_size_y);
 
419
        const dim3 grid(divUp(cols, block.x * Policy::patch_size_x), divUp(rows, block.y * Policy::patch_size_y));
 
420
 
 
421
        reduce<Reductor, Policy::block_size_x * Policy::block_size_y, Policy::patch_size_x, Policy::patch_size_y><<<grid, block, 0, stream>>>(src, result, mask, rows, cols);
 
422
        CV_CUDEV_SAFE_CALL( cudaGetLastError() );
 
423
 
 
424
        if (stream == 0)
 
425
            CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
 
426
    }
 
427
 
 
428
    // callers
 
429
 
 
430
    template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
 
431
    __host__ void sum(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
 
432
    {
 
433
        typedef typename PtrTraits<SrcPtr>::value_type src_type;
 
434
        typedef typename VecTraits<ResType>::elem_type res_elem_type;
 
435
 
 
436
        reduce<SumReductor<src_type, ResType>, Policy>(src, (res_elem_type*) result, mask, rows, cols, stream);
 
437
    }
 
438
 
 
439
    template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
 
440
    __host__ void minVal(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
 
441
    {
 
442
        typedef typename PtrTraits<SrcPtr>::value_type src_type;
 
443
 
 
444
        reduce<MinMaxReductor<minop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
 
445
    }
 
446
 
 
447
    template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
 
448
    __host__ void maxVal(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
 
449
    {
 
450
        typedef typename PtrTraits<SrcPtr>::value_type src_type;
 
451
 
 
452
        reduce<MinMaxReductor<maxop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
 
453
    }
 
454
 
 
455
    template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
 
456
    __host__ void minMaxVal(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
 
457
    {
 
458
        typedef typename PtrTraits<SrcPtr>::value_type src_type;
 
459
 
 
460
        reduce<MinMaxReductor<both, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
 
461
    }
 
462
}
 
463
 
 
464
}}
 
465
 
 
466
#endif