TVM 部署到 Adreno GPU 全指南：OpenCL 纹理增强、OpenCLML 加速与 Android 端到端部署实战

TVM 部署到 Adreno GPU 全指南OpenCL 纹理增强、OpenCLML 加速与 Android 端到端部署实战【免费下载链接】tvmOpen deep learning compiler stack for cpu, gpu and specialized accelerators项目地址: https://gitcode.com/gh_mirrors/tvm7/tvm本指南系统讲解 Apache TVM 如何在 Qualcomm Adreno GPU 上完成深度学习模型的加速部署一方面介绍 TVM 为 Adreno 优化的原生 OpenCL 后端纹理内存、Adreno 友好布局与专用调度另一方面介绍基于 Qualcomm OpenCLML SDK 的算子加速集成。读者将掌握从开发环境搭建、RPC 调优、tvmc/Python 编译、Android 目标部署到精度切换float16 / float16_acc32的完整实战路径。引言TVM 与 Adreno GPUAdreno 是 Qualcomm 设计的一系列图形处理器GPUIP 核心广泛应用于其 SoC 中。它负责加速复杂几何图形的渲染以较低功耗提供高性能图形体验和丰富的用户交互。TVM 对 Adreno GPU 的深度学习加速支持分为两条路径原生 OpenCL 后端TVM 对 OpenCL 后端做了 Adreno 定向增强引入了纹理内存texture memory的使用和 Adreno 友好的数据布局并配套全新的调度schedule策略OpenCLML 后端OpenCLML 是 Qualcomm 发布的 SDK为大部分深度学习算子提供内核加速库TVM 以 BYOCBring Your Own Codegen方式将其集成。本指南依次讲解三大部分OpenCL 后端增强对应文档opencl_enhancements一节、OpenCLML 简介about_openclml以及构建与部署全流程build_deploy。TVM OpenCL 后端的 Adreno 增强TVM 的 OpenCL 后端针对 Adreno 特性做了三项增强纹理内存texture memory使用Adreno 友好的激活activation数据布局为配合上述特性而新增的调度策略。为什么使用纹理Adreno 在纹理处理上具有独特优势这也是 TVM 为 Adreno 引入纹理支持的原因专用 L1 缓存纹理处理器Texture Processor, TP拥有独立的只读 L1 缓存用于缓存从 L2 缓存取回的纹理操作数据这是使用纹理的最主要原因内建图像边界处理纹理对图像边界的处理是内建的无需额外代码丰富的格式组合纹理支持多种图像格式与数据类型组合并具备自动格式转换能力。总体上与基于 OpenCL buffer 的方案相比纹理方案可以带来显著的性能提升。目标字符串与内核形态对比在常规 OpenCL 目标下我们指定targetopencl生成的卷积内核使用__global指针即 OpenCL buffer 对象__kernel void tvmgen_default_fused_nn_conv2d_kernel0(__global float* restrict p0, __global double* restrict p1, __global float* restrict conv2d_nhwc) { // body..当启用纹理增强将目标改为targetopencl -deviceadreno后生成的内核改为使用纹理支撑的 OpenCL image 对象__kernel void tvmgen_default_fused_nn_conv2d_kernel0(__write_only image2d_t pad_temp_global_texture, __read_only image2d_t p0) { // body..image2d_t是 OpenCL 内建类型表示二维图像对象并提供若干附加函数。使用image2d_t时一次可读取 4 个元素从而更高效地利用硬件。从仓库源码看这一能力由多个层面协同支撑存储标注 passsrc/relay/transforms/annotate_texture_storage.cc 中的StorageInfo负责分析 Relay 图产出表达式到存储作用域storage scope的映射其中accept_textures_记录可接受纹理的算子配合PlanDevice算法完成内存作用域标注为纹理代码生成提供前提算子层实现python/tvm/topi/adreno/ 下提供了conv2d_nchw.py、conv2d_nhwc.py、conv2d_winograd_common.py、depthwise_conv2d_nchw.py、depthwise_conv2d_nhwc.py、reduction.py等 Adreno 专用 compute/schedule以及 python/tvm/topi/adreno/utils.py 中诸如split_to_chunks、pack_input运行时打包输入、将通道维度扩展为 4 的倍数等工具函数目标定义python/tvm/target/target.py 中定义了adreno目标工厂函数其选项即[-deviceadreno, -model%s % model]测试佐证tests/python/relay/opencl_texture/test_conv2d_nchw_texture.py、tests/python/relay/opencl_texture/test_conv2d_nhwc_texture.py 等测试覆盖了纹理路径的卷积用例。关于内核源码的生成与查看方式详见后文「高级用法生成源码检查」。关于 OpenCLMLOpenCLML 是 Qualcomm 发布的 SDK提供加速的深度学习算子。这些算子以扩展cl_qcom_ml_ops的形式暴露于标准 OpenCL 规范之上作为 TVM 的 BYOCBring Your Own Codegen方案集成。OpenCLML 算子与原生 OpenCL 使用相同的 context并可入队到与原生 OpenCL 相同的 command queue。TVM 正是利用这一点避免在回退到原生 OpenCL 时发生上下文切换context switching开销。在仓库中OpenCLML 的集成包含三块编译期 codegensrc/relay/backend/contrib/clml/codegen.cc 中的CLMLJSONSerializer将clml.conv2d、clml.pad_conv2d、clml.conv2d_transpose、clml.batch_norm、clml.dense1d、clml.dense2d、clml.pad、clml.concat等算子序列化为 CLML 子图描述运行时src/runtime/contrib/clml/ 下的clml_runtime.cc、clml_memory_planner.cc、clml_utils.cc实现 CLML 模块的加载、内存规划与算子执行构建配置cmake/modules/contrib/CLML.cmake 中USE_CLML控制 CLML codegen 与运行时模块的编译USE_CLML_GRAPH_EXECUTOR则启用 CLML graph runtime并在启用时自动set(USE_OPENCL ON)将 OpenCL 作为 CLML 的回退路径代码注释即写明 Enable OpenCL as fallback to CLML。TVM for Adreno构建与部署总览Adreno 是通过 ADB 连接到主机的远程目标设备在其上部署编译模型需要在主机与目标端分别使用相应工具。TVM 提供了友好的命令行工具也提供了面向开发者的 Python API覆盖自动调优、编译与部署等环节。整体部署管线包含以下阶段模型导入Model import从 TensorFlow、PyTorch、ONNX 等主流框架导入模型转换为 TVM 的 relay Module 格式也可以直接用 TVM 算子清单手工构建 relay Module。此阶段产出的 TVM 模块是与目标无关的图表示。自动调优Auto Tuning针对具体目标调优 TVM 生成的 kernel。调优需要目标设备可用对 Adreno 这类 Android 上的远程目标通过 RPC 进行通信后文详述。调优并非编译模型的必要步骤但它是让 TVM kernel 达到最佳性能的关键。编译Compilation针对具体目标编译模型。若上一步完成了调优编译会使用调优日志tuning log生成性能最优的 kernel最终产出包含 kernel 动态库mod.so、json 格式图定义mod.json与 TVM 特有格式参数二进制文件mod.params的编译产物。目标端部署/测试运行Deploy在目标端运行编译产物既可通过 RPC 从 Python 环境部署也可使用针对 Android 交叉编译的原生工具。此阶段可在 Android 目标上运行模型并验证正确性与性能。应用集成Application Integration在应用中集成 TVM 编译模型涉及从 AndroidC 原生环境或 JNI调用 TVM runtime 设置输入、获取输出。高级用法Advanced Usage面向高级用户例如查看生成源码、修改模块精度等。以下各节将逐项展开这些阶段。开发环境搭建自动DockerTVM 提供了预构建的 docker 容器环境内含全部前置依赖可快速上手。更想自主控制依赖的用户可参考后面的手动搭建章节。Docker 方式唯一的前置条件是主机上装有 docker。构建 adreno 镜像的命令如下./docker/build.sh ci_adreno docker tag tvm.ci_adreno ci_adreno随后用下面命令同时构建主机端与目标端工具./tests/scripts/ci.py adreno -i如需带 OpenCLML SDK 构建 TVM需先导出 SDK 路径再执行构建export ADRENO_OPENCLPath to OpenCLML SDK ./tests/scripts/ci.py adreno -i编译成功后即进入 docker shell。构建会生成两个目录build-adreno主机端 TVM 编译器构建build-adreno-targetAndroid 目标端组件包含libtvm_runtime.soTVM runtime 库tvm_rpcRPC 运行时环境工具rtvm独立的原生部署工具。由于 docker 环境与主机共享 Android 设备主机上的 adb 版本需要为1.0.41docker 内使用同一版本。可在 docker 内用adb devices检查设备userci-adreno-fpeqs:~$ adb devices List of devices attached aaaabbbb device cccddddd device自动化脚本 tests/scripts/setup-adreno-env.sh 是这套 CI 流程的核心它定义了def_environment导出TVM_TRACKER_HOST、TVM_TRACKER_PORT、RPC_DEVICE_KEYandroid、RPC_TARGETadreno、TVM_NDK_CC指向aarch64-linux-android28-clang等变量并支持-e tracker、-e device、-e query三种模式见后文 RPC 章节。开发环境搭建手动手动构建需要分别构建主机端与目标端组件。主机端Host Compiler构建mkdir -p build cd build cp ../cmake/config.cmake . # Enable RPC capability to communicate to remote device. echo set\(USE_RPC ON\) config.cmake # We use graph executor for any host(x86) side verification of the model. echo set\(USE_GRAPH_EXECUTOR ON\) config.cmake # Enable backtrace if possible for more debug information on any crash. echo set\(USE_LIBBACKTRACE AUTO\) config.cmake # The target_host will be llvm. echo set\(USE_LLVM ON\) config.cmake若需编译 OpenCLML 支持追加如下配置export ADRENO_OPENCLPath to OpenCLML SDK echo set\(USE_CLML ${ADRENO_OPENCL}\) config.cmake然后构建cmake .. make最后导出 Python 路径并验证export PYTHONPATH$TVM_HOME/python:${PYTHONPATH} python3 -c import tvm # Verify tvm python package目标端Android Target构建目标端构建需要先安装 Android NDK以及获取 adb 工具的 Android Debug Bridge。然后按如下配置mkdir -p build-adreno cd build-adreno cp ../cmake/config.cmake . # Enable OpenCL backend. echo set\(USE_OPENCL ON\) config.cmake # Enable RPC functionality. echo set\(USE_RPC ON\) config.cmake # Build tvm_rpc tool that runs on target device. echo set\(USE_CPP_RPC ON\) config.cmake # Build native rtvm deploy tool. echo set\(USE_CPP_RTVM ON\) config.cmake # We use graph executor for deploying on devices like Android. echo set\(USE_GRAPH_EXECUTOR ON\) config.cmake # Backtrace enablement if possible. echo set\(USE_LIBBACKTRACE AUTO\) config.cmake # Adreno supports 32bit alignment for OpenCL allocations rather 64bit. echo set\(USE_KALLOC_ALIGNMENT 32\) config.cmake # Android build related defines. echo set\(ANDROID_ABI arm64-v8a\) config.cmake echo set\(ANDROID_PLATFORM android-28\) config.cmake echo set\(MACHINE_NAME aarch64-linux-gnu\) config.cmake同样地要开启 OpenCLML 支持需追加export ADRENO_OPENCLPath to OpenCLML SDK echo set\(USE_CLML ${ADRENO_OPENCL}\) config.cmake echo set\(USE_CLML_GRAPH_EXECUTOR ${ADRENO_OPENCL}\) config.cmakeAndroid 目标构建依赖环境变量ANDROID_NDK_HOME。随后执行cmake -DCMAKE_TOOLCHAIN_FILE${ANDROID_NDK_HOME}/build/cmake/android.toolchain.cmake \ -DANDROID_ABIarm64-v8a \ -DANDROID_PLATFORMandroid-28 \ -DCMAKE_SYSTEM_VERSION1 \ -DCMAKE_FIND_ROOT_PATH${ADRENO_OPENCL} \ -DCMAKE_FIND_ROOT_PATH_MODE_PROGRAMNEVER \ -DCMAKE_FIND_ROOT_PATH_MODE_LIBRARYONLY \ -DCMAKE_CXX_COMPILER${ANDROID_NDK_HOME}/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang \ -DCMAKE_C_COMPILER${ANDROID_NDK_HOME}/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang \ -DMACHINE_NAMEaarch64-linux-gnu .. make tvm_runtime tvm_rpc rtvm关键 CMake 配置项一览配置项作用默认值见 cmake/config.cmakeUSE_OPENCL启用 OpenCL 后端OFFUSE_OPENCL_ENABLE_HOST_PTR允许 host 指针配合 CLML 使用OFFUSE_RPC与远程设备通信的 RPC 能力—USE_CPP_RPC构建运行在目标端的tvm_rpc工具—USE_CPP_RTVM构建原生rtvm部署工具—USE_GRAPH_EXECUTOR图执行器Android 部署用—USE_LIBBACKTRACE崩溃回溯支持AUTO时自动探测—USE_KALLOC_ALIGNMENTOpenCL 分配对齐字节数Adreno 用 32 位对齐—USE_CLML编译 CLML codegen 与运行时模块OFFcmake/config.cmakeUSE_CLML_GRAPH_EXECUTOR启用 CLML graph runtime值为 SDK 路径或ON/OFF启用时自动打开 OpenCL 作为回退OFFANDROID_ABI/ANDROID_PLATFORM/MACHINE_NAMEAndroid 构建三元组文档示例为arm64-v8a/android-28/aarch64-linux-gnu—从 cmake/modules/contrib/CLML.cmake 的实现可见USE_CLML非ON时会从${USE_CLML}/CL/cl_qcom_ml_ops.h读取CL_QCOM_ML_OPS_H_MAJOR_VERSION并定义TVM_CLML_VERSIONUSE_CLML_GRAPH_EXECUTOR则会在指定路径下通过find_library查找libOpenCL并把lib64/libOpenCL.so与lib64/libOpenCL_system.so链接进运行时。RPC 设置RPCRemote Procedure Call让远程目标可通过 TCP/IP 被访问是自动调优的前提——调优需要在真实设备上运行自动生成的 kernel并用机器学习方法优化。同时RPC 也用于从主机端通过 Python 或tvmc向远程设备部署模型。RPC 设置包含两类组件TVM Tracker主机端守护进程管理远程设备并服务于主机端应用。应用可连接 tracker 获取远程设备句柄进行通信TVM RPC运行在远程设备本例为 Android上的原生应用向主机上的 TVM Tracker 注册自己。自动化 RPC 设置Docker 环境启动 tracker监听 9190 端口./tests/scripts/ci.py adreno -i # Launch a new shell on the adreno docker source tests/scripts/setup-adreno-env.sh -e tracker -p 9190在序列号为abcdefgh的 Android 设备上运行 TVM RPC./tests/scripts/ci.py adreno -i # Launch a new shell on adreno docker. source tests/scripts/setup-adreno-env.sh -e device -p 9190 -d abcdefgh在任意其他 docker 终端查询 RPC 设置详情./tests/scripts/ci.py adreno -i # Launch a new shell on adreno docker. source tests/scripts/setup-adreno-env.sh -e query -p 9190对照 tests/scripts/setup-adreno-env.sh 的实现各模式行为如下tracker运行python3 -m tvm.exec.rpc_tracker --host 0.0.0.0 --port RPC_PORTdevice利用adb push将build-adreno-target/tvm_rpc与libtvm_runtime.so推送到设备的/data/local/tmp/tvm_ci-USER执行adb reverse/adb forward端口转发再在设备上以--tracker127.0.0.1:RPC_PORT --keyandroid启动tvm_rpc server默认监听 5000-l参数可改并导出CLML_PROFILING1、CLML_IS_TUNING_RUN1、CLML_TUNING_CACHEclml.bin等 CLML 环境变量query运行python3 -m tvm.exec.query_rpc_tracker --port RPC_PORT列出已注册设备。手动 RPC 设置更细粒度的手动 RPC 环境搭建可参考 TVM 官方「如何在 Adreno 上部署模型」教程deploy_model_on_adreno。RPC 设置完成后主机侧得到 rpc-tracker127.0.0.1即 rpc 主机与 rpc 端口9190。命令行工具tvmc 全流程tvmc是 TVM 的命令行工具可完成模型导入、自动调优、编译以及通过 RPC 部署选项丰富。模型导入与调优下面命令从任意框架导入模型并自动调优。示例使用 Keras 模型借助 RPC 调优并生成调优日志keras-resnet50.logpython3 -m tvm.driver.tvmc tune --targetopencl -deviceadreno \ --target-hostllvm -mtripleaarch64-linux-gnu \ resnet50.h5 -o \ keras-resnet50.log \ --early-stopping 0 --repeat 30 --rpc-key android \ --rpc-tracker 127.0.0.1:9190 --trials 1024 \ --tuning-records keras-resnet50-records.log --tuner xgb关键参数--target指定opencl -deviceadreno纹理增强目标--target-host使用llvm -mtripleaarch64-linux-gnu与目标端架构一致--rpc-key android、--rpc-tracker 127.0.0.1:9190对接前述 RPC 环境--trials控制搜索轮数--tuner xgb选择 XGBoost 调优器--early-stopping 0关闭提前停止。模型编译编译模型并产出 TVM 编译产物python3 -m tvm.driver.tvmc compile \ --cross-compiler ${ANDROID_NDK_HOME}/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang \ --targetopencl, llvm --target-llvm-mtriple aarch64-linux-gnu --target-opencl-device adreno \ --tuning-records keras-resnet50.log -o keras-resnet50.tar resnet50.h5启用 OpenCLML 卸载时需在 target 中追加clml。调优日志对 OpenCLML 路径同样有效——因为任何未走 OpenCLML 路径的算子都会回退到 OpenCL回退算子会复用该调优日志python3 -m tvm.driver.tvmc compile \ --cross-compiler ${ANDROID_NDK_HOME}/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang \ --targetopencl, clml, llvm --desired-layout NCHW --target-llvm-mtriple aarch64-linux-gnu --target-opencl-device adreno \ --tuning-records keras-resnet50.log -o keras-resnet50.tar resnet50.h5编译成功后产出keras-resnet50.tar这是一个压缩归档内含 kernel 动态库mod.so、图 jsonmod.json与参数字节文件mod.params。目标端部署与运行编译模型既可通过 RPC 运行也可原生部署。RPC 方式python3 -m tvm.driver.tvmc run --devicecl keras-resnet50.tar \ --rpc-key android --rpc-tracker 127.0.0.1:9190 --print-timetvmc run提供多种输入初始化模式fill、random等。tvmc部署通常用于在目标上快速验证编译模型。生产环境一般使用原生部署环境Android JNI 或 C 原生环境此时需要调用交叉编译的tvm_runtime接口来部署编译产物TVMPackage。TVM 提供了独立原生工具rtvm构建产物位于build-adreno-target下见 apps/cpp_rtvm可直接在 ADB shell 中原生部署运行模型。应用集成细节见后文。Python 接口TVM 提供基于tvmc抽象的高层接口也提供底层的 Relay API。tvmc 接口使用tvmcPython 接口时先加载模型得到TVMCModelTVMCModel用于自动调优并产出 tuning cache编译过程结合TVMCModel与可选的tuning cache 生成TVMCPackageTVMCPackage可保存到文件系统也可直接用于目标端部署运行。保存的TVMCPackage也可交给rtvm工具做原生部署。更多 API 细节参见 tvmc 文档。Relay 接口Relay API 提供更底层的编译接口。与tvmc类似它提供多种前端 API 将模型转换为 relayModule。Module用于各类变换精度转换、CLML 卸载、其他自定义变换变换后的Module再用于自动调优最后用relay.build生成库模块。从库模块可以导出 mod.so、mod.params、mod.json 等编译产物并据此创建 graph runtime 在目标设备上部署运行。另外TVM 还支持通过 TVM4J 提供 Java 接口见 jvm 目录。应用集成TVM 编译产物为模块动态库mod.so、图 jsonmod.json与参数mod.paramsTVMPackage归档中也包含这些内容。对 Android 应用集成而言C/C 接口通常已足够。TVM 原生暴露c_runtime_api见 include/tvm/runtime/c_runtime_api.h用于加载并运行编译模块更简化的封装可参考 apps/cpp_rtvm/tvm_runner.h 的TVMRunner接口。高级用法本节介绍使用 Adreno 目标时的一些高级用法与附加信息。生成源码检查除标准编译产物mod.so、mod.json、mod.params外还可以从库句柄导出 OpenCL kernel 源码、CLML 卸载子图等。TVM 编译输出组织为一个 TVM 模块其中导入多个子模块。按 json 格式导出 CLML 子图# Look for clml typed module imported. clml_modules list(filter(lambda mod: mod.type_key clml, lib.get_lib().imported_modules)) # Loop through all clml sub graphs and dump the json formatted CLML sub graphs. for cmod in clml_modules: print(CLML Src:, cmod.get_source())类似地导出 OpenCL kernel 源码# Similarly we can dump open kernel source too as shown below # Look for opencl typed module imported. opencl_modules list(filter(lambda mod: mod.type_key opencl, lib.get_lib().imported_modules)) # Now dump kernel source for each OpenCL targetted sub graph. for omod in opencl_modules: print(OpenCL Src:, omod.get_source())mod.type_key对应模块类型clml/opencl运行时侧则由 src/runtime/contrib/clml/clml_runtime.cc 等实现承载。精度选择为特定负载选择正确的精度能显著提升效率使精度与速度的初始平衡向任务优先级倾斜。TVM 提供float16、float16_acc32混合精度与float32标准三种选择。float16为充分利用 GPU 半精度计算与内存管理的优势可将浮点模型转换为半精度模型。较低精度会提升性能但也可能造成精度损失。转换需在任意前端生成 relay Module 后立即调用 Adreno 专用变换 APIfrom tvm.driver.tvmc.transform import apply_graph_transforms mod apply_graph_transforms( mod, { mixed_precision: True, mixed_precision_ops: [nn.conv2d, nn.dense], mixed_precision_calculation_type: float16, mixed_precision_acc_type: float16, }, )tvm.driver.tvmc.transform.apply_graph_transforms是ToMixedPrecisionpass 的简化封装用于获得所需精度。随后可按任意习惯方式编译with tvm.transform.PassContext(opt_level3): lib relay.build( mod, target_hosttarget_host, targettarget, paramsparams )使用tvmcPython 接口时以下参数启用 float16 精度转换mixed_precision True, mixed_precision_ops [nn.conv2d, nn.dense], mixed_precision_calculation_type float16, mixed_precision_acc_type float16tvmc命令行接口对应选项--mixed-precision --mixed-precision-ops nn.conv2d nn.dense --mixed-precision-calculation-type float16 --mixed-precision-acc-type float16float16_acc32混合精度ToMixedPrecisionpass 会遍历网络将网络切分为处理 float 或 float16 数据类型的算子簇。簇由三类算子定义始终转换为 float16 数据类型的算子若后续紧跟已转换的簇则可以转换的算子永不转换为 float16 数据类型的算子。该清单定义于ToMixedPrecision的实现见 python/tvm/relay/transform/mixed_precision.py且可由用户覆盖。ToMixedPrecision将 FP32 Relay 图转换为 FP16 版本累积数据类型为 FP16 或 FP32。该变换有助于减小模型体积——在 FP16_acc16 情形下权重体积约减半。使用方式同样被简化为一次调用from tvm.driver.tvmc.transform import apply_graph_transforms mod apply_graph_transforms( mod, { mixed_precision: True, mixed_precision_ops: [nn.conv2d, nn.dense], mixed_precision_calculation_type: float16, mixed_precision_acc_type: float32, }, )编译方式与 float16 相同with tvm.transform.PassContext(opt_level3): lib relay.build( mod, target_hosttarget_host, targettarget, paramsparams )tvmcPython 接口参数mixed_precision True, mixed_precision_ops [nn.conv2d, nn.dense], mixed_precision_calculation_type float16, mixed_precision_acc_type float32tvmc命令行接口选项--mixed-precision --mixed-precision-ops nn.conv2d nn.dense --mixed-precision-calculation-type float16 --mixed-precision-acc-type float32从实现上看python/tvm/driver/tvmc/transform.py 的convert_to_mixed_precision基于generate_mixed_precision_rule(acc_dtype)生成转换规则区分MIXED_PRECISION_ALWAYS等类别apply_graph_transformspython/tvm/driver/tvmc/transform.py在mixed_precision为真时调用它完成图变换。延伸阅读围绕本指南涉及的各个环节仓库内还提供了可继续深入的材料构建脚本tests/scripts/task_config_build_adreno.sh、tests/scripts/task_build_adreno_bins.sh、tests/scripts/task_python_adreno.shCI 镜像定义docker/Dockerfile.ci_adrenoCLML 集成测试tests/python/contrib/test_clml/test_compiler.py、test_network.py、test_ops.py与基础设施infrastructure.pyOpenCL 纹理/Adreno 代码生成测试tests/python/codegen/test_target_texture_codegen_opencl.py、tests/python/relay/opencl_texture/原生部署示例apps/cpp_rtvmmain.cc、tvm_runner.cc等以及 apps/cpp_rpc。需要说明的是Adreno 属于远程目标无论是自动调优还是 RPC 部署都要求主机能通过 ADB 访问 Android 设备且主机与设备间的 adb/NDK 版本应保持一致这是本指南所有命令生效的前提。【免费下载链接】tvmOpen deep learning compiler stack for cpu, gpu and specialized accelerators项目地址: https://gitcode.com/gh_mirrors/tvm7/tvm创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考