Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Xiaomi
Mace
提交
b77d113c
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,发现更多精彩内容 >>
提交
b77d113c
编写于
3月 22, 2018
作者:
L
Liangliang He
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Diasble OpenCL symbol loading error checking and fix kernels syntax error
上级
50f5ab5e
变更
4
显示空白变更内容
内联
并排
Showing
4 changed file
with
17 addition
and
17 deletion
+17
-17
mace/core/runtime/opencl/opencl_wrapper.cc
mace/core/runtime/opencl/opencl_wrapper.cc
+9
-10
mace/kernels/opencl/cl/common.h
mace/kernels/opencl/cl/common.h
+4
-4
mace/kernels/opencl/cl/fully_connected.cl
mace/kernels/opencl/cl/fully_connected.cl
+3
-2
mace/kernels/opencl/cl/softmax.cl
mace/kernels/opencl/cl/softmax.cl
+1
-1
未找到文件。
mace/core/runtime/opencl/opencl_wrapper.cc
浏览文件 @
b77d113c
...
@@ -294,9 +294,8 @@ void *OpenCLLibraryImpl::LoadFromPath(const std::string &path) {
...
@@ -294,9 +294,8 @@ void *OpenCLLibraryImpl::LoadFromPath(const std::string &path) {
do { \
do { \
void *ptr = dlsym(handle, #func); \
void *ptr = dlsym(handle, #func); \
if (ptr == nullptr) { \
if (ptr == nullptr) { \
LOG(ERROR) << "Failed to load " << #func << " from " << path; \
VLOG(1) << "Failed to load " << #func << " from " << path; \
dlclose(handle); \
continue; \
return nullptr; \
} \
} \
func = reinterpret_cast<func##Func>(ptr); \
func = reinterpret_cast<func##Func>(ptr); \
VLOG(2) << "Loaded " << #func << " from " << path; \
VLOG(2) << "Loaded " << #func << " from " << path; \
...
...
mace/kernels/opencl/cl/common.h
浏览文件 @
b77d113c
...
@@ -28,19 +28,19 @@ inline DATA_TYPE4 do_activation(DATA_TYPE4 in,
...
@@ -28,19 +28,19 @@ inline DATA_TYPE4 do_activation(DATA_TYPE4 in,
__private
const
float
relux_max_limit
)
{
__private
const
float
relux_max_limit
)
{
DATA_TYPE4
out
;
DATA_TYPE4
out
;
#ifdef USE_RELU
#ifdef USE_RELU
out
=
fmax
(
in
,
0
);
out
=
fmax
(
in
,
(
DATA_TYPE
)
0
);
#endif
#endif
#ifdef USE_RELUX
#ifdef USE_RELUX
out
=
clamp
(
in
,
0
,
relux_max_limit
);
out
=
clamp
(
in
,
(
DATA_TYPE4
)
0
,
relux_max_limit
);
#endif
#endif
#ifdef USE_PRELU
#ifdef USE_PRELU
out
=
select
(
prelu_alpha
*
in
,
in
,
in
>=
0
);
out
=
select
(
prelu_alpha
*
in
,
in
,
in
>=
(
DATA_TYPE
)
0
);
#endif
#endif
#ifdef USE_TANH
#ifdef USE_TANH
out
=
tanh
(
in
);
out
=
tanh
(
in
);
#endif
#endif
#ifdef USE_SIGMOID
#ifdef USE_SIGMOID
out
=
native_recip
(
1
.
0
+
native_exp
(
-
in
));
out
=
native_recip
(
(
DATA_TYPE
)
1
+
native_exp
(
-
in
));
#endif
#endif
return
out
;
return
out
;
}
}
...
...
mace/kernels/opencl/cl/fully_connected.cl
浏览文件 @
b77d113c
...
@@ -81,7 +81,7 @@ __kernel void fully_connected_width(__read_only image2d_t input,
...
@@ -81,7 +81,7 @@ __kernel void fully_connected_width(__read_only image2d_t input,
int2 input_coord, weight_coord;
int2 input_coord, weight_coord;
DATA_TYPE4 in, w;
DATA_TYPE4 in, w;
DATA_TYPE sum = 0
.0
;
DATA_TYPE sum = 0;
input_coord = (int2)(0, mul24(batch_idx, input_height));
input_coord = (int2)(0, mul24(batch_idx, input_height));
...
@@ -107,7 +107,8 @@ __kernel void fully_connected_width(__read_only image2d_t input,
...
@@ -107,7 +107,8 @@ __kernel void fully_connected_width(__read_only image2d_t input,
input_coord.y++;
input_coord.y++;
}
}
const short inter_out_offset = mad24(get_local_id(1), 4, get_local_id(0));
const short inter_out_offset = mad24((short)get_local_id(1), (short)4,
(short)get_local_id(0));
const short local_width_blk_size = (short)get_local_size(1);
const short local_width_blk_size = (short)get_local_size(1);
const short local_size = mul24((short)get_local_size(0),
const short local_size = mul24((short)get_local_size(0),
local_width_blk_size);
local_width_blk_size);
...
...
mace/kernels/opencl/cl/softmax.cl
浏览文件 @
b77d113c
...
@@ -12,7 +12,7 @@ __kernel void softmax(__read_only image2d_t input,
...
@@ -12,7 +12,7 @@ __kernel void softmax(__read_only image2d_t input,
int
pos
=
width_idx
;
int
pos
=
width_idx
;
DATA_TYPE
max_value
=
-FLT_MAX
;
DATA_TYPE
max_value
=
-FLT_MAX
;
DATA_TYPE
sum
=
0
.0
;
DATA_TYPE
sum
=
0
;
DATA_TYPE4
data
;
DATA_TYPE4
data
;
for
(
short
i
=
0
; i < chan_blks; ++i) {
for
(
short
i
=
0
; i < chan_blks; ++i) {
data
=
READ_IMAGET
(
input,
SAMPLER,
(
int2
)(
pos,
hb_idx
))
;
data
=
READ_IMAGET
(
input,
SAMPLER,
(
int2
)(
pos,
hb_idx
))
;
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录