Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Greenplum
Opencv
提交
e12e4798
O
Opencv
项目概览
Greenplum
/
Opencv
11 个月 前同步成功
通知
7
Star
0
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
DevOps
流水线
流水线任务
计划
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
O
Opencv
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
DevOps
DevOps
流水线
流水线任务
计划
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
流水线任务
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
提交
e12e4798
编写于
10月 18, 2010
作者:
V
Vladislav Vinogradov
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
added gpu version of magnitude, magnitudeSqr, phase, cartToPolar, polarToCart
上级
a594f01a
变更
4
隐藏空白更改
内联
并排
Showing
4 changed file
with
443 addition
and
39 deletion
+443
-39
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/include/opencv2/gpu/gpu.hpp
+30
-6
modules/gpu/src/arithm.cpp
modules/gpu/src/arithm.cpp
+111
-31
modules/gpu/src/cuda/mathfunc.cu
modules/gpu/src/cuda/mathfunc.cu
+212
-0
tests/gpu/src/arithm.cpp
tests/gpu/src/arithm.cpp
+90
-2
未找到文件。
modules/gpu/include/opencv2/gpu/gpu.hpp
浏览文件 @
e12e4798
...
...
@@ -452,20 +452,44 @@ namespace cv
//! supports only CV_32FC1 type
CV_EXPORTS
void
log
(
const
GpuMat
&
a
,
GpuMat
&
b
);
//! computes magnitude of each (x(i), y(i)) vector
//! supports only CV_32FC1 type
CV_EXPORTS
void
magnitude
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
);
//! computes magnitude of complex (x(i).re, x(i).im) vector
//! supports only CV_32FC2 type
CV_EXPORTS
void
magnitude
(
const
GpuMat
&
x
,
GpuMat
&
magnitude
);
//! computes squared magnitude of each (x(i), y(i)) vector
//! supports only CV_32FC1 type
CV_EXPORTS
void
magnitudeSqr
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
);
//! computes squared magnitude of complex (x(i).re, x(i).im) vector
//! supports only CV_32FC2 type
CV_EXPORTS
void
magnitudeSqr
(
const
GpuMat
&
x
,
GpuMat
&
magnitude
);
//! computes magnitude of each (x(i), y(i)) vector
//! supports only floating-point source
CV_EXPORTS
void
magnitude
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
);
//! Acync version
CV_EXPORTS
void
magnitude
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
,
const
Stream
&
stream
);
//! computes squared magnitude of each (x(i), y(i)) vector
//! supports only floating-point source
CV_EXPORTS
void
magnitudeSqr
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
);
//! Acync version
CV_EXPORTS
void
magnitudeSqr
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
,
const
Stream
&
stream
);
//! computes angle (angle(i)) of each (x(i), y(i)) vector
//! supports only floating-point source
CV_EXPORTS
void
phase
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
angle
,
bool
angleInDegrees
=
false
);
//! Acync version
CV_EXPORTS
void
phase
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
angle
,
bool
angleInDegrees
,
const
Stream
&
stream
);
//! converts Cartesian coordinates to polar
//! supports only floating-point source
CV_EXPORTS
void
cartToPolar
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
,
GpuMat
&
angle
,
bool
angleInDegrees
=
false
);
//! Acync version
CV_EXPORTS
void
cartToPolar
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
magnitude
,
GpuMat
&
angle
,
bool
angleInDegrees
,
const
Stream
&
stream
);
//! converts polar coordinates to Cartesian
//! supports only floating-point source
CV_EXPORTS
void
polarToCart
(
const
GpuMat
&
magnitude
,
const
GpuMat
&
angle
,
GpuMat
&
x
,
GpuMat
&
y
,
bool
angleInDegrees
=
false
);
//! Acync version
CV_EXPORTS
void
polarToCart
(
const
GpuMat
&
magnitude
,
const
GpuMat
&
angle
,
GpuMat
&
x
,
GpuMat
&
y
,
bool
angleInDegrees
,
const
Stream
&
stream
);
////////////////////////////// Image processing //////////////////////////////
//! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.
...
...
modules/gpu/src/arithm.cpp
浏览文件 @
e12e4798
...
...
@@ -69,10 +69,18 @@ void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); }
void
cv
::
gpu
::
LUT
(
const
GpuMat
&
,
const
Mat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
exp
(
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
log
(
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
const
Stream
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
const
Stream
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
phase
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
bool
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
phase
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
bool
,
const
Stream
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
cartToPolar
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
GpuMat
&
,
bool
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
cartToPolar
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
GpuMat
&
,
bool
,
const
Stream
&
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
polarToCart
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
GpuMat
&
,
bool
)
{
throw_nogpu
();
}
void
cv
::
gpu
::
polarToCart
(
const
GpuMat
&
,
const
GpuMat
&
,
GpuMat
&
,
GpuMat
&
,
bool
,
const
Stream
&
)
{
throw_nogpu
();
}
#else
/* !defined (HAVE_CUDA) */
...
...
@@ -611,56 +619,128 @@ void cv::gpu::log(const GpuMat& src, GpuMat& dst)
}
////////////////////////////////////////////////////////////////////////
// magnitude
// NPP magnitide
namespace
{
typedef
NppStatus
(
*
nppMagnitude_t
)(
const
Npp32fc
*
pSrc
,
int
nSrcStep
,
Npp32f
*
pDst
,
int
nDstStep
,
NppiSize
oSizeROI
);
inline
void
npp_magnitude
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
nppMagnitude_t
func
)
{
CV_Assert
(
src
.
type
()
==
CV_32FC2
);
dst
.
create
(
src
.
size
(),
CV_32FC1
);
NppiSize
sz
;
sz
.
width
=
src
.
cols
;
sz
.
height
=
src
.
rows
;
nppSafeCall
(
func
(
src
.
ptr
<
Npp32fc
>
(),
src
.
step
,
dst
.
ptr
<
Npp32f
>
(),
dst
.
step
,
sz
)
);
}
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
src
,
GpuMat
&
dst
)
{
CV_Assert
(
src
.
type
()
==
CV_32FC2
);
::
npp_magnitude
(
src
,
dst
,
nppiMagnitude_32fc32f_C1R
);
}
dst
.
create
(
src
.
size
(),
CV_32FC1
);
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
src
,
GpuMat
&
dst
)
{
::
npp_magnitude
(
src
,
dst
,
nppiMagnitudeSqr_32fc32f_C1R
);
}
NppiSize
sz
;
sz
.
width
=
src
.
cols
;
sz
.
height
=
src
.
rows
;
////////////////////////////////////////////////////////////////////////
// Polar <-> Cart
nppSafeCall
(
nppiMagnitude_32fc32f_C1R
(
src
.
ptr
<
Npp32fc
>
(),
src
.
step
,
dst
.
ptr
<
Npp32f
>
(),
dst
.
step
,
sz
)
);
namespace
cv
{
namespace
gpu
{
namespace
mathfunc
{
void
cartToPolar_gpu
(
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
const
DevMem2Df
&
mag
,
bool
magSqr
,
const
DevMem2Df
&
angle
,
bool
angleInDegrees
,
cudaStream_t
stream
);
void
polarToCart_gpu
(
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
bool
angleInDegrees
,
cudaStream_t
stream
);
}}}
namespace
{
inline
void
cartToPolar_caller
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
mag
,
bool
magSqr
,
GpuMat
&
angle
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
CV_DbgAssert
(
x
.
size
()
==
y
.
size
()
&&
x
.
type
()
==
y
.
type
());
CV_Assert
(
x
.
depth
()
==
CV_32F
);
mag
.
create
(
x
.
size
(),
x
.
type
());
angle
.
create
(
x
.
size
(),
x
.
type
());
GpuMat
x1cn
=
x
.
reshape
(
1
);
GpuMat
y1cn
=
y
.
reshape
(
1
);
GpuMat
mag1cn
=
mag
.
reshape
(
1
);
GpuMat
angle1cn
=
angle
.
reshape
(
1
);
mathfunc
::
cartToPolar_gpu
(
x1cn
,
y1cn
,
mag1cn
,
magSqr
,
angle1cn
,
angleInDegrees
,
stream
);
}
inline
void
polarToCart_caller
(
const
GpuMat
&
mag
,
const
GpuMat
&
angle
,
GpuMat
&
x
,
GpuMat
&
y
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
CV_DbgAssert
((
mag
.
empty
()
||
mag
.
size
()
==
angle
.
size
())
&&
mag
.
type
()
==
angle
.
type
());
CV_Assert
(
mag
.
depth
()
==
CV_32F
);
x
.
create
(
mag
.
size
(),
mag
.
type
());
y
.
create
(
mag
.
size
(),
mag
.
type
());
GpuMat
mag1cn
=
mag
.
reshape
(
1
);
GpuMat
angle1cn
=
angle
.
reshape
(
1
);
GpuMat
x1cn
=
x
.
reshape
(
1
);
GpuMat
y1cn
=
y
.
reshape
(
1
);
mathfunc
::
polarToCart_gpu
(
mag1cn
,
angle1cn
,
x1cn
,
y1cn
,
angleInDegrees
,
stream
);
}
}
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
src1
,
const
GpuMat
&
src2
,
GpuMat
&
dst
)
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
dst
)
{
CV_DbgAssert
(
src1
.
type
()
==
src2
.
type
()
&&
src1
.
size
()
==
src2
.
size
()
);
CV_Assert
(
src1
.
type
()
==
CV_32FC1
);
::
cartToPolar_caller
(
x
,
y
,
dst
,
false
,
GpuMat
(),
false
,
0
);
}
GpuMat
src
(
src1
.
size
(),
CV_32FC2
);
GpuMat
srcs
[]
=
{
src1
,
src2
};
cv
::
gpu
::
merge
(
srcs
,
2
,
src
);
void
cv
::
gpu
::
magnitude
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
dst
,
const
Stream
&
stream
)
{
::
cartToPolar_caller
(
x
,
y
,
dst
,
false
,
GpuMat
(),
false
,
StreamAccessor
::
getStream
(
stream
));
}
cv
::
gpu
::
magnitude
(
src
,
dst
);
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
dst
)
{
::
cartToPolar_caller
(
x
,
y
,
dst
,
true
,
GpuMat
(),
false
,
0
);
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
src
,
GpuMat
&
dst
)
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
dst
,
const
Stream
&
stream
)
{
CV_Assert
(
src
.
type
()
==
CV_32FC2
);
::
cartToPolar_caller
(
x
,
y
,
dst
,
true
,
GpuMat
(),
false
,
StreamAccessor
::
getStream
(
stream
));
}
dst
.
create
(
src
.
size
(),
CV_32FC1
);
void
cv
::
gpu
::
phase
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
angle
,
bool
angleInDegrees
)
{
::
cartToPolar_caller
(
x
,
y
,
GpuMat
(),
false
,
angle
,
angleInDegrees
,
0
);
}
NppiSize
sz
;
sz
.
width
=
src
.
cols
;
sz
.
height
=
src
.
rows
;
void
cv
::
gpu
::
phase
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
angle
,
bool
angleInDegrees
,
const
Stream
&
stream
)
{
::
cartToPolar_caller
(
x
,
y
,
GpuMat
(),
false
,
angle
,
angleInDegrees
,
StreamAccessor
::
getStream
(
stream
));
}
nppSafeCall
(
nppiMagnitudeSqr_32fc32f_C1R
(
src
.
ptr
<
Npp32fc
>
(),
src
.
step
,
dst
.
ptr
<
Npp32f
>
(),
dst
.
step
,
sz
)
);
void
cv
::
gpu
::
cartToPolar
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
mag
,
GpuMat
&
angle
,
bool
angleInDegrees
)
{
::
cartToPolar_caller
(
x
,
y
,
mag
,
false
,
angle
,
angleInDegrees
,
0
);
}
void
cv
::
gpu
::
magnitudeSqr
(
const
GpuMat
&
src1
,
const
GpuMat
&
src2
,
GpuMat
&
dst
)
void
cv
::
gpu
::
cartToPolar
(
const
GpuMat
&
x
,
const
GpuMat
&
y
,
GpuMat
&
mag
,
GpuMat
&
angle
,
bool
angleInDegrees
,
const
Stream
&
stream
)
{
CV_DbgAssert
(
src1
.
type
()
==
src2
.
type
()
&&
src1
.
size
()
==
src2
.
size
(
));
CV_Assert
(
src1
.
type
()
==
CV_32FC1
);
::
cartToPolar_caller
(
x
,
y
,
mag
,
false
,
angle
,
angleInDegrees
,
StreamAccessor
::
getStream
(
stream
));
}
GpuMat
src
(
src1
.
size
(),
CV_32FC2
);
GpuMat
srcs
[]
=
{
src1
,
src2
};
cv
::
gpu
::
merge
(
srcs
,
2
,
src
);
void
cv
::
gpu
::
polarToCart
(
const
GpuMat
&
magnitude
,
const
GpuMat
&
angle
,
GpuMat
&
x
,
GpuMat
&
y
,
bool
angleInDegrees
)
{
::
polarToCart_caller
(
magnitude
,
angle
,
x
,
y
,
angleInDegrees
,
0
);
}
cv
::
gpu
::
magnitudeSqr
(
src
,
dst
);
void
cv
::
gpu
::
polarToCart
(
const
GpuMat
&
magnitude
,
const
GpuMat
&
angle
,
GpuMat
&
x
,
GpuMat
&
y
,
bool
angleInDegrees
,
const
Stream
&
stream
)
{
::
polarToCart_caller
(
magnitude
,
angle
,
x
,
y
,
angleInDegrees
,
StreamAccessor
::
getStream
(
stream
));
}
#endif
/* !defined (HAVE_CUDA) */
\ No newline at end of file
#endif
/* !defined (HAVE_CUDA) */
modules/gpu/src/cuda/mathfunc.cu
0 → 100644
浏览文件 @
e12e4798
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "cuda_shared.hpp"
using
namespace
cv
::
gpu
;
#ifndef CV_PI
#define CV_PI 3.1415926535897932384626433832795f
#endif
namespace
mathfunc_krnls
{
struct
Nothing
{
static
__device__
void
calc
(
int
,
int
,
float
,
float
,
float
*
,
size_t
,
float
)
{
}
};
struct
Magnitude
{
static
__device__
void
calc
(
int
x
,
int
y
,
float
x_data
,
float
y_data
,
float
*
dst
,
size_t
dst_step
,
float
)
{
dst
[
y
*
dst_step
+
x
]
=
sqrtf
(
x_data
*
x_data
+
y_data
*
y_data
);
}
};
struct
MagnitudeSqr
{
static
__device__
void
calc
(
int
x
,
int
y
,
float
x_data
,
float
y_data
,
float
*
dst
,
size_t
dst_step
,
float
)
{
dst
[
y
*
dst_step
+
x
]
=
x_data
*
x_data
+
y_data
*
y_data
;
}
};
struct
Atan2
{
static
__device__
void
calc
(
int
x
,
int
y
,
float
x_data
,
float
y_data
,
float
*
dst
,
size_t
dst_step
,
float
scale
)
{
dst
[
y
*
dst_step
+
x
]
=
scale
*
atan2f
(
y_data
,
x_data
);
}
};
template
<
typename
Mag
,
typename
Angle
>
__global__
void
cartToPolar
(
const
float
*
xptr
,
size_t
x_step
,
const
float
*
yptr
,
size_t
y_step
,
float
*
mag
,
size_t
mag_step
,
float
*
angle
,
size_t
angle_step
,
float
scale
,
int
width
,
int
height
)
{
const
int
x
=
blockDim
.
x
*
blockIdx
.
x
+
threadIdx
.
x
;
const
int
y
=
blockDim
.
y
*
blockIdx
.
y
+
threadIdx
.
y
;
if
(
x
<
width
&&
y
<
height
)
{
float
x_data
=
xptr
[
y
*
x_step
+
x
];
float
y_data
=
yptr
[
y
*
y_step
+
x
];
Mag
::
calc
(
x
,
y
,
x_data
,
y_data
,
mag
,
mag_step
,
scale
);
Angle
::
calc
(
x
,
y
,
x_data
,
y_data
,
angle
,
angle_step
,
scale
);
}
}
struct
NonEmptyMag
{
static
__device__
float
get
(
const
float
*
mag
,
size_t
mag_step
,
int
x
,
int
y
)
{
return
mag
[
y
*
mag_step
+
x
];
}
};
struct
EmptyMag
{
static
__device__
float
get
(
const
float
*
,
size_t
,
int
,
int
)
{
return
1.0
f
;
}
};
template
<
typename
Mag
>
__global__
void
polarToCart
(
const
float
*
mag
,
size_t
mag_step
,
const
float
*
angle
,
size_t
angle_step
,
float
scale
,
float
*
xptr
,
size_t
x_step
,
float
*
yptr
,
size_t
y_step
,
int
width
,
int
height
)
{
const
int
x
=
blockDim
.
x
*
blockIdx
.
x
+
threadIdx
.
x
;
const
int
y
=
blockDim
.
y
*
blockIdx
.
y
+
threadIdx
.
y
;
if
(
x
<
width
&&
y
<
height
)
{
float
mag_data
=
Mag
::
get
(
mag
,
mag_step
,
x
,
y
);
float
angle_data
=
angle
[
y
*
angle_step
+
x
];
float
sin_a
,
cos_a
;
sincosf
(
scale
*
angle_data
,
&
sin_a
,
&
cos_a
);
xptr
[
y
*
x_step
+
x
]
=
mag_data
*
cos_a
;
yptr
[
y
*
y_step
+
x
]
=
mag_data
*
sin_a
;
}
}
}
namespace
cv
{
namespace
gpu
{
namespace
mathfunc
{
template
<
typename
Mag
,
typename
Angle
>
void
cartToPolar_caller
(
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
dim3
threads
(
16
,
16
,
1
);
dim3
grid
(
1
,
1
,
1
);
grid
.
x
=
divUp
(
x
.
cols
,
threads
.
x
);
grid
.
y
=
divUp
(
x
.
rows
,
threads
.
y
);
const
float
scale
=
angleInDegrees
?
(
float
)(
180.0
f
/
CV_PI
)
:
1.
f
;
mathfunc_krnls
::
cartToPolar
<
Mag
,
Angle
><<<
grid
,
threads
,
0
,
stream
>>>
(
x
.
ptr
,
x
.
step
/
sizeof
(
float
),
y
.
ptr
,
y
.
step
/
sizeof
(
float
),
mag
.
ptr
,
mag
.
step
/
sizeof
(
float
),
angle
.
ptr
,
angle
.
step
/
sizeof
(
float
),
scale
,
x
.
cols
,
x
.
rows
);
if
(
stream
==
0
)
cudaSafeCall
(
cudaThreadSynchronize
()
);
}
void
cartToPolar_gpu
(
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
const
DevMem2Df
&
mag
,
bool
magSqr
,
const
DevMem2Df
&
angle
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
typedef
void
(
*
caller_t
)(
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
bool
angleInDegrees
,
cudaStream_t
stream
);
static
const
caller_t
callers
[
2
][
2
][
2
]
=
{
{
{
cartToPolar_caller
<
mathfunc_krnls
::
Magnitude
,
mathfunc_krnls
::
Atan2
>
,
cartToPolar_caller
<
mathfunc_krnls
::
Magnitude
,
mathfunc_krnls
::
Nothing
>
},
{
cartToPolar_caller
<
mathfunc_krnls
::
MagnitudeSqr
,
mathfunc_krnls
::
Atan2
>
,
cartToPolar_caller
<
mathfunc_krnls
::
MagnitudeSqr
,
mathfunc_krnls
::
Nothing
>
,
}
},
{
{
cartToPolar_caller
<
mathfunc_krnls
::
Nothing
,
mathfunc_krnls
::
Atan2
>
,
cartToPolar_caller
<
mathfunc_krnls
::
Nothing
,
mathfunc_krnls
::
Nothing
>
},
{
cartToPolar_caller
<
mathfunc_krnls
::
Nothing
,
mathfunc_krnls
::
Atan2
>
,
cartToPolar_caller
<
mathfunc_krnls
::
Nothing
,
mathfunc_krnls
::
Nothing
>
,
}
}
};
callers
[
mag
.
ptr
==
0
][
magSqr
][
angle
.
ptr
==
0
](
x
,
y
,
mag
,
angle
,
angleInDegrees
,
stream
);
}
template
<
typename
Mag
>
void
polarToCart_caller
(
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
dim3
threads
(
16
,
16
,
1
);
dim3
grid
(
1
,
1
,
1
);
grid
.
x
=
divUp
(
mag
.
cols
,
threads
.
x
);
grid
.
y
=
divUp
(
mag
.
rows
,
threads
.
y
);
const
float
scale
=
angleInDegrees
?
(
float
)(
CV_PI
/
180.0
f
)
:
1.0
f
;
mathfunc_krnls
::
polarToCart
<
Mag
><<<
grid
,
threads
,
0
,
stream
>>>
(
mag
.
ptr
,
mag
.
step
/
sizeof
(
float
),
angle
.
ptr
,
angle
.
step
/
sizeof
(
float
),
scale
,
x
.
ptr
,
x
.
step
/
sizeof
(
float
),
y
.
ptr
,
y
.
step
/
sizeof
(
float
),
mag
.
cols
,
mag
.
rows
);
if
(
stream
==
0
)
cudaSafeCall
(
cudaThreadSynchronize
()
);
}
void
polarToCart_gpu
(
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
bool
angleInDegrees
,
cudaStream_t
stream
)
{
typedef
void
(
*
caller_t
)(
const
DevMem2Df
&
mag
,
const
DevMem2Df
&
angle
,
const
DevMem2Df
&
x
,
const
DevMem2Df
&
y
,
bool
angleInDegrees
,
cudaStream_t
stream
);
static
const
caller_t
callers
[
2
]
=
{
polarToCart_caller
<
mathfunc_krnls
::
NonEmptyMag
>
,
polarToCart_caller
<
mathfunc_krnls
::
EmptyMag
>
};
callers
[
mag
.
ptr
==
0
](
mag
,
angle
,
x
,
y
,
angleInDegrees
,
stream
);
}
}}}
tests/gpu/src/arithm.cpp
浏览文件 @
e12e4798
...
...
@@ -81,7 +81,7 @@ int CV_GpuArithmTest::CheckNorm(const Mat& m1, const Mat& m2)
{
double
ret
=
norm
(
m1
,
m2
,
NORM_INF
);
if
(
ret
<
std
::
numeric_limits
<
double
>::
epsilon
()
)
if
(
ret
<
1e-5
)
return
CvTS
::
OK
;
ts
->
printf
(
CvTS
::
LOG
,
"
\n
Norm: %f
\n
"
,
ret
);
...
...
@@ -99,7 +99,7 @@ int CV_GpuArithmTest::CheckNorm(double d1, double d2)
{
double
ret
=
::
fabs
(
d1
-
d2
);
if
(
ret
<
std
::
numeric_limits
<
double
>::
epsilon
()
)
if
(
ret
<
1e-5
)
return
CvTS
::
OK
;
ts
->
printf
(
CvTS
::
LOG
,
"
\n
Norm: %f
\n
"
,
ret
);
...
...
@@ -605,6 +605,91 @@ struct CV_GpuNppImageMagnitudeTest : public CV_GpuArithmTest
}
};
////////////////////////////////////////////////////////////////////////////////
// phase
struct
CV_GpuNppImagePhaseTest
:
public
CV_GpuArithmTest
{
CV_GpuNppImagePhaseTest
()
:
CV_GpuArithmTest
(
"GPU-NppImagePhase"
,
"phase"
)
{}
int
test
(
const
Mat
&
mat1
,
const
Mat
&
mat2
)
{
if
(
mat1
.
type
()
!=
CV_32FC1
)
{
ts
->
printf
(
CvTS
::
LOG
,
"
\n
Unsupported type
\n
"
);
return
CvTS
::
OK
;
}
cv
::
Mat
cpuRes
;
cv
::
phase
(
mat1
,
mat2
,
cpuRes
);
GpuMat
gpu1
(
mat1
);
GpuMat
gpu2
(
mat2
);
GpuMat
gpuRes
;
cv
::
gpu
::
phase
(
gpu1
,
gpu2
,
gpuRes
);
return
CheckNorm
(
cpuRes
,
gpuRes
);
}
};
////////////////////////////////////////////////////////////////////////////////
// cartToPolar
struct
CV_GpuNppImageCartToPolarTest
:
public
CV_GpuArithmTest
{
CV_GpuNppImageCartToPolarTest
()
:
CV_GpuArithmTest
(
"GPU-NppImageCartToPolar"
,
"cartToPolar"
)
{}
int
test
(
const
Mat
&
mat1
,
const
Mat
&
mat2
)
{
if
(
mat1
.
type
()
!=
CV_32FC1
)
{
ts
->
printf
(
CvTS
::
LOG
,
"
\n
Unsupported type
\n
"
);
return
CvTS
::
OK
;
}
cv
::
Mat
cpuMag
,
cpuAngle
;
cv
::
cartToPolar
(
mat1
,
mat2
,
cpuMag
,
cpuAngle
);
GpuMat
gpu1
(
mat1
);
GpuMat
gpu2
(
mat2
);
GpuMat
gpuMag
,
gpuAngle
;
cv
::
gpu
::
cartToPolar
(
gpu1
,
gpu2
,
gpuMag
,
gpuAngle
);
int
magRes
=
CheckNorm
(
cpuMag
,
gpuMag
);
int
angleRes
=
CheckNorm
(
cpuAngle
,
gpuAngle
);
return
magRes
==
CvTS
::
OK
&&
angleRes
==
CvTS
::
OK
?
CvTS
::
OK
:
CvTS
::
FAIL_GENERIC
;
}
};
////////////////////////////////////////////////////////////////////////////////
// polarToCart
struct
CV_GpuNppImagePolarToCartTest
:
public
CV_GpuArithmTest
{
CV_GpuNppImagePolarToCartTest
()
:
CV_GpuArithmTest
(
"GPU-NppImagePolarToCart"
,
"polarToCart"
)
{}
int
test
(
const
Mat
&
mat1
,
const
Mat
&
mat2
)
{
if
(
mat1
.
type
()
!=
CV_32FC1
)
{
ts
->
printf
(
CvTS
::
LOG
,
"
\n
Unsupported type
\n
"
);
return
CvTS
::
OK
;
}
cv
::
Mat
cpuX
,
cpuY
;
cv
::
polarToCart
(
mat1
,
mat2
,
cpuX
,
cpuY
);
GpuMat
gpu1
(
mat1
);
GpuMat
gpu2
(
mat2
);
GpuMat
gpuX
,
gpuY
;
cv
::
gpu
::
polarToCart
(
gpu1
,
gpu2
,
gpuX
,
gpuY
);
int
xRes
=
CheckNorm
(
cpuX
,
gpuX
);
int
yRes
=
CheckNorm
(
cpuY
,
gpuY
);
return
xRes
==
CvTS
::
OK
&&
yRes
==
CvTS
::
OK
?
CvTS
::
OK
:
CvTS
::
FAIL_GENERIC
;
}
};
/////////////////////////////////////////////////////////////////////////////
/////////////////// tests registration /////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////
...
...
@@ -629,3 +714,6 @@ CV_GpuNppImageLUTTest CV_GpuNppImageLUT_test;
CV_GpuNppImageExpTest
CV_GpuNppImageExp_test
;
CV_GpuNppImageLogTest
CV_GpuNppImageLog_test
;
CV_GpuNppImageMagnitudeTest
CV_GpuNppImageMagnitude_test
;
CV_GpuNppImagePhaseTest
CV_GpuNppImagePhase_test
;
CV_GpuNppImageCartToPolarTest
CV_GpuNppImageCartToPolar_test
;
CV_GpuNppImagePolarToCartTest
CV_GpuNppImagePolarToCart_test
;
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录