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 #ifndef HAVE_OPENCV_CUDEV
46
47 #error "opencv_cudev is required"
48
49 #else
50
51 #include "opencv2/cudev.hpp"
52
53 using namespace cv::cudev;
54
55 void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int);
56 void divMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
57 void divMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
58
59 namespace
60 {
61 template <typename T, typename D> struct DivOp : binary_function<T, T, D>
62 {
operator ()__anonf89046690111::DivOp63 __device__ __forceinline__ D operator ()(T a, T b) const
64 {
65 return b != 0 ? saturate_cast<D>(a / b) : 0;
66 }
67 };
68 template <typename T> struct DivOp<T, float> : binary_function<T, T, float>
69 {
operator ()__anonf89046690111::DivOp70 __device__ __forceinline__ float operator ()(T a, T b) const
71 {
72 return b != 0 ? static_cast<float>(a) / b : 0.0f;
73 }
74 };
75 template <typename T> struct DivOp<T, double> : binary_function<T, T, double>
76 {
operator ()__anonf89046690111::DivOp77 __device__ __forceinline__ double operator ()(T a, T b) const
78 {
79 return b != 0 ? static_cast<double>(a) / b : 0.0;
80 }
81 };
82
83 template <typename T, typename S, typename D> struct DivScaleOp : binary_function<T, T, D>
84 {
85 S scale;
86
operator ()__anonf89046690111::DivScaleOp87 __device__ __forceinline__ D operator ()(T a, T b) const
88 {
89 return b != 0 ? saturate_cast<D>(scale * a / b) : 0;
90 }
91 };
92
93 template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
94 {
95 };
96 template <> struct TransformPolicy<double> : DefaultTransformPolicy
97 {
98 enum {
99 shift = 1
100 };
101 };
102
103 template <typename T, typename S, typename D>
divMatImpl(const GpuMat & src1,const GpuMat & src2,const GpuMat & dst,double scale,Stream & stream)104 void divMatImpl(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream)
105 {
106 if (scale == 1)
107 {
108 DivOp<T, D> op;
109 gridTransformBinary_< TransformPolicy<S> >(globPtr<T>(src1), globPtr<T>(src2), globPtr<D>(dst), op, stream);
110 }
111 else
112 {
113 DivScaleOp<T, S, D> op;
114 op.scale = static_cast<S>(scale);
115 gridTransformBinary_< TransformPolicy<S> >(globPtr<T>(src1), globPtr<T>(src2), globPtr<D>(dst), op, stream);
116 }
117 }
118 }
119
divMat(const GpuMat & src1,const GpuMat & src2,GpuMat & dst,const GpuMat &,double scale,Stream & stream,int)120 void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int)
121 {
122 typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream);
123 static const func_t funcs[7][7] =
124 {
125 {
126 divMatImpl<uchar, float, uchar>,
127 divMatImpl<uchar, float, schar>,
128 divMatImpl<uchar, float, ushort>,
129 divMatImpl<uchar, float, short>,
130 divMatImpl<uchar, float, int>,
131 divMatImpl<uchar, float, float>,
132 divMatImpl<uchar, double, double>
133 },
134 {
135 divMatImpl<schar, float, uchar>,
136 divMatImpl<schar, float, schar>,
137 divMatImpl<schar, float, ushort>,
138 divMatImpl<schar, float, short>,
139 divMatImpl<schar, float, int>,
140 divMatImpl<schar, float, float>,
141 divMatImpl<schar, double, double>
142 },
143 {
144 0 /*divMatImpl<ushort, float, uchar>*/,
145 0 /*divMatImpl<ushort, float, schar>*/,
146 divMatImpl<ushort, float, ushort>,
147 divMatImpl<ushort, float, short>,
148 divMatImpl<ushort, float, int>,
149 divMatImpl<ushort, float, float>,
150 divMatImpl<ushort, double, double>
151 },
152 {
153 0 /*divMatImpl<short, float, uchar>*/,
154 0 /*divMatImpl<short, float, schar>*/,
155 divMatImpl<short, float, ushort>,
156 divMatImpl<short, float, short>,
157 divMatImpl<short, float, int>,
158 divMatImpl<short, float, float>,
159 divMatImpl<short, double, double>
160 },
161 {
162 0 /*divMatImpl<int, float, uchar>*/,
163 0 /*divMatImpl<int, float, schar>*/,
164 0 /*divMatImpl<int, float, ushort>*/,
165 0 /*divMatImpl<int, float, short>*/,
166 divMatImpl<int, float, int>,
167 divMatImpl<int, float, float>,
168 divMatImpl<int, double, double>
169 },
170 {
171 0 /*divMatImpl<float, float, uchar>*/,
172 0 /*divMatImpl<float, float, schar>*/,
173 0 /*divMatImpl<float, float, ushort>*/,
174 0 /*divMatImpl<float, float, short>*/,
175 0 /*divMatImpl<float, float, int>*/,
176 divMatImpl<float, float, float>,
177 divMatImpl<float, double, double>
178 },
179 {
180 0 /*divMatImpl<double, double, uchar>*/,
181 0 /*divMatImpl<double, double, schar>*/,
182 0 /*divMatImpl<double, double, ushort>*/,
183 0 /*divMatImpl<double, double, short>*/,
184 0 /*divMatImpl<double, double, int>*/,
185 0 /*divMatImpl<double, double, float>*/,
186 divMatImpl<double, double, double>
187 }
188 };
189
190 const int sdepth = src1.depth();
191 const int ddepth = dst.depth();
192
193 CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F );
194
195 GpuMat src1_ = src1.reshape(1);
196 GpuMat src2_ = src2.reshape(1);
197 GpuMat dst_ = dst.reshape(1);
198
199 const func_t func = funcs[sdepth][ddepth];
200
201 if (!func)
202 CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
203
204 func(src1_, src2_, dst_, scale, stream);
205 }
206
207 namespace
208 {
209 template <typename T>
210 struct DivOpSpecial : binary_function<T, float, T>
211 {
operator ()__anonf89046690311::DivOpSpecial212 __device__ __forceinline__ T operator ()(const T& a, float b) const
213 {
214 typedef typename VecTraits<T>::elem_type elem_type;
215
216 T res = VecTraits<T>::all(0);
217
218 if (b != 0)
219 {
220 b = 1.0f / b;
221 res.x = saturate_cast<elem_type>(a.x * b);
222 res.y = saturate_cast<elem_type>(a.y * b);
223 res.z = saturate_cast<elem_type>(a.z * b);
224 res.w = saturate_cast<elem_type>(a.w * b);
225 }
226
227 return res;
228 }
229 };
230 }
231
divMat_8uc4_32f(const GpuMat & src1,const GpuMat & src2,GpuMat & dst,Stream & stream)232 void divMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
233 {
234 gridTransformBinary(globPtr<uchar4>(src1), globPtr<float>(src2), globPtr<uchar4>(dst), DivOpSpecial<uchar4>(), stream);
235 }
236
divMat_16sc4_32f(const GpuMat & src1,const GpuMat & src2,GpuMat & dst,Stream & stream)237 void divMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
238 {
239 gridTransformBinary(globPtr<short4>(src1), globPtr<float>(src2), globPtr<short4>(dst), DivOpSpecial<short4>(), stream);
240 }
241
242 #endif
243