Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Xiaomi
Mace
提交
c5bc6a5a
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,发现更多精彩内容 >>
提交
c5bc6a5a
编写于
3月 01, 2018
作者:
L
liuqi
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Convolution support arbitrary padding value.
上级
123f4938
变更
15
隐藏空白更改
内联
并排
Showing
15 changed file
with
133 addition
and
83 deletion
+133
-83
mace/kernels/conv_2d.h
mace/kernels/conv_2d.h
+21
-12
mace/kernels/depthwise_conv2d.h
mace/kernels/depthwise_conv2d.h
+24
-15
mace/kernels/opencl/conv_2d_opencl.cc
mace/kernels/opencl/conv_2d_opencl.cc
+8
-6
mace/kernels/opencl/depthwise_conv_opencl.cc
mace/kernels/opencl/depthwise_conv_opencl.cc
+8
-6
mace/kernels/opencl/pooling_opencl.cc
mace/kernels/opencl/pooling_opencl.cc
+9
-7
mace/kernels/opencl/winograd_transform.cc
mace/kernels/opencl/winograd_transform.cc
+8
-6
mace/kernels/pooling.h
mace/kernels/pooling.h
+24
-16
mace/kernels/winograd_transform.h
mace/kernels/winograd_transform.h
+12
-7
mace/ops/addn.h
mace/ops/addn.h
+6
-1
mace/ops/conv_2d.h
mace/ops/conv_2d.h
+2
-1
mace/ops/conv_pool_2d_base.h
mace/ops/conv_pool_2d_base.h
+4
-2
mace/ops/depthwise_conv2d.h
mace/ops/depthwise_conv2d.h
+2
-1
mace/ops/fused_conv_2d.h
mace/ops/fused_conv_2d.h
+2
-1
mace/ops/pooling.h
mace/ops/pooling.h
+1
-1
mace/ops/winograd_transform.h
mace/ops/winograd_transform.h
+2
-1
未找到文件。
mace/kernels/conv_2d.h
浏览文件 @
c5bc6a5a
...
...
@@ -178,12 +178,14 @@ void Conv2dKernelFunc(const T *input_ptr, // batch start
struct
Conv2dFunctorBase
{
Conv2dFunctorBase
(
const
int
*
strides
,
const
Padding
&
paddings
,
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
strides_
(
strides
),
padding_type_
(
padding_type
),
paddings_
(
paddings
),
dilations_
(
dilations
),
activation_
(
activation
),
...
...
@@ -191,7 +193,8 @@ struct Conv2dFunctorBase {
prelu_alpha_
(
prelu_alpha
)
{}
const
int
*
strides_
;
// [stride_h, stride_w]
const
Padding
paddings_
;
const
Padding
padding_type_
;
std
::
vector
<
int
>
paddings_
;
const
int
*
dilations_
;
// [dilation_h, dilation_w]
const
ActivationType
activation_
;
const
float
relux_max_limit_
;
...
...
@@ -201,12 +204,14 @@ struct Conv2dFunctorBase {
template
<
DeviceType
D
,
typename
T
>
struct
Conv2dFunctor
:
Conv2dFunctorBase
{
Conv2dFunctor
(
const
int
*
strides
,
const
Padding
&
paddings
,
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
Conv2dFunctorBase
(
strides
,
padding_type
,
paddings
,
dilations
,
activation
,
...
...
@@ -223,10 +228,12 @@ struct Conv2dFunctor : Conv2dFunctorBase {
MACE_CHECK_NOTNULL
(
output
);
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
paddings_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
output
->
Resize
(
output_shape
);
int
batch
=
output
->
dim
(
0
);
...
...
@@ -253,13 +260,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
;
}
...
...
@@ -625,12 +632,14 @@ void Conv2dFunctor<DeviceType::NEON, float>::operator()(const Tensor *input,
template
<
typename
T
>
struct
Conv2dFunctor
<
DeviceType
::
OPENCL
,
T
>
:
Conv2dFunctorBase
{
Conv2dFunctor
(
const
int
*
strides
,
const
Padding
&
paddings
,
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
Conv2dFunctorBase
(
strides
,
padding_type
,
paddings
,
dilations
,
activation
,
...
...
mace/kernels/depthwise_conv2d.h
浏览文件 @
c5bc6a5a
...
...
@@ -237,20 +237,23 @@ void DepthwiseConv2dNoOOBCheckKernel(const T *input_ptr,
struct
DepthwiseConv2dFunctorBase
{
DepthwiseConv2dFunctorBase
(
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
strides_
(
strides
),
padding_
(
padding
),
padding_type_
(
padding_type
),
paddings_
(
paddings
),
dilations_
(
dilations
),
activation_
(
activation
),
relux_max_limit_
(
relux_max_limit
),
prelu_alpha_
(
prelu_alpha
)
{}
const
int
*
strides_
;
// [stride_h, stride_w]
const
Padding
padding_
;
const
Padding
padding_type_
;
std
::
vector
<
int
>
paddings_
;
const
int
*
dilations_
;
// [dilation_h, dilation_w]
const
ActivationType
activation_
;
const
float
relux_max_limit_
;
...
...
@@ -260,13 +263,15 @@ struct DepthwiseConv2dFunctorBase {
template
<
DeviceType
D
,
typename
T
>
struct
DepthwiseConv2dFunctor
:
public
DepthwiseConv2dFunctorBase
{
DepthwiseConv2dFunctor
(
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
DepthwiseConv2dFunctorBase
(
strides
,
padding
,
padding_type
,
paddings
,
dilations
,
activation
,
relux_max_limit
,
...
...
@@ -289,10 +294,12 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
fake_filter_shape
[
3
]
=
1
;
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
auto
input_shape
=
fake_filter_shape
;
output
->
Resize
(
output_shape
);
...
...
@@ -322,10 +329,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
;
...
...
@@ -413,13 +420,15 @@ template <typename T>
struct
DepthwiseConv2dFunctor
<
DeviceType
::
OPENCL
,
T
>
:
DepthwiseConv2dFunctorBase
{
DepthwiseConv2dFunctor
(
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
,
const
ActivationType
activation
,
const
float
relux_max_limit
,
const
float
prelu_alpha
)
:
DepthwiseConv2dFunctorBase
(
strides
,
padding
,
padding_type
,
paddings
,
dilations
,
activation
,
relux_max_limit
,
...
...
mace/kernels/opencl/conv_2d_opencl.cc
浏览文件 @
c5bc6a5a
...
...
@@ -80,10 +80,12 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
}
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
paddings_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter
->
shape
().
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
std
::
vector
<
size_t
>
output_image_shape
;
CalImage2DShape
(
output_shape
,
BufferType
::
IN_OUT_CHANNEL
,
output_image_shape
);
...
...
@@ -93,11 +95,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
浏览文件 @
c5bc6a5a
...
...
@@ -140,7 +140,7 @@ void DepthwiseConv2dFunctor<DeviceType::OPENCL, T>::operator()(
<<
" is not implemented yet, using slow version"
;
// TODO(heliangliang) The CPU/NEON kernel should map the buffer
DepthwiseConv2dFunctor
<
DeviceType
::
CPU
,
float
>
(
strides_
,
padding_
,
dilations_
,
activation_
,
relux_max_limit_
,
strides_
,
padding_
type_
,
paddings_
,
dilations_
,
activation_
,
relux_max_limit_
,
prelu_alpha_
)(
input
,
filter
,
bias
,
output
,
future
);
return
;
}
...
...
@@ -153,16 +153,18 @@ void DepthwiseConv2dFunctor<DeviceType::OPENCL, T>::operator()(
fake_filter_shape
[
3
]
=
1
;
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
fake_filter_shape
.
data
(),
dilations_
,
strides_
,
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
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
浏览文件 @
c5bc6a5a
...
...
@@ -18,16 +18,18 @@ void PoolingFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
MACE_CHECK
(
dilations_
[
0
]
==
1
&&
dilations_
[
1
]
==
1
)
<<
"Pooling opencl kernel not support dilation yet"
;
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
std
::
vector
<
index_t
>
filter_shape
=
{
kernels_
[
0
],
kernels_
[
1
],
input
->
dim
(
3
),
input
->
dim
(
3
)
};
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
std
::
vector
<
size_t
>
output_image_shape
;
CalImage2DShape
(
output_shape
,
BufferType
::
IN_OUT_CHANNEL
,
output_image_shape
);
...
...
@@ -64,8 +66,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
浏览文件 @
c5bc6a5a
...
...
@@ -17,10 +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
};
std
::
vector
<
int
>
paddings
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
.
data
(),
strides_
.
data
(),
paddings_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
.
data
(),
strides_
.
data
(),
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
const
index_t
round_h
=
(
output_shape
[
1
]
+
1
)
/
2
;
const
index_t
round_w
=
(
output_shape
[
2
]
+
1
)
/
2
;
...
...
@@ -50,8 +52,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
浏览文件 @
c5bc6a5a
...
...
@@ -24,18 +24,21 @@ struct PoolingFunctorBase {
PoolingFunctorBase
(
const
PoolingType
pooling_type
,
const
int
*
kernels
,
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
)
:
pooling_type_
(
pooling_type
),
kernels_
(
kernels
),
strides_
(
strides
),
padding_
(
padding
),
padding_type_
(
padding_type
),
paddings_
(
paddings
),
dilations_
(
dilations
)
{}
const
PoolingType
pooling_type_
;
const
int
*
kernels_
;
const
int
*
strides_
;
const
Padding
padding_
;
const
Padding
padding_type_
;
std
::
vector
<
int
>
paddings_
;
const
int
*
dilations_
;
};
...
...
@@ -44,27 +47,31 @@ struct PoolingFunctor : PoolingFunctorBase {
PoolingFunctor
(
const
PoolingType
pooling_type
,
const
int
*
kernels
,
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
)
:
PoolingFunctorBase
(
pooling_type
,
kernels
,
strides
,
padding
,
dilations
)
{}
strides
,
padding
_type
,
paddings
,
dilations
)
{}
void
operator
()(
const
Tensor
*
input_tensor
,
Tensor
*
output_tensor
,
StatsFuture
*
future
)
{
std
::
vector
<
index_t
>
output_shape
(
4
);
std
::
vector
<
int
>
paddings
(
2
);
std
::
vector
<
index_t
>
filter_shape
=
{
kernels_
[
0
],
kernels_
[
1
],
input_tensor
->
dim
(
3
),
input_tensor
->
dim
(
3
)
};
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_
,
output_shape
.
data
(),
paddings
.
data
());
if
(
paddings_
.
empty
())
{
paddings_
.
resize
(
2
);
kernels
::
CalcNHWCPaddingAndOutputSize
(
input_tensor
->
shape
().
data
(),
filter_shape
.
data
(),
dilations_
,
strides_
,
this
->
padding_type_
,
output_shape
.
data
(),
paddings_
.
data
());
}
output_tensor
->
Resize
(
output_shape
);
Tensor
::
MappingGuard
in_guard
(
input_tensor
);
...
...
@@ -92,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)
...
...
@@ -163,11 +170,12 @@ struct PoolingFunctor<DeviceType::OPENCL, T> : PoolingFunctorBase {
PoolingFunctor
(
const
PoolingType
pooling_type
,
const
int
*
kernels
,
const
int
*
strides
,
const
Padding
padding
,
const
Padding
padding_type
,
const
std
::
vector
<
int
>
&
paddings
,
const
int
*
dilations
)
:
PoolingFunctorBase
(
pooling_type
,
kernels
,
strides
,
padding
,
dilations
)
{}
strides
,
padding
_type
,
paddings
,
dilations
)
{}
void
operator
()(
const
Tensor
*
input_tensor
,
Tensor
*
output_tensor
,
StatsFuture
*
future
);
...
...
mace/kernels/winograd_transform.h
浏览文件 @
c5bc6a5a
...
...
@@ -15,18 +15,22 @@ namespace mace {
namespace
kernels
{
struct
WinogradTransformFunctorBase
{
WinogradTransformFunctorBase
(
const
Padding
&
paddings
)
:
strides_
({
1
,
1
}),
dilations_
({
1
,
1
}),
paddings_
(
paddings
)
{}
WinogradTransformFunctorBase
(
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
)
:
strides_
({
1
,
1
}),
dilations_
({
1
,
1
}),
padding_type_
(
padding_type
),
paddings_
(
paddings
)
{}
const
std
::
vector
<
int
>
strides_
;
// [stride_h, stride_w]
const
std
::
vector
<
int
>
dilations_
;
// [dilation_h, dilation_w]
Padding
paddings_
;
Padding
padding_type_
;
std
::
vector
<
int
>
paddings_
;
};
template
<
DeviceType
D
,
typename
T
>
struct
WinogradTransformFunctor
:
WinogradTransformFunctorBase
{
WinogradTransformFunctor
(
const
Padding
&
paddings
)
:
WinogradTransformFunctorBase
(
paddings
)
{}
WinogradTransformFunctor
(
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
)
:
WinogradTransformFunctorBase
(
padding_type
,
paddings
)
{}
void
operator
()(
const
Tensor
*
input
,
Tensor
*
output
,
...
...
@@ -38,8 +42,9 @@ struct WinogradTransformFunctor : WinogradTransformFunctorBase {
template
<
typename
T
>
struct
WinogradTransformFunctor
<
DeviceType
::
OPENCL
,
T
>
:
WinogradTransformFunctorBase
{
WinogradTransformFunctor
(
const
Padding
&
paddings
)
:
WinogradTransformFunctorBase
(
paddings
)
{}
WinogradTransformFunctor
(
const
Padding
&
padding_type
,
const
std
::
vector
<
int
>
&
paddings
)
:
WinogradTransformFunctorBase
(
padding_type
,
paddings
)
{}
void
operator
()(
const
Tensor
*
input
,
Tensor
*
output
,
...
...
mace/ops/addn.h
浏览文件 @
c5bc6a5a
...
...
@@ -26,7 +26,12 @@ class AddNOp : public Operator<D, T> {
for
(
int
i
=
1
;
i
<
n
;
++
i
)
{
inputs
[
i
]
=
this
->
Input
(
i
);
MACE_CHECK
(
inputs
[
0
]
->
dim_size
()
==
inputs
[
i
]
->
dim_size
());
MACE_CHECK
(
inputs
[
0
]
->
size
()
==
inputs
[
i
]
->
size
());
MACE_CHECK
(
inputs
[
0
]
->
size
()
==
inputs
[
i
]
->
size
())
<<
"Input 0: "
<<
MakeString
(
inputs
[
0
]
->
shape
())
<<
", size: "
<<
inputs
[
0
]
->
size
()
<<
". Input "
<<
i
<<
": "
<<
MakeString
(
inputs
[
i
]
->
shape
())
<<
", size: "
<<
inputs
[
i
]
->
size
();
}
functor_
(
inputs
,
output_tensor
,
future
);
...
...
mace/ops/conv_2d.h
浏览文件 @
c5bc6a5a
...
...
@@ -19,7 +19,8 @@ class Conv2dOp : public ConvPool2dOpBase<D, T> {
Conv2dOp
(
const
OperatorDef
&
op_def
,
Workspace
*
ws
)
:
ConvPool2dOpBase
<
D
,
T
>
(
op_def
,
ws
),
functor_
(
this
->
strides_
.
data
(),
this
->
padding_
,
this
->
padding_type_
,
this
->
paddings_
,
this
->
dilations_
.
data
(),
kernels
::
ActivationType
::
NOOP
,
0.0
f
,
...
...
mace/ops/conv_pool_2d_base.h
浏览文件 @
c5bc6a5a
...
...
@@ -16,14 +16,16 @@ class ConvPool2dOpBase : public Operator<D, T> {
ConvPool2dOpBase
(
const
OperatorDef
&
op_def
,
Workspace
*
ws
)
:
Operator
<
D
,
T
>
(
op_def
,
ws
),
strides_
(
OperatorBase
::
GetRepeatedArgument
<
int
>
(
"strides"
)),
padding_
(
static_cast
<
Padding
>
(
OperatorBase
::
GetSingleArgument
<
int
>
(
padding_
type_
(
static_cast
<
Padding
>
(
OperatorBase
::
GetSingleArgument
<
int
>
(
"padding"
,
static_cast
<
int
>
(
SAME
)))),
paddings_
(
OperatorBase
::
GetRepeatedArgument
<
int
>
(
"padding_values"
)),
dilations_
(
OperatorBase
::
GetRepeatedArgument
<
int
>
(
"dilations"
,
{
1
,
1
}))
{}
protected:
std
::
vector
<
int
>
strides_
;
Padding
padding_
;
Padding
padding_type_
;
std
::
vector
<
int
>
paddings_
;
std
::
vector
<
int
>
dilations_
;
};
...
...
mace/ops/depthwise_conv2d.h
浏览文件 @
c5bc6a5a
...
...
@@ -20,7 +20,8 @@ class DepthwiseConv2dOp : public ConvPool2dOpBase<D, T> {
DepthwiseConv2dOp
(
const
OperatorDef
&
op_def
,
Workspace
*
ws
)
:
ConvPool2dOpBase
<
D
,
T
>
(
op_def
,
ws
),
functor_
(
this
->
strides_
.
data
(),
this
->
padding_
,
this
->
padding_type_
,
this
->
paddings_
,
this
->
dilations_
.
data
(),
kernels
::
StringToActivationType
(
OperatorBase
::
GetSingleArgument
<
std
::
string
>
(
"activation"
,
...
...
mace/ops/fused_conv_2d.h
浏览文件 @
c5bc6a5a
...
...
@@ -19,7 +19,8 @@ class FusedConv2dOp : public ConvPool2dOpBase<D, T> {
FusedConv2dOp
(
const
OperatorDef
&
op_def
,
Workspace
*
ws
)
:
ConvPool2dOpBase
<
D
,
T
>
(
op_def
,
ws
),
functor_
(
this
->
strides_
.
data
(),
this
->
padding_
,
this
->
padding_type_
,
this
->
paddings_
,
this
->
dilations_
.
data
(),
kernels
::
StringToActivationType
(
OperatorBase
::
GetSingleArgument
<
std
::
string
>
(
"activation"
,
...
...
mace/ops/pooling.h
浏览文件 @
c5bc6a5a
...
...
@@ -21,7 +21,7 @@ class PoolingOp : public ConvPool2dOpBase<D, T> {
static_cast
<
PoolingType
>
(
OperatorBase
::
GetSingleArgument
<
int
>
(
"pooling_type"
,
static_cast
<
int
>
(
AVG
)))),
functor_
(
pooling_type_
,
kernels_
.
data
(),
this
->
strides_
.
data
(),
this
->
padding_
,
this
->
dilations_
.
data
()){};
this
->
padding_
type_
,
this
->
paddings_
,
this
->
dilations_
.
data
()){};
bool
Run
(
StatsFuture
*
future
)
override
{
const
Tensor
*
input
=
this
->
Input
(
INPUT
);
...
...
mace/ops/winograd_transform.h
浏览文件 @
c5bc6a5a
...
...
@@ -18,7 +18,8 @@ class WinogradTransformOp : public Operator<D, T> {
WinogradTransformOp
(
const
OperatorDef
&
op_def
,
Workspace
*
ws
)
:
Operator
<
D
,
T
>
(
op_def
,
ws
),
functor_
(
static_cast
<
Padding
>
(
OperatorBase
::
GetSingleArgument
<
int
>
(
"padding"
,
static_cast
<
int
>
(
VALID
))))
{}
"padding"
,
static_cast
<
int
>
(
VALID
))),
OperatorBase
::
GetRepeatedArgument
<
int
>
(
"padding_values"
))
{}
bool
Run
(
StatsFuture
*
future
)
override
{
const
Tensor
*
input_tensor
=
this
->
Input
(
INPUT
);
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录