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
5e02b204
Commit
5e02b204
authored
Nov 10, 2013
by
Ilya Lavrenov
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
added RGB -> XYZ conversion to ocl::cvtColor
parent
33ae6420
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
149 additions
and
29 deletions
+149
-29
color.cpp
modules/ocl/src/color.cpp
+65
-6
cvt_color.cl
modules/ocl/src/opencl/cvt_color.cl
+37
-0
test_color.cpp
modules/ocl/test/test_color.cpp
+47
-23
No files found.
modules/ocl/src/color.cpp
View file @
5e02b204
...
...
@@ -50,7 +50,8 @@
using
namespace
cv
;
using
namespace
cv
::
ocl
;
static
void
fromRGB_caller
(
const
oclMat
&
src
,
oclMat
&
dst
,
int
bidx
,
const
std
::
string
&
kernelName
)
static
void
fromRGB_caller
(
const
oclMat
&
src
,
oclMat
&
dst
,
int
bidx
,
const
std
::
string
&
kernelName
,
const
oclMat
&
data
=
oclMat
())
{
int
src_offset
=
src
.
offset
/
src
.
elemSize1
(),
src_step
=
src
.
step1
();
int
dst_offset
=
dst
.
offset
/
dst
.
elemSize1
(),
dst_step
=
dst
.
step1
();
...
...
@@ -68,6 +69,9 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::
args
.
push_back
(
make_pair
(
sizeof
(
cl_int
)
,
(
void
*
)
&
src_offset
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_int
)
,
(
void
*
)
&
dst_offset
));
if
(
!
data
.
empty
())
args
.
push_back
(
make_pair
(
sizeof
(
cl_mem
)
,
(
void
*
)
&
data
.
data
));
size_t
gt
[
3
]
=
{
dst
.
cols
,
dst
.
rows
,
1
},
lt
[
3
]
=
{
16
,
16
,
1
};
openCLExecuteKernel
(
src
.
clCxt
,
&
cvt_color
,
kernelName
.
c_str
(),
gt
,
lt
,
args
,
-
1
,
-
1
,
build_options
.
c_str
());
}
...
...
@@ -112,10 +116,10 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
*/
case
CV_BGR2GRAY
:
case
CV_BGRA2GRAY
:
case
CV_RGB2GRAY
:
case
CV_BGR2GRAY
:
case
CV_RGBA2GRAY
:
case
CV_BGRA2GRAY
:
{
CV_Assert
(
scn
==
3
||
scn
==
4
);
bidx
=
code
==
CV_BGR2GRAY
||
code
==
CV_BGRA2GRAY
?
0
:
2
;
...
...
@@ -190,9 +194,64 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
/*
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
case CV_BGR2YCrCb: case CV_RGB2YCrCb:
case CV_BGR2XYZ: case CV_RGB2XYZ:
case CV_XYZ2BGR: case CV_XYZ2RGB:
*/
case
CV_BGR2XYZ
:
case
CV_RGB2XYZ
:
{
CV_Assert
(
scn
==
3
||
scn
==
4
);
bidx
=
code
==
CV_BGR2XYZ
?
0
:
2
;
dst
.
create
(
sz
,
CV_MAKE_TYPE
(
depth
,
3
));
void
*
pdata
=
NULL
;
if
(
depth
==
CV_32F
)
{
float
coeffs
[]
=
{
0.412453
f
,
0.357580
f
,
0.180423
f
,
0.212671
f
,
0.715160
f
,
0.072169
f
,
0.019334
f
,
0.119193
f
,
0.950227
f
};
if
(
bidx
==
0
)
{
std
::
swap
(
coeffs
[
0
],
coeffs
[
2
]);
std
::
swap
(
coeffs
[
3
],
coeffs
[
5
]);
std
::
swap
(
coeffs
[
6
],
coeffs
[
8
]);
}
pdata
=
coeffs
;
}
else
{
int
coeffs
[]
=
{
1689
,
1465
,
739
,
871
,
2929
,
296
,
79
,
488
,
3892
};
if
(
bidx
==
0
)
{
std
::
swap
(
coeffs
[
0
],
coeffs
[
2
]);
std
::
swap
(
coeffs
[
3
],
coeffs
[
5
]);
std
::
swap
(
coeffs
[
6
],
coeffs
[
8
]);
}
pdata
=
coeffs
;
}
oclMat
oclCoeffs
(
1
,
9
,
depth
==
CV_32F
?
CV_32F
:
CV_32S
,
pdata
);
fromRGB_caller
(
src
,
dst
,
bidx
,
"RGB2XYZ"
,
oclCoeffs
);
break
;
}
case
CV_XYZ2BGR
:
case
CV_XYZ2RGB
:
{
if
(
dcn
<=
0
)
dcn
=
3
;
CV_Assert
(
scn
==
3
&&
(
dcn
==
3
||
dcn
==
4
));
bidx
=
code
==
CV_XYZ2BGR
?
0
:
2
;
dst
.
create
(
sz
,
CV_MAKE_TYPE
(
depth
,
dcn
));
toRGB_caller
(
src
,
dst
,
bidx
,
"XYZ2RGB"
);
break
;
}
/*
case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:
case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:
case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:
...
...
modules/ocl/src/opencl/cvt_color.cl
View file @
5e02b204
...
...
@@ -48,6 +48,7 @@
#
if
defined
(
DEPTH_0
)
#
define
DATA_TYPE
uchar
#
define
COEFF_TYPE
int
#
define
MAX_NUM
255
#
define
HALF_MAX
128
#
define
SAT_CAST
(
num
)
convert_uchar_sat_rte
(
num
)
...
...
@@ -55,6 +56,7 @@
#
if
defined
(
DEPTH_2
)
#
define
DATA_TYPE
ushort
#
define
COEFF_TYPE
int
#
define
MAX_NUM
65535
#
define
HALF_MAX
32768
#
define
SAT_CAST
(
num
)
convert_ushort_sat_rte
(
num
)
...
...
@@ -62,6 +64,7 @@
#
if
defined
(
DEPTH_5
)
#
define
DATA_TYPE
float
#
define
COEFF_TYPE
float
#
define
MAX_NUM
1.0f
#
define
HALF_MAX
0.5f
#
define
SAT_CAST
(
num
)
(
num
)
...
...
@@ -331,3 +334,37 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int chan
dst[dst_idx
+
3]
=
MAX_NUM
;
}
}
/////////////////////////////////////
RGB
<->
XYZ
//////////////////////////////////////
#
pragma
OPENCL
EXTENSION
cl_amd_printf:enable
__kernel
void
RGB2XYZ
(
int
cols,
int
rows,
int
src_step,
int
dst_step,
int
bidx,
__global
const
DATA_TYPE*
src,
__global
DATA_TYPE*
dst,
int
src_offset,
int
dst_offset,
__global
COEFF_TYPE
*
coeffs
)
{
int
x
=
get_global_id
(
0
)
;
int
y
=
get_global_id
(
1
)
;
if
(
y
<
rows
&&
x
<
cols
)
{
x
<<=
2
;
int
src_idx
=
mad24
(
y,
src_step,
src_offset
+
x
)
;
int
dst_idx
=
mad24
(
y,
dst_step,
dst_offset
+
x
)
;
DATA_TYPE
r
=
src[src_idx],
g
=
src[src_idx
+
1],
b
=
src[src_idx
+
2]
;
#
ifdef
DEPTH_5
DATA_TYPE
x
=
r
*
coeffs[0]
+
g
*
coeffs[1]
+
b
*
coeffs[2]
;
DATA_TYPE
y
=
r
*
coeffs[3]
+
g
*
coeffs[4]
+
b
*
coeffs[5]
;
DATA_TYPE
z
=
r
*
coeffs[6]
+
g
*
coeffs[7]
+
b
*
coeffs[8]
;
#
else
DATA_TYPE
x
=
CV_DESCALE
(
r
*
coeffs[0]
+
g
*
coeffs[1]
+
b
*
coeffs[2],
xyz_shift
)
;
DATA_TYPE
y
=
CV_DESCALE
(
r
*
coeffs[3]
+
g
*
coeffs[4]
+
b
*
coeffs[5],
xyz_shift
)
;
DATA_TYPE
z
=
CV_DESCALE
(
r
*
coeffs[6]
+
g
*
coeffs[7]
+
b
*
coeffs[8],
xyz_shift
)
;
#
endif
dst[dst_idx]
=
SAT_CAST
(
x
)
;
dst[dst_idx
+
1]
=
SAT_CAST
(
y
)
;
dst[dst_idx
+
2]
=
SAT_CAST
(
z
)
;
}
}
modules/ocl/test/test_color.cpp
View file @
5e02b204
...
...
@@ -47,6 +47,7 @@
#ifdef HAVE_OPENCL
using
namespace
testing
;
using
namespace
cv
;
///////////////////////////////////////////////////////////////////////////////////////////////////////
// cvtColor
...
...
@@ -57,20 +58,20 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
bool
use_roi
;
// src mat
cv
::
Mat
src1
;
cv
::
Mat
dst1
;
Mat
src
;
Mat
dst
;
// src mat with roi
cv
::
Mat
src1
_roi
;
cv
::
Mat
dst1
_roi
;
Mat
src
_roi
;
Mat
dst
_roi
;
// ocl dst mat for testing
cv
::
ocl
::
oclMat
gsrc1
_whole
;
cv
::
ocl
::
oclMat
gdst1
_whole
;
ocl
::
oclMat
gsrc
_whole
;
ocl
::
oclMat
gdst
_whole
;
// ocl mat with roi
cv
::
ocl
::
oclMat
gsrc1
_roi
;
cv
::
ocl
::
oclMat
gdst1
_roi
;
ocl
::
oclMat
gsrc
_roi
;
ocl
::
oclMat
gdst
_roi
;
virtual
void
SetUp
()
{
...
...
@@ -85,19 +86,23 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
Size
roiSize
=
randomSize
(
1
,
MAX_VALUE
);
Border
srcBorder
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
src
1
,
src1
_roi
,
roiSize
,
srcBorder
,
srcType
,
2
,
100
);
randomSubMat
(
src
,
src
_roi
,
roiSize
,
srcBorder
,
srcType
,
2
,
100
);
Border
dst
1
Border
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
dst
1
,
dst1_roi
,
roiSize
,
dst1
Border
,
dstType
,
5
,
16
);
Border
dstBorder
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
dst
,
dst_roi
,
roiSize
,
dst
Border
,
dstType
,
5
,
16
);
generateOclMat
(
gsrc
1_whole
,
gsrc1_roi
,
src1
,
roiSize
,
srcBorder
);
generateOclMat
(
gdst
1_whole
,
gdst1_roi
,
dst1
,
roiSize
,
dst1
Border
);
generateOclMat
(
gsrc
_whole
,
gsrc_roi
,
src
,
roiSize
,
srcBorder
);
generateOclMat
(
gdst
_whole
,
gdst_roi
,
dst
,
roiSize
,
dst
Border
);
}
void
Near
(
double
threshold
=
1e-3
)
{
EXPECT_MAT_NEAR
(
dst1
,
gdst1_whole
,
threshold
);
EXPECT_MAT_NEAR
(
dst1_roi
,
gdst1_roi
,
threshold
);
Mat
whole
,
roi
;
gdst_whole
.
download
(
whole
);
gdst_roi
.
download
(
roi
);
EXPECT_MAT_NEAR
(
dst
,
whole
,
threshold
);
EXPECT_MAT_NEAR
(
dst_roi
,
roi
,
threshold
);
}
void
doTest
(
int
channelsIn
,
int
channelsOut
,
int
code
)
...
...
@@ -106,15 +111,15 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
{
random_roi
(
channelsIn
,
channelsOut
);
cv
::
cvtColor
(
src1_roi
,
dst1
_roi
,
code
,
channelsOut
);
cv
::
ocl
::
cvtColor
(
gsrc1_roi
,
gdst1
_roi
,
code
,
channelsOut
);
cv
tColor
(
src_roi
,
dst
_roi
,
code
,
channelsOut
);
ocl
::
cvtColor
(
gsrc_roi
,
gdst
_roi
,
code
,
channelsOut
);
Near
();
}
}
};
#define CVTCODE(name)
cv::
COLOR_ ## name
#define CVTCODE(name) COLOR_ ## name
// RGB <-> Gray
...
...
@@ -224,6 +229,25 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA)
doTest
(
3
,
4
,
CVTCODE
(
YCrCb2BGR
));
}
// RGB <-> XYZ
OCL_TEST_P
(
CvtColor
,
RGB2XYZ
)
{
doTest
(
3
,
3
,
CVTCODE
(
RGB2XYZ
));
}
OCL_TEST_P
(
CvtColor
,
BGR2XYZ
)
{
doTest
(
3
,
3
,
CVTCODE
(
BGR2XYZ
));
}
OCL_TEST_P
(
CvtColor
,
RGBA2XYZ
)
{
doTest
(
4
,
3
,
CVTCODE
(
RGB2XYZ
));
}
OCL_TEST_P
(
CvtColor
,
BGRA2XYZ
)
{
doTest
(
4
,
3
,
CVTCODE
(
BGR2XYZ
));
}
// YUV -> RGBA_NV12
struct
CvtColor_YUV420
:
...
...
@@ -238,13 +262,13 @@ struct CvtColor_YUV420 :
roiSize
.
width
*=
2
;
roiSize
.
height
*=
3
;
Border
srcBorder
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
src
1
,
src1
_roi
,
roiSize
,
srcBorder
,
srcType
,
2
,
100
);
randomSubMat
(
src
,
src
_roi
,
roiSize
,
srcBorder
,
srcType
,
2
,
100
);
Border
dst
1
Border
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
dst
1
,
dst1_roi
,
roiSize
,
dst1
Border
,
dstType
,
5
,
16
);
Border
dstBorder
=
randomBorder
(
0
,
use_roi
?
MAX_VALUE
:
0
);
randomSubMat
(
dst
,
dst_roi
,
roiSize
,
dst
Border
,
dstType
,
5
,
16
);
generateOclMat
(
gsrc
1_whole
,
gsrc1_roi
,
src1
,
roiSize
,
srcBorder
);
generateOclMat
(
gdst
1_whole
,
gdst1_roi
,
dst1
,
roiSize
,
dst1
Border
);
generateOclMat
(
gsrc
_whole
,
gsrc_roi
,
src
,
roiSize
,
srcBorder
);
generateOclMat
(
gdst
_whole
,
gdst_roi
,
dst
,
roiSize
,
dst
Border
);
}
};
...
...
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