部署到 Adreno™ GPU

介绍​

Adreno™ 是由高通开发并用于许多 SoC 的图形处理单元（GPU）半导体 IP 核系列。

Adreno™ GPU 能够加快复杂几何图形的渲染，在供给高功能图形和丰富的用户体验的同时具有很低的功耗。

TVM 运用 TVM 的原生 OpenCL 后端 和 OpenCLML 后端以支撑加快 Adreno™ GPU 上的深度学习。TVM 的原生 OpenCL 后端经过结合纹路内存运用和 Adreno™ 友爱布局来改进 Adreno™ 。 OpenCLML 是由高通发布的 SDK ，供给了大多数深度学习运算符的内核加快库。

本攻略展示以下方面的不同设计

• OpenCL 后端增强
• 关于 OpenCLML
• 构建与布置

OpenCL 后端增强​

TVM 的 OpenCL 后端已被增强以运用 Adreno™ 特征功用，如

• 纹路内存运用。
• Adreno™ 友爱的激活布局。
• 全新的调度以加快上述功用。

Adreno™ 的一个优势是对纹路的奇妙处理。目前，TVM 能够经过对 Adreno™ 的纹路支撑取得益处。下图显现了 Adreno™ A5x 架构。

图1 用于 OpenCL Adreno™ A5x 架构的高层次概览

来历： Qualcomm Adreno™ GPU 的 OpenCL 优化和最佳实践

运用纹路的原因：

• 纹路处理器（TP）具有专用的 L1 缓存，它是只读缓存，并存储了从 2 级缓存（L2）中提取的用于纹路操作的数据（主要原因）
• 图画边界的处理已内置。
• 支撑多种图画格局和数据类型组合，支撑主动格局转化

整体而言，与根据 OpenCL 缓冲区的解决方案相比，运用纹路能够取得明显的功能提升。

