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_TO_VEC_HPP__
47 #define __OPENCV_CUDEV_GRID_REDUCE_TO_VEC_HPP__
48
49 #include "../common.hpp"
50 #include "../util/vec_traits.hpp"
51 #include "../util/limits.hpp"
52 #include "../util/saturate_cast.hpp"
53 #include "../ptr2d/traits.hpp"
54 #include "../ptr2d/gpumat.hpp"
55 #include "../ptr2d/mask.hpp"
56 #include "../functional/functional.hpp"
57 #include "detail/reduce_to_column.hpp"
58 #include "detail/reduce_to_row.hpp"
59
60 namespace cv { namespace cudev {
61
62 //! @addtogroup cudev
63 //! @{
64
65 template <typename T> struct Sum : plus<T>
66 {
67 typedef T work_type;
68
69 template <typename U> struct rebind
70 {
71 typedef Sum<U> other;
72 };
73
initialValuecv::cudev::Sum74 __device__ __forceinline__ static T initialValue()
75 {
76 return VecTraits<T>::all(0);
77 }
78
resultcv::cudev::Sum79 __device__ __forceinline__ static T result(T r, int)
80 {
81 return r;
82 }
83 };
84
85 template <typename T> struct Avg : plus<T>
86 {
87 typedef T work_type;
88
89 template <typename U> struct rebind
90 {
91 typedef Avg<U> other;
92 };
93
initialValuecv::cudev::Avg94 __device__ __forceinline__ static T initialValue()
95 {
96 return VecTraits<T>::all(0);
97 }
98
resultcv::cudev::Avg99 __device__ __forceinline__ static T result(T r, float sz)
100 {
101 return saturate_cast<T>(r / sz);
102 }
103 };
104
105 template <typename T> struct Min : minimum<T>
106 {
107 typedef T work_type;
108
109 template <typename U> struct rebind
110 {
111 typedef Min<U> other;
112 };
113
initialValuecv::cudev::Min114 __device__ __forceinline__ static T initialValue()
115 {
116 return VecTraits<T>::all(numeric_limits<typename VecTraits<T>::elem_type>::max());
117 }
118
resultcv::cudev::Min119 __device__ __forceinline__ static T result(T r, int)
120 {
121 return r;
122 }
123 };
124
125 template <typename T> struct Max : maximum<T>
126 {
127 typedef T work_type;
128
129 template <typename U> struct rebind
130 {
131 typedef Max<U> other;
132 };
133
initialValuecv::cudev::Max134 __device__ __forceinline__ static T initialValue()
135 {
136 return VecTraits<T>::all(-numeric_limits<typename VecTraits<T>::elem_type>::max());
137 }
138
resultcv::cudev::Max139 __device__ __forceinline__ static T result(T r, int)
140 {
141 return r;
142 }
143 };
144
145 template <class Reductor, class SrcPtr, typename ResType, class MaskPtr>
gridReduceToRow(const SrcPtr & src,GpuMat_<ResType> & dst,const MaskPtr & mask,Stream & stream=Stream::Null ())146 __host__ void gridReduceToRow(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
147 {
148 const int rows = getRows(src);
149 const int cols = getCols(src);
150
151 CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
152
153 dst.create(1, cols);
154
155 grid_reduce_to_vec_detail::reduceToRow<Reductor>(shrinkPtr(src),
156 dst[0],
157 shrinkPtr(mask),
158 rows, cols,
159 StreamAccessor::getStream(stream));
160 }
161
162 template <class Reductor, class SrcPtr, typename ResType>
gridReduceToRow(const SrcPtr & src,GpuMat_<ResType> & dst,Stream & stream=Stream::Null ())163 __host__ void gridReduceToRow(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
164 {
165 const int rows = getRows(src);
166 const int cols = getCols(src);
167
168 dst.create(1, cols);
169
170 grid_reduce_to_vec_detail::reduceToRow<Reductor>(shrinkPtr(src),
171 dst[0],
172 WithOutMask(),
173 rows, cols,
174 StreamAccessor::getStream(stream));
175 }
176
177 template <class Reductor, class Policy, class SrcPtr, typename ResType, class MaskPtr>
gridReduceToColumn_(const SrcPtr & src,GpuMat_<ResType> & dst,const MaskPtr & mask,Stream & stream=Stream::Null ())178 __host__ void gridReduceToColumn_(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
179 {
180 const int rows = getRows(src);
181 const int cols = getCols(src);
182
183 CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
184
185 dst.create(1, rows);
186
187 grid_reduce_to_vec_detail::reduceToColumn<Reductor, Policy>(shrinkPtr(src),
188 dst[0],
189 shrinkPtr(mask),
190 rows, cols,
191 StreamAccessor::getStream(stream));
192 }
193
194 template <class Reductor, class Policy, class SrcPtr, typename ResType>
gridReduceToColumn_(const SrcPtr & src,GpuMat_<ResType> & dst,Stream & stream=Stream::Null ())195 __host__ void gridReduceToColumn_(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
196 {
197 const int rows = getRows(src);
198 const int cols = getCols(src);
199
200 dst.create(1, rows);
201
202 grid_reduce_to_vec_detail::reduceToColumn<Reductor, Policy>(shrinkPtr(src),
203 dst[0],
204 WithOutMask(),
205 rows, cols,
206 StreamAccessor::getStream(stream));
207 }
208
209 // default policy
210
211 struct DefaultReduceToVecPolicy
212 {
213 enum {
214 block_size_x = 32,
215 block_size_y = 8
216 };
217 };
218
219 template <class Reductor, class SrcPtr, typename ResType, class MaskPtr>
gridReduceToColumn(const SrcPtr & src,GpuMat_<ResType> & dst,const MaskPtr & mask,Stream & stream=Stream::Null ())220 __host__ void gridReduceToColumn(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
221 {
222 gridReduceToColumn_<Reductor, DefaultReduceToVecPolicy>(src, dst, mask, stream);
223 }
224
225 template <class Reductor, class SrcPtr, typename ResType>
gridReduceToColumn(const SrcPtr & src,GpuMat_<ResType> & dst,Stream & stream=Stream::Null ())226 __host__ void gridReduceToColumn(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
227 {
228 gridReduceToColumn_<Reductor, DefaultReduceToVecPolicy>(src, dst, stream);
229 }
230
231 //! @}
232
233 }}
234
235 #endif
236