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 GPU

Adreno 是 Qualcomm 设计的一系列图形处理器（GPU）IP 核心，广泛应用于其 SoC 中。它负责加速复杂几何图形的渲染，以较低功耗提供高性能图形体验和丰富的用户交互。

TVM 对 Adreno GPU 的深度学习加速支持分为两条路径：

原生 OpenCL 后端：TVM 对 OpenCL 后端做了 Adreno 定向增强，引入了纹理内存（texture memory）的使用和 Adreno 友好的数据布局，并配套全新的调度（schedule）策略；
OpenCLML 后端：OpenCLML 是 Qualcomm 发布的 SDK，为大部分深度学习算子提供内核加速库，TVM 以 BYOC（Bring 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 目标下，我们指定target="opencl"，生成的卷积内核使用__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..

当启用纹理增强，将目标改为target="opencl -device=adreno"后，生成的内核改为使用纹理支撑的 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 个元素，从而更高效地利用硬件。

从仓库源码看，这一能力由多个层面协同支撑：

存储标注 pass：src/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目标工厂函数，其选项即["-device=adreno", "-model=%s" % model]；
测试佐证：tests/python/relay/opencl_texture/test_conv2d_nchw_texture.py、tests/python/relay/opencl_texture/test_conv2d_nhwc_texture.py 等测试覆盖了纹理路径的卷积用例。

关于内核源码的生成与查看方式，详见后文「高级用法：生成源码检查」。

关于 OpenCLML

OpenCLML 是 Qualcomm 发布的 SDK，提供加速的深度学习算子。这些算子以扩展cl_qcom_ml_ops的形式暴露于标准 OpenCL 规范之上，作为 TVM 的 BYOC（Bring Your Own Codegen）方案集成。

OpenCLML 算子与原生 OpenCL 使用相同的 context，并可入队到与原生 OpenCL 相同的 command queue。TVM 正是利用这一点，避免在回退到原生 OpenCL 时发生上下文切换（context switching）开销。

在仓库中，OpenCLML 的集成包含三块：

编译期 codegen：src/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 编译模型，涉及从 Android（C++ 原生环境或 JNI）调用 TVM runtime 设置输入、获取输出。

高级用法（Advanced Usage）：面向高级用户，例如查看生成源码、修改模块精度等。

以下各节将逐项展开这些阶段。

开发环境搭建：自动（Docker）

TVM 提供了预构建的 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_OPENCL=<Path to OpenCLML SDK> ./tests/scripts/ci.py adreno -i

编译成功后即进入 docker shell。构建会生成两个目录：

build-adreno：主机端 TVM 编译器构建；
build-adreno-target：Android 目标端组件，包含
- libtvm_runtime.so：TVM runtime 库；
- tvm_rpc：RPC 运行时环境工具；
- rtvm：独立的原生部署工具。

由于 docker 环境与主机共享 Android 设备，主机上的 adb 版本需要为1.0.41（docker 内使用同一版本）。可在 docker 内用adb devices检查设备：

user@ci-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_KEY="android"、RPC_TARGET="adreno"、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_OPENCL=<Path 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_OPENCL=<Path to OpenCLML SDK> echo set\(USE_CLML "${ADRENO_OPENCL}"\) >> config.cmake echo set\(USE_CLML_GRAPH_EXECUTOR "${ADRENO_OPENCL}"\) >> config.cmake

Android 目标构建依赖环境变量ANDROID_NDK_HOME。随后执行：

cmake -DCMAKE_TOOLCHAIN_FILE="${ANDROID_NDK_HOME}/build/cmake/android.toolchain.cmake" \ -DANDROID_ABI=arm64-v8a \ -DANDROID_PLATFORM=android-28 \ -DCMAKE_SYSTEM_VERSION=1 \ -DCMAKE_FIND_ROOT_PATH="${ADRENO_OPENCL}" \ -DCMAKE_FIND_ROOT_PATH_MODE_PROGRAM=NEVER \ -DCMAKE_FIND_ROOT_PATH_MODE_LIBRARY=ONLY \ -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_NAME="aarch64-linux-gnu" .. make tvm_runtime tvm_rpc rtvm

关键 CMake 配置项一览

