Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Greenplum
Opencv
提交
d9de8409
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,发现更多精彩内容 >>
提交
d9de8409
编写于
4月 13, 2013
作者:
P
Peng Xiao
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Allow more input query/train types for ocl::bfmatcher
RadiusMatch for HammingDist cannot pass yet.
上级
1db20099
变更
3
隐藏空白更改
内联
并排
Showing
3 changed file
with
30 addition
and
100 deletion
+30
-100
modules/ocl/src/brute_force_matcher.cpp
modules/ocl/src/brute_force_matcher.cpp
+2
-72
modules/ocl/src/opencl/brute_force_match.cl
modules/ocl/src/opencl/brute_force_match.cl
+27
-23
modules/ocl/test/test_brute_force_matcher.cpp
modules/ocl/test/test_brute_force_matcher.cpp
+1
-5
未找到文件。
modules/ocl/src/brute_force_matcher.cpp
浏览文件 @
d9de8409
...
...
@@ -77,7 +77,6 @@ template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ >
void
matchUnrolledCached
(
const
oclMat
&
query
,
const
oclMat
&
train
,
const
oclMat
&
/*mask*/
,
const
oclMat
&
trainIdx
,
const
oclMat
&
distance
,
int
distType
)
{
assert
(
query
.
type
()
==
CV_32F
);
cv
::
ocl
::
Context
*
ctx
=
query
.
clCxt
;
size_t
globalSize
[]
=
{(
query
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
size_t
localSize
[]
=
{
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
...
...
@@ -121,7 +120,6 @@ template < int BLOCK_SIZE/*, typename Mask*/ >
void
match
(
const
oclMat
&
query
,
const
oclMat
&
train
,
const
oclMat
&
/*mask*/
,
const
oclMat
&
trainIdx
,
const
oclMat
&
distance
,
int
distType
)
{
assert
(
query
.
type
()
==
CV_32F
);
cv
::
ocl
::
Context
*
ctx
=
query
.
clCxt
;
size_t
globalSize
[]
=
{(
query
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
size_t
localSize
[]
=
{
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
...
...
@@ -164,7 +162,6 @@ template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ >
void
matchUnrolledCached
(
const
oclMat
&
query
,
const
oclMat
&
train
,
float
maxDistance
,
const
oclMat
&
/*mask*/
,
const
oclMat
&
trainIdx
,
const
oclMat
&
distance
,
const
oclMat
&
nMatches
,
int
distType
)
{
assert
(
query
.
type
()
==
CV_32F
);
cv
::
ocl
::
Context
*
ctx
=
query
.
clCxt
;
size_t
globalSize
[]
=
{(
train
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
(
query
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
1
};
size_t
localSize
[]
=
{
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
...
...
@@ -207,7 +204,6 @@ template < int BLOCK_SIZE/*, typename Mask*/ >
void
radius_match
(
const
oclMat
&
query
,
const
oclMat
&
train
,
float
maxDistance
,
const
oclMat
&
/*mask*/
,
const
oclMat
&
trainIdx
,
const
oclMat
&
distance
,
const
oclMat
&
nMatches
,
int
distType
)
{
assert
(
query
.
type
()
==
CV_32F
);
cv
::
ocl
::
Context
*
ctx
=
query
.
clCxt
;
size_t
globalSize
[]
=
{(
train
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
(
query
.
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE
,
1
};
size_t
localSize
[]
=
{
BLOCK_SIZE
,
BLOCK_SIZE
,
1
};
...
...
@@ -566,17 +562,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchSingle(const oclMat &query, const
if
(
query
.
empty
()
||
train
.
empty
())
return
;
// match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
int
callType
=
query
.
depth
();
if
(
callType
!=
5
)
CV_Error
(
CV_UNSUPPORTED_FORMAT_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
if
((
distType
==
0
&&
callType
==
1
)
||
(
distType
==
1
&&
callType
!=
5
)
||
(
distType
==
2
&&
(
callType
!=
0
||
callType
!=
2
||
callType
!=
4
)))
{
CV_Error
(
CV_UNSUPPORTED_DEPTH_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
}
CV_Assert
(
query
.
channels
()
==
1
&&
query
.
depth
()
<
CV_64F
);
CV_Assert
(
train
.
cols
==
query
.
cols
&&
train
.
type
()
==
query
.
type
());
...
...
@@ -687,17 +672,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchCollection(const oclMat &query, c
if
(
query
.
empty
()
||
trainCollection
.
empty
())
return
;
// match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
int
callType
=
query
.
depth
();
if
(
callType
!=
5
)
CV_Error
(
CV_UNSUPPORTED_FORMAT_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
if
((
distType
==
0
&&
callType
==
1
)
||
(
distType
==
1
&&
callType
!=
5
)
||
(
distType
==
2
&&
(
callType
!=
0
||
callType
!=
2
||
callType
!=
4
)))
{
CV_Error
(
CV_UNSUPPORTED_DEPTH_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
}
CV_Assert
(
query
.
channels
()
==
1
&&
query
.
depth
()
<
CV_64F
);
const
int
nQuery
=
query
.
rows
;
...
...
@@ -706,7 +680,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchCollection(const oclMat &query, c
ensureSizeIsEnough
(
1
,
nQuery
,
CV_32S
,
imgIdx
);
ensureSizeIsEnough
(
1
,
nQuery
,
CV_32F
,
distance
);
matchDispatcher
(
query
,
(
const
oclMat
*
)
trainCollection
.
ptr
(),
trainCollection
.
cols
,
masks
,
trainIdx
,
imgIdx
,
distance
,
distType
);
return
;
...
...
@@ -778,18 +751,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatchSingle(const oclMat &query, co
if
(
query
.
empty
()
||
train
.
empty
())
return
;
// match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
int
callType
=
query
.
depth
();
if
(
callType
!=
5
)
CV_Error
(
CV_UNSUPPORTED_FORMAT_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
if
((
distType
==
0
&&
callType
==
1
)
||
(
distType
==
1
&&
callType
!=
5
)
||
(
distType
==
2
&&
(
callType
!=
0
||
callType
!=
2
||
callType
!=
4
)))
{
CV_Error
(
CV_UNSUPPORTED_DEPTH_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
}
CV_Assert
(
query
.
channels
()
==
1
&&
query
.
depth
()
<
CV_64F
);
CV_Assert
(
train
.
type
()
==
query
.
type
()
&&
train
.
cols
==
query
.
cols
);
...
...
@@ -886,26 +847,7 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatch2Collection(const oclMat &quer
typedef
void
(
*
caller_t
)(
const
oclMat
&
query
,
const
oclMat
&
trains
,
const
oclMat
&
masks
,
const
oclMat
&
trainIdx
,
const
oclMat
&
imgIdx
,
const
oclMat
&
distance
);
#if 0
static const caller_t callers[3][6] =
{
{
ocl_match2L1_gpu<unsigned char>, 0/*match2L1_gpu<signed char>*/,
ocl_match2L1_gpu<unsigned short>, ocl_match2L1_gpu<short>,
ocl_match2L1_gpu<int>, ocl_match2L1_gpu<float>
},
{
0/*match2L2_gpu<unsigned char>*/, 0/*match2L2_gpu<signed char>*/,
0/*match2L2_gpu<unsigned short>*/, 0/*match2L2_gpu<short>*/,
0/*match2L2_gpu<int>*/, ocl_match2L2_gpu<float>
},
{
ocl_match2Hamming_gpu<unsigned char>, 0/*match2Hamming_gpu<signed char>*/,
ocl_match2Hamming_gpu<unsigned short>, 0/*match2Hamming_gpu<short>*/,
ocl_match2Hamming_gpu<int>, 0/*match2Hamming_gpu<float>*/
}
};
#endif
CV_Assert
(
query
.
channels
()
==
1
&&
query
.
depth
()
<
CV_64F
);
const
int
nQuery
=
query
.
rows
;
...
...
@@ -1051,23 +993,11 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatch(const oclMat &query, vector<
// radiusMatchSingle
void
cv
::
ocl
::
BruteForceMatcher_OCL_base
::
radiusMatchSingle
(
const
oclMat
&
query
,
const
oclMat
&
train
,
oclMat
&
trainIdx
,
oclMat
&
distance
,
oclMat
&
nMatches
,
float
maxDistance
,
const
oclMat
&
mask
)
oclMat
&
trainIdx
,
oclMat
&
distance
,
oclMat
&
nMatches
,
float
maxDistance
,
const
oclMat
&
mask
)
{
if
(
query
.
empty
()
||
train
.
empty
())
return
;
// match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
int
callType
=
query
.
depth
();
if
(
callType
!=
5
)
CV_Error
(
CV_UNSUPPORTED_FORMAT_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
if
((
distType
==
0
&&
callType
==
1
)
||
(
distType
==
1
&&
callType
!=
5
)
||
(
distType
==
2
&&
(
callType
!=
0
||
callType
!=
2
||
callType
!=
4
)))
{
CV_Error
(
CV_UNSUPPORTED_DEPTH_ERR
,
"BruteForceMatch OpenCL only support float type query!
\n
"
);
}
const
int
nQuery
=
query
.
rows
;
const
int
nTrain
=
train
.
rows
;
...
...
modules/ocl/src/opencl/brute_force_match.cl
浏览文件 @
d9de8409
...
...
@@ -85,14 +85,17 @@ int bit1Count(int x)
typedef
int
value_type
;
typedef
int
result_type
;
#
endif
#
define
DIST_RES
(
x
)
(
x
)
#
elif
(
DIST_TYPE
==
1
)
//
L2Dist
#
define
DIST
(
x,
y
)
(((
x
)
-
(
y
))
*
((
x
)
-
(
y
)))
typedef
float
value_type
;
typedef
float
result_type
;
#
define
DIST_RES
(
x
)
sqrt
(
x
)
#
elif
(
DIST_TYPE
==
2
)
//
Hamming
#
define
DIST
(
x,
y
)
bit1Count
(
((
x
)
^
(
y
)
)
#
define
DIST
(
x,
y
)
bit1Count
(
(
x
)
^
(
y
)
)
typedef
int
value_type
;
typedef
int
result_type
;
#
define
DIST_RES
(
x
)
(
x
)
#
endif
result_type
reduce_block
(
...
...
@@ -107,10 +110,10 @@ result_type reduce_block(
for
(
int
j
=
0
; j < BLOCK_SIZE ; j++)
{
result
+=
DIST
(
s_query[lidy
*
BLOCK_SIZE
+
j],
s_query[lidy
*
BLOCK_SIZE
+
j],
s_train[j
*
BLOCK_SIZE
+
lidx]
)
;
}
return
result
;
return
DIST_RES
(
result
)
;
}
result_type
reduce_multi_block
(
...
...
@@ -126,10 +129,10 @@ result_type reduce_multi_block(
for
(
int
j
=
0
; j < BLOCK_SIZE ; j++)
{
result
+=
DIST
(
s_query[lidy
*
MAX_DESC_LEN
+
block_index
*
BLOCK_SIZE
+
j],
s_query[lidy
*
MAX_DESC_LEN
+
block_index
*
BLOCK_SIZE
+
j],
s_train[j
*
BLOCK_SIZE
+
lidx]
)
;
}
return
result
;
return
DIST_RES
(
result
)
;
}
/*
2dim
launch,
global
size:
dim0
is
(
query
rows
+
BLOCK_SIZE
-
1
)
/
BLOCK_SIZE
*
BLOCK_SIZE,
dim1
is
BLOCK_SIZE
...
...
@@ -153,8 +156,8 @@ __kernel void BruteForceMatch_UnrollMatch(
const
int
lidy
=
get_local_id
(
1
)
;
const
int
groupidx
=
get_group_id
(
0
)
;
__local
value_type
*s_query
=
sharebuffer
;
__local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
MAX_DESC_LEN
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
MAX_DESC_LEN
;
int
queryIdx
=
groupidx
*
BLOCK_SIZE
+
lidy
;
//
load
the
query
into
local
memory.
...
...
@@ -251,8 +254,8 @@ __kernel void BruteForceMatch_Match(
float
myBestDistance
=
MAX_FLOAT
;
int
myBestTrainIdx
=
-1
;
__local
value_type
*s_query
=
sharebuffer
;
__local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
//
loop
for
(
int
t
=
0
; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
...
...
@@ -345,8 +348,8 @@ __kernel void BruteForceMatch_RadiusUnrollMatch(
const
int
queryIdx
=
groupidy
*
BLOCK_SIZE
+
lidy
;
const
int
trainIdx
=
groupidx
*
BLOCK_SIZE
+
lidx
;
__local
value_type
*s_query
=
sharebuffer
;
__local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
result_type
result
=
0
;
for
(
int
i
=
0
; i < MAX_DESC_LEN / BLOCK_SIZE ; ++i)
...
...
@@ -365,7 +368,8 @@ __kernel void BruteForceMatch_RadiusUnrollMatch(
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
}
if
(
queryIdx
<
query_rows
&&
trainIdx
<
train_rows
&&
result
<
maxDistance/*
&&
mask
(
queryIdx,
trainIdx
)
*/
)
if
(
queryIdx
<
query_rows
&&
trainIdx
<
train_rows
&&
convert_float
(
result
)
<
maxDistance/*
&&
mask
(
queryIdx,
trainIdx
)
*/
)
{
unsigned
int
ind
=
atom_inc
(
nMatches
+
queryIdx/*,
(
unsigned
int
)
-1*/
)
;
...
...
@@ -405,8 +409,8 @@ __kernel void BruteForceMatch_RadiusMatch(
const
int
queryIdx
=
groupidy
*
BLOCK_SIZE
+
lidy
;
const
int
trainIdx
=
groupidx
*
BLOCK_SIZE
+
lidx
;
__local
value_type
*s_query
=
sharebuffer
;
__local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
result_type
result
=
0
;
for
(
int
i
=
0
; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; ++i)
...
...
@@ -425,7 +429,8 @@ __kernel void BruteForceMatch_RadiusMatch(
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
}
if
(
queryIdx
<
query_rows
&&
trainIdx
<
train_rows
&&
result
<
maxDistance/*
&&
mask
(
queryIdx,
trainIdx
)
*/
)
if
(
queryIdx
<
query_rows
&&
trainIdx
<
train_rows
&&
convert_float
(
result
)
<
maxDistance/*
&&
mask
(
queryIdx,
trainIdx
)
*/
)
{
unsigned
int
ind
=
atom_inc
(
nMatches
+
queryIdx
)
;
...
...
@@ -458,8 +463,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
const
int
groupidx
=
get_group_id
(
0
)
;
const
int
queryIdx
=
groupidx
*
BLOCK_SIZE
+
lidy
;
local
value_type
*s_query
=
sharebuffer
;
local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
MAX_DESC_LEN
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
MAX_DESC_LEN
;
//
load
the
query
into
local
memory.
for
(
int
i
=
0
; i < MAX_DESC_LEN / BLOCK_SIZE; i ++)
...
...
@@ -480,7 +485,6 @@ __kernel void BruteForceMatch_knnUnrollMatch(
result_type
result
=
0
;
for
(
int
i
=
0
; i < MAX_DESC_LEN / BLOCK_SIZE ; i++)
{
const
int
loadX
=
lidx
+
i
*
BLOCK_SIZE
;
//load
a
BLOCK_SIZE
*
BLOCK_SIZE
block
into
local
train.
const
int
loadx
=
lidx
+
i
*
BLOCK_SIZE
;
s_train[lidx
*
BLOCK_SIZE
+
lidy]
=
loadx
<
train_cols
?
train[min
(
t
*
BLOCK_SIZE
+
lidy,
train_rows
-
1
)
*
(
step
/
sizeof
(
float
))
+
loadx]
:
0
;
...
...
@@ -514,8 +518,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
local
float
*s_distance
=
(
local
float
*
)
sharebuffer
;
local
int
*s_trainIdx
=
(
local
int
*
)(
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
)
;
__
local
float
*s_distance
=
(
local
float
*
)
sharebuffer
;
__
local
int
*s_trainIdx
=
(
local
int
*
)(
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
)
;
//
find
BestMatch
s_distance
+=
lidy
*
BLOCK_SIZE
;
...
...
@@ -604,8 +608,8 @@ __kernel void BruteForceMatch_knnMatch(
const
int
groupidx
=
get_group_id
(
0
)
;
const
int
queryIdx
=
groupidx
*
BLOCK_SIZE
+
lidy
;
local
value_type
*s_query
=
sharebuffer
;
local
value_type
*s_train
=
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
__local
value_type
*s_query
=
(
__local
value_type
*
)
sharebuffer
;
__local
value_type
*s_train
=
(
__local
value_type
*
)
sharebuffer
+
BLOCK_SIZE
*
BLOCK_SIZE
;
float
myBestDistance1
=
MAX_FLOAT
;
float
myBestDistance2
=
MAX_FLOAT
;
...
...
@@ -766,4 +770,4 @@ kernel void BruteForceMatch_findBestMatch(
)
{
/*
Todo
*/
}
\ No newline at end of file
}
modules/ocl/test/test_brute_force_matcher.cpp
浏览文件 @
d9de8409
...
...
@@ -158,11 +158,7 @@ namespace
TEST_P
(
BruteForceMatcher
,
RadiusMatch_Single
)
{
float
radius
;
if
(
distType
==
cv
::
ocl
::
BruteForceMatcher_OCL_base
::
L2Dist
)
radius
=
1.
f
/
countFactor
/
countFactor
;
else
radius
=
1.
f
/
countFactor
;
float
radius
=
1.
f
/
countFactor
;
cv
::
ocl
::
BruteForceMatcher_OCL_base
matcher
(
distType
);
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录