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>
swap(T & p1,T & p2)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>
divUp(T a,T b)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 {
assignfunctorAddValues68 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 }
reducefunctorAddValues73 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 {
assignfunctorMinValues83 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 }
reducefunctorMinValues88 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 {
assignfunctorMaxValues98 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 }
reducefunctorMaxValues103 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>
subReduce(Tdata threadElem)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_
156