Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
as350144
Mace
提交
62216ad1
Mace
项目概览
as350144
/
Mace
与 Fork 源项目一致
Fork自
Xiaomi / Mace
通知
2
Star
1
Fork
1
代码
文件
提交
分支
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,发现更多精彩内容 >>
提交
62216ad1
编写于
11月 13, 2017
作者:
L
liuqi
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Finish add2 opencl kernel and update tf converter.
上级
47eece0b
变更
8
隐藏空白更改
内联
并排
Showing
8 changed file
with
235 addition
and
59 deletion
+235
-59
mace/kernels/addn.h
mace/kernels/addn.h
+16
-7
mace/kernels/neon/addn_neon.cc
mace/kernels/neon/addn_neon.cc
+8
-5
mace/kernels/opencl/addn.cc
mace/kernels/opencl/addn.cc
+54
-0
mace/kernels/opencl/cl/addn.cl
mace/kernels/opencl/cl/addn.cl
+17
-0
mace/ops/addn.cc
mace/ops/addn.cc
+2
-0
mace/ops/addn.h
mace/ops/addn.h
+4
-7
mace/ops/addn_test.cc
mace/ops/addn_test.cc
+87
-10
mace/python/tools/tf_converter_lib.py
mace/python/tools/tf_converter_lib.py
+47
-30
未找到文件。
mace/kernels/addn.h
浏览文件 @
62216ad1
...
...
@@ -10,22 +10,31 @@
namespace
mace
{
namespace
kernels
{
template
<
DeviceType
D
,
typename
T
>
template
<
DeviceType
D
,
typename
T
>
struct
AddNFunctor
{
void
operator
()(
const
vector
<
const
T
*>
&
inputs
,
T
*
output
,
index_t
size
)
{
memset
(
output
,
0
,
size
*
sizeof
(
T
));
int
n
=
inputs
.
size
();
void
operator
()(
std
::
vector
<
const
Tensor
*>
&
input_tensors
,
Tensor
*
output_tensor
)
{
Tensor
::
MappingGuard
output_map
(
output_tensor
);
index_t
size
=
input_tensors
[
0
]
->
size
();
T
*
output_ptr
=
output_tensor
->
mutable_data
<
T
>
();
memset
(
output_ptr
,
0
,
size
*
sizeof
(
T
));
int
n
=
input_tensors
.
size
();
for
(
int
i
=
0
;
i
<
n
;
++
i
)
{
Tensor
::
MappingGuard
input_map
(
input_tensors
[
i
]);
const
T
*
input_ptr
=
input_tensors
[
i
]
->
data
<
T
>
();
for
(
index_t
j
=
0
;
j
<
size
;
++
j
)
{
output
[
j
]
+=
inputs
[
i
]
[
j
];
output
_ptr
[
j
]
+=
input_ptr
[
j
];
}
}
}
};
template
<
>
template
<
>
void
AddNFunctor
<
DeviceType
::
NEON
,
float
>::
operator
()(
const
vector
<
const
float
*>
&
inputs
,
float
*
output
,
index_t
size
);
std
::
vector
<
const
Tensor
*>
&
input_tensors
,
Tensor
*
output_tensor
);
template
<
>
void
AddNFunctor
<
DeviceType
::
OPENCL
,
float
>::
operator
()(
std
::
vector
<
const
Tensor
*>
&
inputs
,
Tensor
*
output
);
}
// namespace kernels
}
// namespace mace
...
...
mace/kernels/neon/addn_neon.cc
浏览文件 @
62216ad1
...
...
@@ -10,10 +10,12 @@ namespace kernels {
template
<
>
void
AddNFunctor
<
DeviceType
::
NEON
,
float
>::
operator
()(
const
vector
<
const
float
*>
&
inputs
,
float
*
output
,
index_t
size
)
{
std
::
vector
<
const
Tensor
*>
&
input_tensors
,
Tensor
*
output_tensor
)
{
// TODO: neon mem copy
memset
(
output
,
0
,
size
*
sizeof
(
float
));
int
n
=
inputs
.
size
();
index_t
size
=
output_tensor
->
size
();
float
*
output_ptr
=
output_tensor
->
mutable_data
<
float
>
();
memset
(
output_ptr
,
0
,
size
*
sizeof
(
float
));
int
n
=
input_tensors
.
size
();
int64_t
cost
=
size
*
n
;
int64_t
groups
=
1
;
if
(
cost
>
kCostPerGroup
)
{
...
...
@@ -27,8 +29,9 @@ void AddNFunctor<DeviceType::NEON, float>::operator()(
int
nn
=
count
>>
2
;
int
remain
=
count
-
(
nn
<<
2
);
for
(
int64_t
j
=
0
;
j
<
n
;
++
j
)
{
const
float
*
inptr
=
inputs
[
j
]
+
i
;
float
*
outptr
=
output
+
i
;
const
float
*
input_base
=
input_tensors
[
j
]
->
data
<
float
>
();
const
float
*
inptr
=
input_base
+
i
;
float
*
outptr
=
output_ptr
+
i
;
for
(
int
k
=
0
;
k
<
nn
;
++
k
)
{
float32x4_t
_inptr
=
vld1q_f32
(
inptr
);
float32x4_t
_outptr
=
vld1q_f32
(
outptr
);
...
...
mace/kernels/opencl/addn.cc
0 → 100644
浏览文件 @
62216ad1
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#include "mace/kernels/addn.h"
#include "mace/core/runtime/opencl/opencl_runtime.h"
namespace
mace
{
namespace
kernels
{
static
void
Add2
(
const
Tensor
*
input0
,
const
Tensor
*
input1
,
Tensor
*
output
)
{
index_t
element_size
=
input0
->
NumElements
();
index_t
blocks
=
(
element_size
+
3
)
/
4
;
const
uint32_t
gws
=
blocks
;
auto
runtime
=
OpenCLRuntime
::
Get
();
auto
program
=
runtime
->
program
();
auto
addn_kernel
=
cl
::
Kernel
(
program
,
"add2"
);
const
uint32_t
lws
=
runtime
->
GetKernelMaxWorkGroupSize
(
addn_kernel
);
uint32_t
idx
=
0
;
addn_kernel
.
setArg
(
idx
++
,
*
(
static_cast
<
const
cl
::
Buffer
*>
(
input0
->
buffer
())));
addn_kernel
.
setArg
(
idx
++
,
*
(
static_cast
<
const
cl
::
Buffer
*>
(
input1
->
buffer
())));
addn_kernel
.
setArg
(
idx
++
,
static_cast
<
int32_t
>
(
element_size
));
addn_kernel
.
setArg
(
idx
++
,
*
(
static_cast
<
cl
::
Buffer
*>
(
output
->
buffer
())));
cl_int
error
=
runtime
->
command_queue
().
enqueueNDRangeKernel
(
addn_kernel
,
cl
::
NullRange
,
cl
::
NDRange
(
gws
),
cl
::
NDRange
(
lws
));
MACE_CHECK
(
error
==
CL_SUCCESS
);
}
template
<
>
void
AddNFunctor
<
DeviceType
::
OPENCL
,
float
>::
operator
()(
std
::
vector
<
const
Tensor
*>
&
input_tensors
,
Tensor
*
output_tensor
)
{
if
(
input_tensors
.
empty
()
||
input_tensors
.
front
()
==
nullptr
)
{
return
;
}
size_t
size
=
input_tensors
.
size
();
switch
(
size
)
{
case
2
:
Add2
(
input_tensors
[
0
],
input_tensors
[
1
],
output_tensor
);
break
;
default:
MACE_NOT_IMPLEMENTED
;
}
};
}
// namespace kernels
}
// namespace mace
mace/kernels/opencl/cl/addn.cl
0 → 100644
浏览文件 @
62216ad1
__kernel
void
add2
(
__global
const
float
*input0,
__global
const
float
*input1,
__private
const
int
size,
__global
float
*output
)
{
int
idx
=
get_global_id
(
0
)
;
if
(
idx
+
4
>
size
)
{
for
(
; idx < size; ++idx) {
*
(
output+idx
)
=
*
(
input0+idx
)
+
*
(
input1+idx
)
;
}
}
else
{
float4
in_data0
=
vload4
(
idx,
input0
)
;
float4
in_data1
=
vload4
(
idx,
input1
)
;
vstore4
(
in_data0+in_data1,
idx,
output
)
;
}
}
mace/ops/addn.cc
浏览文件 @
62216ad1
...
...
@@ -12,4 +12,6 @@ REGISTER_CPU_OPERATOR(AddN, AddNOp<DeviceType::CPU, float>);
REGISTER_NEON_OPERATOR
(
AddN
,
AddNOp
<
DeviceType
::
NEON
,
float
>
);
#endif // __ARM_NEON
REGISTER_OPENCL_OPERATOR
(
AddN
,
AddNOp
<
DeviceType
::
OPENCL
,
float
>
);
}
// namespace mace
mace/ops/addn.h
浏览文件 @
62216ad1
...
...
@@ -10,7 +10,7 @@
namespace
mace
{
template
<
DeviceType
D
,
class
T
>
template
<
DeviceType
D
,
class
T
>
class
AddNOp
:
public
Operator
<
D
,
T
>
{
public:
AddNOp
(
const
OperatorDef
&
operator_def
,
Workspace
*
ws
)
...
...
@@ -19,16 +19,13 @@ class AddNOp : public Operator<D, T> {
bool
Run
()
override
{
Tensor
*
output_tensor
=
this
->
outputs_
[
0
];
output_tensor
->
ResizeLike
(
this
->
inputs_
[
0
]);
T
*
output
=
output_tensor
->
mutable_data
<
T
>
();
index_t
size
=
this
->
inputs_
[
0
]
->
size
();
int
n
=
this
->
inputs_
.
size
();
vector
<
const
T
*>
inputs
(
n
);
vector
<
const
T
ensor
*>
inputs
(
n
,
nullptr
);
for
(
int
i
=
0
;
i
<
n
;
++
i
)
{
const
Tensor
*
input_tensor
=
this
->
inputs_
[
i
];
inputs
[
i
]
=
input_tensor
->
data
<
T
>
();
inputs
[
i
]
=
this
->
inputs_
[
i
];
}
functor_
(
inputs
,
output
,
size
);
functor_
(
inputs
,
output
_tensor
);
return
true
;
}
...
...
mace/ops/addn_test.cc
浏览文件 @
62216ad1
...
...
@@ -9,9 +9,44 @@ namespace mace {
class
AddnOpTest
:
public
OpsTestBase
{};
TEST_F
(
AddnOpTest
,
AddnOp
)
{
template
<
DeviceType
D
>
void
SimpleAdd2
()
{
// Construct graph
auto
&
net
=
test_net
();
OpsTestNet
net
;
OpDefBuilder
(
"AddN"
,
"AddNTest"
)
.
Input
(
"Input1"
)
.
Input
(
"Input2"
)
.
Output
(
"Output"
)
.
Finalize
(
net
.
NewOperatorDef
());
// Add input data
net
.
AddInputFromArray
<
D
,
float
>
(
"Input1"
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
});
net
.
AddInputFromArray
<
D
,
float
>
(
"Input2"
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
});
// Run
net
.
RunOp
(
D
);
auto
expected
=
CreateTensor
<
float
>
({
1
,
1
,
2
,
3
},
{
2
,
4
,
6
,
8
,
10
,
12
});
ExpectTensorNear
<
float
>
(
*
expected
,
*
net
.
GetOutput
(
"Output"
),
1e-5
);
}
TEST_F
(
AddnOpTest
,
CPUSimpleAdd2
)
{
SimpleAdd2
<
DeviceType
::
CPU
>
();
}
TEST_F
(
AddnOpTest
,
NEONSimpleAdd2
)
{
SimpleAdd2
<
DeviceType
::
NEON
>
();
}
TEST_F
(
AddnOpTest
,
OPENCLSimpleAdd2
)
{
SimpleAdd2
<
DeviceType
::
OPENCL
>
();
}
template
<
DeviceType
D
>
void
SimpleAdd3
()
{
// Construct graph
OpsTestNet
net
;
OpDefBuilder
(
"AddN"
,
"AddNTest"
)
.
Input
(
"Input1"
)
.
Input
(
"Input2"
)
...
...
@@ -20,20 +55,62 @@ TEST_F(AddnOpTest, AddnOp) {
.
Finalize
(
net
.
NewOperatorDef
());
// Add input data
net
.
Add
RandomInput
<
DeviceType
::
CPU
,
float
>
(
"Input1"
,
{
1
,
2
,
3
,
4
});
net
.
Add
RandomInput
<
DeviceType
::
CPU
,
float
>
(
"Input2"
,
{
1
,
2
,
3
,
4
});
net
.
Add
RandomInput
<
DeviceType
::
CPU
,
float
>
(
"Input3"
,
{
1
,
2
,
3
,
4
});
net
.
Add
InputFromArray
<
D
,
float
>
(
"Input1"
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
});
net
.
Add
InputFromArray
<
D
,
float
>
(
"Input2"
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
});
net
.
Add
InputFromArray
<
D
,
float
>
(
"Input3"
,
{
1
,
1
,
2
,
3
},
{
1
,
2
,
3
,
4
,
5
,
6
});
// Run
net
.
RunOp
();
net
.
RunOp
(
D
);
auto
expected
=
CreateTensor
<
float
>
({
1
,
1
,
2
,
3
},
{
3
,
6
,
9
,
12
,
15
,
18
});
ExpectTensorNear
<
float
>
(
*
expected
,
*
net
.
GetOutput
(
"Output"
),
1e-5
);
}
TEST_F
(
AddnOpTest
,
CPUSimpleAdd3
)
{
SimpleAdd3
<
DeviceType
::
CPU
>
();
}
Tensor
expected
;
expected
.
Copy
(
*
net
.
GetOutput
(
"Output"
));
TEST_F
(
AddnOpTest
,
NEONSimpleAdd3
)
{
SimpleAdd3
<
DeviceType
::
NEON
>
();
}
template
<
DeviceType
D
>
void
RandomTest
()
{
// Construct graph
OpsTestNet
net
;
OpDefBuilder
(
"AddN"
,
"AddNTest"
)
.
Input
(
"Input1"
)
.
Input
(
"Input2"
)
.
Output
(
"Output"
)
.
Finalize
(
net
.
NewOperatorDef
());
// Add input data
net
.
AddRandomInput
<
D
,
float
>
(
"Input1"
,
{
1
,
2
,
3
,
4
});
net
.
AddRandomInput
<
D
,
float
>
(
"Input2"
,
{
1
,
2
,
3
,
4
});
// Check
net
.
RunOp
(
DeviceType
::
NEON
);
net
.
RunOp
(
D
);
Tensor
result
;
result
.
Copy
(
*
net
.
GetOutput
(
"Output"
));
// Run
net
.
RunOp
();
ExpectTensorNear
<
float
>
(
*
net
.
GetOutput
(
"Output"
),
result
,
1e-5
);
}
TEST_F
(
AddnOpTest
,
CPURandom
)
{
RandomTest
<
DeviceType
::
CPU
>
();
}
TEST_F
(
AddnOpTest
,
NEONRandom
)
{
RandomTest
<
DeviceType
::
NEON
>
();
}
ExpectTensorNear
<
float
>
(
expected
,
*
net
.
GetOutput
(
"Output"
),
0.01
);
TEST_F
(
AddnOpTest
,
OPENCLRandom
)
{
RandomTest
<
DeviceType
::
OPENCL
>
();
}
}
// namespace mace
mace/python/tools/tf_converter_lib.py
浏览文件 @
62216ad1
...
...
@@ -12,6 +12,33 @@ pooling_type_mode = {
'MaxPool'
:
2
}
def
convert_tensor
(
op
,
tensor
):
tf_tensor
=
op
.
outputs
[
0
].
eval
()
tensor
.
name
=
op
.
outputs
[
0
].
name
shape
=
list
(
tf_tensor
.
shape
)
if
(
op
.
name
.
find
(
'pointwise_kernel'
)
!=
-
1
or
op
.
name
.
find
(
'depthwise_kernel'
)
!=
-
1
or
op
.
name
.
endswith
(
'weights'
)
or
op
.
name
.
endswith
(
'kernel'
))
\
and
op
.
outputs
[
0
].
consumers
()[
0
].
type
.
find
(
'Conv'
)
!=
-
1
:
if
op
.
outputs
[
0
].
consumers
()[
0
].
get_attr
(
'data_format'
)
==
'NCHW'
:
tf_tensor
=
np
.
transpose
(
tf_tensor
,
axes
=
(
3
,
2
,
0
,
1
))
shape
=
[
shape
[
3
],
shape
[
2
],
shape
[
0
],
shape
[
1
]]
# print (tensor.name, shape)
tensor
.
dims
.
extend
(
shape
)
tf_dt
=
op
.
get_attr
(
'dtype'
)
if
tf_dt
==
tf
.
float32
:
tensor
.
data_type
=
mace_pb2
.
DT_FLOAT
tensor
.
float_data
.
extend
(
tf_tensor
.
astype
(
float
).
flat
)
elif
tf_dt
==
tf
.
int32
:
tensor
.
data_type
=
mace_pb2
.
DT_INT32
tensor
.
int32_data
.
extend
(
tf_tensor
.
astype
(
np
.
int32
).
flat
)
else
:
raise
Exception
(
"Not supported tensor type: "
+
tf_dt
.
name
)
def
get_input_tensor
(
op
,
index
):
input_tensor
=
op
.
inputs
[
index
]
if
input_tensor
.
op
.
type
==
'Reshape'
:
...
...
@@ -24,26 +51,11 @@ def convert_ops(unresolved_ops, net_def):
first_op
=
unresolved_ops
[
0
]
if
first_op
.
type
==
'Placeholder'
or
first_op
.
type
==
'Reshape'
:
if
first_op
.
type
in
[
'Placeholder'
,
'Reshape'
,
'Identity'
]
:
pass
elif
first_op
.
type
==
'Const'
:
tf_tensor
=
first_op
.
outputs
[
0
].
eval
()
tensor
=
net_def
.
tensors
.
add
()
tensor
.
name
=
first_op
.
outputs
[
0
].
name
# TODO: support other type than float
tensor
.
data_type
=
mace_pb2
.
DT_FLOAT
shape
=
list
(
tf_tensor
.
shape
)
if
(
first_op
.
name
.
find
(
'pointwise_kernel'
)
!=
-
1
or
first_op
.
name
.
find
(
'depthwise_kernel'
)
!=
-
1
or
first_op
.
name
.
endswith
(
'weights'
)
or
first_op
.
name
.
endswith
(
'kernel'
))
\
and
first_op
.
outputs
[
0
].
consumers
()[
0
].
type
.
find
(
'Conv'
)
!=
-
1
:
tf_tensor
=
np
.
transpose
(
tf_tensor
,
axes
=
(
3
,
2
,
0
,
1
))
shape
=
[
shape
[
3
],
shape
[
2
],
shape
[
0
],
shape
[
1
]]
# print (tensor.name, shape)
tensor
.
dims
.
extend
(
shape
)
tensor
.
float_data
.
extend
(
tf_tensor
.
astype
(
float
).
flat
)
convert_tensor
(
first_op
,
tensor
)
elif
first_op
.
type
==
'Conv2D'
or
first_op
.
type
==
'DepthwiseConv2dNative'
:
op_def
=
net_def
.
op
.
add
()
op_def
.
name
=
first_op
.
name
...
...
@@ -61,9 +73,7 @@ def convert_ops(unresolved_ops, net_def):
strides_arg
.
ints
.
extend
(
first_op
.
get_attr
(
'strides'
)[
2
:])
data_format_arg
=
op_def
.
arg
.
add
()
data_format_arg
.
name
=
'data_format'
data_format_arg
.
s
=
first_op
.
get_attr
(
'data_format'
)
if
first_op
.
get_attr
(
'data_format'
)
!=
'NCHW'
:
raise
Exception
(
'only support NCHW now'
)
data_format_arg
.
s
=
'NCHW'
if
ops_count
>=
2
and
unresolved_ops
[
1
].
type
==
'BiasAdd'
:
bias_add_op
=
unresolved_ops
[
1
]
...
...
@@ -78,7 +88,8 @@ def convert_ops(unresolved_ops, net_def):
sub_op
=
unresolved_ops
[
5
]
add_1_op
=
unresolved_ops
[
6
]
# print (mul_op.type, mul_2_op.type, mul_1_op.type, sub_op.type)
if
mul_op
.
type
!=
'Mul'
or
mul_2_op
.
type
!=
'Mul'
or
mul_1_op
.
type
!=
'Mul'
or
sub_op
.
type
!=
'Sub'
or
add_1_op
.
type
!=
'Add'
:
if
mul_op
.
type
!=
'Mul'
or
mul_2_op
.
type
!=
'Mul'
or
\
mul_1_op
.
type
!=
'Mul'
or
sub_op
.
type
!=
'Sub'
or
add_1_op
.
type
!=
'Add'
:
raise
Exception
(
'Invalid BatchNorm Op'
)
input_name
=
get_input_tensor
(
mul_1_op
,
0
).
name
...
...
@@ -104,12 +115,6 @@ def convert_ops(unresolved_ops, net_def):
max_limit_arg
=
op_def
.
arg
.
add
()
max_limit_arg
.
name
=
'max_limit'
max_limit_arg
.
f
=
6
elif
first_op
.
type
==
'Relu'
:
op_def
=
net_def
.
op
.
add
()
op_def
.
name
=
first_op
.
name
op_def
.
type
=
first_op
.
type
op_def
.
input
.
extend
([
input
.
name
for
input
in
first_op
.
inputs
])
op_def
.
output
.
extend
([
output
.
name
for
output
in
first_op
.
outputs
])
elif
first_op
.
type
==
'AvgPool'
or
first_op
.
type
==
'MaxPool'
:
op_def
=
net_def
.
op
.
add
()
op_def
.
name
=
first_op
.
name
...
...
@@ -130,9 +135,19 @@ def convert_ops(unresolved_ops, net_def):
kernels_arg
.
ints
.
extend
(
first_op
.
get_attr
(
'ksize'
)[
2
:])
data_format_arg
=
op_def
.
arg
.
add
()
data_format_arg
.
name
=
'data_format'
data_format_arg
.
s
=
first_op
.
get_attr
(
'data_format'
)
if
first_op
.
get_attr
(
'data_format'
)
!=
'NCHW'
:
raise
Exception
(
'only support NCHW now'
)
data_format_arg
.
s
=
'NCHW'
elif
first_op
.
type
==
'Add'
:
op_def
=
net_def
.
op
.
add
()
op_def
.
name
=
first_op
.
name
op_def
.
type
=
"AddN"
op_def
.
input
.
extend
([
input
.
name
for
input
in
first_op
.
inputs
])
op_def
.
output
.
extend
([
output
.
name
for
output
in
first_op
.
outputs
])
elif
first_op
.
type
in
[
'Relu'
,
'ResizeBilinear'
,
'SpaceToBatchND'
,
'BatchToSpaceND'
]:
op_def
=
net_def
.
op
.
add
()
op_def
.
name
=
first_op
.
name
op_def
.
type
=
first_op
.
type
op_def
.
input
.
extend
([
input
.
name
for
input
in
first_op
.
inputs
])
op_def
.
output
.
extend
([
output
.
name
for
output
in
first_op
.
outputs
])
else
:
raise
Exception
(
'Unknown Op: '
+
first_op
.
name
)
pass
...
...
@@ -152,4 +167,6 @@ def convert_to_mace_pb(input_graph_def):
while
len
(
unresolved_ops
)
>
0
:
convert_ops
(
unresolved_ops
,
net_def
)
print
"Done."
return
net_def
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录