Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Greenplum
Opencv
提交
7b025474
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,发现更多精彩内容 >>
提交
7b025474
编写于
8月 02, 2012
作者:
V
Vladislav Vinogradov
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
added debayer to gpu::cvtColor
上级
c0d3adef
变更
3
隐藏空白更改
内联
并排
Showing
3 changed file
with
414 addition
and
4 deletion
+414
-4
modules/gpu/src/color.cpp
modules/gpu/src/color.cpp
+54
-4
modules/gpu/src/cuda/debayer.cu
modules/gpu/src/cuda/debayer.cu
+208
-0
modules/gpu/test/test_color.cpp
modules/gpu/test/test_color.cpp
+152
-0
未找到文件。
modules/gpu/src/color.cpp
浏览文件 @
7b025474
...
...
@@ -54,6 +54,15 @@ void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nog
#else
/* !defined (HAVE_CUDA) */
#include <cvt_colot_internal.h>
namespace
cv
{
namespace
gpu
{
namespace
device
{
template
<
typename
T
,
int
cn
>
void
Bayer2BGR_gpu
(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
}
}}
using
namespace
::
cv
::
gpu
::
device
;
namespace
...
...
@@ -1302,6 +1311,47 @@ namespace
nppSafeCall
(
nppiAlphaPremul_16u_AC4R
(
src
.
ptr
<
Npp16u
>
(),
static_cast
<
int
>
(
src
.
step
),
dst
.
ptr
<
Npp16u
>
(),
static_cast
<
int
>
(
dst
.
step
),
oSizeROI
)
);
#endif
}
void
bayer_to_bgr
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
dcn
,
bool
blue_last
,
bool
start_with_green
,
Stream
&
stream
)
{
typedef
void
(
*
func_t
)(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
static
const
func_t
funcs
[
3
][
4
]
=
{
{
0
,
0
,
Bayer2BGR_gpu
<
uchar
,
3
>
,
Bayer2BGR_gpu
<
uchar
,
4
>
},
{
0
,
0
,
0
,
0
},
{
0
,
0
,
Bayer2BGR_gpu
<
ushort
,
3
>
,
Bayer2BGR_gpu
<
ushort
,
4
>
}
};
if
(
dcn
<=
0
)
dcn
=
3
;
CV_Assert
(
src
.
type
()
==
CV_8UC1
||
src
.
type
()
==
CV_16UC1
);
CV_Assert
(
src
.
rows
>
2
&&
src
.
cols
>
2
);
CV_Assert
(
dcn
==
3
||
dcn
==
4
);
dst
.
create
(
src
.
size
(),
CV_MAKETYPE
(
src
.
depth
(),
dcn
));
funcs
[
src
.
depth
()][
dcn
-
1
](
src
,
dst
,
blue_last
,
start_with_green
,
StreamAccessor
::
getStream
(
stream
));
}
void
bayerBG_to_bgr
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
dcn
,
Stream
&
stream
)
{
bayer_to_bgr
(
src
,
dst
,
dcn
,
false
,
false
,
stream
);
}
void
bayerGB_to_bgr
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
dcn
,
Stream
&
stream
)
{
bayer_to_bgr
(
src
,
dst
,
dcn
,
false
,
true
,
stream
);
}
void
bayerRG_to_bgr
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
dcn
,
Stream
&
stream
)
{
bayer_to_bgr
(
src
,
dst
,
dcn
,
true
,
false
,
stream
);
}
void
bayerGR_to_bgr
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
dcn
,
Stream
&
stream
)
{
bayer_to_bgr
(
src
,
dst
,
dcn
,
true
,
true
,
stream
);
}
}
void
cv
::
gpu
::
cvtColor
(
const
GpuMat
&
src
,
GpuMat
&
dst
,
int
code
,
int
dcn
,
Stream
&
stream
)
...
...
@@ -1366,10 +1416,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
bgr_to_lab
,
// CV_BGR2Lab =44
rgb_to_lab
,
// CV_RGB2Lab =45
0
,
// CV_BayerBG2BGR =46
0
,
// CV_BayerGB2BGR =47
0
,
// CV_BayerRG2BGR =48
0
,
// CV_BayerGR2BGR =49
bayerBG_to_bgr
,
// CV_BayerBG2BGR =46
bayerGB_to_bgr
,
// CV_BayerGB2BGR =47
bayerRG_to_bgr
,
// CV_BayerRG2BGR =48
bayerGR_to_bgr
,
// CV_BayerGR2BGR =49
bgr_to_luv
,
// CV_BGR2Luv =50
rgb_to_luv
,
// CV_RGB2Luv =51
...
...
modules/gpu/src/cuda/debayer.cu
0 → 100644
浏览文件 @
7b025474
/*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 <opencv2/gpu/device/common.hpp>
#include <opencv2/gpu/device/vec_traits.hpp>
namespace
cv
{
namespace
gpu
{
namespace
device
{
template
<
class
SrcPtr
,
typename
T
>
__global__
void
Bayer2BGR
(
const
SrcPtr
src
,
PtrStep_
<
T
>
dst
,
const
int
width
,
const
int
height
,
const
bool
glob_blue_last
,
const
bool
glob_start_with_green
)
{
const
int
tx
=
blockIdx
.
x
*
blockDim
.
x
+
threadIdx
.
x
;
const
int
y
=
blockIdx
.
y
*
blockDim
.
y
+
threadIdx
.
y
;
if
(
y
>=
height
)
return
;
const
bool
blue_last
=
(
y
&
1
)
?
!
glob_blue_last
:
glob_blue_last
;
const
bool
start_with_green
=
(
y
&
1
)
?
!
glob_start_with_green
:
glob_start_with_green
;
int
x
=
tx
*
2
;
if
(
start_with_green
)
{
--
x
;
if
(
tx
==
0
)
{
const
int
t0
=
(
src
(
y
,
1
)
+
src
(
y
+
2
,
1
)
+
1
)
>>
1
;
const
int
t1
=
(
src
(
y
+
1
,
0
)
+
src
(
y
+
1
,
2
)
+
1
)
>>
1
;
T
res
;
res
.
x
=
blue_last
?
t0
:
t1
;
res
.
y
=
src
(
y
+
1
,
1
);
res
.
z
=
blue_last
?
t1
:
t0
;
dst
(
y
+
1
,
0
)
=
dst
(
y
+
1
,
1
)
=
res
;
if
(
y
==
0
)
{
dst
(
0
,
0
)
=
dst
(
0
,
1
)
=
res
;
}
else
if
(
y
==
height
-
1
)
{
dst
(
height
+
1
,
0
)
=
dst
(
height
+
1
,
1
)
=
res
;
}
}
}
if
(
x
>=
0
&&
x
<=
width
-
2
)
{
const
int
t0
=
(
src
(
y
,
x
)
+
src
(
y
,
x
+
2
)
+
src
(
y
+
2
,
x
)
+
src
(
y
+
2
,
x
+
2
)
+
2
)
>>
2
;
const
int
t1
=
(
src
(
y
,
x
+
1
)
+
src
(
y
+
1
,
x
)
+
src
(
y
+
1
,
x
+
2
)
+
src
(
y
+
2
,
x
+
1
)
+
2
)
>>
2
;
const
int
t2
=
(
src
(
y
,
x
+
2
)
+
src
(
y
+
2
,
x
+
2
)
+
1
)
>>
1
;
const
int
t3
=
(
src
(
y
+
1
,
x
+
1
)
+
src
(
y
+
1
,
x
+
3
)
+
1
)
>>
1
;
T
res1
,
res2
;
if
(
blue_last
)
{
res1
.
x
=
t0
;
res1
.
y
=
t1
;
res1
.
z
=
src
(
y
+
1
,
x
+
1
);
res2
.
x
=
t2
;
res2
.
y
=
src
(
y
+
1
,
x
+
2
);
res2
.
z
=
t3
;
}
else
{
res1
.
x
=
src
(
y
+
1
,
x
+
1
);
res1
.
y
=
t1
;
res1
.
z
=
t0
;
res2
.
x
=
t3
;
res2
.
y
=
src
(
y
+
1
,
x
+
2
);
res2
.
z
=
t2
;
}
dst
(
y
+
1
,
x
+
1
)
=
res1
;
dst
(
y
+
1
,
x
+
2
)
=
res2
;
if
(
y
==
0
)
{
dst
(
0
,
x
+
1
)
=
res1
;
dst
(
0
,
x
+
2
)
=
res2
;
if
(
x
==
0
)
{
dst
(
0
,
0
)
=
res1
;
}
else
if
(
x
==
width
-
2
)
{
dst
(
0
,
width
+
1
)
=
res2
;
}
}
else
if
(
y
==
height
-
1
)
{
dst
(
height
+
1
,
x
+
1
)
=
res1
;
dst
(
height
+
1
,
x
+
2
)
=
res2
;
if
(
x
==
0
)
{
dst
(
height
+
1
,
0
)
=
res1
;
}
else
if
(
x
==
width
-
2
)
{
dst
(
height
+
1
,
width
+
1
)
=
res2
;
}
}
if
(
x
==
0
)
{
dst
(
y
+
1
,
0
)
=
res1
;
}
else
if
(
x
==
width
-
2
)
{
dst
(
y
+
1
,
width
+
1
)
=
res2
;
}
}
else
if
(
x
==
width
-
1
)
{
const
int
t0
=
(
src
(
y
,
x
)
+
src
(
y
,
x
+
2
)
+
src
(
y
+
2
,
x
)
+
src
(
y
+
2
,
x
+
2
)
+
2
)
>>
2
;
const
int
t1
=
(
src
(
y
,
x
+
1
)
+
src
(
y
+
1
,
x
)
+
src
(
y
+
1
,
x
+
2
)
+
src
(
y
+
2
,
x
+
1
)
+
2
)
>>
2
;
T
res
;
res
.
x
=
blue_last
?
t0
:
src
(
y
+
1
,
x
+
1
);
res
.
y
=
t1
;
res
.
z
=
blue_last
?
src
(
y
+
1
,
x
+
1
)
:
t0
;
dst
(
y
+
1
,
x
+
1
)
=
dst
(
y
+
1
,
x
+
2
)
=
res
;
if
(
y
==
0
)
{
dst
(
0
,
x
+
1
)
=
dst
(
0
,
x
+
2
)
=
res
;
}
else
if
(
y
==
height
-
1
)
{
dst
(
height
+
1
,
x
+
1
)
=
dst
(
height
+
1
,
x
+
2
)
=
res
;
}
}
}
template
<
typename
T
,
int
cn
>
void
Bayer2BGR_gpu
(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
)
{
typedef
typename
TypeVec
<
T
,
cn
>::
vec_type
dst_t
;
const
int
width
=
src
.
cols
-
2
;
const
int
height
=
src
.
rows
-
2
;
const
dim3
total
(
divUp
(
width
,
2
),
height
);
const
dim3
block
(
32
,
8
);
const
dim3
grid
(
divUp
(
total
.
x
,
block
.
x
),
divUp
(
total
.
y
,
block
.
y
));
Bayer2BGR
<
PtrStep_
<
T
>
,
dst_t
><<<
grid
,
block
,
0
,
stream
>>>
((
DevMem2D_
<
T
>
)
src
,
(
DevMem2D_
<
dst_t
>
)
dst
,
width
,
height
,
blue_last
,
start_with_green
);
cudaSafeCall
(
cudaGetLastError
()
);
if
(
stream
==
0
)
cudaSafeCall
(
cudaDeviceSynchronize
()
);
}
template
void
Bayer2BGR_gpu
<
uchar
,
3
>(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
template
void
Bayer2BGR_gpu
<
uchar
,
4
>(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
template
void
Bayer2BGR_gpu
<
ushort
,
3
>(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
template
void
Bayer2BGR_gpu
<
ushort
,
4
>(
DevMem2Db
src
,
DevMem2Db
dst
,
bool
blue_last
,
bool
start_with_green
,
cudaStream_t
stream
);
}
}}
modules/gpu/test/test_color.cpp
浏览文件 @
7b025474
...
...
@@ -1744,6 +1744,158 @@ TEST_P(CvtColor, RGBA2mRGBA)
}
}
TEST_P
(
CvtColor
,
BayerBG2BGR
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerBG2BGR
);
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerBG2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst
,
0
);
}
TEST_P
(
CvtColor
,
BayerBG2BGR4
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerBG2BGR
,
4
);
ASSERT_EQ
(
4
,
dst
.
channels
());
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerBG2BGR
);
cv
::
Mat
dst4
(
dst
);
cv
::
Mat
dst3
;
cv
::
cvtColor
(
dst4
,
dst3
,
cv
::
COLOR_BGRA2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst3
,
0
);
}
TEST_P
(
CvtColor
,
BayerGB2BGR
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerGB2BGR
);
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerGB2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst
,
0
);
}
TEST_P
(
CvtColor
,
BayerGB2BGR4
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerGB2BGR
,
4
);
ASSERT_EQ
(
4
,
dst
.
channels
());
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerGB2BGR
);
cv
::
Mat
dst4
(
dst
);
cv
::
Mat
dst3
;
cv
::
cvtColor
(
dst4
,
dst3
,
cv
::
COLOR_BGRA2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst3
,
0
);
}
TEST_P
(
CvtColor
,
BayerRG2BGR
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerRG2BGR
);
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerRG2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst
,
0
);
}
TEST_P
(
CvtColor
,
BayerRG2BGR4
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerRG2BGR
,
4
);
ASSERT_EQ
(
4
,
dst
.
channels
());
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerRG2BGR
);
cv
::
Mat
dst4
(
dst
);
cv
::
Mat
dst3
;
cv
::
cvtColor
(
dst4
,
dst3
,
cv
::
COLOR_BGRA2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst3
,
0
);
}
TEST_P
(
CvtColor
,
BayerGR2BGR
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerGR2BGR
);
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerGR2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst
,
0
);
}
TEST_P
(
CvtColor
,
BayerGR2BGR4
)
{
if
(
depth
!=
CV_8U
&&
depth
!=
CV_16U
)
return
;
cv
::
Mat
src
=
randomMat
(
size
,
CV_8UC1
);
cv
::
gpu
::
GpuMat
dst
;
cv
::
gpu
::
cvtColor
(
loadMat
(
src
,
useRoi
),
dst
,
cv
::
COLOR_BayerGR2BGR
,
4
);
ASSERT_EQ
(
4
,
dst
.
channels
());
cv
::
Mat
dst_gold
;
cv
::
cvtColor
(
src
,
dst_gold
,
cv
::
COLOR_BayerGR2BGR
);
cv
::
Mat
dst4
(
dst
);
cv
::
Mat
dst3
;
cv
::
cvtColor
(
dst4
,
dst3
,
cv
::
COLOR_BGRA2BGR
);
EXPECT_MAT_NEAR
(
dst_gold
,
dst3
,
0
);
}
INSTANTIATE_TEST_CASE_P
(
GPU_ImgProc
,
CvtColor
,
testing
::
Combine
(
ALL_DEVICES
,
DIFFERENT_SIZES
,
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录