Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
慢慢CG
Mace
提交
00b70758
Mace
项目概览
慢慢CG
/
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看板
提交
00b70758
编写于
3月 10, 2018
作者:
L
liuqi
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Fix code style problem.
上级
b9d8e771
变更
3
隐藏空白更改
内联
并排
Showing
3 changed file
with
16 addition
and
10 deletion
+16
-10
mace/kernels/opencl/cl/slice.cl
mace/kernels/opencl/cl/slice.cl
+4
-2
mace/kernels/opencl/slice.cc
mace/kernels/opencl/slice.cc
+10
-7
mace/kernels/slice.h
mace/kernels/slice.h
+2
-1
未找到文件。
mace/kernels/opencl/cl/slice.cl
浏览文件 @
00b70758
...
...
@@ -8,6 +8,8 @@ __kernel void slice(__read_only image2d_t input,
const
int
width
=
get_global_size
(
1
)
;
const
int
hb_idx
=
get_global_id
(
2
)
;
DATA_TYPE4
data
=
READ_IMAGET
(
input,
SAMPLER,
(
int2
)(
mad24
(
chan_blk_idx
+
chan_blk_offset,
width,
width_idx
)
,
hb_idx
))
;
WRITE_IMAGET
(
output,
(
int2
)(
mad24
(
chan_blk_idx,
width,
width_idx
)
,
hb_idx
)
,
data
)
;
(
int2
)(
mad24
(
chan_blk_idx
+
chan_blk_offset,
width,
width_idx
)
,
hb_idx
))
;
WRITE_IMAGET
(
output,
(
int2
)(
mad24
(
chan_blk_idx,
width,
width_idx
)
,
hb_idx
)
,
data
)
;
}
mace/kernels/opencl/slice.cc
浏览文件 @
00b70758
...
...
@@ -11,13 +11,15 @@ namespace mace {
namespace
kernels
{
template
<
typename
T
>
void
SliceFunctor
<
DeviceType
::
OPENCL
,
T
>::
operator
()(
const
Tensor
*
input
,
const
std
::
vector
<
Tensor
*>
&
output_list
,
StatsFuture
*
future
)
{
void
SliceFunctor
<
DeviceType
::
OPENCL
,
T
>::
operator
()(
const
Tensor
*
input
,
const
std
::
vector
<
Tensor
*>
&
output_list
,
StatsFuture
*
future
)
{
const
index_t
input_channels
=
input
->
dim
(
3
);
const
size_t
outputs_count
=
output_list
.
size
();
const
index_t
output_channels
=
input_channels
/
outputs_count
;
MACE_CHECK
(
output_channels
%
4
==
0
)
<<
"output channels of slice op must be divisible by 4"
;
MACE_CHECK
(
output_channels
%
4
==
0
)
<<
"output channels of slice op must be divisible by 4"
;
std
::
vector
<
index_t
>
output_shape
({
input
->
dim
(
0
),
input
->
dim
(
1
),
input
->
dim
(
2
),
output_channels
});
...
...
@@ -33,7 +35,8 @@ void SliceFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
std
::
string
kernel_name
=
MACE_OBFUSCATE_SYMBOL
(
"slice"
);
built_options
.
emplace
(
"-Dslice="
+
kernel_name
);
built_options
.
emplace
(
"-DDATA_TYPE="
+
DtToCLDt
(
DataTypeToEnum
<
T
>::
value
));
built_options
.
emplace
(
"-DCMD_DATA_TYPE="
+
DtToCLCMDDt
(
DataTypeToEnum
<
T
>::
value
));
built_options
.
emplace
(
"-DCMD_DATA_TYPE="
+
DtToCLCMDDt
(
DataTypeToEnum
<
T
>::
value
));
kernel_
=
runtime
->
BuildKernel
(
"slice"
,
kernel_name
,
built_options
);
}
const
index_t
channel_blk
=
RoundUpDiv4
(
output_channels
);
...
...
@@ -53,9 +56,9 @@ void SliceFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
<<
outputs_count
;
for
(
int
i
=
0
;
i
<
outputs_count
;
++
i
)
{
uint32_t
idx
=
0
;
kernel_
.
setArg
(
idx
++
,
*
(
static_cast
<
const
cl
::
Image2D
*>
(
input
->
opencl_image
()
)));
kernel_
.
setArg
(
idx
++
,
*
(
input
->
opencl_image
(
)));
kernel_
.
setArg
(
idx
++
,
static_cast
<
int32_t
>
(
channel_blk
*
i
));
kernel_
.
setArg
(
idx
++
,
*
(
static_cast
<
cl
::
Image2D
*>
(
output_list
[
i
]
->
opencl_image
()
)));
kernel_
.
setArg
(
idx
++
,
*
(
output_list
[
i
]
->
opencl_image
(
)));
TuningOrRun3DKernel
(
kernel_
,
ss
.
str
(),
gws
,
lws
,
future
);
}
...
...
mace/kernels/slice.h
浏览文件 @
00b70758
...
...
@@ -41,7 +41,8 @@ struct SliceFunctor {
int
output_idx
=
outer_idx
*
output_channels
;
for
(
size_t
i
=
0
;
i
<
outputs_count
;
++
i
)
{
if
(
DataTypeCanUseMemcpy
(
DataTypeToEnum
<
T
>::
v
()))
{
memcpy
(
output_ptrs
[
i
]
+
output_idx
,
input_ptr
+
input_idx
,
output_channels
*
sizeof
(
T
));
memcpy
(
output_ptrs
[
i
]
+
output_idx
,
input_ptr
+
input_idx
,
output_channels
*
sizeof
(
T
));
}
else
{
for
(
index_t
k
=
0
;
k
<
output_channels
;
++
k
)
{
*
(
output_ptrs
[
i
]
+
output_idx
+
k
)
=
*
(
input_ptr
+
input_idx
+
k
);
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录