1
/*M///////////////////////////////////////////////////////////////////////////////////////
3
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
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.
11
// For Open Source Computer Vision Library
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.
17
// Redistribution and use in source and binary forms, with or without modification,
18
// are permitted provided that the following conditions are met:
20
// * Redistribution's of source code must retain the above copyright notice,
21
// this list of conditions and the following disclaimer.
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.
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.
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.
43
#include "opencv2/opencv_modules.hpp"
45
#ifndef HAVE_OPENCV_CUDEV
47
#error "opencv_cudev is required"
51
#include "opencv2/cudev.hpp"
53
using namespace cv::cudev;
55
void minMaxMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int op);
57
void minMaxScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int op);
59
///////////////////////////////////////////////////////////////////////
64
template <template <typename> class Op, typename T>
65
void minMaxMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
67
gridTransformBinary(globPtr<T>(src1), globPtr<T>(src2), globPtr<T>(dst), Op<T>(), stream);
70
struct MinOp2 : binary_function<uint, uint, uint>
72
__device__ __forceinline__ uint operator ()(uint a, uint b) const
78
struct MaxOp2 : binary_function<uint, uint, uint>
80
__device__ __forceinline__ uint operator ()(uint a, uint b) const
87
void minMaxMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
89
const int vcols = src1.cols >> 1;
91
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
92
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
93
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
95
gridTransformBinary(src1_, src2_, dst_, Op2(), stream);
98
struct MinOp4 : binary_function<uint, uint, uint>
100
__device__ __forceinline__ uint operator ()(uint a, uint b) const
106
struct MaxOp4 : binary_function<uint, uint, uint>
108
__device__ __forceinline__ uint operator ()(uint a, uint b) const
115
void minMaxMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
117
const int vcols = src1.cols >> 2;
119
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
120
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
121
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
123
gridTransformBinary(src1_, src2_, dst_, Op4(), stream);
127
void minMaxMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int op)
129
typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
130
static const func_t funcs_v1[2][7] =
133
minMaxMat_v1<minimum, uchar>,
134
minMaxMat_v1<minimum, schar>,
135
minMaxMat_v1<minimum, ushort>,
136
minMaxMat_v1<minimum, short>,
137
minMaxMat_v1<minimum, int>,
138
minMaxMat_v1<minimum, float>,
139
minMaxMat_v1<minimum, double>
142
minMaxMat_v1<maximum, uchar>,
143
minMaxMat_v1<maximum, schar>,
144
minMaxMat_v1<maximum, ushort>,
145
minMaxMat_v1<maximum, short>,
146
minMaxMat_v1<maximum, int>,
147
minMaxMat_v1<maximum, float>,
148
minMaxMat_v1<maximum, double>
152
static const func_t funcs_v2[2] =
154
minMaxMat_v2<MinOp2>, minMaxMat_v2<MaxOp2>
157
static const func_t funcs_v4[2] =
159
minMaxMat_v4<MinOp4>, minMaxMat_v4<MaxOp4>
162
const int depth = src1.depth();
164
CV_DbgAssert( depth <= CV_64F );
166
GpuMat src1_ = src1.reshape(1);
167
GpuMat src2_ = src2.reshape(1);
168
GpuMat dst_ = dst.reshape(1);
170
if (depth == CV_8U || depth == CV_16U)
172
const intptr_t src1ptr = reinterpret_cast<intptr_t>(src1_.data);
173
const intptr_t src2ptr = reinterpret_cast<intptr_t>(src2_.data);
174
const intptr_t dstptr = reinterpret_cast<intptr_t>(dst_.data);
176
const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0;
180
if (depth == CV_8U && (src1_.cols & 3) == 0)
182
funcs_v4[op](src1_, src2_, dst_, stream);
185
else if (depth == CV_16U && (src1_.cols & 1) == 0)
187
funcs_v2[op](src1_, src2_, dst_, stream);
193
const func_t func = funcs_v1[op][depth];
195
func(src1_, src2_, dst_, stream);
198
///////////////////////////////////////////////////////////////////////
203
template <template <typename> class Op, typename T>
204
void minMaxScalar(const GpuMat& src, double value, GpuMat& dst, Stream& stream)
206
gridTransformUnary(globPtr<T>(src), globPtr<T>(dst), bind2nd(Op<T>(), cv::saturate_cast<T>(value)), stream);
210
void minMaxScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int op)
212
typedef void (*func_t)(const GpuMat& src, double value, GpuMat& dst, Stream& stream);
213
static const func_t funcs[2][7] =
216
minMaxScalar<minimum, uchar>,
217
minMaxScalar<minimum, schar>,
218
minMaxScalar<minimum, ushort>,
219
minMaxScalar<minimum, short>,
220
minMaxScalar<minimum, int>,
221
minMaxScalar<minimum, float>,
222
minMaxScalar<minimum, double>
225
minMaxScalar<maximum, uchar>,
226
minMaxScalar<maximum, schar>,
227
minMaxScalar<maximum, ushort>,
228
minMaxScalar<maximum, short>,
229
minMaxScalar<maximum, int>,
230
minMaxScalar<maximum, float>,
231
minMaxScalar<maximum, double>
235
const int depth = src.depth();
237
CV_DbgAssert( depth <= CV_64F );
238
CV_DbgAssert( src.channels() == 1 );
240
funcs[op][depth](src, value[0], dst, stream);