Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Xiaomi
Mace
提交
c1d3bcdc
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,发现更多精彩内容 >>
提交
c1d3bcdc
编写于
4月 04, 2018
作者:
Y
yejianwu
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
move macro to inline function for check_out_of_range_for_image2d
上级
ebb979dc
变更
23
隐藏空白更改
内联
并排
Showing
23 changed file
with
22 addition
and
116 deletion
+22
-116
mace/kernels/opencl/cl/activation.cl
mace/kernels/opencl/cl/activation.cl
+1
-2
mace/kernels/opencl/cl/addn.cl
mace/kernels/opencl/cl/addn.cl
+0
-2
mace/kernels/opencl/cl/batch_norm.cl
mace/kernels/opencl/cl/batch_norm.cl
+1
-2
mace/kernels/opencl/cl/bias_add.cl
mace/kernels/opencl/cl/bias_add.cl
+1
-2
mace/kernels/opencl/cl/buffer_to_image.cl
mace/kernels/opencl/cl/buffer_to_image.cl
+1
-15
mace/kernels/opencl/cl/channel_shuffle.cl
mace/kernels/opencl/cl/channel_shuffle.cl
+1
-3
mace/kernels/opencl/cl/common.h
mace/kernels/opencl/cl/common.h
+2
-0
mace/kernels/opencl/cl/concat.cl
mace/kernels/opencl/cl/concat.cl
+2
-4
mace/kernels/opencl/cl/conv_2d.cl
mace/kernels/opencl/cl/conv_2d.cl
+0
-8
mace/kernels/opencl/cl/conv_2d_1x1.cl
mace/kernels/opencl/cl/conv_2d_1x1.cl
+0
-8
mace/kernels/opencl/cl/conv_2d_3x3.cl
mace/kernels/opencl/cl/conv_2d_3x3.cl
+0
-10
mace/kernels/opencl/cl/cwise.cl
mace/kernels/opencl/cl/cwise.cl
+1
-2
mace/kernels/opencl/cl/depth_to_space.cl
mace/kernels/opencl/cl/depth_to_space.cl
+2
-4
mace/kernels/opencl/cl/depthwise_conv2d.cl
mace/kernels/opencl/cl/depthwise_conv2d.cl
+0
-16
mace/kernels/opencl/cl/eltwise.cl
mace/kernels/opencl/cl/eltwise.cl
+1
-2
mace/kernels/opencl/cl/fully_connected.cl
mace/kernels/opencl/cl/fully_connected.cl
+2
-4
mace/kernels/opencl/cl/matmul.cl
mace/kernels/opencl/cl/matmul.cl
+0
-8
mace/kernels/opencl/cl/pooling.cl
mace/kernels/opencl/cl/pooling.cl
+0
-2
mace/kernels/opencl/cl/resize_bilinear.cl
mace/kernels/opencl/cl/resize_bilinear.cl
+1
-2
mace/kernels/opencl/cl/slice.cl
mace/kernels/opencl/cl/slice.cl
+1
-4
mace/kernels/opencl/cl/softmax.cl
mace/kernels/opencl/cl/softmax.cl
+1
-2
mace/kernels/opencl/cl/space_to_batch.cl
mace/kernels/opencl/cl/space_to_batch.cl
+2
-4
mace/kernels/opencl/cl/winograd_transform.cl
mace/kernels/opencl/cl/winograd_transform.cl
+2
-10
未找到文件。
mace/kernels/opencl/cl/activation.cl
浏览文件 @
c1d3bcdc
...
...
@@ -31,8 +31,7 @@ __kernel void activation(KERNEL_ERROR_PARAMS
DATA_TYPE4
out
=
do_activation
(
in,
relux_max_limit
)
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/addn.cl
浏览文件 @
c1d3bcdc
...
...
@@ -32,9 +32,7 @@ __kernel void addn(KERNEL_ERROR_PARAMS
out
=
out
+
in3
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
w,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
w,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/batch_norm.cl
浏览文件 @
c1d3bcdc
...
...
@@ -49,8 +49,7 @@ __kernel void batch_norm(KERNEL_ERROR_PARAMS
out
=
do_activation
(
out,
relux_max_limit
)
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/bias_add.cl
浏览文件 @
c1d3bcdc
...
...
@@ -24,8 +24,7 @@ __kernel void bias_add(KERNEL_ERROR_PARAMS
DATA_TYPE4
bias_value
=
READ_IMAGET
(
bias,
SAMPLER,
(
int2
)(
ch_blk,
0
))
;
DATA_TYPE4
out
=
in
+
bias_value
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/buffer_to_image.cl
浏览文件 @
c1d3bcdc
...
...
@@ -50,9 +50,7 @@ __kernel void filter_buffer_to_image(KERNEL_ERROR_PARAMS
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -160,9 +158,7 @@ __kernel void dw_filter_buffer_to_image(KERNEL_ERROR_PARAMS
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -205,9 +201,7 @@ __kernel void in_out_buffer_to_image(KERNEL_ERROR_PARAMS
values = vload4(0, input + offset);
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -284,9 +278,7 @@ __kernel void arg_buffer_to_image(KERNEL_ERROR_PARAMS
values = vload4(0, input + offset);
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -364,9 +356,7 @@ __kernel void in_out_height_buffer_to_image(KERNEL_ERROR_PARAMS
values.x = *(input + offset);
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -448,9 +438,7 @@ __kernel void in_out_width_buffer_to_image(KERNEL_ERROR_PARAMS
values.x = *(input + offset);
}
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, w, h, kernel_error);
#endif
WRITE_IMAGET(output, coord, values);
}
...
...
@@ -525,9 +513,7 @@ __kernel void winograd_filter_buffer_to_image(KERNEL_ERROR_PARAMS
int2 coord = (int2)(w, h);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, coord.x, coord.y + out_channels * 15, kernel_error);
#endif
check_out_of_range_for_image2d(output, coord.x, coord.y + out_channels * 15, kernel_error);
#pragma unroll
for (short i = 0; i < 4; ++i) {
...
...
mace/kernels/opencl/cl/channel_shuffle.cl
浏览文件 @
c1d3bcdc
...
...
@@ -51,9 +51,7 @@ __kernel void channel_shuffle(KERNEL_ERROR_PARAMS
int
out_x
=
mad24
(
mad24
(
group_chan_blk_idx,
groups,
g_blk
)
,
width,
width_idx
)
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x
+
groups_blks_width
*
3
,
hb_idx,
kernel_error
)
;
#
endif
check_out_of_range_for_image2d
(
output,
out_x
+
groups_blks_width
*
3
,
hb_idx,
kernel_error
)
;
WRITE_IMAGET
(
output,
(
int2
)(
out_x,
hb_idx
)
,
out_chan_data0
)
;
out_x
+=
groups_blks_width
;
...
...
mace/kernels/opencl/cl/common.h
浏览文件 @
c1d3bcdc
...
...
@@ -78,10 +78,12 @@ inline void check_out_of_range_for_image2d(__write_only image2d_t image,
__private
const
int
x
,
__private
const
int
y
,
global
char
*
kernel_error
)
{
#ifdef OUT_OF_RANGE_CHECK
int2
image_dim
=
get_image_dim
(
image
);
if
(
x
>=
image_dim
.
x
||
y
>=
image_dim
.
y
)
{
*
kernel_error
=
'1'
;
}
#endif
}
#endif // MACE_KERNELS_OPENCL_CL_COMMON_H_
mace/kernels/opencl/cl/concat.cl
浏览文件 @
c1d3bcdc
...
...
@@ -82,9 +82,8 @@ __kernel void concat_channel(KERNEL_ERROR_PARAMS
const int pos = mad24(chan_blk_idx, width, width_idx);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(pos, hb_idx), data);
}
...
...
@@ -115,9 +114,8 @@ __kernel void concat_channel_multi(KERNEL_ERROR_PARAMS
const int pos = mad24(chan_blk_idx + chan_blk_offset, width, width_idx);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(pos, hb_idx), data);
}
...
...
mace/kernels/opencl/cl/conv_2d.cl
浏览文件 @
c1d3bcdc
...
...
@@ -129,30 +129,22 @@ __kernel void conv_2d(KERNEL_ERROR_PARAMS
const
int
out_x_base
=
mul24
(
out_ch_blk,
out_width
)
;
int
w
=
out_w_blk
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out0
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out1
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out2
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out3
)
;
}
mace/kernels/opencl/cl/conv_2d_1x1.cl
浏览文件 @
c1d3bcdc
...
...
@@ -105,29 +105,21 @@ __kernel void conv_2d_1x1(KERNEL_ERROR_PARAMS
const
int
out_x_base
=
mul24
(
out_ch_blk,
width
)
;
int
out_x_idx
=
out_w_blk
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
out_x_idx,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
out_x_idx,
out_hb
)
,
out0
)
;
out_x_idx
+=
out_w_blks
;
if
(
out_x_idx
>=
width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
out_x_idx,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
out_x_idx,
out_hb
)
,
out1
)
;
out_x_idx
+=
out_w_blks
;
if
(
out_x_idx
>=
width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
out_x_idx,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
out_x_idx,
out_hb
)
,
out2
)
;
out_x_idx
+=
out_w_blks
;
if
(
out_x_idx
>=
width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
out_x_idx,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
out_x_idx,
out_hb
)
,
out3
)
;
}
mace/kernels/opencl/cl/conv_2d_3x3.cl
浏览文件 @
c1d3bcdc
...
...
@@ -136,45 +136,35 @@ __kernel void conv_2d_3x3(KERNEL_ERROR_PARAMS
const
int
out_x_base
=
mul24
(
out_ch_blk,
out_width
)
;
int
w
=
out_w_blk
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out0
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out1
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out2
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out3
)
;
w
+=
out_w_blks
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out4
)
;
...
...
mace/kernels/opencl/cl/cwise.cl
浏览文件 @
c1d3bcdc
...
...
@@ -44,8 +44,7 @@ __kernel void cwise(KERNEL_ERROR_PARAMS
out.w
=
fabs
(
in0.w
)
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
w,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
w,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/depth_to_space.cl
浏览文件 @
c1d3bcdc
...
...
@@ -37,9 +37,8 @@ __kernel void depth_to_space(KERNEL_ERROR_PARAMS
const int in_pos = mad24(in_d, input_width, in_w);
DATA_TYPE4 in_data = READ_IMAGET(input, SAMPLER, (int2)(in_pos, in_hb));
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_pos, out_hb, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_pos, out_hb), in_data);
}
...
...
@@ -80,8 +79,7 @@ __kernel void space_to_depth(KERNEL_ERROR_PARAMS
const
int
out_pos
=
mad24
(
out_d,
output_width,
out_w
)
;
DATA_TYPE4
in_data
=
READ_IMAGET
(
input,
SAMPLER,
(
int2
)(
in_pos,
hb
))
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_pos,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_pos,
out_hb
)
,
in_data
)
;
}
mace/kernels/opencl/cl/depthwise_conv2d.cl
浏览文件 @
c1d3bcdc
...
...
@@ -123,30 +123,22 @@ __kernel void depthwise_conv2d(KERNEL_ERROR_PARAMS
const short out_x_base = mul24(out_ch_blk, out_width);
short w = out_w_blk;
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0);
w += out_w_blks;
if (w >= out_width) return;
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1);
w += out_w_blks;
if (w >= out_width) return;
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2);
w += out_w_blks;
if (w >= out_width) return;
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3);
}
...
...
@@ -261,29 +253,21 @@ __kernel void depthwise_conv2d_s1(KERNEL_ERROR_PARAMS
const
short
out_x_base
=
mul24
(
out_ch_blk,
out_width
)
;
short
w
=
out_w_blk
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out0
)
;
w
+=
1
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out1
)
;
w
+=
1
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out2
)
;
w
+=
1
;
if
(
w
>=
out_width
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_x_base
+
w,
out_hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_x_base
+
w,
out_hb
)
,
out3
)
;
}
mace/kernels/opencl/cl/eltwise.cl
浏览文件 @
c1d3bcdc
...
...
@@ -37,8 +37,7 @@ __kernel void eltwise(KERNEL_ERROR_PARAMS
out
=
in0
-
in1
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
w,
hb,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
w,
hb
)
,
out
)
;
}
mace/kernels/opencl/cl/fully_connected.cl
浏览文件 @
c1d3bcdc
...
...
@@ -60,9 +60,8 @@ __kernel void fully_connected(KERNEL_ERROR_PARAMS
result = do_activation(result, relux_max_limit);
#endif
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_blk_idx, batch_idx, kernel_error);
#endif
WRITE_IMAGET(output, (int2)(out_blk_idx, batch_idx), result);
}
...
...
@@ -154,9 +153,8 @@ __kernel void fully_connected_width(KERNEL_ERROR_PARAMS
result
=
do_activation
(
result,
relux_max_limit
)
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_blk_idx,
batch_idx,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_blk_idx,
batch_idx
)
,
result
)
;
}
}
mace/kernels/opencl/cl/matmul.cl
浏览文件 @
c1d3bcdc
...
...
@@ -48,26 +48,18 @@ __kernel void matmul(KERNEL_ERROR_PARAMS
c3
+=
(
DATA_TYPE4
)(
dot
(
a0,
b3
)
,
dot
(
a1,
b3
)
,
dot
(
a2,
b3
)
,
dot
(
a3,
b3
))
;
}
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
C,
gx,
gy,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
C,
(
int2
)(
gx,
gy
)
,
c0
)
;
if
((
gx
+
1
)
>=
N
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
C,
gx
+
1
,
gy,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
C,
(
int2
)(
gx
+
1
,
gy
)
,
c1
)
;
if
((
gx
+
2
)
>=
N
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
C,
gx
+
2
,
gy,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
C,
(
int2
)(
gx
+
2
,
gy
)
,
c2
)
;
if
((
gx
+
3
)
>=
N
)
return
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
C,
gx
+
3
,
gy,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
C,
(
int2
)(
gx
+
3
,
gy
)
,
c3
)
;
}
mace/kernels/opencl/cl/pooling.cl
浏览文件 @
c1d3bcdc
...
...
@@ -96,8 +96,6 @@ __kernel void pooling(KERNEL_ERROR_PARAMS
#
endif
const
int
pos
=
mad24
(
out_chan_idx,
out_width,
out_width_idx
)
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
out_hb_idx,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
out_hb_idx
)
,
res
)
;
}
mace/kernels/opencl/cl/resize_bilinear.cl
浏览文件 @
c1d3bcdc
...
...
@@ -58,9 +58,8 @@ __kernel void resize_bilinear_nocache(KERNEL_ERROR_PARAMS
const
int
out_w_offset
=
mul24
(
ch_blk,
out_width
)
;
const
int
out_h_offset
=
mul24
(
b,
out_height
)
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
out_w_offset
+
w,
out_h_offset
+
h,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
out_w_offset
+
w,
out_h_offset
+
h
)
,
out
)
;
}
mace/kernels/opencl/cl/slice.cl
浏览文件 @
c1d3bcdc
...
...
@@ -24,9 +24,6 @@ __kernel void slice(KERNEL_ERROR_PARAMS
width,
width_idx
)
,
hb_idx
))
;
const
int
pos
=
mad24
(
chan_blk_idx,
width,
width_idx
)
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
hb_idx,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb_idx
)
,
data
)
;
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb_idx
)
,
data
)
;
}
mace/kernels/opencl/cl/softmax.cl
浏览文件 @
c1d3bcdc
...
...
@@ -85,8 +85,7 @@ __kernel void softmax(KERNEL_ERROR_PARAMS
data
=
native_exp
(
data
)
/
sum
;
}
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
pos,
hb_idx,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
pos,
hb_idx
)
,
data
)
;
}
mace/kernels/opencl/cl/space_to_batch.cl
浏览文件 @
c1d3bcdc
...
...
@@ -46,9 +46,8 @@ __kernel void space_to_batch(KERNEL_ERROR_PARAMS
int2 batch_coord = (int2)(mul24(chan_idx, batch_width) + batch_w_idx, batch_hb_idx);
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(batch_data, batch_coord.x, batch_coord.y, kernel_error);
#endif
WRITE_IMAGET(batch_data, batch_coord, value);
}
...
...
@@ -94,9 +93,8 @@ __kernel void batch_to_space(KERNEL_ERROR_PARAMS
int2
space_coord
=
(
int2
)(
mul24
(
chan_idx,
space_width
)
+
space_w_idx,
space_b_idx
*
space_height
+
space_h_idx
)
;
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
space_data,
space_coord.x,
space_coord.y,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
space_data,
space_coord,
value
)
;
}
}
mace/kernels/opencl/cl/winograd_transform.cl
浏览文件 @
c1d3bcdc
...
...
@@ -94,9 +94,8 @@ __kernel void winograd_transform_2x2(KERNEL_ERROR_PARAMS
input3[2] = tv3[2] - tv3[1];
input3[3] = tv3[1] - tv3[3];
#ifdef OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d(output, out_width_idx, chan_blk_idx + chan_blk_idx * 15, kernel_error);
#endif
#pragma unroll
for (short i = 0; i < 4; ++i) {
WRITE_IMAGET(output, (int2)(out_width_idx, chan_blk_idx), input0[i]);
...
...
@@ -213,30 +212,23 @@ __kernel void winograd_inverse_transform_2x2(KERNEL_ERROR_PARAMS
in1[1]
=
do_activation
(
in1[1],
relux_max_limit
)
;
#
endif
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
coord_x,
coord_y,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
coord_x,
coord_y
)
,
in0[0]
)
;
t
=
0
;
if
(
out_width_idx
+
1
<
out_width
)
{
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
coord_x
+
1
,
coord_y,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
coord_x
+
1
,
coord_y
)
,
in0[1]
)
;
t
+=
1
;
}
if
(
out_height_idx
+
1
<
out_height
)
{
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
coord_x,
coord_y
+
1
,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
coord_x,
coord_y
+
1
)
,
in1[0]
)
;
t
+=
1
;
}
if
(
t
==
2
)
{
#
ifdef
OUT_OF_RANGE_CHECK
check_out_of_range_for_image2d
(
output,
coord_x
+
1
,
coord_y
+
1
,
kernel_error
)
;
#
endif
WRITE_IMAGET
(
output,
(
int2
)(
coord_x
+
1
,
coord_y
+
1
)
,
in1[1]
)
;
}
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录