配置项	作用	默认值（见 cmake/config.cmake）
`USE_OPENCL`	启用 OpenCL 后端	`OFF`
`USE_OPENCL_ENABLE_HOST_PTR`	允许 host 指针，配合 CLML 使用	`OFF`
`USE_RPC`	与远程设备通信的 RPC 能力	—
`USE_CPP_RPC`	构建运行在目标端的`tvm_rpc`工具	—
`USE_CPP_RTVM`	构建原生`rtvm`部署工具	—
`USE_GRAPH_EXECUTOR`	图执行器（Android 部署用）	—
`USE_LIBBACKTRACE`	崩溃回溯支持，`AUTO`时自动探测	—
`USE_KALLOC_ALIGNMENT`	OpenCL 分配对齐字节数，Adreno 用 32 位对齐	—
`USE_CLML`	编译 CLML codegen 与运行时模块	`OFF`（cmake/config.cmake）
`USE_CLML_GRAPH_EXECUTOR`	启用 CLML graph runtime，值为 SDK 路径或`ON`/`OFF`，启用时自动打开 OpenCL 作为回退	`OFF`
`ANDROID_ABI`/`ANDROID_PLATFORM`/`MACHINE_NAME`	Android 构建三元组，文档示例为`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_VERSION；USE_CLML_GRAPH_EXECUTOR则会在指定路径下通过find_library查找libOpenCL，并把lib64/libOpenCL.so与lib64/libOpenCL_system.so链接进运行时。

RPC 设置

RPC（Remote 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_PORT>；
device：利用adb push将build-adreno-target/tvm_rpc与libtvm_runtime.so推送到设备的/data/local/tmp/tvm_ci-<USER>，执行adb reverse/adb forward端口转发，再在设备上以--tracker=127.0.0.1:<RPC_PORT> --key=android启动tvm_rpc server（默认监听 5000，-l参数可改），并导出CLML_PROFILING=1、CLML_IS_TUNING_RUN=1、CLML_TUNING_CACHE=clml.bin等 CLML 环境变量；
query：运行python3 -m tvm.exec.query_rpc_tracker --port <RPC_PORT>列出已注册设备。

手动 RPC 设置

更细粒度的手动 RPC 环境搭建可参考 TVM 官方「如何在 Adreno 上部署模型」教程（deploy_model_on_adreno）。

RPC 设置完成后，主机侧得到 rpc-tracker（127.0.0.1，即 rpc 主机）与 rpc 端口9190。

命令行工具：tvmc 全流程

tvmc是 TVM 的命令行工具，可完成模型导入、自动调优、编译以及通过 RPC 部署，选项丰富。

模型导入与调优

下面命令从任意框架导入模型并自动调优。示例使用 Keras 模型，借助 RPC 调优并生成调优日志keras-resnet50.log：

python3 -m tvm.driver.tvmc tune --target="opencl -device=adreno" \ --target-host="llvm -mtriple=aarch64-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 -device=adreno（纹理增强目标）；--target-host使用llvm -mtriple=aarch64-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 \ --target="opencl, 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 \ --target="opencl, 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）、图 json（mod.json）与参数字节文件（mod.params）。

目标端部署与运行

编译模型既可通过 RPC 运行，也可原生部署。RPC 方式：

python3 -m tvm.driver.tvmc run --device="cl" keras-resnet50.tar \ --rpc-key android --rpc-tracker 127.0.0.1:9190 --print-time

tvmc 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 接口时，先加载模型得到TVMCModel；TVMCModel用于自动调优并产出 tuning cache；编译过程结合TVMCModel与（可选的）tuning cache 生成TVMCPackage；TVMCPackage可保存到文件系统，也可直接用于目标端部署运行。

保存的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）、图 json（mod.json）与参数（mod.params），TVMPackage归档中也包含这些内容。对 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 专用变换 API：

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": "float16", }, )

tvm.driver.tvmc.transform.apply_graph_transforms是ToMixedPrecisionpass 的简化封装，用于获得所需精度。随后可按任意习惯方式编译：

with tvm.transform.PassContext(opt_level=3): lib = relay.build( mod, target_host=target_host, target=target, params=params )

使用tvmcPython 接口时，以下参数启用 float16 精度转换：

mixed_precision = True, mixed_precision_ops = ["nn.conv2d", "nn.dense"], mixed_precision_calculation_type = "float16", mixed_precision_acc_type = "float16"

tvmc命令行接口对应选项：

--mixed-precision --mixed-precision-ops nn.conv2d nn.dense --mixed-precision-calculation-type float16 --mixed-precision-acc-type float16

float16_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_level=3): lib = relay.build( mod, target_host=target_host, target=target, params=params )

tvmcPython 接口参数：

mixed_precision = True, mixed_precision_ops = ["nn.conv2d", "nn.dense"], mixed_precision_calculation_type = "float16", mixed_precision_acc_type = "float32"

tvmc命令行接口选项：

--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_transforms（python/tvm/driver/tvmc/transform.py）在mixed_precision为真时调用它完成图变换。

企业官网建设流程全解析