Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Greenplum
Opencv
提交
5c0dabf1
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,发现更多精彩内容 >>
提交
5c0dabf1
编写于
12月 12, 2013
作者:
I
Ilya Lavrenov
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
added cv::equalizeHist to T-API
上级
f3e2bfea
变更
4
隐藏空白更改
内联
并排
Showing
4 changed file
with
193 addition
and
6 deletion
+193
-6
modules/imgproc/src/histogram.cpp
modules/imgproc/src/histogram.cpp
+70
-5
modules/imgproc/src/opencl/histogram.cl
modules/imgproc/src/opencl/histogram.cl
+120
-0
modules/imgproc/test/ocl/test_imgproc.cpp
modules/imgproc/test/ocl/test_imgproc.cpp
+1
-1
modules/imgproc/test/test_imgproc_umat.cpp
modules/imgproc/test/test_imgproc_umat.cpp
+2
-0
未找到文件。
modules/imgproc/src/histogram.cpp
浏览文件 @
5c0dabf1
...
...
@@ -38,7 +38,9 @@
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#include "opencl_kernels.hpp"
namespace
cv
{
...
...
@@ -3129,17 +3131,80 @@ CV_IMPL void cvEqualizeHist( const CvArr* srcarr, CvArr* dstarr )
cv
::
equalizeHist
(
cv
::
cvarrToMat
(
srcarr
),
cv
::
cvarrToMat
(
dstarr
));
}
namespace
cv
{
enum
{
BINS
=
256
};
static
bool
ocl_calcHist
(
InputArray
_src
,
OutputArray
_hist
)
{
int
compunits
=
ocl
::
Device
::
getDefault
().
maxComputeUnits
();
size_t
wgs
=
ocl
::
Device
::
getDefault
().
maxWorkGroupSize
();
ocl
::
Kernel
k1
(
"calculate_histogram"
,
ocl
::
imgproc
::
histogram_oclsrc
,
format
(
"-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d"
,
BINS
,
compunits
,
wgs
));
if
(
k1
.
empty
())
return
false
;
_hist
.
create
(
1
,
BINS
,
CV_32SC1
);
UMat
src
=
_src
.
getUMat
(),
hist
=
_hist
.
getUMat
(),
ghist
(
1
,
BINS
*
compunits
,
CV_32SC1
);
k1
.
args
(
ocl
::
KernelArg
::
ReadOnly
(
src
),
ocl
::
KernelArg
::
PtrWriteOnly
(
ghist
),
(
int
)
src
.
total
());
size_t
globalsize
=
compunits
*
wgs
;
if
(
!
k1
.
run
(
1
,
&
globalsize
,
&
wgs
,
false
))
return
false
;
ocl
::
Kernel
k2
(
"merge_histogram"
,
ocl
::
imgproc
::
histogram_oclsrc
,
format
(
"-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d"
,
BINS
,
compunits
,
(
int
)
wgs
));
if
(
k2
.
empty
())
return
false
;
k2
.
args
(
ocl
::
KernelArg
::
PtrReadOnly
(
ghist
),
ocl
::
KernelArg
::
PtrWriteOnly
(
hist
));
return
k2
.
run
(
1
,
&
wgs
,
&
wgs
,
false
);
}
static
bool
ocl_equalizeHist
(
InputArray
_src
,
OutputArray
_dst
)
{
size_t
wgs
=
ocl
::
Device
::
getDefault
().
maxWorkGroupSize
();
// calculation of histogram
UMat
hist
;
if
(
!
ocl_calcHist
(
_src
,
hist
))
return
false
;
UMat
lut
(
1
,
256
,
CV_8UC1
);
ocl
::
Kernel
k
(
"calcLUT"
,
ocl
::
imgproc
::
histogram_oclsrc
,
format
(
"-D BINS=%d -D HISTS_COUNT=1 -D WGS=%d"
,
BINS
,
(
int
)
wgs
));
k
.
args
(
ocl
::
KernelArg
::
PtrWriteOnly
(
lut
),
ocl
::
KernelArg
::
PtrReadOnly
(
hist
),
(
int
)
_src
.
total
());
// calculation of LUT
if
(
!
k
.
run
(
1
,
&
wgs
,
&
wgs
,
false
))
return
false
;
// execute LUT transparently
LUT
(
_src
,
lut
,
_dst
);
return
true
;
}
}
void
cv
::
equalizeHist
(
InputArray
_src
,
OutputArray
_dst
)
{
Mat
src
=
_src
.
getMat
();
CV_Assert
(
src
.
type
()
==
CV_8UC1
);
CV_Assert
(
_src
.
type
()
==
CV_8UC1
);
_dst
.
create
(
src
.
size
(),
src
.
type
()
);
Mat
dst
=
_dst
.
getMat
()
;
if
(
_src
.
empty
())
return
;
if
(
src
.
empty
(
))
if
(
ocl
::
useOpenCL
()
&&
_dst
.
isUMat
()
&&
ocl_equalizeHist
(
_src
,
_dst
))
return
;
Mat
src
=
_src
.
getMat
();
_dst
.
create
(
src
.
size
(),
src
.
type
()
);
Mat
dst
=
_dst
.
getMat
();
Mutex
histogramLockInstance
;
const
int
hist_sz
=
EqualizeHistCalcHist_Invoker
::
HIST_SZ
;
...
...
modules/imgproc/src/opencl/histogram.cl
0 → 100644
浏览文件 @
5c0dabf1
//
License
Agreement
//
For
Open
Source
Computer
Vision
Library
//
//
Copyright
(
C
)
2010-2012,
Institute
Of
Software
Chinese
Academy
Of
Science,
all
rights
reserved.
//
Copyright
(
C
)
2010-2012,
Advanced
Micro
Devices,
Inc.,
all
rights
reserved.
//
Copyright
(
C
)
2010-2012,
Multicoreware,
Inc.,
all
rights
reserved.
//
Third
party
copyrights
are
property
of
their
respective
owners.
//
//
@Authors
//
Niko
Li,
newlife20080214@gmail.com
//
Jia
Haipeng,
jiahaipeng95@gmail.com
//
Xu
Pang,
pangxu010@163.com
//
Wenju
He,
wenju@multicorewareinc.com
//
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.
//
//
__kernel
void
calculate_histogram
(
__global
const
uchar
*
src,
int
src_step,
int
src_offset,
int
src_rows,
int
src_cols,
__global
uchar
*
hist,
int
total
)
{
int
lid
=
get_local_id
(
0
)
;
int
id
=
get_global_id
(
0
)
;
int
gid
=
get_group_id
(
0
)
;
__local
int
localhist[BINS]
;
for
(
int
i
=
lid
; i < BINS; i += WGS)
localhist[i]
=
0
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
for
(
int
grain
=
HISTS_COUNT
*
WGS
; id < total; id += grain)
{
int
src_index
=
mad24
(
id
/
src_cols,
src_step,
src_offset
+
id
%
src_cols
)
;
atomic_inc
(
localhist
+
(
int
)
src[src_index]
)
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
for
(
int
i
=
lid
; i < BINS; i += WGS)
*
(
__global
int
*
)(
hist
+
mad24
(
gid,
BINS
*
(
int
)
sizeof
(
int
)
,
i
*
(
int
)
sizeof
(
int
)))
=
localhist[i]
;
}
__kernel
void
merge_histogram
(
__global
const
int
*
ghist,
__global
int
*
hist
)
{
int
lid
=
get_local_id
(
0
)
;
#
pragma
unroll
for
(
int
i
=
lid
; i < BINS; i += WGS)
hist[i]
=
0
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
#
pragma
unroll
for
(
int
i
=
0
; i < HISTS_COUNT; ++i)
{
#
pragma
unroll
for
(
int
j
=
lid
; j < BINS; j += WGS)
hist[j]
+=
ghist[mad24
(
i,
BINS,
j
)
]
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
}
}
__kernel
void
calcLUT
(
__global
uchar
*
dst,
__constant
int
*
hist,
int
total
)
{
int
lid
=
get_local_id
(
0
)
;
__local
int
sumhist[BINS]
;
__local
float
scale
;
sumhist[lid]
=
hist[lid]
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
lid
==
0
)
{
int
sum
=
0
,
i
=
0
;
while
(
!sumhist[i]
)
++i
;
if
(
total
==
sumhist[i]
)
{
scale
=
1
;
for
(
int
j
=
0
; j < BINS; ++j)
sumhist[i]
=
i
;
}
else
{
scale
=
255.f
/
(
total
-
sumhist[i]
)
;
for
(
sumhist[i++]
=
0
; i < BINS; i++)
{
sum
+=
sumhist[i]
;
sumhist[i]
=
sum
;
}
}
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
#
pragma
unroll
for
(
int
i
=
lid
; i < BINS; i += WGS)
dst[i]=
convert_uchar_sat_rte
(
convert_float
(
sumhist[i]
)
*
scale
)
;
}
modules/imgproc/test/ocl/test_imgproc.cpp
浏览文件 @
5c0dabf1
...
...
@@ -184,7 +184,7 @@ OCL_TEST_P(EqualizeHist, Mat)
OCL_OFF
(
cv
::
equalizeHist
(
src_roi
,
dst_roi
));
OCL_ON
(
cv
::
equalizeHist
(
usrc_roi
,
udst_roi
));
Near
(
1
.1
);
Near
(
1
);
}
}
...
...
modules/imgproc/test/test_imgproc_umat.cpp
浏览文件 @
5c0dabf1
...
...
@@ -76,6 +76,8 @@ protected:
destroyWindow("equalized gray");
#endif
ts
->
set_failed_test_info
(
cvtest
::
TS
::
OK
);
(
void
)
uresult
.
getMat
(
ACCESS_READ
);
}
};
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录