Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Greenplum
Opencv
提交
572cfc99
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,发现更多精彩内容 >>
提交
572cfc99
编写于
6月 11, 2013
作者:
R
Roman Donchenko
提交者:
OpenCV Buildbot
6月 11, 2013
浏览文件
操作
浏览文件
下载
差异文件
Merge pull request #973 from pengx17:2.4_oclclahe
上级
3af21cad
1d8cd3a7
变更
5
隐藏空白更改
内联
并排
Showing
5 changed file
with
570 addition
and
1 deletion
+570
-1
modules/ocl/include/opencv2/ocl/ocl.hpp
modules/ocl/include/opencv2/ocl/ocl.hpp
+17
-0
modules/ocl/perf/perf_imgproc.cpp
modules/ocl/perf/perf_imgproc.cpp
+48
-1
modules/ocl/src/imgproc.cpp
modules/ocl/src/imgproc.cpp
+185
-0
modules/ocl/src/opencl/imgproc_clahe.cl
modules/ocl/src/opencl/imgproc_clahe.cl
+275
-0
modules/ocl/test/test_imgproc.cpp
modules/ocl/test/test_imgproc.cpp
+45
-0
未找到文件。
modules/ocl/include/opencv2/ocl/ocl.hpp
浏览文件 @
572cfc99
...
...
@@ -483,6 +483,23 @@ namespace cv
CV_EXPORTS
void
calcHist
(
const
oclMat
&
mat_src
,
oclMat
&
mat_hist
);
//! only 8UC1 and 256 bins is supported now
CV_EXPORTS
void
equalizeHist
(
const
oclMat
&
mat_src
,
oclMat
&
mat_dst
);
//! only 8UC1 is supported now
class
CV_EXPORTS
CLAHE
{
public:
virtual
void
apply
(
const
oclMat
&
src
,
oclMat
&
dst
)
=
0
;
virtual
void
setClipLimit
(
double
clipLimit
)
=
0
;
virtual
double
getClipLimit
()
const
=
0
;
virtual
void
setTilesGridSize
(
Size
tileGridSize
)
=
0
;
virtual
Size
getTilesGridSize
()
const
=
0
;
virtual
void
collectGarbage
()
=
0
;
};
CV_EXPORTS
Ptr
<
cv
::
ocl
::
CLAHE
>
createCLAHE
(
double
clipLimit
=
40.0
,
Size
tileGridSize
=
Size
(
8
,
8
));
//! bilateralFilter
// supports 8UC1 8UC4
CV_EXPORTS
void
bilateralFilter
(
const
oclMat
&
src
,
oclMat
&
dst
,
int
d
,
double
sigmaColor
,
double
sigmaSpave
,
int
borderType
=
BORDER_DEFAULT
);
...
...
modules/ocl/perf/perf_imgproc.cpp
浏览文件 @
572cfc99
...
...
@@ -921,4 +921,51 @@ PERFTEST(remap)
}
}
}
\ No newline at end of file
}
///////////// CLAHE ////////////////////////
PERFTEST
(
CLAHE
)
{
Mat
src
,
dst
,
ocl_dst
;
cv
::
ocl
::
oclMat
d_src
,
d_dst
;
int
all_type
[]
=
{
CV_8UC1
};
std
::
string
type_name
[]
=
{
"CV_8UC1"
};
double
clipLimit
=
40.0
;
cv
::
Ptr
<
cv
::
CLAHE
>
clahe
=
cv
::
createCLAHE
(
clipLimit
);
cv
::
Ptr
<
cv
::
ocl
::
CLAHE
>
d_clahe
=
cv
::
ocl
::
createCLAHE
(
clipLimit
);
for
(
int
size
=
Min_Size
;
size
<=
Max_Size
;
size
*=
Multiple
)
{
for
(
size_t
j
=
0
;
j
<
sizeof
(
all_type
)
/
sizeof
(
int
);
j
++
)
{
SUBTEST
<<
size
<<
'x'
<<
size
<<
"; "
<<
type_name
[
j
]
;
gen
(
src
,
size
,
size
,
all_type
[
j
],
0
,
256
);
CPU_ON
;
clahe
->
apply
(
src
,
dst
);
CPU_OFF
;
d_src
.
upload
(
src
);
WARMUP_ON
;
d_clahe
->
apply
(
d_src
,
d_dst
);
WARMUP_OFF
;
ocl_dst
=
d_dst
;
TestSystem
::
instance
().
ExpectedMatNear
(
dst
,
ocl_dst
,
1.0
);
GPU_ON
;
d_clahe
->
apply
(
d_src
,
d_dst
);
GPU_OFF
;
GPU_FULL_ON
;
d_src
.
upload
(
src
);
d_clahe
->
apply
(
d_src
,
d_dst
);
d_dst
.
download
(
dst
);
GPU_FULL_OFF
;
}
}
}
modules/ocl/src/imgproc.cpp
浏览文件 @
572cfc99
...
...
@@ -25,6 +25,7 @@
// Xu Pang, pangxu010@163.com
// Wu Zailong, bullet@yeah.net
// Wenju He, wenju@multicorewareinc.com
// Sen Liu, swjtuls1987@126.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
...
...
@@ -80,6 +81,7 @@ namespace cv
extern
const
char
*
imgproc_calcHarris
;
extern
const
char
*
imgproc_calcMinEigenVal
;
extern
const
char
*
imgproc_convolve
;
extern
const
char
*
imgproc_clahe
;
////////////////////////////////////OpenCL call wrappers////////////////////////////
template
<
typename
T
>
struct
index_and_sizeof
;
...
...
@@ -1511,6 +1513,189 @@ namespace cv
openCLExecuteKernel
(
clCxt
,
&
imgproc_histogram
,
kernelName
,
globalThreads
,
localThreads
,
args
,
-
1
,
-
1
);
LUT
(
mat_src
,
lut
,
mat_dst
);
}
////////////////////////////////////////////////////////////////////////
// CLAHE
namespace
clahe
{
inline
int
divUp
(
int
total
,
int
grain
)
{
return
(
total
+
grain
-
1
)
/
grain
*
grain
;
}
static
void
calcLut
(
const
oclMat
&
src
,
oclMat
&
dst
,
const
int
tilesX
,
const
int
tilesY
,
const
cv
::
Size
tileSize
,
const
int
clipLimit
,
const
float
lutScale
)
{
cl_int2
tile_size
;
tile_size
.
s
[
0
]
=
tileSize
.
width
;
tile_size
.
s
[
1
]
=
tileSize
.
height
;
std
::
vector
<
pair
<
size_t
,
const
void
*>
>
args
;
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
src
.
data
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
dst
.
data
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
src
.
step
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
dst
.
step
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int2
),
(
void
*
)
&
tile_size
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
tilesX
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
clipLimit
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_float
),
(
void
*
)
&
lutScale
));
String
kernelName
=
"calcLut"
;
size_t
localThreads
[
3
]
=
{
32
,
8
,
1
};
size_t
globalThreads
[
3
]
=
{
tilesX
*
localThreads
[
0
],
tilesY
*
localThreads
[
1
],
1
};
bool
is_cpu
=
queryDeviceInfo
<
IS_CPU_DEVICE
,
bool
>
();
if
(
is_cpu
)
{
openCLExecuteKernel
(
Context
::
getContext
(),
&
imgproc_clahe
,
kernelName
,
globalThreads
,
localThreads
,
args
,
-
1
,
-
1
,
(
char
*
)
" -D CPU"
);
}
else
{
cl_kernel
kernel
=
openCLGetKernelFromSource
(
Context
::
getContext
(),
&
imgproc_clahe
,
kernelName
);
int
wave_size
=
queryDeviceInfo
<
WAVEFRONT_SIZE
,
int
>
(
kernel
);
openCLSafeCall
(
clReleaseKernel
(
kernel
));
static
char
opt
[
20
]
=
{
0
};
sprintf
(
opt
,
" -D WAVE_SIZE=%d"
,
wave_size
);
openCLExecuteKernel
(
Context
::
getContext
(),
&
imgproc_clahe
,
kernelName
,
globalThreads
,
localThreads
,
args
,
-
1
,
-
1
,
opt
);
}
}
static
void
transform
(
const
oclMat
&
src
,
oclMat
&
dst
,
const
oclMat
&
lut
,
const
int
tilesX
,
const
int
tilesY
,
const
cv
::
Size
tileSize
)
{
cl_int2
tile_size
;
tile_size
.
s
[
0
]
=
tileSize
.
width
;
tile_size
.
s
[
1
]
=
tileSize
.
height
;
std
::
vector
<
pair
<
size_t
,
const
void
*>
>
args
;
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
src
.
data
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
dst
.
data
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
lut
.
data
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
src
.
step
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
dst
.
step
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
lut
.
step
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
src
.
cols
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
src
.
rows
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int2
),
(
void
*
)
&
tile_size
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
tilesX
));
args
.
push_back
(
std
::
make_pair
(
sizeof
(
cl_int
),
(
void
*
)
&
tilesY
));
String
kernelName
=
"transform"
;
size_t
localThreads
[
3
]
=
{
32
,
8
,
1
};
size_t
globalThreads
[
3
]
=
{
divUp
(
src
.
cols
,
localThreads
[
0
]),
divUp
(
src
.
rows
,
localThreads
[
1
]),
1
};
openCLExecuteKernel
(
Context
::
getContext
(),
&
imgproc_clahe
,
kernelName
,
globalThreads
,
localThreads
,
args
,
-
1
,
-
1
);
}
}
namespace
{
class
CLAHE_Impl
:
public
cv
::
ocl
::
CLAHE
{
public:
CLAHE_Impl
(
double
clipLimit
=
40.0
,
int
tilesX
=
8
,
int
tilesY
=
8
);
cv
::
AlgorithmInfo
*
info
()
const
;
void
apply
(
const
oclMat
&
src
,
oclMat
&
dst
);
void
setClipLimit
(
double
clipLimit
);
double
getClipLimit
()
const
;
void
setTilesGridSize
(
cv
::
Size
tileGridSize
);
cv
::
Size
getTilesGridSize
()
const
;
void
collectGarbage
();
private:
double
clipLimit_
;
int
tilesX_
;
int
tilesY_
;
oclMat
srcExt_
;
oclMat
lut_
;
};
CLAHE_Impl
::
CLAHE_Impl
(
double
clipLimit
,
int
tilesX
,
int
tilesY
)
:
clipLimit_
(
clipLimit
),
tilesX_
(
tilesX
),
tilesY_
(
tilesY
)
{
}
void
CLAHE_Impl
::
apply
(
const
oclMat
&
src
,
oclMat
&
dst
)
{
CV_Assert
(
src
.
type
()
==
CV_8UC1
);
dst
.
create
(
src
.
size
(),
src
.
type
()
);
const
int
histSize
=
256
;
ensureSizeIsEnough
(
tilesX_
*
tilesY_
,
histSize
,
CV_8UC1
,
lut_
);
cv
::
Size
tileSize
;
oclMat
srcForLut
;
if
(
src
.
cols
%
tilesX_
==
0
&&
src
.
rows
%
tilesY_
==
0
)
{
tileSize
=
cv
::
Size
(
src
.
cols
/
tilesX_
,
src
.
rows
/
tilesY_
);
srcForLut
=
src
;
}
else
{
cv
::
ocl
::
copyMakeBorder
(
src
,
srcExt_
,
0
,
tilesY_
-
(
src
.
rows
%
tilesY_
),
0
,
tilesX_
-
(
src
.
cols
%
tilesX_
),
cv
::
BORDER_REFLECT_101
,
cv
::
Scalar
());
tileSize
=
cv
::
Size
(
srcExt_
.
cols
/
tilesX_
,
srcExt_
.
rows
/
tilesY_
);
srcForLut
=
srcExt_
;
}
const
int
tileSizeTotal
=
tileSize
.
area
();
const
float
lutScale
=
static_cast
<
float
>
(
histSize
-
1
)
/
tileSizeTotal
;
int
clipLimit
=
0
;
if
(
clipLimit_
>
0.0
)
{
clipLimit
=
static_cast
<
int
>
(
clipLimit_
*
tileSizeTotal
/
histSize
);
clipLimit
=
std
::
max
(
clipLimit
,
1
);
}
clahe
::
calcLut
(
srcForLut
,
lut_
,
tilesX_
,
tilesY_
,
tileSize
,
clipLimit
,
lutScale
);
//finish();
clahe
::
transform
(
src
,
dst
,
lut_
,
tilesX_
,
tilesY_
,
tileSize
);
}
void
CLAHE_Impl
::
setClipLimit
(
double
clipLimit
)
{
clipLimit_
=
clipLimit
;
}
double
CLAHE_Impl
::
getClipLimit
()
const
{
return
clipLimit_
;
}
void
CLAHE_Impl
::
setTilesGridSize
(
cv
::
Size
tileGridSize
)
{
tilesX_
=
tileGridSize
.
width
;
tilesY_
=
tileGridSize
.
height
;
}
cv
::
Size
CLAHE_Impl
::
getTilesGridSize
()
const
{
return
cv
::
Size
(
tilesX_
,
tilesY_
);
}
void
CLAHE_Impl
::
collectGarbage
()
{
srcExt_
.
release
();
lut_
.
release
();
}
}
cv
::
Ptr
<
cv
::
ocl
::
CLAHE
>
createCLAHE
(
double
clipLimit
,
cv
::
Size
tileGridSize
)
{
return
new
CLAHE_Impl
(
clipLimit
,
tileGridSize
.
width
,
tileGridSize
.
height
);
}
//////////////////////////////////bilateralFilter////////////////////////////////////////////////////
static
void
oclbilateralFilter_8u
(
const
oclMat
&
src
,
oclMat
&
dst
,
int
d
,
...
...
modules/ocl/src/opencl/imgproc_clahe.cl
0 → 100644
浏览文件 @
572cfc99
/*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
)
2010-2012,
Multicoreware,
Inc.,
all
rights
reserved.
//
Copyright
(
C
)
2010-2012,
Advanced
Micro
Devices,
Inc.,
all
rights
reserved.
//
Third
party
copyrights
are
property
of
their
respective
owners.
//
//
@Authors
//
Sen
Liu,
swjtuls1987@126.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
oclMaterials
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*/
#
ifndef
WAVE_SIZE
#
define
WAVE_SIZE
1
#
endif
int
calc_lut
(
__local
int*
smem,
int
val,
int
tid
)
{
smem[tid]
=
val
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
==
0
)
{
for
(
int
i
=
1
; i < 256; ++i)
{
smem[i]
+=
smem[i
-
1]
;
}
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
return
smem[tid]
;
}
#
ifdef
CPU
void
reduce
(
volatile
__local
int*
smem,
int
val,
int
tid
)
{
smem[tid]
=
val
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
128
)
{
smem[tid]
=
val
+=
smem[tid
+
128]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
64
)
{
smem[tid]
=
val
+=
smem[tid
+
64]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
32
)
{
smem[tid]
+=
smem[tid
+
32]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
16
)
{
smem[tid]
+=
smem[tid
+
16]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
8
)
{
smem[tid]
+=
smem[tid
+
8]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
4
)
{
smem[tid]
+=
smem[tid
+
4]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
2
)
{
smem[tid]
+=
smem[tid
+
2]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
1
)
{
smem[256]
=
smem[tid]
+
smem[tid
+
1]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
}
#
else
void
reduce
(
__local
volatile
int*
smem,
int
val,
int
tid
)
{
smem[tid]
=
val
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
128
)
{
smem[tid]
=
val
+=
smem[tid
+
128]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
64
)
{
smem[tid]
=
val
+=
smem[tid
+
64]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
32
)
{
smem[tid]
+=
smem[tid
+
32]
;
#
if
WAVE_SIZE
<
32
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
16
)
{
#
endif
smem[tid]
+=
smem[tid
+
16]
;
#
if
WAVE_SIZE
<
16
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
tid
<
8
)
{
#
endif
smem[tid]
+=
smem[tid
+
8]
;
smem[tid]
+=
smem[tid
+
4]
;
smem[tid]
+=
smem[tid
+
2]
;
smem[tid]
+=
smem[tid
+
1]
;
}
}
#
endif
__kernel
void
calcLut
(
__global
__const
uchar
*
src,
__global
uchar
*
lut,
const
int
srcStep,
const
int
dstStep,
const
int2
tileSize,
const
int
tilesX,
const
int
clipLimit,
const
float
lutScale
)
{
__local
int
smem[512]
;
const
int
tx
=
get_group_id
(
0
)
;
const
int
ty
=
get_group_id
(
1
)
;
const
unsigned
int
tid
=
get_local_id
(
1
)
*
get_local_size
(
0
)
+
get_local_id
(
0
)
;
smem[tid]
=
0
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
for
(
int
i
=
get_local_id
(
1
)
; i < tileSize.y; i += get_local_size(1))
{
__global
const
uchar*
srcPtr
=
src
+
mad24
(
ty
*
tileSize.y
+
i,
srcStep,
tx
*
tileSize.x
)
;
for
(
int
j
=
get_local_id
(
0
)
; j < tileSize.x; j += get_local_size(0))
{
const
int
data
=
srcPtr[j]
;
atomic_inc
(
&smem[data]
)
;
}
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
int
tHistVal
=
smem[tid]
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
clipLimit
>
0
)
{
//
clip
histogram
bar
int
clipped
=
0
;
if
(
tHistVal
>
clipLimit
)
{
clipped
=
tHistVal
-
clipLimit
;
tHistVal
=
clipLimit
;
}
//
find
number
of
overall
clipped
samples
reduce
(
smem,
clipped,
tid
)
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
#
ifdef
CPU
clipped
=
smem[256]
;
#
else
clipped
=
smem[0]
;
#
endif
//
broadcast
evaluated
value
__local
int
totalClipped
;
if
(
tid
==
0
)
totalClipped
=
clipped
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
//
redistribute
clipped
samples
evenly
int
redistBatch
=
totalClipped
/
256
;
tHistVal
+=
redistBatch
;
int
residual
=
totalClipped
-
redistBatch
*
256
;
if
(
tid
<
residual
)
++tHistVal
;
}
const
int
lutVal
=
calc_lut
(
smem,
tHistVal,
tid
)
;
uint
ires
=
(
uint
)
convert_int_rte
(
lutScale
*
lutVal
)
;
lut[
(
ty
*
tilesX
+
tx
)
*
dstStep
+
tid]
=
convert_uchar
(
clamp
(
ires,
(
uint
)
0
,
(
uint
)
255
))
;
}
__kernel
void
transform
(
__global
__const
uchar
*
src,
__global
uchar
*
dst,
__global
uchar
*
lut,
const
int
srcStep,
const
int
dstStep,
const
int
lutStep,
const
int
cols,
const
int
rows,
const
int2
tileSize,
const
int
tilesX,
const
int
tilesY
)
{
const
int
x
=
get_global_id
(
0
)
;
const
int
y
=
get_global_id
(
1
)
;
if
(
x
>=
cols
||
y
>=
rows
)
return
;
const
float
tyf
=
(
convert_float
(
y
)
/
tileSize.y
)
-
0.5f
;
int
ty1
=
convert_int_rtn
(
tyf
)
;
int
ty2
=
ty1
+
1
;
const
float
ya
=
tyf
-
ty1
;
ty1
=
max
(
ty1,
0
)
;
ty2
=
min
(
ty2,
tilesY
-
1
)
;
const
float
txf
=
(
convert_float
(
x
)
/
tileSize.x
)
-
0.5f
;
int
tx1
=
convert_int_rtn
(
txf
)
;
int
tx2
=
tx1
+
1
;
const
float
xa
=
txf
-
tx1
;
tx1
=
max
(
tx1,
0
)
;
tx2
=
min
(
tx2,
tilesX
-
1
)
;
const
int
srcVal
=
src[mad24
(
y,
srcStep,
x
)
]
;
float
res
=
0
;
res
+=
lut[mad24
(
ty1
*
tilesX
+
tx1,
lutStep,
srcVal
)
]
*
((
1.0f
-
xa
)
*
(
1.0f
-
ya
))
;
res
+=
lut[mad24
(
ty1
*
tilesX
+
tx2,
lutStep,
srcVal
)
]
*
((
xa
)
*
(
1.0f
-
ya
))
;
res
+=
lut[mad24
(
ty2
*
tilesX
+
tx1,
lutStep,
srcVal
)
]
*
((
1.0f
-
xa
)
*
(
ya
))
;
res
+=
lut[mad24
(
ty2
*
tilesX
+
tx2,
lutStep,
srcVal
)
]
*
((
xa
)
*
(
ya
))
;
uint
ires
=
(
uint
)
convert_int_rte
(
res
)
;
dst[mad24
(
y,
dstStep,
x
)
]
=
convert_uchar
(
clamp
(
ires,
(
uint
)
0
,
(
uint
)
255
))
;
}
modules/ocl/test/test_imgproc.cpp
浏览文件 @
572cfc99
...
...
@@ -23,6 +23,7 @@
// Rock Li, Rock.Li@amd.com
// Wu Zailong, bullet@yeah.net
// Xu Pang, pangxu010@163.com
// Sen Liu, swjtuls1987@126.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
...
...
@@ -1393,6 +1394,46 @@ TEST_P(calcHist, Mat)
EXPECT_MAT_NEAR
(
dst_hist
,
cpu_hist
,
0.0
);
}
}
///////////////////////////////////////////////////////////////////////////////////////////////////////
// CLAHE
namespace
{
IMPLEMENT_PARAM_CLASS
(
ClipLimit
,
double
)
}
PARAM_TEST_CASE
(
CLAHE
,
cv
::
Size
,
ClipLimit
)
{
cv
::
Size
size
;
double
clipLimit
;
cv
::
Mat
src
;
cv
::
Mat
dst_gold
;
cv
::
ocl
::
oclMat
g_src
;
cv
::
ocl
::
oclMat
g_dst
;
virtual
void
SetUp
()
{
size
=
GET_PARAM
(
0
);
clipLimit
=
GET_PARAM
(
1
);
cv
::
RNG
&
rng
=
TS
::
ptr
()
->
get_rng
();
src
=
randomMat
(
rng
,
size
,
CV_8UC1
,
0
,
256
,
false
);
g_src
.
upload
(
src
);
}
};
TEST_P
(
CLAHE
,
Accuracy
)
{
cv
::
Ptr
<
cv
::
ocl
::
CLAHE
>
clahe
=
cv
::
ocl
::
createCLAHE
(
clipLimit
);
clahe
->
apply
(
g_src
,
g_dst
);
cv
::
Mat
dst
(
g_dst
);
cv
::
Ptr
<
cv
::
CLAHE
>
clahe_gold
=
cv
::
createCLAHE
(
clipLimit
);
clahe_gold
->
apply
(
src
,
dst_gold
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst
,
1.0
);
}
///////////////////////////Convolve//////////////////////////////////
PARAM_TEST_CASE
(
ConvolveTestBase
,
MatType
,
bool
)
...
...
@@ -1643,6 +1684,10 @@ INSTANTIATE_TEST_CASE_P(histTestBase, calcHist, Combine(
ONE_TYPE
(
CV_32SC1
)
//no use
));
INSTANTIATE_TEST_CASE_P
(
ImgProc
,
CLAHE
,
Combine
(
Values
(
cv
::
Size
(
128
,
128
),
cv
::
Size
(
113
,
113
),
cv
::
Size
(
1300
,
1300
)),
Values
(
0.0
,
40.0
)));
//INSTANTIATE_TEST_CASE_P(ConvolveTestBase, Convolve, Combine(
// Values(CV_32FC1, CV_32FC1),
// Values(false))); // Values(false) is the reserved parameter
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录