← 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/cudalegacy/src/cuda/NCVAlg.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/cudalegacy/src/cuda/NCVAlg.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
// 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
#ifndef _ncv_alg_hpp_
 
44
#define _ncv_alg_hpp_
 
45
 
 
46
#include "opencv2/cudalegacy/NCV.hpp"
 
47
 
 
48
 
 
49
template <class T>
 
50
static void swap(T &p1, T &p2)
 
51
{
 
52
    T tmp = p1;
 
53
    p1 = p2;
 
54
    p2 = tmp;
 
55
}
 
56
 
 
57
 
 
58
template<typename T>
 
59
static T divUp(T a, T b)
 
60
{
 
61
    return (a + b - 1) / b;
 
62
}
 
63
 
 
64
 
 
65
template<typename T>
 
66
struct functorAddValues
 
67
{
 
68
    static __device__ __inline__ void assign(volatile T *dst, volatile T *src)
 
69
    {
 
70
        //Works only for integral types. If you see compiler error here, then you have to specify how to copy your object as a set of integral fields.
 
71
        *dst = *src;
 
72
    }
 
73
    static __device__ __inline__ void reduce(volatile T &in1out, const volatile T &in2)
 
74
    {
 
75
        in1out += in2;
 
76
    }
 
77
};
 
78
 
 
79
 
 
80
template<typename T>
 
81
struct functorMinValues
 
82
{
 
83
    static __device__ __inline__ void assign(volatile T *dst, volatile T *src)
 
84
    {
 
85
        //Works only for integral types. If you see compiler error here, then you have to specify how to copy your object as a set of integral fields.
 
86
        *dst = *src;
 
87
    }
 
88
    static __device__ __inline__ void reduce(volatile T &in1out, const volatile T &in2)
 
89
    {
 
90
        in1out = in1out > in2 ? in2 : in1out;
 
91
    }
 
92
};
 
93
 
 
94
 
 
95
template<typename T>
 
96
struct functorMaxValues
 
97
{
 
98
    static __device__ __inline__ void assign(volatile T *dst, volatile T *src)
 
99
    {
 
100
        //Works only for integral types. If you see compiler error here, then you have to specify how to copy your object as a set of integral fields.
 
101
        *dst = *src;
 
102
    }
 
103
    static __device__ __inline__ void reduce(volatile T &in1out, const volatile T &in2)
 
104
    {
 
105
        in1out = in1out > in2 ? in1out : in2;
 
106
    }
 
107
};
 
108
 
 
109
 
 
110
template<typename Tdata, class Tfunc, Ncv32u nThreads>
 
111
static __device__ Tdata subReduce(Tdata threadElem)
 
112
{
 
113
    Tfunc functor;
 
114
 
 
115
    __shared__ Tdata _reduceArr[nThreads];
 
116
    volatile Tdata *reduceArr = _reduceArr;
 
117
    functor.assign(reduceArr + threadIdx.x, &threadElem);
 
118
    __syncthreads();
 
119
 
 
120
    if (nThreads >= 256 && threadIdx.x < 128)
 
121
    {
 
122
        functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 128]);
 
123
    }
 
124
    __syncthreads();
 
125
 
 
126
    if (nThreads >= 128 && threadIdx.x < 64)
 
127
    {
 
128
        functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 64]);
 
129
    }
 
130
    __syncthreads();
 
131
 
 
132
    if (threadIdx.x < 32)
 
133
    {
 
134
        if (nThreads >= 64)
 
135
        {
 
136
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 32]);
 
137
        }
 
138
        if (nThreads >= 32 && threadIdx.x < 16)
 
139
        {
 
140
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 16]);
 
141
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 8]);
 
142
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 4]);
 
143
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 2]);
 
144
            functor.reduce(reduceArr[threadIdx.x], reduceArr[threadIdx.x + 1]);
 
145
        }
 
146
    }
 
147
 
 
148
    __syncthreads();
 
149
    Tdata reduceRes;
 
150
    functor.assign(&reduceRes, reduceArr);
 
151
    return reduceRes;
 
152
}
 
153
 
 
154
 
 
155
#endif //_ncv_alg_hpp_