Toybrick

关于自定义elu算子的几个问题

nanami

新手上路

积分
42
楼主
发表于 2019-11-4 10:54:23    查看: 10164|回复: 3 | [复制链接]    打印 | 只看该作者
本帖最后由 nanami 于 2019-11-4 11:07 编辑

问题描述:
因为我的tensorflow模型需要elu算子,rknn好像不支持,所以我在尝试自定义一个算子。因为想使用NPU加速elu的运算,所以我想写一下OpenVX的kernel。
我参考了rknn-toolkit-V1.2.1里面的example/custom_op/resize_area/,遇到了几个问题:
1)在vx_kernel_initializer函数里面,如何设置globalWorkSize?我做过OpenCL,我的理解是这个应该根据input tensor的total_bytes动态计算,是否正确?我看到rknn_kernel_resizearea.c里面把globalWorkSize设置成固定的值64x64了,这个是否正确?
目前我是这么写的:
shaderParam.workdim = 1;
shaderParam.globalWorkSize[0]   = in_tensor_attr.total_bytes / sizeof(float);
这样写是否正确?
2)关于VXC_ReadImage2DArray API的使用我参考了文档Rockchip_Developer_Guide_RKNN_Toolkit_Custom_OP_CN.pdf,但是还是没搞懂。我想每次读取4个float32做转化,Coord坐标参数该如何设置?image的坐标是如何对应到tensor数据的?
我看到rknn_kernel_resizearea.vx里面这样读取数据:
int4 dst_coord = (int4)(0, get_global_id(0), get_global_id(1), 0);
dst_coord的第一个和第四个值为什么是0?
3)如果我的模型使用混合量化,没有对elu层做量化,那么vx kernel处理的是不是float32类型的数据?
4)我在PC上把带自定义算子的模型转化成rknn模型文件之后,拷贝到rk3399pro设备上加载,遇到了以下错误:
E RKNNAPI: rknn_init,  recv(MsgLoadAck) fail, -9(ERROR_PIPE) != 368!
rknn_init fail! ret=-3
如何debug这个错误?rk3399pro的固件是1.5版本,rknn API版本是 1.2.1。
5)自定义算子在NPU上的执行方式和内置算子是否相同?会不会带来额外的性能开销?
问题比较多,见谅。



回复

使用道具 举报

jefferyzhang

版主

积分
13580
沙发
发表于 2019-11-4 11:27:36 | 只看该作者
1、2 太具体的问题我没能力回答,可以自行google学习openVX或者等其他高手回答。
3. 是的,是float32的类型
4. 你可以先运行下demo的自定义算子程序看下能否正常执行,也可以用c语言写自定义op尝试
5. 如果你写的是openVX程序,大概率是会运行到NPU上;如果你写的是c版程序,是一定会运行在CPU上。
至于哪个快,不一定NPU都会快于CPU,NPU的IP核如果没有支持你定义的特定运算的话,实际通用运算速度还不如CPU。
回复

使用道具 举报

nanami

新手上路

积分
42
板凳
 楼主| 发表于 2019-11-6 22:08:33 | 只看该作者
谢谢解答,能否把1和2转给你们NPU部门的人?因为rknn api 1.2.1提供的custom_op demo代码功能不完整,Custom_OP的文档也没有解释清楚如何使用VXC_ReadImage2DArray API读取float32数据,这个不是OpenVX的API,是你们自己写的,文档不全,demo也不完整,我没法掌握用法。
回复

使用道具 举报

linpanda

注册会员

积分
135
地板
发表于 2019-12-10 09:05:44 | 只看该作者
您好。我现在也需要elu这个op,但是没有opencl的开发经验,写的很懵圈,不知道可不可以分享一下代码。感谢!
回复

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

产品中心 购买渠道 开源社区 Wiki教程 资料下载 关于Toybrick


快速回复 返回顶部 返回列表