一般，咱们将方针指定为target="opencl"，以生成如下所示的内核的常规OpenCL方针。

   __kernel void tvmgen_default_fused_nn_conv2d_kernel0(__global float* restrict p0, __global double* restrict p1, __global float* restrict conv2d_nhwc) {
   // body..

上述 OpenCL 内核界说有__global float*类型指针，它们实质上是 OpenCLbuffer方针。

经过修正方针界说为target="opencl -device=adreno"启用根据纹路的增强后，咱们能够看到生成的内核运用纹路支撑的 OpenCL 图画方针，如下所示。

   __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个元素能够更有效地运用硬件。

有关生成和检查内核源的更多具体信息，请参阅高档用法。

关于 OpenCLML​

OpenCLML 是高通发布的 SDK ，供给了更快的深度学习运算符。这些运算符作为规范 OpenCL 规范的扩展cl_qcom_ml_ops公开。有关更多具体信息，请参见Accelerate your models with our OpenCL ML SDK。

OpenCLML 已集成到 TVM ,作为BYOC解决方案。 OpenCLML 运算符能够运用相同的上下文，并能够加入到相同被原生 OpenCL 运用的指令行列中。咱们运用了这一点防止在回退到原生 OpenCL 时的上下文切换。

Adreno™ 下的 TVM​

本节供给有关构建和布置模型到 Adreno™ 方针机的办法阐明。 Adreno™ 是经过 ADB 衔接与主机衔接的长途方针。在这里布置已编译的模型需求在主机和方针上运用一些东西。

TVM 供给了简单的、用户友爱的指令行东西以及专为开发者设计的 Python API 接口，可用于各种进程，如主动调整、构建和布置。

上图展示了以下各个阶段的通用流程：

• 导入模型：在此阶段，咱们从知名结构（如 TensorFlow、PyTorch、ONNX 等）导入模型。该阶段将给定的模型转化为 TVM 的 relay 模块格局。或许也能够经过运用 TVM 的操作库手动构建 relay 模块。此处生成的 TVM 模块的是独立于图的表明方法的方针。
• 主动调整：在此阶段，咱们调整特定于方针的 TVM 生成的内核。主动调整进程需求方针设备的有效性，在如 Android 设备上的 Adreno™ 这样的长途方针中，咱们运用 RPC 设置进行通信。本攻略的后续部分将具体介绍 Android 设备的 RPC 设置。主动调整不是模型编译的必需进程，但关于取得 TVM 生成的内核的最佳功能是必要的。
• 编译：在此阶段，咱们为特定方针编译模型。鉴于咱们在前一阶段主动调整了模块， TVM 编译运用调整日志以生成功能最佳的内核。 TVM 编译进程中产生包括内核的共享库、以 json 格局界说的图和以 TVM 特定格局的二进制参数文件。
• 布置（或测试运转）到方针：在此阶段，咱们在方针上运转 TVM 编译输出。布置能够从 Python 环境中运用 RPC 设置。也能够运用 TVM 的本地东西进行，该东西能够为 Android 进行本地二进制穿插编译。在此阶段，咱们能够在 Android 方针上运转已编译的模型，并对输出的正确性和功能方面进行单元测试。
• 运用集成：本阶段触及将 TVM 编译的模型集成到运用程序中。咱们在这里讨论怎么从 Android（cpp 本地环境或 JNI）中设置输入和获取输出的 TVM 运转时接口。
• 高档用法：本部分涵盖了高档用户感兴趣的主题，如检查生成的源代码、更改模块的精度等。

此教程以下各节将涵盖上述内容。

• 开发环境
• RPC 设置
• 指令行东西
• Python 接口
• 运用集成
• 高档用法

主动开发环境设置​

TVM 供给了一个预界说的 Docker 容器环境，其间包括了所有入门所需的先决条件。假如您想要更多地控制依靠联系，请参阅手动环境设置。

关于 Docker 设置，先决条件仅仅主机上有 Docker 东西。

以下指令能够构建 Adreno 的 Docker 镜像：

./docker/build.sh ci_adreno
docker tag tvm.ci_adreno ci_adreno

现在，咱们能够运用以下指令构建主机和方针东西：

./tests/scripts/ci.py adreno -i

要运用 OpenCLML SDK 构建TVM，在构建时需求导出 OpenCLML SDK，如下所示：

export ADRENO_OPENCL=<OpenCLML SDK途径>
./tests/scripts/ci.py adreno -i

成功编译后，您将进入 Docker shell。构建将生成两个文件夹：

• build-adreno：主机端TVM编译器构建。

• build-adreno-target：包括 Android 方针组件。

  • libtvm_runtime.so：TVM 运转时库
  • tvm_rpc：rpc 运转时环境东西
  • rtvm：独立的原生东西

在运用Docker环境时， Android 设备与主机共享，所以在主机上需求装置 ADB 版别为1.0.41，由于 Docker 运用相同的版别。

您还能够在 Docker 环境中检查 ADB 设备的可用性：

user@ci-adreno-fpeqs:~$ adb devices
List of devices attached
aaaabbbb	device
ccccdddd	device

手动开发环境设置​

手动构建进程需求构建主机和方针组件。

以下指令将装备主机编译：

mkdir -p build
cd build
cp ../cmake/config.cmake .
# 启用 RPC 功用以与长途设备通信。
echo set(USE_RPC ON) >> config.cmake
# 咱们在主机（x86）上运用图履行器验证模型。
echo set(USE_GRAPH_EXECUTOR ON) >> config.cmake
# 启用溢出时的回溯以获取更多调试信息。
echo set(USE_LIBBACKTRACE AUTO) >> config.cmake
# 方针主机将是 llvm。
echo set(USE_LLVM ON) >> config.cmake

此外，咱们能够推送以下装备条目，以运用 OpenCLML 支撑进行编译：

export ADRENO_OPENCL=<OpenCLML SDK途径>
echo set(USE_CLML ${ADRENO_OPENCL}) >> config.cmake

现在咱们能够像下面这样构建：

cmake ..
make

终究，咱们能够导出 Python 途径：

export PYTHONPATH=$TVM_HOME/python:${PYTHONPATH}
python3 -c "import tvm" # 验证tvm Python包

现在，咱们能够运用以下装备来装备和构建方针组件。方针构建需求装置 Android NDK。

• 在此处阅读有关Android NDK 装置的文档：developer.android.com/ndk
• 要获取 adb 东西的拜访权限，您能够在此处检查Android Debug Bridge 装置：developer.android.com/studio/comm…

mkdir -p build-adreno
cd build-adreno
cp ../cmake/config.cmake .
# 启用 OpenCL 后端。
echo set(USE_OPENCL ON) >> config.cmake
# 启用 RPC 功用。
echo set(USE_RPC ON) >> config.cmake
# 构建在方针设备上运转的 tvm_rpc 东西。
echo set(USE_CPP_RPC ON) >> config.cmake
# 构建本机 rtvm 布置东西。
echo set(USE_CPP_RTVM ON) >> config.cmake
# 咱们在像 Android 这样的设备上运用图履行器。
echo set(USE_GRAPH_EXECUTOR ON) >> config.cmake
# 在或许的情况下启用回溯。
echo set(USE_LIBBACKTRACE AUTO) >> config.cmake
# Adreno 支撑 OpenCL 分配的32位对齐而不是64位。
echo set(USE_KALLOC_ALIGNMENT 32) >> config.cmake
# Android 构建相关界说。
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=<OpenCLML SDK途径>
echo set(USE_CLML "${ADRENO_OPENCL}") >> config.cmake
echo set(USE_CLML_GRAPH_EXECUTOR "${ADRENO_OPENCL}") >> config.cmake

关于 Android 方针构建，ANDROID_NDK_HOME是一个依靠项，咱们应该在环境变量中设置相同的依靠项。以下指令将构建 Adreno™ 方针组件：

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

RPC 设置​

RPC 设置允许经过 TCP/IP 网络接口拜访长途方针。RPC 设置关于主动调整阶段是必不可少的，由于调整触及在实在设备上运转主动生成的内核，并经过运用机器学习办法对其进行优化。请参阅运用模板和 AutoTVM 进行主动调整以获取有关 AutoTVM 的具体信息。

RPC 设置也可用于经过 python 接口或来自主机设备的tvmc东西将已编译的模型布置到长途设备。

RPC 设置有多个组件，如下所示。

TVM Tracker: TVM Tracker 是主机端的看护程序，办理长途设备并为主机端运用程序供给服务。运用程序能够衔接到此跟踪器并获取长途设备句柄以进行通信。

TVM RPC: TVM RPC 是在长途设备上运转的本地运用程序（本情况下是 Android ），并向在主机上运转的 TVM Tracker 注册自身。

因而，关于根据 RPC 的设置，咱们将在主机和方针设备上运转上述组件。以下部分解说了怎么手动设置相同的内容，以及怎么在运用主动化东西的 Docker 环境中设置相同的内容。

**主动化 RPC 设置:**此处，咱们将解说怎么在 Docker 环境中设置 RPC 。

以下指令在 Docker 环境中发动 tracker，其间 tracker 监听端口 9190。

./tests/scripts/ci.py adreno -i # 在 anreno docker 上发动一个新的 shell
source tests/scripts/setup-adreno-env.sh -e tracker -p 9190

现在，以下指令能够在具有 IDabcdefgh的长途 Android 设备上运转 TVM RPC 。

./tests/scripts/ci.py adreno -i # 在adreno docker上发动一个新的shell。
source tests/scripts/setup-adreno-env.sh -e device -p 9190 -d abcdefgh

此外，以下指令可用于在任何其他 Docker 终端上查询 RPC 设置具体信息。

./tests/scripts/ci.py adreno -i # 在 adreno docker 上发动一个新的 shell。
source tests/scripts/setup-adreno-env.sh -e query -p 9190

**手动 RPC 设置：**请参阅教程在 Adreno 上布置模型了解手动 RPC 环境设置。

这些 RPC 设置完成后，咱们在主机127.0.0.1（rpc-tracker）上具有 rpc-tracker，端口9190（rpc-port）可用。

.. _commandline_interface:

指令行东西​

此处咱们运用指令行东西描绘整个编译进程。 TVM 具有指令行实用程序tvmc，用于履行模型导入、主动调整、编译和 rpc 布置。tvmc有许多选项可供探究和测验。

**模型导入和调整：**运用以下指令从任何结构导入模型并对其进行主动调整。此处咱们运用来自 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

模型编译:

运用以下指令编译模型并生成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 卸载时，咱们需求如下添加方针clml。调整日志关于 OpenCLML 卸载相同有效，由于 OpenCL 途径是没有经过 OpenCLML 途径的任何运算符的回退选项。调整日志将用于这类运算符。

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" --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 shared lib（ mod.so ）、graph json（ mod.json ）和参数二进制（ mod.params ）。

在方针上布置和运转:

能够经过 RPC 方法和本地布置方法实现在 Android 方针上运转编译后的模型。

咱们能够运用下面的 tvmc 指令经过 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供给了更多选项，用于以各种形式（如 fill 、 random 等）初始化输入。

根据布置的tvmc一般是经过 RPC 设置从长途主机快速验证在方针上编译的模型的一种方法。

一般，生产环境运用原生布置环境，如 Android JNI 或 CPP 原生环境。此处咱们需求运用穿插编译的tvm_runtime接口来布置 tvm 编译输出，即TVMPackage。

TVM 有一个名为rtvm的独立东西，用于在 ADB shell 上本地布置和运转模型。构建进程会在 build-adreno-target 下生成此东西。有关有关此东西的更多具体信息，请参阅rtvm。

在将其集成到现有 Android 运用程序中时， TVM 有多种挑选。关于 JNI 或 CPP 原生，咱们能够运用C Runtime API。 您还能够参阅rtvm的简化接口” ref=”nofollow noopener noreferrer”>TVMRunner。

Python 接口​

该节解说怎么运用 Python 接口导入、主动调整、编译和运转模型的进程。 TVM 经过tvmc笼统供给高档接口，以及初级的 Relay API 。咱们将具体讨论这两者。

TVMC 接口：

在运用tvmcPython接口时，咱们首要加载一个生成TVMCModel的模型。TVMCModel将用于主动调整以生成调整缓存。编译进程运用TVMCModel和调整缓存（可选）来生成TVMCPackage。现在TVMCPackage将保存到文件系统，或许能够用于布置和在方针设备上运转。

请参阅相应的教程怎么运用 TVMC 在 Adreno 上布置预训练模型。

保存的TVMCPackage也能够用于运用rtvm东西进行原生布置。

此外，请参阅运用 TVMC 编译和优化模型文档，了解有关 API 接口的更多具体信息。

Relay 接口：

Relay API 接口供给了对 TVM 编译器接口的初级 API 拜访。与tvmc接口类似，Relay API 接口供给了各种前端 API，用于将模型转化为 RelayModule。RelayModule将用于所有种类的转化，如精度转化、CLML 卸载和其他自界说转化（假如有）。生成的 Module 也将用于主动调整。终究，咱们运用relay.buildAPI 生成库模块。从这个库模块，咱们能够导出编译产物，如模块共享库（mod.so）、参数（mod.params）和 json 图（mod.json）。这个库模块将用于创立图形运转时以在方针设备上布置和运转。

请参阅教程怎么在 Adreno 上布置预训练模型，其间逐渐解说了相同的内容。

此外，TVM 还经过TVM4J支撑 Java 接口。

运用程序集成​

TVM 编译输出以模块共享库（mod.so）、图形 json （mod.json）和参数（mod.params）的方法表明。 TVMPackage 的存档表明也包括相同的内容。

一般关于任何 Android 运用程序集成来说，根据 CPP/C 的接口就足够了。

TVM 原生地公开了c_runtime_api，用于加载 TVM 编译的模块并运转相同的模块。

或许，用户还能够参阅cpp_rtvm中的TVMRunner接口，以取得相同的进一步简化的版别。

.. _advanced_usage:

高档用法​

本节具体介绍在 TVM 上 运用 Adreno™ 方针机时的一些高档用法和其他信息。

生成源码检查

除了规范的 tvm 编译产物（kernel 库 mod.so 、图形 mod.json 和参数 mod.params）之外，咱们还能够从 lib handle 生成 opencl kernel 源码、clml 卸载图等。TVM 编译的输出被组织为一个 TVM 模块，其间包括许多其他导入的 TVM 模块。

下面的代码段能够将 CLML子 图以 json 格局转储。

# 查找 "clml" 类型的导入模块。
clml_modules = list(filter(lambda mod: mod.type_key == "clml", lib.get_lib().imported_modules))
# 循环遍历所有 clml 子图并转储以 json 格局的 CLML 子图。
for cmod in clml_modules:
   print("CLML Src:", cmod.get_source())

类似地，下面的代码段能够从编译的 TVM 模块中提取 opencl

# 类似地，咱们能够转存敞开的内核源码如下所示
# 寻觅 "opencl" 类型模块导入.
opencl_modules = list(filter(lambda mod: mod.type_key == "opencl", lib.get_lib().imported_modules))
# 现在为每一个 Open CL 的方针子图转储内核源码
for omod in opencl_modules:
    print("OpenCL Src:", omod.get_source())

精度​

挑选特定作业负载的正确精度能够极大地进步解决方案的效率，将精度和速度的初始平衡转移到问题的首要一侧。

咱们能够挑选float16，float16_acc32（混合精度），float32（规范）。

Float16​

为了充分运用 GPU 硬件功用并运用半精度核算和内存办理的优势，咱们能够将具有浮点运算的原始模型转化为运用半精度运算的模型。挑选较低的精度将积极地影响模型的功能，但也或许导致模型精度下降。

要履行转化，您需求在任何一个前端生成 Relay 模块的同时调用 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是对ToMixedPrecision传递的简化 API ，为了获取所需的精度。

然后咱们能够以任意便利的方法编译咱们的模型：

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 (混合精度)​

ToMixedPrecision进程遍历网络并将网络拆分为处理 float 或 float16 数据类型的操作簇。这些簇由三种类型的操作界说：

• 总是转化为 float16 数据类型的操作
• 假如后面跟着转化簇，能够转化的操作
• 从不转化为 float16 数据类型的操作

此列表在 ToMixedPrecision 的实现relay/transform/mixed_precision.py中界说，用户能够掩盖。

ToMixedPrecision办法是将 FP32 的 Relay 图转化为 FP16 版别（运用 FP16 或 FP32 累积数据类型）的进程。进行此转化关于减小模型巨细很有用，由于它将权重的期望巨细折半（FP16_acc16情况）。

ToMixedPrecision进程的运用简化为以下代码：

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

tvm.driver.tvmc.transform.apply_graph_transforms是对ToMixedPrecision进程的简化 API ，以获取所需的精度。

然后咱们能够以任何便利的方法编译咱们的模型：

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

在运用tvmcPython 接口时，下面的参数启用到 float16_acc32 的精度转化：

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