Skip to content
Projects
Groups
Snippets
Help
Loading...
Sign in / Register
Toggle navigation
O
opencv
Project
Project
Details
Activity
Cycle Analytics
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Charts
Issues
0
Issues
0
List
Board
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Charts
Packages
Packages
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Charts
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
submodule
opencv
Commits
ef9f6905
Commit
ef9f6905
authored
Oct 30, 2013
by
Andrey Pavlenko
Committed by
OpenCV Buildbot
Oct 30, 2013
Browse files
Options
Browse Files
Download
Plain Diff
Merge pull request #1709 from ilya-lavrenov:ocl_cornerharris
parents
5795bb32
bf6b5ee9
Hide whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
186 additions
and
164 deletions
+186
-164
imgproc.cpp
modules/ocl/src/imgproc.cpp
+8
-6
imgproc_calcHarris.cl
modules/ocl/src/opencl/imgproc_calcHarris.cl
+75
-75
imgproc_calcMinEigenVal.cl
modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl
+65
-69
test_imgproc.cpp
modules/ocl/test/test_imgproc.cpp
+35
-11
utility.cpp
modules/ocl/test/utility.cpp
+3
-3
No files found.
modules/ocl/src/imgproc.cpp
View file @
ef9f6905
...
...
@@ -991,6 +991,7 @@ namespace cv
args
.
push_back
(
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
dst
.
cols
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
dst
.
step
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_float
)
,
(
void
*
)
&
k
));
openCLExecuteKernel
(
dst
.
clCxt
,
source
,
kernelName
,
gt
,
lt
,
args
,
-
1
,
-
1
,
buildOptions
.
c_str
());
}
...
...
@@ -1006,15 +1007,15 @@ namespace cv
{
if
(
!
src
.
clCxt
->
supportsFeature
(
FEATURE_CL_DOUBLE
)
&&
src
.
depth
()
==
CV_64F
)
{
CV_Error
(
CV_OpenCLDoubleNotSupported
,
"Select device doesn't support double"
);
CV_Error
(
CV_OpenCLDoubleNotSupported
,
"Select
ed
device doesn't support double"
);
return
;
}
CV_Assert
(
src
.
cols
>=
blockSize
/
2
&&
src
.
rows
>=
blockSize
/
2
);
CV_Assert
(
borderType
==
cv
::
BORDER_CONSTANT
||
borderType
==
cv
::
BORDER_REFLECT101
||
borderType
==
cv
::
BORDER_REPLICATE
||
borderType
==
cv
::
BORDER_REFLECT
);
extractCovData
(
src
,
dx
,
dy
,
blockSize
,
ksize
,
borderType
);
dst
.
create
(
src
.
size
(),
CV_32F
);
dst
.
create
(
src
.
size
(),
CV_32F
C1
);
corner_ocl
(
&
imgproc_calcHarris
,
"calcHarris"
,
blockSize
,
static_cast
<
float
>
(
k
),
dx
,
dy
,
dst
,
borderType
);
}
...
...
@@ -1028,12 +1029,13 @@ namespace cv
{
if
(
!
src
.
clCxt
->
supportsFeature
(
FEATURE_CL_DOUBLE
)
&&
src
.
depth
()
==
CV_64F
)
{
CV_Error
(
CV_OpenCLDoubleNotSupported
,
"
select device do
n't support double"
);
CV_Error
(
CV_OpenCLDoubleNotSupported
,
"
Selected device does
n't support double"
);
return
;
}
CV_Assert
(
src
.
cols
>=
blockSize
/
2
&&
src
.
rows
>=
blockSize
/
2
);
CV_Assert
(
borderType
==
cv
::
BORDER_CONSTANT
||
borderType
==
cv
::
BORDER_REFLECT101
||
borderType
==
cv
::
BORDER_REPLICATE
||
borderType
==
cv
::
BORDER_REFLECT
);
CV_Assert
(
borderType
==
cv
::
BORDER_CONSTANT
||
borderType
==
cv
::
BORDER_REFLECT101
||
borderType
==
cv
::
BORDER_REPLICATE
||
borderType
==
cv
::
BORDER_REFLECT
);
extractCovData
(
src
,
dx
,
dy
,
blockSize
,
ksize
,
borderType
);
dst
.
create
(
src
.
size
(),
CV_32F
);
...
...
modules/ocl/src/opencl/imgproc_calcHarris.cl
View file @
ef9f6905
...
...
@@ -43,60 +43,64 @@
//
//M*/
#
if
defined
(
DOUBLE_SUPPORT
)
#
pragma
OPENCL
EXTENSION
cl_khr_fp64:enable
#
endif
///////////////////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////Macro
for
border
type////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////////////////
#
ifdef
BORDER_REPLICATE
//BORDER_REPLICATE:
aaaaaa|abcdefgh|hhhhhhh
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
(
l_edge
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
(
r_edge
)
-1
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
(
t_edge
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
(
b_edge
)
-1
:
(
addr
))
#
endif
#
ifdef
BORDER_CONSTANT
#
elif
defined
BORDER_REPLICATE
#
define
EXTRAPOLATE
(
x,
maxV
)
\
{
\
x
=
max
(
min
(
x,
maxV
-
1
)
,
0
)
; \
}
#
elif
defined
BORDER_WRAP
#
define
EXTRAPOLATE
(
x,
maxV
)
\
{
\
if
(
x
<
0
)
\
x
-=
((
x
-
maxV
+
1
)
/
maxV
)
*
maxV
; \
if
(
x
>=
maxV
)
\
x
%=
maxV
; \
}
#
elif
defined
(
BORDER_REFLECT
)
|
| defined(BORDER_REFLECT101)
#define EXTRAPOLATE_(x, maxV, delta) \
{ \
if (maxV == 1) \
x = 0; \
else \
do \
{ \
if ( x < 0 ) \
x = -x - 1 + delta; \
else \
x = maxV - 1 - (x - maxV) - delta; \
} \
while (x >= maxV |
|
x
<
0
)
; \
}
#
ifdef
BORDER_REFLECT
//BORDER_REFLECT:
fedcba|abcdefgh|hgfedcb
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
-
(
i
)
-1
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
-
(
i
)
-1+
((
r_edge
)
<<1
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
-
(
i
)
-1
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
-
(
i
)
-1+
((
b_edge
)
<<1
)
:
(
addr
))
#
endif
#
ifdef
BORDER_REFLECT101
//BORDER_REFLECT101:
gfedcb|abcdefgh|gfedcba
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
-
(
i
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
-
(
i
)
-2+
((
r_edge
)
<<1
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
-
(
i
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
-
(
i
)
-2+
((
b_edge
)
<<1
)
:
(
addr
))
#
define
EXTRAPOLATE
(
x,
maxV
)
EXTRAPOLATE_
(
x,
maxV,
0
)
#
else
#
define
EXTRAPOLATE
(
x,
maxV
)
EXTRAPOLATE_
(
x,
maxV,
1
)
#
endif
#
ifdef
BORDER_WRAP
//BORDER_WRAP:
cdefgh|abcdefgh|abcdefg
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
(
i
)
+
(
r_edge
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
(
i
)
-
(
r_edge
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
(
i
)
+
(
b_edge
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
(
i
)
-
(
b_edge
)
:
(
addr
))
#
else
#
error
No
extrapolation
method
#
endif
#
define
THREADS
256
#
define
ELEM
(
i,
l_edge,
r_edge,
elem1,
elem2
)
(
i
)
>=
(
l_edge
)
&&
(
i
)
<
(
r_edge
)
?
(
elem1
)
:
(
elem2
)
///////////////////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////calcHarris////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////////////
__kernel
void
calcHarris
(
__global
const
float
*Dx,__global
const
float
*Dy,
__global
float
*dst,
int
dx_offset,
int
dx_whole_rows,
int
dx_whole_cols,
int
dx_step
,
int
dy_offset,
int
dy_whole_rows,
int
dy_whole_cols,
int
dy
_step,
int
dst_offset,
int
dst_rows,
int
dst_cols,
int
dst
_step,
float
k
)
__kernel
void
calcHarris
(
__global
const
float
*Dx,
__global
const
float
*Dy,
__global
float
*dst
,
int
dx_offset,
int
dx_whole_rows,
int
dx_whole_cols,
int
dx
_step,
int
dy_offset,
int
dy_whole_rows,
int
dy_whole_cols,
int
dy
_step,
int
dst_offset,
int
dst_rows,
int
dst_cols,
int
dst_step,
float
k
)
{
int
col
=
get_local_id
(
0
)
;
const
int
gX
=
get_group_id
(
0
)
;
const
int
gY
=
get_group_id
(
1
)
;
const
int
glx
=
get_global_id
(
0
)
;
const
int
gly
=
get_global_id
(
1
)
;
int
gX
=
get_group_id
(
0
)
;
int
gY
=
get_group_id
(
1
)
;
int
glx
=
get_global_id
(
0
)
;
int
gly
=
get_global_id
(
1
)
;
int
dx_x_off
=
(
dx_offset
%
dx_step
)
>>
2
;
int
dx_y_off
=
dx_offset
/
dx_step
;
...
...
@@ -112,41 +116,38 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
int
dst_startX
=
gX
*
(
THREADS-ksX+1
)
+
dst_x_off
;
int
dst_startY
=
(
gY
<<
1
)
+
dst_y_off
;
float
dx_data[ksY+1],dy_data[ksY+1],data[3][ksY+1]
;
float
dx_data[ksY+1],dy_data[ksY+1],
data[3][ksY+1]
;
__local
float
temp[6][THREADS]
;
#
ifdef
BORDER_CONSTANT
bool
dx_con,dy_con
;
float
dx_s,dy_s
;
for
(
int
i=0
; i < ksY+1; i++)
float
dx_s,
dy_s
;
for
(
int
i=0
; i < ksY+1; i++)
{
dx_con
=
dx_startX+col
>=
0
&&
dx_startX+col
<
dx_whole_cols
&&
dx_startY+i
>=
0
&&
dx_startY+i
<
dx_whole_rows
;
dx_s
=
Dx[
(
dx_startY+i
)
*
(
dx_step>>2
)
+
(
dx_startX+col
)
]
;
dx_data[i]
=
dx_con
?
dx_s
:
0.0
;
dy_con
=
dy_startX+col
>=
0
&&
dy_startX+col
<
dy_whole_cols
&&
dy_startY+i
>=
0
&&
dy_startY+i
<
dy_whole_rows
;
dy_s
=
Dy[
(
dy_startY+i
)
*
(
dy_step>>2
)
+
(
dy_startX+col
)
]
;
dy_data[i]
=
dy_con
?
dy_s
:
0.0
;
data[0][i]
=
dx_data[i]
*
dx_data[i]
;
data[1][i]
=
dx_data[i]
*
dy_data[i]
;
data[2][i]
=
dy_data[i]
*
dy_data[i]
;
}
#
else
int
clamped_col
=
min
(
dst_cols,
col
)
;
for
(
int
i=0
; i < ksY+1; i++)
for
(
int
i=0
; i < ksY+1; i++)
{
int
dx_selected_row
;
int
dx_selected_col
;
dx_selected_row
=
ADDR_H
(
dx_startY+i,
0
,
dx_whole_rows
)
;
dx_selected_row
=
ADDR_B
(
dx_startY+i,
dx_whole_rows,
dx_selected_row
)
;
dx_selected_col
=
ADDR_L
(
dx_startX+clamped_col,
0
,
dx_whole_cols
)
;
dx_selected_col
=
ADDR_R
(
dx_startX+clamped_col,
dx_whole_cols,
dx_selected_col
)
;
int
dx_selected_row
=
dx_startY+i,
dx_selected_col
=
dx_startX+clamped_col
;
EXTRAPOLATE
(
dx_selected_row,
dx_whole_rows
)
EXTRAPOLATE
(
dx_selected_col,
dx_whole_cols
)
dx_data[i]
=
Dx[dx_selected_row
*
(
dx_step>>2
)
+
dx_selected_col]
;
int
dy_selected_row
;
int
dy_selected_col
;
dy_selected_row
=
ADDR_H
(
dy_startY+i,
0
,
dy_whole_rows
)
;
dy_selected_row
=
ADDR_B
(
dy_startY+i,
dy_whole_rows,
dy_selected_row
)
;
dy_selected_col
=
ADDR_L
(
dy_startX+clamped_col,
0
,
dy_whole_cols
)
;
dy_selected_col
=
ADDR_R
(
dy_startX+clamped_col,
dy_whole_cols,
dy_selected_col
)
;
int
dy_selected_row
=
dy_startY+i,
dy_selected_col
=
dy_startX+clamped_col
;
EXTRAPOLATE
(
dy_selected_row,
dy_whole_rows
)
EXTRAPOLATE
(
dy_selected_col,
dy_whole_cols
)
dy_data[i]
=
Dy[dy_selected_row
*
(
dy_step>>2
)
+
dy_selected_col]
;
data[0][i]
=
dx_data[i]
*
dx_data[i]
;
...
...
@@ -155,45 +156,44 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
}
#
endif
float
sum0
=
0.0
,
sum1
=
0.0
,
sum2
=
0.0
;
for
(
int
i=1
; i < ksY; i++)
for
(
int
i=1
; i < ksY; i++)
{
sum0
+=
(
data[0][i]
)
;
sum1
+=
(
data[1][i]
)
;
sum2
+=
(
data[2][i]
)
;
sum0
+=
data[0][i]
;
sum1
+=
data[1][i]
;
sum2
+=
data[2][i]
;
}
float
sum01,sum02,sum11,sum12,sum21,sum22
;
sum01
=
sum0
+
(
data[0][0]
)
;
sum02
=
sum0
+
(
data[0][ksY]
)
;
float
sum01
=
sum0
+
data[0][0]
;
float
sum02
=
sum0
+
data[0][ksY]
;
temp[0][col]
=
sum01
;
temp[1][col]
=
sum02
;
sum11
=
sum1
+
(
data[1][0]
)
;
sum12
=
sum1
+
(
data[1][ksY]
)
;
float
sum11
=
sum1
+
data[1][0]
;
float
sum12
=
sum1
+
data[1][ksY]
;
temp[2][col]
=
sum11
;
temp[3][col]
=
sum12
;
sum21
=
sum2
+
(
data[2][0]
)
;
sum22
=
sum2
+
(
data[2][ksY]
)
;
float
sum21
=
sum2
+
data[2][0]
;
float
sum22
=
sum2
+
data[2][ksY]
;
temp[4][col]
=
sum21
;
temp[5][col]
=
sum22
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
col
<
(
THREADS-
(
ksX-1
)))
if
(
col
<
(
THREADS-
(
ksX
-
1
)))
{
col
+=
anX
;
int
posX
=
dst_startX
-
dst_x_off
+
col
-
anX
;
int
posY
=
(
gly
<<
1
)
;
int
till
=
(
ksX
+
1
)
%2
;
float
tmp_sum[6]={
0.0
,
0.0
,
0.0
,
0.0
,
0.0
,
0.0
}
;
for
(
int
k=0
; k<6; k++)
for
(
int
i=-anX
; i<=anX - till; i++)
{
float
tmp_sum[6]
=
{
0.0
,
0.0
,
0.0
,
0.0
,
0.0
,
0.0
}
;
for
(
int
k=0
; k<6; k++)
for
(
int
i=-anX
; i<=anX - till; i++)
tmp_sum[k]
+=
temp[k][col+i]
;
}
if
(
posX
<
dst_cols
&&
(
posY
)
<
dst_rows
)
if
(
posX
<
dst_cols
&&
(
posY
)
<
dst_rows
)
{
dst[
(
dst_startY+0
)
*
(
dst_step>>2
)
+
dst_startX
+
col
-
anX]
=
tmp_sum[0]
*
tmp_sum[4]
-
tmp_sum[2]
*
tmp_sum[2]
-
k
*
(
tmp_sum[0]
+
tmp_sum[4]
)
*
(
tmp_sum[0]
+
tmp_sum[4]
)
;
}
if
(
posX
<
dst_cols
&&
(
posY
+
1
)
<
dst_rows
)
if
(
posX
<
dst_cols
&&
(
posY
+
1
)
<
dst_rows
)
{
dst[
(
dst_startY+1
)
*
(
dst_step>>2
)
+
dst_startX
+
col
-
anX]
=
tmp_sum[1]
*
tmp_sum[5]
-
tmp_sum[3]
*
tmp_sum[3]
-
k
*
(
tmp_sum[1]
+
tmp_sum[5]
)
*
(
tmp_sum[1]
+
tmp_sum[5]
)
;
...
...
modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl
View file @
ef9f6905
...
...
@@ -43,60 +43,63 @@
//
//M*/
#
if
defined
(
DOUBLE_SUPPORT
)
#
pragma
OPENCL
EXTENSION
cl_khr_fp64:enable
#
endif
///////////////////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////Macro
for
border
type////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////////////////
#
ifdef
BORDER_REPLICATE
//BORDER_REPLICATE:
aaaaaa|abcdefgh|hhhhhhh
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
(
l_edge
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
(
r_edge
)
-1
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
(
t_edge
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
(
b_edge
)
-1
:
(
addr
))
#
endif
#
ifdef
BORDER_CONSTANT
#
elif
defined
BORDER_REPLICATE
#
define
EXTRAPOLATE
(
x,
maxV
)
\
{
\
x
=
max
(
min
(
x,
maxV
-
1
)
,
0
)
; \
}
#
elif
defined
BORDER_WRAP
#
define
EXTRAPOLATE
(
x,
maxV
)
\
{
\
if
(
x
<
0
)
\
x
-=
((
x
-
maxV
+
1
)
/
maxV
)
*
maxV
; \
if
(
x
>=
maxV
)
\
x
%=
maxV
; \
}
#
elif
defined
(
BORDER_REFLECT
)
|
| defined(BORDER_REFLECT101)
#define EXTRAPOLATE_(x, maxV, delta) \
{ \
if (maxV == 1) \
x = 0; \
else \
do \
{ \
if ( x < 0 ) \
x = -x - 1 + delta; \
else \
x = maxV - 1 - (x - maxV) - delta; \
} \
while (x >= maxV |
|
x
<
0
)
; \
}
#
ifdef
BORDER_REFLECT
//BORDER_REFLECT:
fedcba|abcdefgh|hgfedcb
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
-
(
i
)
-1
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
-
(
i
)
-1+
((
r_edge
)
<<1
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
-
(
i
)
-1
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
-
(
i
)
-1+
((
b_edge
)
<<1
)
:
(
addr
))
#
endif
#
ifdef
BORDER_REFLECT101
//BORDER_REFLECT101:
gfedcb|abcdefgh|gfedcba
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
-
(
i
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
-
(
i
)
-2+
((
r_edge
)
<<1
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
-
(
i
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
-
(
i
)
-2+
((
b_edge
)
<<1
)
:
(
addr
))
#
define
EXTRAPOLATE
(
x,
maxV
)
EXTRAPOLATE_
(
x,
maxV,
0
)
#
else
#
define
EXTRAPOLATE
(
x,
maxV
)
EXTRAPOLATE_
(
x,
maxV,
1
)
#
endif
#
ifdef
BORDER_WRAP
//BORDER_WRAP:
cdefgh|abcdefgh|abcdefg
#
define
ADDR_L
(
i,
l_edge,
r_edge
)
((
i
)
<
(
l_edge
)
?
(
i
)
+
(
r_edge
)
:
(
i
))
#
define
ADDR_R
(
i,
r_edge,
addr
)
((
i
)
>=
(
r_edge
)
?
(
i
)
-
(
r_edge
)
:
(
addr
))
#
define
ADDR_H
(
i,
t_edge,
b_edge
)
((
i
)
<
(
t_edge
)
?
(
i
)
+
(
b_edge
)
:
(
i
))
#
define
ADDR_B
(
i,
b_edge,
addr
)
((
i
)
>=
(
b_edge
)
?
(
i
)
-
(
b_edge
)
:
(
addr
))
#
else
#
error
No
extrapolation
method
#
endif
#
define
THREADS
256
#
define
ELEM
(
i,
l_edge,
r_edge,
elem1,
elem2
)
(
i
)
>=
(
l_edge
)
&&
(
i
)
<
(
r_edge
)
?
(
elem1
)
:
(
elem2
)
///////////////////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////calcHarris////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////////////
__kernel
void
calcMinEigenVal
(
__global
const
float
*Dx,__global
const
float
*Dy,
__global
float
*dst,
int
dx_offset,
int
dx_whole_rows,
int
dx_whole_cols,
int
dx_step,
int
dy_offset,
int
dy_whole_rows,
int
dy_whole_cols,
int
dy_step,
int
dst_offset,
int
dst_rows,
int
dst_cols,
int
dst_step,
float
k
)
int
dst_offset,
int
dst_rows,
int
dst_cols,
int
dst_step,
float
k
)
{
int
col
=
get_local_id
(
0
)
;
const
int
gX
=
get_group_id
(
0
)
;
const
int
gY
=
get_group_id
(
1
)
;
const
int
glx
=
get_global_id
(
0
)
;
const
int
gly
=
get_global_id
(
1
)
;
int
gX
=
get_group_id
(
0
)
;
int
gY
=
get_group_id
(
1
)
;
int
glx
=
get_global_id
(
0
)
;
int
gly
=
get_global_id
(
1
)
;
int
dx_x_off
=
(
dx_offset
%
dx_step
)
>>
2
;
int
dx_y_off
=
dx_offset
/
dx_step
;
...
...
@@ -112,12 +115,13 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
int
dst_startX
=
gX
*
(
THREADS-ksX+1
)
+
dst_x_off
;
int
dst_startY
=
(
gY
<<
1
)
+
dst_y_off
;
float
dx_data[ksY+1],
dy_data[ksY+1],
data[3][ksY+1]
;
float
dx_data[ksY+1],
dy_data[ksY+1],
data[3][ksY+1]
;
__local
float
temp[6][THREADS]
;
#
ifdef
BORDER_CONSTANT
bool
dx_con,dy_con
;
float
dx_s,dy_s
;
for
(
int
i=0
; i < ksY+1; i++)
bool
dx_con,
dy_con
;
float
dx_s,
dy_s
;
for
(
int
i=0
; i < ksY+1; i++)
{
dx_con
=
dx_startX+col
>=
0
&&
dx_startX+col
<
dx_whole_cols
&&
dx_startY+i
>=
0
&&
dx_startY+i
<
dx_whole_rows
;
dx_s
=
Dx[
(
dx_startY+i
)
*
(
dx_step>>2
)
+
(
dx_startX+col
)
]
;
...
...
@@ -131,23 +135,16 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
}
#
else
int
clamped_col
=
min
(
dst_cols,
col
)
;
for
(
int
i=0
; i < ksY+1; i++)
for
(
int
i=0
; i < ksY+1; i++)
{
int
dx_selected_row
;
int
dx_selected_col
;
dx_selected_row
=
ADDR_H
(
dx_startY+i,
0
,
dx_whole_rows
)
;
dx_selected_row
=
ADDR_B
(
dx_startY+i,
dx_whole_rows,
dx_selected_row
)
;
dx_selected_col
=
ADDR_L
(
dx_startX+clamped_col,
0
,
dx_whole_cols
)
;
dx_selected_col
=
ADDR_R
(
dx_startX+clamped_col,
dx_whole_cols,
dx_selected_col
)
;
int
dx_selected_row
=
dx_startY+i,
dx_selected_col
=
dx_startX+clamped_col
;
EXTRAPOLATE
(
dx_selected_row,
dx_whole_rows
)
EXTRAPOLATE
(
dx_selected_col,
dx_whole_cols
)
dx_data[i]
=
Dx[dx_selected_row
*
(
dx_step>>2
)
+
dx_selected_col]
;
int
dy_selected_row
;
int
dy_selected_col
;
dy_selected_row
=
ADDR_H
(
dy_startY+i,
0
,
dy_whole_rows
)
;
dy_selected_row
=
ADDR_B
(
dy_startY+i,
dy_whole_rows,
dy_selected_row
)
;
dy_selected_col
=
ADDR_L
(
dy_startX+clamped_col,
0
,
dy_whole_cols
)
;
dy_selected_col
=
ADDR_R
(
dy_startX+clamped_col,
dy_whole_cols,
dy_selected_col
)
;
int
dy_selected_row
=
dy_startY+i,
dy_selected_col
=
dy_startX+clamped_col
;
EXTRAPOLATE
(
dy_selected_row,
dy_whole_rows
)
EXTRAPOLATE
(
dy_selected_col,
dy_whole_cols
)
dy_data[i]
=
Dy[dy_selected_row
*
(
dy_step>>2
)
+
dy_selected_col]
;
data[0][i]
=
dx_data[i]
*
dx_data[i]
;
...
...
@@ -156,38 +153,37 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
}
#
endif
float
sum0
=
0.0
,
sum1
=
0.0
,
sum2
=
0.0
;
for
(
int
i=1
; i < ksY; i++)
for
(
int
i=1
; i < ksY; i++)
{
sum0
+=
(
data[0][i]
)
;
sum1
+=
(
data[1][i]
)
;
sum2
+=
(
data[2][i]
)
;
}
float
sum01,sum02,sum11,sum12,sum21,sum22
;
sum01
=
sum0
+
(
data[0][0]
)
;
sum02
=
sum0
+
(
data[0][ksY]
)
;
float
sum01
=
sum0
+
(
data[0][0]
)
;
float
sum02
=
sum0
+
(
data[0][ksY]
)
;
temp[0][col]
=
sum01
;
temp[1][col]
=
sum02
;
sum11
=
sum1
+
(
data[1][0]
)
;
sum12
=
sum1
+
(
data[1][ksY]
)
;
float
sum11
=
sum1
+
(
data[1][0]
)
;
float
sum12
=
sum1
+
(
data[1][ksY]
)
;
temp[2][col]
=
sum11
;
temp[3][col]
=
sum12
;
sum21
=
sum2
+
(
data[2][0]
)
;
sum22
=
sum2
+
(
data[2][ksY]
)
;
float
sum21
=
sum2
+
(
data[2][0]
)
;
float
sum22
=
sum2
+
(
data[2][ksY]
)
;
temp[4][col]
=
sum21
;
temp[5][col]
=
sum22
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
col
<
(
THREADS-
(
ksX-1
)))
{
col
+=
anX
;
int
posX
=
dst_startX
-
dst_x_off
+
col
-
anX
;
int
posY
=
(
gly
<<
1
)
;
int
till
=
(
ksX
+
1
)
%2
;
float
tmp_sum[6]={
0.0
,
0.0
,
0.0
,
0.0
,
0.0
,
0.0
}
;
for
(
int
k=0
; k<6; k++)
for
(
int
i=-anX
; i<=anX - till; i++)
{
float
tmp_sum[6]
=
{
0.0
,
0.0
,
0.0
,
0.0
,
0.0
,
0.0
}
;
for
(
int
k=0
; k<6; k++)
for
(
int
i=-anX
; i<=anX - till; i++)
tmp_sum[k]
+=
temp[k][col+i]
;
}
if
(
posX
<
dst_cols
&&
(
posY
)
<
dst_rows
)
{
...
...
@@ -196,7 +192,7 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
float
c
=
tmp_sum[4]
*
0.5f
;
dst[
(
dst_startY+0
)
*
(
dst_step>>2
)
+
dst_startX
+
col
-
anX]
=
(
float
)((
a+c
)
-
sqrt
((
a-c
)
*
(
a-c
)
+
b*b
))
;
}
if
(
posX
<
dst_cols
&&
(
posY
+
1
)
<
dst_rows
)
if
(
posX
<
dst_cols
&&
(
posY
+
1
)
<
dst_rows
)
{
float
a
=
tmp_sum[1]
*
0.5f
;
float
b
=
tmp_sum[3]
;
...
...
modules/ocl/test/test_imgproc.cpp
View file @
ef9f6905
...
...
@@ -80,7 +80,7 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType,
useRoi
=
GET_PARAM
(
3
);
}
void
random_roi
()
v
irtual
v
oid
random_roi
()
{
Size
roiSize
=
randomSize
(
1
,
MAX_VALUE
);
Border
srcBorder
=
randomBorder
(
0
,
useRoi
?
MAX_VALUE
:
0
);
...
...
@@ -191,7 +191,31 @@ OCL_TEST_P(EqualizeHist, Mat)
////////////////////////////////cornerMinEigenVal//////////////////////////////////////////
typedef
ImgprocTestBase
CornerMinEigenVal
;
struct
CornerTestBase
:
public
ImgprocTestBase
{
virtual
void
random_roi
()
{
Mat
image
=
readImageType
(
"gpu/stereobm/aloe-L.png"
,
type
);
ASSERT_FALSE
(
image
.
empty
());
Size
roiSize
=
image
.
size
();
Border
srcBorder
=
randomBorder
(
0
,
useRoi
?
MAX_VALUE
:
0
);
Size
wholeSize
=
Size
(
roiSize
.
width
+
srcBorder
.
lef
+
srcBorder
.
rig
,
roiSize
.
height
+
srcBorder
.
top
+
srcBorder
.
bot
);
src
=
randomMat
(
wholeSize
,
type
,
-
255
,
255
,
false
);
src_roi
=
src
(
Rect
(
srcBorder
.
lef
,
srcBorder
.
top
,
roiSize
.
width
,
roiSize
.
height
));
image
.
copyTo
(
src_roi
);
Border
dstBorder
=
randomBorder
(
0
,
useRoi
?
MAX_VALUE
:
0
);
randomSubMat
(
dst_whole
,
dst_roi
,
roiSize
,
dstBorder
,
CV_32FC1
,
5
,
16
);
generateOclMat
(
gsrc_whole
,
gsrc_roi
,
src
,
roiSize
,
srcBorder
);
generateOclMat
(
gdst_whole
,
gdst_roi
,
dst_whole
,
roiSize
,
dstBorder
);
}
};
typedef
CornerTestBase
CornerMinEigenVal
;
OCL_TEST_P
(
CornerMinEigenVal
,
Mat
)
{
...
...
@@ -204,13 +228,13 @@ OCL_TEST_P(CornerMinEigenVal, Mat)
cornerMinEigenVal
(
src_roi
,
dst_roi
,
blockSize
,
apertureSize
,
borderType
);
ocl
::
cornerMinEigenVal
(
gsrc_roi
,
gdst_roi
,
blockSize
,
apertureSize
,
borderType
);
Near
(
1.0
);
Near
(
0.02
);
}
}
////////////////////////////////cornerHarris//////////////////////////////////////////
typedef
Imgproc
TestBase
CornerHarris
;
typedef
Corner
TestBase
CornerHarris
;
OCL_TEST_P
(
CornerHarris
,
Mat
)
{
...
...
@@ -219,12 +243,12 @@ OCL_TEST_P(CornerHarris, Mat)
random_roi
();
int
apertureSize
=
3
;
double
k
=
2.0
;
double
k
=
randomDouble
(
0.01
,
0.9
)
;
cornerHarris
(
src_roi
,
dst_roi
,
blockSize
,
apertureSize
,
k
,
borderType
);
ocl
::
cornerHarris
(
gsrc_roi
,
gdst_roi
,
blockSize
,
apertureSize
,
k
,
borderType
);
Near
(
1.0
);
Near
(
0.02
);
}
}
...
...
@@ -484,19 +508,19 @@ INSTANTIATE_TEST_CASE_P(Imgproc, EqualizeHist, Combine(
Bool
()));
INSTANTIATE_TEST_CASE_P
(
Imgproc
,
CornerMinEigenVal
,
Combine
(
Values
(
CV_8UC1
,
CV_32FC1
),
Values
(
3
),
// TODO some fails when blockSize != 3 (for example 5)
Values
((
int
)
BORDER_
REFLECT
,
(
int
)
BORDER_CONSTANT
,
(
int
)
BORDER_REPLICATE
),
// TODO does not work with (int)BORDER_REFLECT101
Values
(
(
MatType
)
CV_8UC1
,
(
MatType
)
CV_32FC1
),
Values
(
3
,
5
),
Values
((
int
)
BORDER_
CONSTANT
,
(
int
)
BORDER_REPLICATE
,
(
int
)
BORDER_REFLECT
,
(
int
)
BORDER_REFLECT101
),
Bool
()));
INSTANTIATE_TEST_CASE_P
(
Imgproc
,
CornerHarris
,
Combine
(
Values
((
MatType
)
CV_8UC1
),
// TODO does not work properly with CV_32FC1
Values
(
3
,
5
),
Values
(
(
int
)
BORDER_REFLECT101
,
(
int
)
BORDER_REFLECT
,
(
int
)
BORDER_CONSTANT
,
(
int
)
BORDER_REPLICATE
),
Values
(
(
int
)
BORDER_CONSTANT
,
(
int
)
BORDER_REPLICATE
,
(
int
)
BORDER_REFLECT
,
(
int
)
BORDER_REFLECT_101
),
Bool
()));
INSTANTIATE_TEST_CASE_P
(
Imgproc
,
Integral
,
Combine
(
Values
((
MatType
)
CV_8UC1
),
// TODO does work with CV_32F, CV_64F
Values
((
MatType
)
CV_8UC1
),
// TODO does
not
work with CV_32F, CV_64F
Values
(
0
),
// not used
Values
(
0
),
// not used
Bool
()));
...
...
modules/ocl/test/utility.cpp
View file @
ef9f6905
...
...
@@ -233,12 +233,12 @@ double checkRectSimilarity(Size sz, std::vector<Rect>& ob1, std::vector<Rect>& o
void
showDiff
(
const
Mat
&
gold
,
const
Mat
&
actual
,
double
eps
,
bool
alwaysShow
)
{
Mat
diff
;
Mat
diff
,
diff_thresh
;
absdiff
(
gold
,
actual
,
diff
);
diff
.
convertTo
(
diff
,
CV_32F
);
threshold
(
diff
,
diff
,
eps
,
255.0
,
cv
::
THRESH_BINARY
);
threshold
(
diff
,
diff
_thresh
,
eps
,
255.0
,
cv
::
THRESH_BINARY
);
if
(
alwaysShow
||
cv
::
countNonZero
(
diff
.
reshape
(
1
))
>
0
)
if
(
alwaysShow
||
cv
::
countNonZero
(
diff
_thresh
.
reshape
(
1
))
>
0
)
{
namedWindow
(
"gold"
,
WINDOW_NORMAL
);
namedWindow
(
"actual"
,
WINDOW_NORMAL
);
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment