Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
毕竟曾有刹那
Mace
提交
405c502a
Mace
项目概览
毕竟曾有刹那
/
Mace
与 Fork 源项目一致
Fork自
Xiaomi / Mace
通知
1
Star
0
Fork
0
代码
文件
提交
分支
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,发现更多精彩内容 >>
提交
405c502a
编写于
5月 10, 2019
作者:
L
liutuo
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
add onnx clip
上级
52a379a2
变更
7
隐藏空白更改
内联
并排
Showing
7 changed file
with
76 addition
and
14 deletion
+76
-14
.gitignore
.gitignore
+1
-0
mace/ops/common/eltwise_type.h
mace/ops/common/eltwise_type.h
+2
-1
mace/ops/eltwise.cc
mace/ops/eltwise.cc
+29
-1
mace/ops/opencl/cl/eltwise.cl
mace/ops/opencl/cl/eltwise.cl
+3
-1
mace/python/tools/converter_tool/base_converter.py
mace/python/tools/converter_tool/base_converter.py
+2
-0
mace/python/tools/converter_tool/onnx_converter.py
mace/python/tools/converter_tool/onnx_converter.py
+15
-1
test/ccunit/mace/ops/eltwise_test.cc
test/ccunit/mace/ops/eltwise_test.cc
+24
-10
未找到文件。
.gitignore
浏览文件 @
405c502a
...
...
@@ -27,3 +27,4 @@ examples/android/macelibrary/src/main/cpp/mace/
examples/android/macelibrary/src/main/cpp/include/
examples/android/macelibrary/src/main/cpp/lib/arm64-v8a/
examples/android/macelibrary/src/main/jniLibs/arm64-v8a/
mace/ops/common/eltwise_type.h
浏览文件 @
405c502a
...
...
@@ -31,7 +31,8 @@ enum EltwiseType {
POW
=
9
,
EQUAL
=
10
,
FLOOR_DIV
=
11
,
NONE
=
12
,
CLIP
=
12
,
NONE
=
13
,
};
}
// namespace ops
...
...
mace/ops/eltwise.cc
浏览文件 @
405c502a
...
...
@@ -376,6 +376,15 @@ inline void TensorBroadcastEltwise(const OpContext *context,
}
}
break
;
case
CLIP
:
for
(
index_t
d
=
start0
;
d
<
end0
;
d
+=
step0
)
{
for
(
index_t
i
=
start1
;
i
<
end1
;
i
+=
step1
)
{
output
[
i
+
d
*
common_size
]
=
std
::
fmaxf
(
coeff
[
0
],
std
::
fminf
(
coeff
[
1
],
input0
[
i
+
d
*
common_size
]));
}
}
break
;
default:
LOG
(
FATAL
)
<<
"Eltwise op not support type "
<<
type
;
}
},
0
,
diff_size
,
1
,
0
,
common_size
,
1
);
...
...
@@ -497,6 +506,11 @@ inline void TensorEltwise(const OpContext *context,
output
[
i
]
=
input0
[
i
]
==
input1
[
i
];
}
break
;
case
CLIP
:
for
(
index_t
i
=
start
;
i
<
end
;
i
+=
step
)
{
output
[
i
]
=
std
::
fmaxf
(
coeff
[
0
],
std
::
fminf
(
coeff
[
1
],
input0
[
i
]));
}
break
;
default:
LOG
(
FATAL
)
<<
"Eltwise op not support type "
<<
type
;
}
},
0
,
size
,
1
);
...
...
@@ -617,7 +631,11 @@ inline void TensorScalarEltwise(const OpContext *context,
for
(
index_t
i
=
start
;
i
<
end
;
i
+=
step
)
{
output
[
i
]
=
input0
[
i
]
==
input1
;
}
break
;
case
CLIP
:
for
(
index_t
i
=
start
;
i
<
end
;
i
+=
step
)
{
output
[
i
]
=
std
::
fmaxf
(
coeff
[
0
],
std
::
fminf
(
coeff
[
1
],
input0
[
i
]));
}
break
;
default:
LOG
(
FATAL
)
<<
"Eltwise op not support type "
<<
type
;
}
...
...
@@ -886,6 +904,11 @@ class EltwiseOp : public Operation {
input1
=
&
scalar_tensor_
;
}
if
(
type_
==
CLIP
)
{
MACE_CHECK
(
coeff_
.
size
()
==
2
&&
coeff_
[
0
]
<
coeff_
[
1
],
"Clip's min/max values are not correct."
);
}
if
(
IsLogicalType
(
type_
))
{
// as we do not have bool-type tensor, we use int type
return
DoEltwise
<
int32_t
>
(
context
,
input0
,
input1
,
output
);
...
...
@@ -1144,6 +1167,11 @@ class EltwiseOp<DeviceType::GPU, T> : public Operation {
Operation
::
GetOptionalArg
<
int
>
(
"type"
,
static_cast
<
int
>
(
ops
::
EltwiseType
::
NONE
)));
std
::
vector
<
float
>
coeff
=
Operation
::
GetRepeatedArgs
<
float
>
(
"coeff"
);
if
(
type
==
ops
::
EltwiseType
::
CLIP
)
{
MACE_CHECK
(
coeff
.
size
()
==
2
&&
coeff
[
0
]
<
coeff
[
1
],
"Clip's min/max values are not correct."
);
}
float
scalar_input
=
Operation
::
GetOptionalArg
<
float
>
(
"scalar_input"
,
1.0
);
int32_t
scalar_input_index
=
Operation
::
GetOptionalArg
<
int32_t
>
(
"scalar_input_index"
,
1
);
...
...
mace/ops/opencl/cl/eltwise.cl
浏览文件 @
405c502a
...
...
@@ -87,13 +87,15 @@ __kernel void eltwise(OUT_OF_RANGE_PARAMS
#else
out = floor(in0 / in1);
#endif
#elif ELTWISE_TYPE == 12
out = fmax(coeff0, fmin(coeff1, in0));
#endif
#if defined(NOT_DIVISIBLE_FOUR) && \
((ELTWISE_TYPE == 3 || ELTWISE_TYPE == 9 || ELTWISE_TYPE == 11) \
|| ((defined(INPUT_SCALAR) || defined(INPUT_TENSOR_BC_CHAN)) && \
(ELTWISE_TYPE == 0 || ELTWISE_TYPE == 1 || ELTWISE_TYPE == 4 || \
ELTWISE_TYPE == 5 |
|
ELTWISE_TYPE
==
8
)
))
ELTWISE_TYPE == 5 || ELTWISE_TYPE == 8
|
|
ELTWISE_TYPE
==
12
)
))
const
int
remain_channel
=
channel
-
4
*
chan_idx
;
if
(
remain_channel
<
4
)
{
switch
(
remain_channel
)
{
...
...
mace/python/tools/converter_tool/base_converter.py
浏览文件 @
405c502a
...
...
@@ -80,6 +80,7 @@ class EltwiseType(Enum):
POW
=
9
EQUAL
=
10
FLOOR_DIV
=
11
CLIP
=
12
class
ReduceType
(
Enum
):
...
...
@@ -273,6 +274,7 @@ class MaceKeyword(object):
mace_exclusive_str
=
'exclusive'
mace_reverse_str
=
'reverse'
mace_const_data_num_arg_str
=
'const_data_num'
mace_coeff_str
=
'coeff'
class
TransformerRule
(
Enum
):
...
...
mace/python/tools/converter_tool/onnx_converter.py
浏览文件 @
405c502a
...
...
@@ -69,7 +69,7 @@ OnnxSupportedOps = [
'BatchNorm'
,
'Cast'
,
# 'Ceil',
#
'Clip',
'Clip'
,
# 'Compress',
'Concat'
,
# 'Constant',
...
...
@@ -300,6 +300,7 @@ class OnnxConverter(base_converter.ConverterInterface):
OnnxOpType
.
Sqrt
.
name
:
EltwiseType
.
POW
,
OnnxOpType
.
Reciprocal
.
name
:
EltwiseType
.
POW
,
OnnxOpType
.
Scale
.
name
:
EltwiseType
.
PROD
,
OnnxOpType
.
Clip
.
name
:
EltwiseType
.
CLIP
,
}
reduce_type
=
{
...
...
@@ -331,6 +332,7 @@ class OnnxConverter(base_converter.ConverterInterface):
OnnxOpType
.
BatchNormalization
.
name
:
self
.
convert_fused_batchnorm
,
OnnxOpType
.
BatchNorm
.
name
:
self
.
convert_fused_batchnorm
,
OnnxOpType
.
Cast
.
name
:
self
.
convert_cast
,
OnnxOpType
.
Clip
.
name
:
self
.
convert_eltwise
,
OnnxOpType
.
Concat
.
name
:
self
.
convert_concat
,
OnnxOpType
.
Conv
.
name
:
self
.
convert_conv2d
,
OnnxOpType
.
ConvTranspose
.
name
:
self
.
convert_deconv
,
...
...
@@ -862,6 +864,18 @@ class OnnxConverter(base_converter.ConverterInterface):
value_arg
=
op
.
arg
.
add
()
value_arg
.
name
=
MaceKeyword
.
mace_scalar_input_str
value_arg
.
f
=
value
elif
node
.
op_type
==
OnnxOpType
.
Clip
.
name
:
if
'min'
in
node
.
attrs
:
min_value
=
node
.
attrs
[
'min'
]
else
:
min_value
=
np
.
finfo
(
np
.
float32
).
min
if
'max'
in
node
.
attrs
:
max_value
=
node
.
attrs
[
'max'
]
else
:
max_value
=
np
.
finfo
(
np
.
float32
).
max
coeff_arg
=
op
.
arg
.
add
()
coeff_arg
.
name
=
MaceKeyword
.
mace_coeff_str
coeff_arg
.
floats
.
extend
([
min_value
,
max_value
])
@
staticmethod
def
copy_node_attr
(
op
,
node
,
attr_name
,
dtype
=
AttributeType
.
INT
,
...
...
test/ccunit/mace/ops/eltwise_test.cc
浏览文件 @
405c502a
...
...
@@ -114,7 +114,8 @@ void SimpleTensorEltwise(const ops::EltwiseType type,
// Add input data
net
.
AddInputFromArray
<
D
,
T
>
(
"Input0"
,
shape0
,
input0
);
net
.
AddInputFromArray
<
D
,
T
>
(
"Input1"
,
shape1
,
input1
);
if
(
shape1
.
size
()
>
0
&&
input1
.
size
()
>
0
)
net
.
AddInputFromArray
<
D
,
T
>
(
"Input1"
,
shape1
,
input1
);
if
(
D
==
DeviceType
::
CPU
)
{
auto
op_builder
=
...
...
@@ -136,7 +137,7 @@ void SimpleTensorEltwise(const ops::EltwiseType type,
net
.
TransformDataFormat
<
D
,
T
>
(
"Input1"
,
DataFormat
::
NHWC
,
"TInput1"
,
DataFormat
::
NCHW
);
op_builder
.
Input
(
"TInput1"
);
}
else
{
}
else
if
(
shape1
.
size
()
>
0
)
{
op_builder
.
Input
(
"Input1"
);
}
op_builder
.
Finalize
(
net
.
NewOperatorDef
());
...
...
@@ -146,14 +147,15 @@ void SimpleTensorEltwise(const ops::EltwiseType type,
net
.
TransformDataFormat
<
D
,
DstType
>
(
"TOutput"
,
DataFormat
::
NCHW
,
"Output"
,
DataFormat
::
NHWC
);
}
else
{
OpDefBuilder
(
"Eltwise"
,
"EltwiseTest"
)
.
Input
(
"Input0"
)
.
Input
(
"Input1"
)
.
AddIntArg
(
"type"
,
static_cast
<
int
>
(
type
))
.
AddFloatsArg
(
"coeff"
,
coeff
)
.
Output
(
"Output"
)
.
Finalize
(
net
.
NewOperatorDef
());
auto
op_builder
=
OpDefBuilder
(
"Eltwise"
,
"EltwiseTest"
)
.
Input
(
"Input0"
)
.
AddIntArg
(
"type"
,
static_cast
<
int
>
(
type
))
.
AddFloatsArg
(
"coeff"
,
coeff
)
.
Output
(
"Output"
);
if
(
input1
.
size
()
>
0
&&
shape1
.
size
()
>
0
)
op_builder
.
Input
(
"Input1"
);
op_builder
.
Finalize
(
net
.
NewOperatorDef
());
// Run
net
.
RunOp
(
D
);
}
...
...
@@ -500,6 +502,10 @@ TEST_F(EltwiseOpTest, CPUSimpleTensorTensor) {
SimpleTensorEltwise
<
DeviceType
::
CPU
,
int32_t
,
int32_t
>
(
ops
::
EltwiseType
::
EQUAL
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
},
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
},
{
1
,
1
,
1
,
1
,
1
,
1
});
SimpleTensorEltwise
<
DeviceType
::
CPU
,
float
,
float
>
(
ops
::
EltwiseType
::
CLIP
,
{
1
,
2
,
1
,
5
},
{
1
,
2
,
3
,
4
,
5
,
1
,
2
,
3
,
4
,
5
},
{},
{},
{
2
,
2
,
3
,
3
,
3
,
2
,
2
,
3
,
3
,
3
},
{
2.0
f
,
3.0
f
});
}
TEST_F
(
EltwiseOpTest
,
GPUSimpleTensorTensor
)
{
SimpleTensorEltwise
<
DeviceType
::
GPU
,
float
,
float
>
(
...
...
@@ -535,6 +541,10 @@ TEST_F(EltwiseOpTest, GPUSimpleTensorTensor) {
ops
::
EltwiseType
::
SQR_DIFF
,
{
1
,
2
,
1
,
5
},
{
1
,
2
,
3
,
4
,
5
,
1
,
2
,
3
,
4
,
5
},
{
1
,
2
,
1
,
5
},
{
1
,
2
,
3
,
4
,
5
,
6
,
7
,
8
,
9
,
10
},
{
0
,
0
,
0
,
0
,
0
,
25
,
25
,
25
,
25
,
25
});
SimpleTensorEltwise
<
DeviceType
::
GPU
,
float
,
float
>
(
ops
::
EltwiseType
::
CLIP
,
{
1
,
2
,
1
,
5
},
{
1
,
2
,
3
,
4
,
5
,
1
,
2
,
3
,
4
,
5
},
{},
{},
{
2
,
2
,
3
,
3
,
3
,
2
,
2
,
3
,
3
,
3
},
{
2.0
f
,
3.0
f
});
}
namespace
{
...
...
@@ -912,6 +922,8 @@ TEST_F(EltwiseOpTest, RandomTensorTensorFloat) {
{
3
,
31
,
37
,
17
});
RandomTensorEltwise
<
float
>
(
ops
::
EltwiseType
::
SQR_DIFF
,
{
3
,
31
,
37
,
17
},
{
3
,
31
,
37
,
17
});
RandomTensorEltwise
<
float
>
(
ops
::
EltwiseType
::
CLIP
,
{
3
,
31
,
37
,
17
},
{
3
,
31
,
37
,
17
},
{
-
0.2
,
0.85
});
}
TEST_F
(
EltwiseOpTest
,
RandomTensorTensorHalf
)
{
...
...
@@ -929,6 +941,8 @@ TEST_F(EltwiseOpTest, RandomTensorTensorHalf) {
{
3
,
31
,
37
,
17
});
RandomTensorEltwise
<
half
>
(
ops
::
EltwiseType
::
SQR_DIFF
,
{
3
,
31
,
37
,
17
},
{
3
,
31
,
37
,
17
});
RandomTensorEltwise
<
half
>
(
ops
::
EltwiseType
::
CLIP
,
{
3
,
31
,
37
,
17
},
{
3
,
31
,
37
,
17
},
{
-
0.2
,
0.85
});
}
TEST_F
(
EltwiseOpTest
,
TensorGeneralBroadcastCPU
)
{
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录