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/cudaarithm.hpp"
52
#include "opencv2/cudev.hpp"
53
#include "opencv2/core/private.cuda.hpp"
56
using namespace cv::cuda;
57
using namespace cv::cudev;
61
template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
64
template <> struct TransformPolicy<double> : DefaultTransformPolicy
72
//////////////////////////////////////////////////////////////////////////////
78
void absMat(const GpuMat& src, const GpuMat& dst, Stream& stream)
80
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), abs_func<T>(), stream);
84
void cv::cuda::abs(InputArray _src, OutputArray _dst, Stream& stream)
86
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
87
static const func_t funcs[] =
98
GpuMat src = getInputMat(_src, stream);
100
CV_Assert( src.depth() <= CV_64F );
102
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
104
funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
106
syncOutput(dst, _dst, stream);
109
//////////////////////////////////////////////////////////////////////////////
114
template <typename T> struct SqrOp : unary_function<T, T>
116
__device__ __forceinline__ T operator ()(T x) const
118
return cudev::saturate_cast<T>(x * x);
122
template <typename T>
123
void sqrMat(const GpuMat& src, const GpuMat& dst, Stream& stream)
125
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), SqrOp<T>(), stream);
129
void cv::cuda::sqr(InputArray _src, OutputArray _dst, Stream& stream)
131
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
132
static const func_t funcs[] =
143
GpuMat src = getInputMat(_src, stream);
145
CV_Assert( src.depth() <= CV_64F );
147
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
149
funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
151
syncOutput(dst, _dst, stream);
154
//////////////////////////////////////////////////////////////////////////////
159
template <typename T>
160
void sqrtMat(const GpuMat& src, const GpuMat& dst, Stream& stream)
162
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), sqrt_func<T>(), stream);
166
void cv::cuda::sqrt(InputArray _src, OutputArray _dst, Stream& stream)
168
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
169
static const func_t funcs[] =
180
GpuMat src = getInputMat(_src, stream);
182
CV_Assert( src.depth() <= CV_64F );
184
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
186
funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
188
syncOutput(dst, _dst, stream);
191
////////////////////////////////////////////////////////////////////////
196
template <typename T> struct ExpOp : unary_function<T, T>
198
__device__ __forceinline__ T operator ()(T x) const
201
return cudev::saturate_cast<T>(f(x));
205
template <typename T>
206
void expMat(const GpuMat& src, const GpuMat& dst, Stream& stream)
208
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), ExpOp<T>(), stream);
212
void cv::cuda::exp(InputArray _src, OutputArray _dst, Stream& stream)
214
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
215
static const func_t funcs[] =
226
GpuMat src = getInputMat(_src, stream);
228
CV_Assert( src.depth() <= CV_64F );
230
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
232
funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
234
syncOutput(dst, _dst, stream);
237
////////////////////////////////////////////////////////////////////////
242
template <typename T>
243
void logMat(const GpuMat& src, const GpuMat& dst, Stream& stream)
245
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), log_func<T>(), stream);
249
void cv::cuda::log(InputArray _src, OutputArray _dst, Stream& stream)
251
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
252
static const func_t funcs[] =
263
GpuMat src = getInputMat(_src, stream);
265
CV_Assert( src.depth() <= CV_64F );
267
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
269
funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
271
syncOutput(dst, _dst, stream);
274
////////////////////////////////////////////////////////////////////////
279
template<typename T, bool Signed = numeric_limits<T>::is_signed> struct PowOp : unary_function<T, T>
283
__device__ __forceinline__ T operator()(T e) const
285
return cudev::saturate_cast<T>(__powf((float)e, power));
288
template<typename T> struct PowOp<T, true> : unary_function<T, T>
292
__device__ __forceinline__ T operator()(T e) const
294
T res = cudev::saturate_cast<T>(__powf((float)e, power));
296
if ((e < 0) && (1 & static_cast<int>(power)))
302
template<> struct PowOp<float> : unary_function<float, float>
306
__device__ __forceinline__ float operator()(float e) const
308
return __powf(::fabs(e), power);
311
template<> struct PowOp<double> : unary_function<double, double>
315
__device__ __forceinline__ double operator()(double e) const
317
return ::pow(::fabs(e), power);
322
void powMat(const GpuMat& src, double power, const GpuMat& dst, Stream& stream)
325
op.power = static_cast<typename LargerType<T, float>::type>(power);
327
gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), op, stream);
331
void cv::cuda::pow(InputArray _src, double power, OutputArray _dst, Stream& stream)
333
typedef void (*func_t)(const GpuMat& src, double power, const GpuMat& dst, Stream& stream);
334
static const func_t funcs[] =
345
GpuMat src = getInputMat(_src, stream);
347
CV_Assert( src.depth() <= CV_64F );
349
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
351
funcs[src.depth()](src.reshape(1), power, dst.reshape(1), stream);
353
syncOutput(dst, _dst, stream);