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

[复制链接] · 发表于 2019-11-4 10:54:23

本帖最后由 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上的执行方式和内置算子是否相同？会不会带来额外的性能开销？
问题比较多，见谅。

只看该作者 · 发表于 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。

只看该作者 · 发表于 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也不完整，我没法掌握用法。

只看该作者 · 发表于 2019-12-10 09:05:44

您好。我现在也需要elu这个op，但是没有opencl的开发经验，写的很懵圈，不知道可不可以分享一下代码。感谢！