Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Xiaomi
Mace
提交
5501056d
Mace
项目概览
Xiaomi
/
Mace
通知
106
Star
40
Fork
27
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
DevOps
流水线
流水线任务
计划
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
Mace
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
DevOps
DevOps
流水线
流水线任务
计划
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
流水线任务
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
提交
5501056d
编写于
3月 01, 2018
作者:
L
liuqi
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Fix arbitrary pad bug.
上级
c5bc6a5a
变更
7
显示空白变更内容
内联
并排
Showing
7 changed file
with
63 addition
and
56 deletion
+63
-56
mace/kernels/conv_2d.h
mace/kernels/conv_2d.h
+10
-9
mace/kernels/depthwise_conv2d.h
mace/kernels/depthwise_conv2d.h
+10
-9
mace/kernels/opencl/conv_2d_opencl.cc
mace/kernels/opencl/conv_2d_opencl.cc
+10
-8
mace/kernels/opencl/depthwise_conv_opencl.cc
mace/kernels/opencl/depthwise_conv_opencl.cc
+7
-6
mace/kernels/opencl/pooling_opencl.cc
mace/kernels/opencl/pooling_opencl.cc
+9
-8
mace/kernels/opencl/winograd_transform.cc
mace/kernels/opencl/winograd_transform.cc
+8
-7
mace/kernels/pooling.h
mace/kernels/pooling.h
+9
-9
未找到文件。
mace/kernels/conv_2d.h
浏览文件 @
5501056d
...
...
@@ -228,11 +228,12 @@ struct Conv2dFunctor : Conv2dFunctorBase {
MACE_CHECK_NOTNULL
(
output
);
std
::
vector
<
index_t
>
output_shape
(
4
);
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
padding_type_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
output
->
Resize
(
output_shape
);
...
...
@@ -260,13 +261,13 @@ struct Conv2dFunctor : Conv2dFunctorBase {
MACE_CHECK
(
batch
==
input_batch
,
"Input/Output batch size mismatch"
);
int
padded_height
=
input_height
+
paddings
_
[
0
];
int
padded_width
=
input_width
+
paddings
_
[
1
];
int
padded_height
=
input_height
+
paddings
[
0
];
int
padded_width
=
input_width
+
paddings
[
1
];
Tensor
padded_input
;
// Keep this alive during kernel execution
if
(
paddings
_
[
0
]
>
0
||
paddings_
[
1
]
>
0
)
{
ConstructNHWCInputWithPadding
(
input
,
paddings
_
.
data
(),
&
padded_input
);
if
(
paddings
[
0
]
>
0
||
paddings
[
1
]
>
0
)
{
ConstructNHWCInputWithPadding
(
input
,
paddings
.
data
(),
&
padded_input
);
input
=
&
padded_input
;
}
...
...
mace/kernels/depthwise_conv2d.h
浏览文件 @
5501056d
...
...
@@ -294,11 +294,12 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
fake_filter_shape
[
3
]
=
1
;
std
::
vector
<
index_t
>
output_shape
(
4
);
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
padding_type_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
auto
input_shape
=
fake_filter_shape
;
output
->
Resize
(
output_shape
);
...
...
@@ -329,10 +330,10 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
MACE_CHECK
(
batch
==
input_batch
,
"Input/Output batch size mismatch"
);
// The left-upper most offset of the padded input
int
paddings_top
=
paddings
_
[
0
]
/
2
;
int
paddings_bottom
=
paddings
_
[
0
]
-
paddings_top
;
int
paddings_left
=
paddings
_
[
1
]
/
2
;
int
paddings_right
=
paddings
_
[
1
]
-
paddings_left
;
int
paddings_top
=
paddings
[
0
]
/
2
;
int
paddings_bottom
=
paddings
[
0
]
-
paddings_top
;
int
paddings_left
=
paddings
[
1
]
/
2
;
int
paddings_right
=
paddings
[
1
]
-
paddings_left
;
int
padded_h_start
=
0
-
paddings_top
;
int
padded_w_start
=
0
-
paddings_left
;
...
...
mace/kernels/opencl/conv_2d_opencl.cc
浏览文件 @
5501056d
...
...
@@ -68,7 +68,8 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
index_t
kernel_h
=
filter
->
dim
(
0
);
index_t
kernel_w
=
filter
->
dim
(
1
);
if
(
!
input
->
is_image
()
||
strides_
[
0
]
!=
strides_
[
1
]
||
strides_
[
0
]
>
2
||
if
(
!
input
->
is_image
()
||
strides_
[
0
]
!=
strides_
[
1
]
||
((
kernel_h
==
1
||
kernel_h
==
3
)
&&
strides_
[
0
]
>
2
)
||
(
dilations_
[
0
]
>
1
&&
(
strides_
[
0
]
>
1
||
kernel_h
==
1
)))
{
LOG
(
WARNING
)
<<
"OpenCL conv2d kernel with "
<<
"filter"
<<
kernel_h
<<
"x"
<<
kernel_w
<<
","
...
...
@@ -80,11 +81,12 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
}
std
::
vector
<
index_t
>
output_shape
(
4
);
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
padding_type_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
std
::
vector
<
size_t
>
output_image_shape
;
...
...
@@ -95,11 +97,11 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
selector
[
kernel_h
-
1
]
!=
nullptr
&&
0
<
strides_
[
0
]
&&
strides_
[
0
]
<
3
)
{
auto
conv2d_func
=
selector
[
kernel_h
-
1
];
conv2d_func
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
_
.
data
(),
dilations_
,
activation_
,
conv2d_func
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
.
data
(),
dilations_
,
activation_
,
relux_max_limit_
,
prelu_alpha_
,
DataTypeToEnum
<
T
>::
value
,
output
,
future
);
}
else
{
Conv2dOpencl
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
_
.
data
(),
dilations_
,
Conv2dOpencl
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
.
data
(),
dilations_
,
activation_
,
relux_max_limit_
,
prelu_alpha_
,
DataTypeToEnum
<
T
>::
value
,
output
,
future
);
}
...
...
mace/kernels/opencl/depthwise_conv_opencl.cc
浏览文件 @
5501056d
...
...
@@ -153,18 +153,19 @@ void DepthwiseConv2dFunctor<DeviceType::OPENCL, T>::operator()(
fake_filter_shape
[
3
]
=
1
;
std
::
vector
<
index_t
>
output_shape
(
4
);
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
padding_type_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
std
::
vector
<
size_t
>
output_image_shape
;
CalImage2DShape
(
output_shape
,
BufferType
::
IN_OUT_CHANNEL
,
output_image_shape
);
output
->
ResizeImage
(
output_shape
,
output_image_shape
);
DepthwiseConv2d
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
_
.
data
(),
dilations_
,
DepthwiseConv2d
(
&
kernel_
,
input
,
filter
,
bias
,
strides_
[
0
],
paddings
.
data
(),
dilations_
,
activation_
,
relux_max_limit_
,
prelu_alpha_
,
DataTypeToEnum
<
T
>::
value
,
output
,
future
);
}
...
...
mace/kernels/opencl/pooling_opencl.cc
浏览文件 @
5501056d
...
...
@@ -23,12 +23,13 @@ void PoolingFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
input
->
dim
(
3
),
input
->
dim
(
3
)
};
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
std
::
vector
<
size_t
>
output_image_shape
;
...
...
@@ -66,8 +67,8 @@ void PoolingFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
kernel_
.
setArg
(
idx
++
,
static_cast
<
int32_t
>
(
input
->
dim
(
1
)));
kernel_
.
setArg
(
idx
++
,
static_cast
<
int32_t
>
(
input
->
dim
(
2
)));
kernel_
.
setArg
(
idx
++
,
static_cast
<
int32_t
>
(
out_height
));
kernel_
.
setArg
(
idx
++
,
paddings
_
[
0
]
/
2
);
kernel_
.
setArg
(
idx
++
,
paddings
_
[
1
]
/
2
);
kernel_
.
setArg
(
idx
++
,
paddings
[
0
]
/
2
);
kernel_
.
setArg
(
idx
++
,
paddings
[
1
]
/
2
);
kernel_
.
setArg
(
idx
++
,
strides_
[
0
]);
kernel_
.
setArg
(
idx
++
,
kernels_
[
0
]);
kernel_
.
setArg
(
idx
++
,
*
(
static_cast
<
cl
::
Image2D
*>
(
output
->
buffer
())));
...
...
mace/kernels/opencl/winograd_transform.cc
浏览文件 @
5501056d
...
...
@@ -17,11 +17,12 @@ void WinogradTransformFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *i
StatsFuture
*
future
)
{
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
index_t
>
filter_shape
=
{
3
,
3
,
input_tensor
->
dim
(
3
),
1
};
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
.
data
(),
strides_
.
data
(),
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
strides_
.
data
(),
padding_type_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
const
index_t
round_h
=
(
output_shape
[
1
]
+
1
)
/
2
;
...
...
@@ -52,8 +53,8 @@ void WinogradTransformFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *i
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
input_tensor
->
dim
(
3
)));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
round_h
*
round_w
));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
round_w
));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
paddings
_
[
0
]
/
2
));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
paddings
_
[
1
]
/
2
));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
paddings
[
0
]
/
2
));
kernel_
.
setArg
(
idx
++
,
static_cast
<
uint32_t
>
(
paddings
[
1
]
/
2
));
}
const
uint32_t
gws
[
2
]
=
{
static_cast
<
uint32_t
>
(
out_width
),
...
...
mace/kernels/pooling.h
浏览文件 @
5501056d
...
...
@@ -64,13 +64,13 @@ struct PoolingFunctor : PoolingFunctorBase {
input_tensor
->
dim
(
3
),
input_tensor
->
dim
(
3
)
};
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
output_shape
.
data
(),
paddings
.
data
());
if
(
!
paddings_
.
empty
())
{
paddings
=
paddings_
;
}
output_tensor
->
Resize
(
output_shape
);
...
...
@@ -99,8 +99,8 @@ struct PoolingFunctor : PoolingFunctorBase {
int
dilation_w
=
dilations_
[
1
];
// The left-upper most offset of the padded input
int
padded_h_start
=
0
-
paddings
_
[
0
]
/
2
;
int
padded_w_start
=
0
-
paddings
_
[
1
]
/
2
;
int
padded_h_start
=
0
-
paddings
[
0
]
/
2
;
int
padded_w_start
=
0
-
paddings
[
1
]
/
2
;
if
(
pooling_type_
==
MAX
)
{
#pragma omp parallel for collapse(4)
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录