人工智能推理引擎模型量化模型优化边缘计算开发工具【免费下载链接】rknn-toolkit2项目地址https://gitcode.com/gh_mirrors/rk/rknn-toolkit2点击查看免费下载导读本文基于 RKNN-Toolkit2 仓库中 rknn_custom_gpu_op_demo 示例系统讲解如何在 RKNPU2 运行时上为 RKNN 模型实现自定义 GPUOpenCL算子。该示例以替代 ONNX 原生ArgMax算子为核心串联起 RKNN-Toolkit2 侧的自定义算子注册与 RKNN Runtime C API 侧的算子实现、构建、部署与精度验证全流程。读完本文你将掌握如何用 RKNN-Toolkit2 将带自定义算子的 ONNX 模型转换为 RKNN 模型、如何用 C/C OpenCL 编写自定义 GPU 算子的插件库、如何在 AArch64 Linux 与 Android 板端构建并运行以及如何利用自定义算子插件库配合 toolkit2 做精度分析。一、示例概览为什么需要自定义 GPU 算子RKNN 模型默认由 NPU神经网络处理器执行但并非所有算子都能被 NPU 直接支持。当模型中含有 NPU 不支持、或者用户希望绕开原生实现自行优化的算子时RKNN 提供了**自定义算子Custom Op**机制允许开发者在RKNN-Toolkit2工具链侧把 ONNX 模型中的某个算子替换为自定义算子并完成量化、编译、导出 RKNN 模型在RKNN Runtime C API板端侧用 CPU、GPUOpenCL或 PyTorch 实现该算子的实际计算逻辑通过插件库注入运行时。本示例rknn_custom_gpu_op_demo是三者中最典型的 GPU 实现它以ArgMax算子为例将 PaddleSeg 的 PP-HumanSegV2-Mobile 人像分割模型中的原生ArgMax替换为自定义 GPU 算子并在板端用 OpenCL kernel 完成通道维求最大值的计算最终输出人像 mask 与原图合成的效果图。示例配套的 RKNN 模型已注册 ArgMax 自定义算子预置于以下平台目录中平台模型路径相对仓库根目录RK3562model/RK3562/pp_human_segv2_custom_argmax.rknnRK3566 / RK3568model/RK3566_RK3568/pp_human_segv2_custom_argmax.rknnRK3576model/RK3576/pp_human_segv2_custom_argmax.rknnRK3588model/RK3588/pp_human_segv2_custom_argmax.rknn测试输入图片为 model/test_image.jpg513×513 人像图。1.1 模型来源PP-HumanSegV2-Mobile示例模型来自 PaddlePaddle 的 PaddleSeg 开源项目PP-HumanSeg 人像分割套件使用的具体模型为PP-HumanSegV2-Mobile192×192 输入。原模型是 Paddle 格式需要先通过 Paddle2onnx 转换为 ONNX 格式再进入 RKNN 转换流程。转换脚本位于 model/convert/test.py其预期输出为result.jpg包含输入图、分割 mask 图、合成叠加图与类别图例四联画效果图可参考 model/convert/result_truth.jpg。1.2 相关文档与配套示例模型转换与自定义算子注册的详细说明见 model/convert/README.md如何在 RKNN-Toolkit2 中替换 ONNX 原生算子可参考 rknn-toolkit2/examples/functions/custom_op/non-onnx_standard/test.py该示例以Sigmoid → cstSigmoid为例演示了 ONNX 图编辑与自定义算子注册同系列还有 CPU 版 rknn_custom_cpu_op_demo 与 PyTorch 版 rknn_custom_pytorch_op_demo可对照学习不同后端实现差异。二、工具链侧用 RKNN-Toolkit2 转换带自定义算子的模型板端运行前需要先用 RKNN-Toolkit2 生成注册了自定义 ArgMax 算子的 RKNN 模型。转换脚本为 model/convert/test.py核心流程如下初始化 RKNN 对象注册自定义算子 ArgMaxrknn.reg_custom_op将 ONNX 模型转换为 RKNN 模型在模拟器中推理验证保存结果图片释放 RKNN 资源。脚本关键代码节选from rknn.api import RKNN class ArgMax: op_type ArgMax # 创建 RKNN 对象 rknn RKNN(verboseTrue) onnx_model pp_humansegv2_mobile.onnx rknn_model pp_human_segv2_custom_argmax.rknn DATASET ./dataset.txt QUANTIZE_ON True rknn.config(mean_values[127.5, 127.5, 127.5], std_values[127.5, 127.5, 127.5], quant_img_RGB2BGRFalse, target_platformrk3576) # 注册自定义 ArgMax 算子 ret rknn.reg_custom_op(ArgMax) # 加载 ONNX 模型输入 x尺寸 1x3x192x192 ret rknn.load_onnx(modelonnx_model, inputs[x], input_size_list[[1,3,192,192]]) # 量化构建 ret rknn.build(do_quantizationQUANTIZE_ON, datasetDATASET) # 导出 RKNN 模型 ret rknn.export_rknn(rknn_model)关键点说明rknn.config的target_platform脚本默认注释为 rk3576转换前必须按实际目标平台修改支持 rk3562、rk3566/rk3568、rk3576、rk3588 等。模型在推理阶段需要注册了同名自定义算子的运行时环境因此转换侧工具链与运行侧板端插件库的算子类型名必须一致此处均为ArgMaxreg_custom_op(ArgMax)传入一个声明了op_type ArgMax的类。在工具链侧只需声明算子类型名用于模型编译与量化工具链会用默认实现完成量化统计真正的计算逻辑在板端插件库中提供load_onnx指定inputs[x]PP-HumanSegV2-Mobile 的输入 tensor 名为x尺寸[1, 3, 192, 192]NCHW与板端模型输入保持一致量化配置mean_values[127.5, 127.5, 127.5]、std_values[127.5, 127.5, 127.5]、quant_img_RGB2BGRFalse即采用 0~255 归一化到 [-1, 1] 的常见预处理约定且保持 RGB 通道顺序不变。关于“替换 ONNX 原生算子”的更多细节在 non-onnx_standard 示例 中可以看到另一种通用做法——先用onnx.load读入模型遍历model.graph.node将目标op_type如Sigmoid改写为自定义类型名如cstSigmoid再保存为新 ONNX 模型后加载转换自定义算子的 Python 类中还可以通过shape_infer与compute方法提供形状推导和参考实现。三、板端实现OpenCL 自定义 GPU 算子插件库板端自定义 GPU 算子的核心是 src/rknn_custom_op_opencl_plugin_lib.cpp它同时被两个目标使用主程序rknn_custom_gpu_op_demo和独立插件库librkcst_argmax.so从 CMakeLists.txt 可见rk_custom_argmax与rknn_custom_gpu_op_demo共享同一源文件。3.1 自定义算子的注册结构rknn_custom_op根据 rknn_custom_op.h 的定义一个自定义算子需要填充以下关键字段字段说明version自定义算子版本号target后端执行设备RKNN_TARGET_TYPE_CPU 1或RKNN_TARGET_TYPE_GPU 2GPU 即 OpenCL 后端op_type自定义算子类型名最长RKNN_MAX_NAME_LEN必须与模型转换时注册的名字一致如ArgMaxcl_kernel_nameOpenCL kernel 函数名cl_kernel_sourceOpenCL kernel 源码字符串若cl_source_size 0则此字段解释为 kernel 文件路径cl_source_sizekernel 源码长度cl_build_optionsOpenCL 构建 clProgram 的编译选项init[可选] 算子初始化回调prepare[可选] 算子 prepare 回调compute[必选] 算子计算回调GPU 算子的核心逻辑compute_native[可选] 原生属性计算回调当前版本暂不支持建议置空destroy[可选] 算子销毁回调GPU 后端的上下文由rknn_gpu_op_context提供包含cl_context、cl_command_queue、cl_kernel三个句柄由运行时在调用compute前创建并注入rknn_custom_op_context。3.2 插件库入口get_rknn_custom_op插件库必须导出get_rknn_custom_op()函数宏RKNN_CUSTOM_OP_EXPORT保证动态库符号默认可见。示例中的实现RKNN_CUSTOM_OP_EXPORT rknn_custom_op* get_rknn_custom_op() { memset(user_op, 0, sizeof(rknn_custom_op)); char op_type[] ArgMax; strcpy(user_op.op_type, op_type); user_op.version 1; user_op.target RKNN_TARGET_TYPE_GPU; // GPU(OpenCL) 后端 user_op.init custom_init_callback_gpu; user_op.compute compute_custom_gpu_op_float32; user_op.destroy destroy_callback_gpu; char kernel_name[] Argmax_channel_float; strcpy(user_op.cl_kernel_name, kernel_name); user_op.cl_kernel_source cl_kernel_source; user_op.cl_source_size strlen(cl_kernel_source); return user_op; }3.3 OpenCL kernel沿通道求 ArgMax示例内置的 kernel 源码Argmax_channel_float完成“沿 channel 维度求最大值所在索引”的计算输入为 float32输出为每个 (i, j) 像素位置对应的通道索引0 为 background、1 为 person__kernel void Argmax_channel_float(__global const float* src_buf, const int inputChannel, __global float* dst_buf) { unsigned int i get_global_id(0); unsigned int j get_global_id(1); unsigned int width get_global_size(0); unsigned int height get_global_size(1); int max_index 0; float max_value src_buf[i*widthj]; for (int c 1; c inputChannel; c) { float value src_buf[c*height*width i*widthj]; if (value max_value) { max_value value; max_index c; } } dst_buf[i*widthj] max_index; }其中二维工作项分别对应该层平面的宽和高每个工作项独立遍历inputChannel个通道比较后输出通道索引。这与语义分割中“取置信度最高的类别标签”的典型后处理需求完全吻合。3.4compute回调导入 DMA-BUF 并调度 kernelGPU 自定义算子的compute回调rknn_custom_op_opencl_plugin_lib.cpp遵循如下模式从op_ctx-gpu_ctx取出运行时注入的cl_context、cl_command_queue、cl_kernel通过 ARM 扩展clImportMemoryARM配合CL_IMPORT_TYPE_DMA_BUF_ARM将输入/输出 tensor 的 fdinputs[0].mem.fd与偏移mem.offset mem.size导入为cl_mem实现零拷贝访问 RKNN 运行时管理的 DMA 内存根据 tensor 属性计算字节数、偏移、通道数与目标宽高outputs[0].attr.dims[3]为宽、dims[2]为高clSetKernelArg设置 kernel 参数clEnqueueNDRangeKernel以二维global_work_size {dst_width, dst_height}调度 kernelclFinish(queue)等待执行完成返回 0 表示成功。init与destroy回调在示例中为占位实现分别打印argmax_init_callback_gpu/Argmax_destroy_callback_gpu用户可在其中完成资源准备与清理。3.5 主程序加载模型并注册自定义算子主程序 src/rknn_api_test_custom_op_opencl.cpp 展示了一套完整的 RKNN Runtime 调用流程与自定义 GPU 算子直接相关的部分包括uint32_t flag 0; // 当算子不受 NPU 支持时优先回退到 GPU 设备执行 flag flag | RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU; int ret rknn_init(ctx, model_path, 0, flag, NULL); ... // 从插件库导出函数获取自定义算子描述 auto gpu_user_op get_rknn_custom_op(); ret rknn_register_custom_ops(ctx, gpu_user_op, 1);RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU定义于 rknn_api.h含义为“当算子不被 NPU 支持时将 GPU 设为优先回退的执行后端”。在混合模型中NPU 算子走 NPU、自定义算子走 GPU此标志必不可少rknn_register_custom_ops(ctx, op, 1)在rknn_init之后将自定义算子注册到该rknn_context。rknn_custom_op.h同时指出compute回调必选其他可选若init回调返回RKNN_WARNING_SKIP_CUSTOM_OP_COMPUTE-14且该 op 类型被 RKNN 原生支持则运行时将改用 RKNN 原生实现。主程序随后依次完成查询 SDK/驱动版本与内存占用rknn_query、查询输入输出 tensor 属性含 layout 自适应NHWC 取dims[1..3]NCHW 取dims[2..3]作为宽高、stb系列库加载与缩放图片、rknn_inputs_set设置输入、rknn_set_core_mask(RKNN_NPU_CORE_ALL)设置 NPU 核掩码、循环rknn_run测时并输出每帧耗时与 FPS、查询RKNN_QUERY_PERF_DETAIL与RKNN_QUERY_PERF_RUN获取性能明细最后取出输出rknn_outputs_getwant_float 1请求 float 输出并把分割 mask 叠加回原图通过stbi_write_png写出test_compose_img.png。四、构建AArch64 Linux 与 Android示例提供两个构建脚本build-linux.sh 与 build-android.sh二者都基于 CMakeLists.txt 组织工程。4.1 AArch64 Linux 构建先导出交叉编译器前缀再执行构建脚本export GCC_COMPILER~/opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu ./build-linux.sh -t rk3588 -a aarch64 -b Release脚本参数说明摘自 build-linux.sh 的用法提示-t target目标平台合法值为rk3562 / rk3566 / rk3568 / rk3576 / rk3588rk356x、rk3566_rk3568等写法也会被归一化到RK3566_RK3568-a arch架构aarch64或armhf-b build_type构建类型Debug或Release缺省为Release。脚本校验GCC_COMPILER可用性要求${GCC_COMPILER}-gcc可执行随后在build/build_TARGET_SOC_linux_arch_build_type目录执行cmake ../.. -DTARGET_SOC... -DCMAKE_SYSTEM_NAMELinux -DCMAKE_SYSTEM_PROCESSORarch -DCMAKE_C_COMPILER${CC} -DCMAKE_CXX_COMPILER${CXX}最后make -j4 make install。4.2 Android 构建先导出 NDK 路径再执行export ANDROID_NDK_PATH~/opts/ndk/android-ndk-r18b ./build-android.sh -t rk3568 -a arm64-v8a -b Release-a arch在 Android 下取值arm64-v8a或armeabi-v7abuild-android.sh 推荐使用NDK r18 / r19其他版本可能导致构建失败内部通过cmake -DANDROID_TOOLCHAINclang -DCMAKE_TOOLCHAIN_FILE$ANDROID_NDK_PATH/build/cmake/android.toolchain.cmake -DANDROID_ABIarch -DANDROID_STLc_static -DANDROID_PLATFORMandroid-24交叉编译。4.3 CMake 工程要点CMakeLists.txt 揭示了工程组织方式运行时头文件与librknnrt.so取自仓库的 runtime/Linux/librknn_apiAndroid 侧按 ABI 子目录选择依赖仓库自带 3rdparty/opencl/libopencl-stubOpenCL 客户端 stub以及3rdparty下的fp16、stb头文件无需外部安装 OpenCL SDK工程产出三个目标OpenCL stub 动态库、可执行程序rknn_custom_gpu_op_demo、自定义算子插件库rk_custom_argmaxCMAKE_INSTALL_PREFIX指向install/rknn_custom_gpu_op_demo_Linux|Androidmake install会把可执行程序、librknnrt.so到lib/、测试图片、对应平台的 RKNN 模型按TARGET_SOC目录以及rk_custom_argmax到lib/统一装配到 install 目录方便整目录推送到板端。五、部署与运行5.1 AArch64 Linux将install/rknn_custom_gpu_op_demo_Linux目录拷贝到设备/userdata/下。瑞芯微 EVB 开发板可用 adb 推送adb push install/rknn_custom_gpu_op_demo_Linux/ /data/若板子开了 sshd也可用 scp 等方式拷贝程序与 RKNN 模型。随后进入目录运行adb shell cd /data/rknn_custom_gpu_op_demo_Linux export LD_LIBRARY_PATH./lib ./rknn_custom_gpu_op_demo model/TARGET_PLATFORM/pp_human_segv2_custom_argmax.rknn model/test_image.jpg 1命令行参数依次为RKNN 模型路径、输入图片路径、循环推理次数loop_count示例传 1。程序运行后会打印 SDK/驱动版本、tensor 属性、每帧耗时与 FPS并写出合成图test_compose_img.png。5.2 Androidadb push install/rknn_custom_gpu_op_demo_Android/ /data/ adb shell cd /data/rknn_custom_gpu_op_demo_Android/若报错library libOpenCL.so not found: needed by main executable说明设备上 OpenCL 运行库未被加载器找到可将系统自带的libOpenCL.so拷到本地再运行cp /vendor/lib64/libOpenCL.so ./lib/ export LD_LIBRARY_PATH./lib ./rknn_custom_gpu_op_demo model/TARGET_PLATFORM/pp_human_segv2_custom_argmax.rknn model/test_image.jpg 1注意不同平台、不同版本的工具链与驱动推理结果可能略有差异。六、插件库与精度分析Accuracy Analysis若想验证自己实现的自定义算子数值是否准确官方推荐流程是将自定义算子单独编译成插件库安装到板端系统目录再用 RKNN-Toolkit2 的精度分析功能进行对比。6.1 编译插件库插件库源码即前文分析的 src/rknn_custom_op_opencl_plugin_lib.cpp。用与主程序相同的构建体系编译仅需要其中的rk_custom_argmax目标./build-linux_TARGET_PLATFORM.sh # 或 ./build-android_TARGET_PLATFORM.sh6.2 安装插件库将插件库推送到运行时的自定义算子插件目录Linux/usr/lib/rknpu/op_pluginsAndroidarm64-v8a/vendor/lib64/。例如 Linux 系统adb push install/rknn_custom_gpu_op_demo_Linux/lib/librkcst_argmax.so /usr/lib/rknpu/op_plugins注意若板端不存在/usr/lib/rknpu/op_plugins目录需要先手动创建。插件库安装到位后RKNN Runtime 在加载含有对应op_type的模型时会自动 dlopen 该插件并调用其get_rknn_custom_op()获取算子实现从而让板端运行及 toolkit2 的精度分析使用用户自写的 GPU 算子。6.3 执行精度分析完成插件部署后即可使用 RKNN-Toolkit2 提供的精度分析accuracy analysis功能对比板端真实自定义 GPU 算子与模拟器/参考实现的输出确认自定义算子的计算精度满足要求。精度分析相关使用方式可参考 rknn-toolkit2/examples/functions/accuracy_analysis/test.py 及配套 README。七、从源码看 GPU 自定义算子的整体数据流综合 rknn_custom_op.h、rknn_api.h 与示例实现一次完整的 GPU 自定义算子推理数据流可以概括为转换阶段工具链PC 端RKNN-Toolkit2 通过reg_custom_op感知模型中的自定义ArgMax算子完成量化与编译生成pp_human_segv2_custom_argmax.rknn算子信息连同cl_kernel_name等元数据写入模型加载阶段板端rknn_init时传入RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU确保模型编译期无法被 NPU 支持的算子回退到 GPU注册阶段get_rknn_custom_op()导出算子描述目标后端 GPU、OpenCL kernel 名与源码、init/compute/destroy 回调rknn_register_custom_ops注入rknn_context执行阶段rknn_run触发推理遇到ArgMax节点时运行时创建 OpenCL context / command queue / program 并调用compute回调回调通过clImportMemoryARM导入输入输出 tensor 的 DMA-BUF 实现零拷贝调度二维Argmax_channel_floatkernel 计算通道索引后clFinish返回验证阶段主程序取出 float 输出将0person区域在原图中置黑生成合成图用户如需更严谨的精度验证则部署插件库后使用 toolkit2 的 accuracy analysis 对比输出。八、常见问题与注意事项算子类型名必须一致转换脚本reg_custom_op(ArgMax)中的op_type、插件库user_op.op_type、以及模型图中的算子名三者必须完全一致否则运行时无法匹配到自定义实现GPU 回退标志不能遗漏若不设置RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU含 NPU 不支持算子的模型在rknn_init阶段可能直接失败OpenCL 环境依赖Linux 侧通过仓库自带 libopencl-stub 链接触发设备上的 OpenCL ICDAndroid 侧若找不到libOpenCL.so需按 5.2 节拷贝设备自带的运行库插件目录权限与存在性Linux 插件目录/usr/lib/rknpu/op_plugins在部分系统上需要手动创建精度差异不同平台、工具链版本与驱动版本下自定义算子的结果可能有轻微差异属正常现象compute_native当前不可用按 rknn_custom_op.h 说明compute_native暂未支持应置空。九、小结rknn_custom_gpu_op_demo提供了一个“从工具链注册、到 OpenCL kernel 编写、再到板端部署验证”的完整闭环范例。开发者可以以此为模板将自己业务中 NPU 不支持的算子或其他需要 GPU 加速的算子按同样流程接入 RKNN 运行时在 toolkit2 侧声明op_type完成模型转换在板端实现get_rknn_custom_op导出结构与compute回调最后通过插件库机制配合精度分析确认正确性。配套的 CPU 版rknn_custom_cpu_op_demo与 PyTorch 版rknn_custom_pytorch_op_demo示例则可提供不同后端实现的横向参考。赞分享人工智能推理引擎模型量化模型优化边缘计算开发工具【免费下载链接】rknn-toolkit2项目地址https://gitcode.com/gh_mirrors/rk/rknn-toolkit2点击查看免费下载相关推荐Wox Shell 插件完全指南在启动器中直接执行命令、管理历史与工作目录Wox Shell 插件完全指南在启动器中直接执行命令、管理历史与工作目录 本指南围绕 Wox 内置系统插件 Shell 展开讲解如何用 触发关键字在启动器人工智能推理引擎模型量化模型优化边缘计算开发工具CANN PTO 自定义算子开发实战以 fused_add_relu_mul 算子融合为例CANN PTO 自定义算子开发实战以 fused_add_relu_mul 算子融合为例 本篇指南围绕 CANN Parallel Tile Operati算子库人工智能CANNRKNN-Toolkit2 自定义 CPU 算子实战以 cstSigmoid 替换 ONNX Sigmoid 的端到端流程RKNPU2 运行时侧RKNN Toolkit2 自定义 CPU 算子实战以 cstSigmoid 替换 ONNX Sigmoid 的端到端流程RKNPU2 运行时侧 导读 本人工智能推理引擎模型量化模型优化边缘计算开发工具上一篇OpenViking Experience Memory 实战指南让 Agent 在执行任务时自动检索并复用历史操作经验下一篇PPTX 视觉素材治理指南为可编辑 PPTX 幻灯片选择与放置图标、图片、SVG、图表与信息图创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 RKNN-Toolkit2 自定义 GPU 算子（OpenCL）实战：以 ArgMax 算子移植 PP-HumanSegV2 为例
人工智能推理引擎模型量化模型优化边缘计算开发工具【免费下载链接】rknn-toolkit2项目地址https://gitcode.com/gh_mirrors/rk/rknn-toolkit2点击查看免费下载导读本文基于 RKNN-Toolkit2 仓库中 rknn_custom_gpu_op_demo 示例系统讲解如何在 RKNPU2 运行时上为 RKNN 模型实现自定义 GPUOpenCL算子。该示例以替代 ONNX 原生ArgMax算子为核心串联起 RKNN-Toolkit2 侧的自定义算子注册与 RKNN Runtime C API 侧的算子实现、构建、部署与精度验证全流程。读完本文你将掌握如何用 RKNN-Toolkit2 将带自定义算子的 ONNX 模型转换为 RKNN 模型、如何用 C/C OpenCL 编写自定义 GPU 算子的插件库、如何在 AArch64 Linux 与 Android 板端构建并运行以及如何利用自定义算子插件库配合 toolkit2 做精度分析。一、示例概览为什么需要自定义 GPU 算子RKNN 模型默认由 NPU神经网络处理器执行但并非所有算子都能被 NPU 直接支持。当模型中含有 NPU 不支持、或者用户希望绕开原生实现自行优化的算子时RKNN 提供了**自定义算子Custom Op**机制允许开发者在RKNN-Toolkit2工具链侧把 ONNX 模型中的某个算子替换为自定义算子并完成量化、编译、导出 RKNN 模型在RKNN Runtime C API板端侧用 CPU、GPUOpenCL或 PyTorch 实现该算子的实际计算逻辑通过插件库注入运行时。本示例rknn_custom_gpu_op_demo是三者中最典型的 GPU 实现它以ArgMax算子为例将 PaddleSeg 的 PP-HumanSegV2-Mobile 人像分割模型中的原生ArgMax替换为自定义 GPU 算子并在板端用 OpenCL kernel 完成通道维求最大值的计算最终输出人像 mask 与原图合成的效果图。示例配套的 RKNN 模型已注册 ArgMax 自定义算子预置于以下平台目录中平台模型路径相对仓库根目录RK3562model/RK3562/pp_human_segv2_custom_argmax.rknnRK3566 / RK3568model/RK3566_RK3568/pp_human_segv2_custom_argmax.rknnRK3576model/RK3576/pp_human_segv2_custom_argmax.rknnRK3588model/RK3588/pp_human_segv2_custom_argmax.rknn测试输入图片为 model/test_image.jpg513×513 人像图。1.1 模型来源PP-HumanSegV2-Mobile示例模型来自 PaddlePaddle 的 PaddleSeg 开源项目PP-HumanSeg 人像分割套件使用的具体模型为PP-HumanSegV2-Mobile192×192 输入。原模型是 Paddle 格式需要先通过 Paddle2onnx 转换为 ONNX 格式再进入 RKNN 转换流程。转换脚本位于 model/convert/test.py其预期输出为result.jpg包含输入图、分割 mask 图、合成叠加图与类别图例四联画效果图可参考 model/convert/result_truth.jpg。1.2 相关文档与配套示例模型转换与自定义算子注册的详细说明见 model/convert/README.md如何在 RKNN-Toolkit2 中替换 ONNX 原生算子可参考 rknn-toolkit2/examples/functions/custom_op/non-onnx_standard/test.py该示例以Sigmoid → cstSigmoid为例演示了 ONNX 图编辑与自定义算子注册同系列还有 CPU 版 rknn_custom_cpu_op_demo 与 PyTorch 版 rknn_custom_pytorch_op_demo可对照学习不同后端实现差异。二、工具链侧用 RKNN-Toolkit2 转换带自定义算子的模型板端运行前需要先用 RKNN-Toolkit2 生成注册了自定义 ArgMax 算子的 RKNN 模型。转换脚本为 model/convert/test.py核心流程如下初始化 RKNN 对象注册自定义算子 ArgMaxrknn.reg_custom_op将 ONNX 模型转换为 RKNN 模型在模拟器中推理验证保存结果图片释放 RKNN 资源。脚本关键代码节选from rknn.api import RKNN class ArgMax: op_type ArgMax # 创建 RKNN 对象 rknn RKNN(verboseTrue) onnx_model pp_humansegv2_mobile.onnx rknn_model pp_human_segv2_custom_argmax.rknn DATASET ./dataset.txt QUANTIZE_ON True rknn.config(mean_values[127.5, 127.5, 127.5], std_values[127.5, 127.5, 127.5], quant_img_RGB2BGRFalse, target_platformrk3576) # 注册自定义 ArgMax 算子 ret rknn.reg_custom_op(ArgMax) # 加载 ONNX 模型输入 x尺寸 1x3x192x192 ret rknn.load_onnx(modelonnx_model, inputs[x], input_size_list[[1,3,192,192]]) # 量化构建 ret rknn.build(do_quantizationQUANTIZE_ON, datasetDATASET) # 导出 RKNN 模型 ret rknn.export_rknn(rknn_model)关键点说明rknn.config的target_platform脚本默认注释为 rk3576转换前必须按实际目标平台修改支持 rk3562、rk3566/rk3568、rk3576、rk3588 等。模型在推理阶段需要注册了同名自定义算子的运行时环境因此转换侧工具链与运行侧板端插件库的算子类型名必须一致此处均为ArgMaxreg_custom_op(ArgMax)传入一个声明了op_type ArgMax的类。在工具链侧只需声明算子类型名用于模型编译与量化工具链会用默认实现完成量化统计真正的计算逻辑在板端插件库中提供load_onnx指定inputs[x]PP-HumanSegV2-Mobile 的输入 tensor 名为x尺寸[1, 3, 192, 192]NCHW与板端模型输入保持一致量化配置mean_values[127.5, 127.5, 127.5]、std_values[127.5, 127.5, 127.5]、quant_img_RGB2BGRFalse即采用 0~255 归一化到 [-1, 1] 的常见预处理约定且保持 RGB 通道顺序不变。关于“替换 ONNX 原生算子”的更多细节在 non-onnx_standard 示例 中可以看到另一种通用做法——先用onnx.load读入模型遍历model.graph.node将目标op_type如Sigmoid改写为自定义类型名如cstSigmoid再保存为新 ONNX 模型后加载转换自定义算子的 Python 类中还可以通过shape_infer与compute方法提供形状推导和参考实现。三、板端实现OpenCL 自定义 GPU 算子插件库板端自定义 GPU 算子的核心是 src/rknn_custom_op_opencl_plugin_lib.cpp它同时被两个目标使用主程序rknn_custom_gpu_op_demo和独立插件库librkcst_argmax.so从 CMakeLists.txt 可见rk_custom_argmax与rknn_custom_gpu_op_demo共享同一源文件。3.1 自定义算子的注册结构rknn_custom_op根据 rknn_custom_op.h 的定义一个自定义算子需要填充以下关键字段字段说明version自定义算子版本号target后端执行设备RKNN_TARGET_TYPE_CPU 1或RKNN_TARGET_TYPE_GPU 2GPU 即 OpenCL 后端op_type自定义算子类型名最长RKNN_MAX_NAME_LEN必须与模型转换时注册的名字一致如ArgMaxcl_kernel_nameOpenCL kernel 函数名cl_kernel_sourceOpenCL kernel 源码字符串若cl_source_size 0则此字段解释为 kernel 文件路径cl_source_sizekernel 源码长度cl_build_optionsOpenCL 构建 clProgram 的编译选项init[可选] 算子初始化回调prepare[可选] 算子 prepare 回调compute[必选] 算子计算回调GPU 算子的核心逻辑compute_native[可选] 原生属性计算回调当前版本暂不支持建议置空destroy[可选] 算子销毁回调GPU 后端的上下文由rknn_gpu_op_context提供包含cl_context、cl_command_queue、cl_kernel三个句柄由运行时在调用compute前创建并注入rknn_custom_op_context。3.2 插件库入口get_rknn_custom_op插件库必须导出get_rknn_custom_op()函数宏RKNN_CUSTOM_OP_EXPORT保证动态库符号默认可见。示例中的实现RKNN_CUSTOM_OP_EXPORT rknn_custom_op* get_rknn_custom_op() { memset(user_op, 0, sizeof(rknn_custom_op)); char op_type[] ArgMax; strcpy(user_op.op_type, op_type); user_op.version 1; user_op.target RKNN_TARGET_TYPE_GPU; // GPU(OpenCL) 后端 user_op.init custom_init_callback_gpu; user_op.compute compute_custom_gpu_op_float32; user_op.destroy destroy_callback_gpu; char kernel_name[] Argmax_channel_float; strcpy(user_op.cl_kernel_name, kernel_name); user_op.cl_kernel_source cl_kernel_source; user_op.cl_source_size strlen(cl_kernel_source); return user_op; }3.3 OpenCL kernel沿通道求 ArgMax示例内置的 kernel 源码Argmax_channel_float完成“沿 channel 维度求最大值所在索引”的计算输入为 float32输出为每个 (i, j) 像素位置对应的通道索引0 为 background、1 为 person__kernel void Argmax_channel_float(__global const float* src_buf, const int inputChannel, __global float* dst_buf) { unsigned int i get_global_id(0); unsigned int j get_global_id(1); unsigned int width get_global_size(0); unsigned int height get_global_size(1); int max_index 0; float max_value src_buf[i*widthj]; for (int c 1; c inputChannel; c) { float value src_buf[c*height*width i*widthj]; if (value max_value) { max_value value; max_index c; } } dst_buf[i*widthj] max_index; }其中二维工作项分别对应该层平面的宽和高每个工作项独立遍历inputChannel个通道比较后输出通道索引。这与语义分割中“取置信度最高的类别标签”的典型后处理需求完全吻合。3.4compute回调导入 DMA-BUF 并调度 kernelGPU 自定义算子的compute回调rknn_custom_op_opencl_plugin_lib.cpp遵循如下模式从op_ctx-gpu_ctx取出运行时注入的cl_context、cl_command_queue、cl_kernel通过 ARM 扩展clImportMemoryARM配合CL_IMPORT_TYPE_DMA_BUF_ARM将输入/输出 tensor 的 fdinputs[0].mem.fd与偏移mem.offset mem.size导入为cl_mem实现零拷贝访问 RKNN 运行时管理的 DMA 内存根据 tensor 属性计算字节数、偏移、通道数与目标宽高outputs[0].attr.dims[3]为宽、dims[2]为高clSetKernelArg设置 kernel 参数clEnqueueNDRangeKernel以二维global_work_size {dst_width, dst_height}调度 kernelclFinish(queue)等待执行完成返回 0 表示成功。init与destroy回调在示例中为占位实现分别打印argmax_init_callback_gpu/Argmax_destroy_callback_gpu用户可在其中完成资源准备与清理。3.5 主程序加载模型并注册自定义算子主程序 src/rknn_api_test_custom_op_opencl.cpp 展示了一套完整的 RKNN Runtime 调用流程与自定义 GPU 算子直接相关的部分包括uint32_t flag 0; // 当算子不受 NPU 支持时优先回退到 GPU 设备执行 flag flag | RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU; int ret rknn_init(ctx, model_path, 0, flag, NULL); ... // 从插件库导出函数获取自定义算子描述 auto gpu_user_op get_rknn_custom_op(); ret rknn_register_custom_ops(ctx, gpu_user_op, 1);RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU定义于 rknn_api.h含义为“当算子不被 NPU 支持时将 GPU 设为优先回退的执行后端”。在混合模型中NPU 算子走 NPU、自定义算子走 GPU此标志必不可少rknn_register_custom_ops(ctx, op, 1)在rknn_init之后将自定义算子注册到该rknn_context。rknn_custom_op.h同时指出compute回调必选其他可选若init回调返回RKNN_WARNING_SKIP_CUSTOM_OP_COMPUTE-14且该 op 类型被 RKNN 原生支持则运行时将改用 RKNN 原生实现。主程序随后依次完成查询 SDK/驱动版本与内存占用rknn_query、查询输入输出 tensor 属性含 layout 自适应NHWC 取dims[1..3]NCHW 取dims[2..3]作为宽高、stb系列库加载与缩放图片、rknn_inputs_set设置输入、rknn_set_core_mask(RKNN_NPU_CORE_ALL)设置 NPU 核掩码、循环rknn_run测时并输出每帧耗时与 FPS、查询RKNN_QUERY_PERF_DETAIL与RKNN_QUERY_PERF_RUN获取性能明细最后取出输出rknn_outputs_getwant_float 1请求 float 输出并把分割 mask 叠加回原图通过stbi_write_png写出test_compose_img.png。四、构建AArch64 Linux 与 Android示例提供两个构建脚本build-linux.sh 与 build-android.sh二者都基于 CMakeLists.txt 组织工程。4.1 AArch64 Linux 构建先导出交叉编译器前缀再执行构建脚本export GCC_COMPILER~/opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu ./build-linux.sh -t rk3588 -a aarch64 -b Release脚本参数说明摘自 build-linux.sh 的用法提示-t target目标平台合法值为rk3562 / rk3566 / rk3568 / rk3576 / rk3588rk356x、rk3566_rk3568等写法也会被归一化到RK3566_RK3568-a arch架构aarch64或armhf-b build_type构建类型Debug或Release缺省为Release。脚本校验GCC_COMPILER可用性要求${GCC_COMPILER}-gcc可执行随后在build/build_TARGET_SOC_linux_arch_build_type目录执行cmake ../.. -DTARGET_SOC... -DCMAKE_SYSTEM_NAMELinux -DCMAKE_SYSTEM_PROCESSORarch -DCMAKE_C_COMPILER${CC} -DCMAKE_CXX_COMPILER${CXX}最后make -j4 make install。4.2 Android 构建先导出 NDK 路径再执行export ANDROID_NDK_PATH~/opts/ndk/android-ndk-r18b ./build-android.sh -t rk3568 -a arm64-v8a -b Release-a arch在 Android 下取值arm64-v8a或armeabi-v7abuild-android.sh 推荐使用NDK r18 / r19其他版本可能导致构建失败内部通过cmake -DANDROID_TOOLCHAINclang -DCMAKE_TOOLCHAIN_FILE$ANDROID_NDK_PATH/build/cmake/android.toolchain.cmake -DANDROID_ABIarch -DANDROID_STLc_static -DANDROID_PLATFORMandroid-24交叉编译。4.3 CMake 工程要点CMakeLists.txt 揭示了工程组织方式运行时头文件与librknnrt.so取自仓库的 runtime/Linux/librknn_apiAndroid 侧按 ABI 子目录选择依赖仓库自带 3rdparty/opencl/libopencl-stubOpenCL 客户端 stub以及3rdparty下的fp16、stb头文件无需外部安装 OpenCL SDK工程产出三个目标OpenCL stub 动态库、可执行程序rknn_custom_gpu_op_demo、自定义算子插件库rk_custom_argmaxCMAKE_INSTALL_PREFIX指向install/rknn_custom_gpu_op_demo_Linux|Androidmake install会把可执行程序、librknnrt.so到lib/、测试图片、对应平台的 RKNN 模型按TARGET_SOC目录以及rk_custom_argmax到lib/统一装配到 install 目录方便整目录推送到板端。五、部署与运行5.1 AArch64 Linux将install/rknn_custom_gpu_op_demo_Linux目录拷贝到设备/userdata/下。瑞芯微 EVB 开发板可用 adb 推送adb push install/rknn_custom_gpu_op_demo_Linux/ /data/若板子开了 sshd也可用 scp 等方式拷贝程序与 RKNN 模型。随后进入目录运行adb shell cd /data/rknn_custom_gpu_op_demo_Linux export LD_LIBRARY_PATH./lib ./rknn_custom_gpu_op_demo model/TARGET_PLATFORM/pp_human_segv2_custom_argmax.rknn model/test_image.jpg 1命令行参数依次为RKNN 模型路径、输入图片路径、循环推理次数loop_count示例传 1。程序运行后会打印 SDK/驱动版本、tensor 属性、每帧耗时与 FPS并写出合成图test_compose_img.png。5.2 Androidadb push install/rknn_custom_gpu_op_demo_Android/ /data/ adb shell cd /data/rknn_custom_gpu_op_demo_Android/若报错library libOpenCL.so not found: needed by main executable说明设备上 OpenCL 运行库未被加载器找到可将系统自带的libOpenCL.so拷到本地再运行cp /vendor/lib64/libOpenCL.so ./lib/ export LD_LIBRARY_PATH./lib ./rknn_custom_gpu_op_demo model/TARGET_PLATFORM/pp_human_segv2_custom_argmax.rknn model/test_image.jpg 1注意不同平台、不同版本的工具链与驱动推理结果可能略有差异。六、插件库与精度分析Accuracy Analysis若想验证自己实现的自定义算子数值是否准确官方推荐流程是将自定义算子单独编译成插件库安装到板端系统目录再用 RKNN-Toolkit2 的精度分析功能进行对比。6.1 编译插件库插件库源码即前文分析的 src/rknn_custom_op_opencl_plugin_lib.cpp。用与主程序相同的构建体系编译仅需要其中的rk_custom_argmax目标./build-linux_TARGET_PLATFORM.sh # 或 ./build-android_TARGET_PLATFORM.sh6.2 安装插件库将插件库推送到运行时的自定义算子插件目录Linux/usr/lib/rknpu/op_pluginsAndroidarm64-v8a/vendor/lib64/。例如 Linux 系统adb push install/rknn_custom_gpu_op_demo_Linux/lib/librkcst_argmax.so /usr/lib/rknpu/op_plugins注意若板端不存在/usr/lib/rknpu/op_plugins目录需要先手动创建。插件库安装到位后RKNN Runtime 在加载含有对应op_type的模型时会自动 dlopen 该插件并调用其get_rknn_custom_op()获取算子实现从而让板端运行及 toolkit2 的精度分析使用用户自写的 GPU 算子。6.3 执行精度分析完成插件部署后即可使用 RKNN-Toolkit2 提供的精度分析accuracy analysis功能对比板端真实自定义 GPU 算子与模拟器/参考实现的输出确认自定义算子的计算精度满足要求。精度分析相关使用方式可参考 rknn-toolkit2/examples/functions/accuracy_analysis/test.py 及配套 README。七、从源码看 GPU 自定义算子的整体数据流综合 rknn_custom_op.h、rknn_api.h 与示例实现一次完整的 GPU 自定义算子推理数据流可以概括为转换阶段工具链PC 端RKNN-Toolkit2 通过reg_custom_op感知模型中的自定义ArgMax算子完成量化与编译生成pp_human_segv2_custom_argmax.rknn算子信息连同cl_kernel_name等元数据写入模型加载阶段板端rknn_init时传入RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU确保模型编译期无法被 NPU 支持的算子回退到 GPU注册阶段get_rknn_custom_op()导出算子描述目标后端 GPU、OpenCL kernel 名与源码、init/compute/destroy 回调rknn_register_custom_ops注入rknn_context执行阶段rknn_run触发推理遇到ArgMax节点时运行时创建 OpenCL context / command queue / program 并调用compute回调回调通过clImportMemoryARM导入输入输出 tensor 的 DMA-BUF 实现零拷贝调度二维Argmax_channel_floatkernel 计算通道索引后clFinish返回验证阶段主程序取出 float 输出将0person区域在原图中置黑生成合成图用户如需更严谨的精度验证则部署插件库后使用 toolkit2 的 accuracy analysis 对比输出。八、常见问题与注意事项算子类型名必须一致转换脚本reg_custom_op(ArgMax)中的op_type、插件库user_op.op_type、以及模型图中的算子名三者必须完全一致否则运行时无法匹配到自定义实现GPU 回退标志不能遗漏若不设置RKNN_FLAG_EXECUTE_FALLBACK_PRIOR_DEVICE_GPU含 NPU 不支持算子的模型在rknn_init阶段可能直接失败OpenCL 环境依赖Linux 侧通过仓库自带 libopencl-stub 链接触发设备上的 OpenCL ICDAndroid 侧若找不到libOpenCL.so需按 5.2 节拷贝设备自带的运行库插件目录权限与存在性Linux 插件目录/usr/lib/rknpu/op_plugins在部分系统上需要手动创建精度差异不同平台、工具链版本与驱动版本下自定义算子的结果可能有轻微差异属正常现象compute_native当前不可用按 rknn_custom_op.h 说明compute_native暂未支持应置空。九、小结rknn_custom_gpu_op_demo提供了一个“从工具链注册、到 OpenCL kernel 编写、再到板端部署验证”的完整闭环范例。开发者可以以此为模板将自己业务中 NPU 不支持的算子或其他需要 GPU 加速的算子按同样流程接入 RKNN 运行时在 toolkit2 侧声明op_type完成模型转换在板端实现get_rknn_custom_op导出结构与compute回调最后通过插件库机制配合精度分析确认正确性。配套的 CPU 版rknn_custom_cpu_op_demo与 PyTorch 版rknn_custom_pytorch_op_demo示例则可提供不同后端实现的横向参考。赞分享人工智能推理引擎模型量化模型优化边缘计算开发工具【免费下载链接】rknn-toolkit2项目地址https://gitcode.com/gh_mirrors/rk/rknn-toolkit2点击查看免费下载相关推荐Wox Shell 插件完全指南在启动器中直接执行命令、管理历史与工作目录Wox Shell 插件完全指南在启动器中直接执行命令、管理历史与工作目录 本指南围绕 Wox 内置系统插件 Shell 展开讲解如何用 触发关键字在启动器人工智能推理引擎模型量化模型优化边缘计算开发工具CANN PTO 自定义算子开发实战以 fused_add_relu_mul 算子融合为例CANN PTO 自定义算子开发实战以 fused_add_relu_mul 算子融合为例 本篇指南围绕 CANN Parallel Tile Operati算子库人工智能CANNRKNN-Toolkit2 自定义 CPU 算子实战以 cstSigmoid 替换 ONNX Sigmoid 的端到端流程RKNPU2 运行时侧RKNN Toolkit2 自定义 CPU 算子实战以 cstSigmoid 替换 ONNX Sigmoid 的端到端流程RKNPU2 运行时侧 导读 本人工智能推理引擎模型量化模型优化边缘计算开发工具上一篇OpenViking Experience Memory 实战指南让 Agent 在执行任务时自动检索并复用历史操作经验下一篇PPTX 视觉素材治理指南为可编辑 PPTX 幻灯片选择与放置图标、图片、SVG、图表与信息图创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考