本篇文章译自英文文档 Cross Compilation and RPC 作者是 Ziheng Jiang，Lianmin Zheng。更多 TVM 中文文档可拜访 →TVM 中文站

本教程介绍了如安在 TVM 中运用 RPC 进行穿插编译和长途设备履行。

使用穿插编译和 RPC，能够完成程序在本地机器编译，在长途设备运转。这个特性在长途设备资源有限时（如在树莓派和移动平台上）尤其有用。本教程将把树莓派作为 CPU 示例，把 Firefly-RK3399 作为 OpenCL 示例进行演示。

在设备上构建 TVM Runtime

首要在长途设备上构建 TVM runtime。

留意 本节和下一节中的所有指令都应在方针设备（例如树莓派）上履行。假定方针设备运转 Linux 体系。

由于在本地机器上只做编译，而长途设备用于运转生成的代码。所以只需在长途设备上构建 TVM runtime。

git clone --recursive https://github.com/apache/tvm tvm
cd tvm
make runtime -j2

成功构建 runtime 后，要在 ~/.bashrc 文件中设置环境变量。能够用 vi ~/.bashrc指令编辑 ~/.bashrc，在这个文件里添加下面这行代码（假定 TVM 目录在 ~/tvm 中）：

export PYTHONPATH=$PYTHONPATH:~/tvm/python

履行 source ~/.bashrc 来更新环境变量。

在设备上设置 RPC 服务器

在长途设备（本例为树莓派）上运转以下指令来发动 RPC 服务器：

python -m tvm.exec.rpc_server --host 0.0.0.0 --port=9090

看到下面这行提示，则表示 RPC 服务器已成功发动。

INFO:root:RPCServer: bind to 0.0.0.0:9090

在本地机器上声明和穿插编译内核

补白 现在回到本地机器（已经用 LLVM 安装了完整的 TVM）。

在本地机器上声明一个简单的内核：

import numpy as np
import tvm
from tvm import te
from tvm import rpc
from tvm.contrib import utils
n = tvm.runtime.convert(1024)
A = te.placeholder((n,), name="A")
B = te.compute((n,), lambda i: A[i] + 1.0, name="B")
s = te.create_schedule(B.op)

然后穿插编译内核。关于树莓派 3B，target 是“llvm -mtriple=armv7l-linux-gnueabihf”，但这儿用的是“llvm”，使得本教程能够在网页构建服务器上运转。请参阅下面的详细说明。

local_demo = True
if local_demo:
    target = "llvm"
else:
    target = "llvm -mtriple=armv7l-linux-gnueabihf"
func = tvm.build(s, [A, B], target=target, name="add_one")
# 将 lib 存储在本地临时文件夹
temp = utils.tempdir()
path = temp.relpath("lib.tar")
func.export_library(path)

补白

要使本教程运转在真实的长途设备上，需要将 local_demo 改为 False，并将 build 中的 target 替换为适合设备的 target 三元组。不同设备的 target 三元组或许不同。例如，关于树莓派 3B，它是 llvm -mtriple=armv7l-linux-gnueabihf；关于 RK3399，它是 llvm -mtriple=aarch64-linux-gnu。

一般，能够在设备上运转 gcc -v 来查询 target，寻觅以 Target 最初的行：（虽然它或许仍然是一个松散的装备。）
除了 -mtriple，还可设置其他编译选项，例如：

  • -mcpu=< cpuname>

指定生成的代码运转的芯片架构。默许状况这是从 target 三元组推断出来的，并自动检测到当前架构。

  • -mattr=a1,+a2,-a3,…

掩盖或操控 target 的指定特点，例如是否启用 SIMD 操作。默许特点集由当前 CPU 设置。要获取可用特点列表，履行：

  llc -mtriple=<your device target triple> -mattr=help

这些选项与 llc 共同。主张设置 target 三元组和功用集，使其包括可用的特定功用，这样我们能够充分使用单板的功用。检查 LLVM 穿插编译攻略获取有关穿插编译特点的详细信息。

经过 RPC 长途运转 CPU 内核

下面将演示如安在长途设备上运转生成的 CPU 内核。首要，从长途设备获取 RPC 会话：

if local_demo:
    remote = rpc.LocalSession()
else:
    # 下面是我的环境，将这个换成你方针设备的 IP 地址
    host = "10.77.1.162"
    port = 9090
    remote = rpc.connect(host, port)

将 lib 上传到长途设备，然后调用设备的本地编译器从头链接它们。其间 func 是一个长途模块对象。

remote.upload(path)
func = remote.load_module("lib.tar")
# 在长途设备上创建数组
dev = remote.cpu()
a = tvm.nd.array(np.random.uniform(size=1024).astype(A.dtype), dev)
b = tvm.nd.array(np.zeros(1024, dtype=A.dtype), dev)
# 这个函数将在长途设备上运转
func(a, b)
np.testing.assert_equal(b.numpy(), a.numpy() + 1)

要想评价内核在长途设备上的性能，避免网络开销很重要。time_evaluator 返回一个长途函数，这个长途函数多次运转 func 函数，并测验每一次在长途设备上运转的成本，然后返回测验的成本（不包括网络开销）。

time_f = func.time_evaluator(func.entry_name, dev, number=10)
cost = time_f(a, b).mean
print("%g secs/op" % cost)

输出结果：

1.369e-07 secs/op

经过 RPC 长途运转 OpenCL 内核

长途 OpenCL 设备的作业流程与上述内容根本相同。能够界说内核、上传文件，然后经过 RPC 运转。

补白
树莓派不支持 OpenCL，下面的代码是在 Firefly-RK3399 上测验的。能够按照 教程 为 RK3399 设置 OS 及 OpenCL 驱动程序。
在 rk3399 板上构建 runtime 也需启用 OpenCL。在 TVM 根目录下履行：

cp cmake/config.cmake .
sed -i "s/USE_OPENCL OFF/USE_OPENCL ON/" config.cmake
make runtime -j4

下面的函数展现了怎么长途运转 OpenCL 内核：

def run_opencl():
    # 留意：这是 rk3399 板的设置。你需要依据你的环境进行修改
    opencl_device_host = "10.77.1.145"
    opencl_device_port = 9090
    target = tvm.target.Target("opencl", host="llvm -mtriple=aarch64-linux-gnu")
    # 为上面的核算声明 "add one" 创建 schedule
    s = te.create_schedule(B.op)
    xo, xi = s[B].split(B.op.axis[0], factor=32)
    s[B].bind(xo, te.thread_axis("blockIdx.x"))
    s[B].bind(xi, te.thread_axis("threadIdx.x"))
    func = tvm.build(s, [A, B], target=target)
    remote = rpc.connect(opencl_device_host, opencl_device_port)
    # 导出并上传
    path = temp.relpath("lib_cl.tar")
    func.export_library(path)
    remote.upload(path)
    func = remote.load_module("lib_cl.tar")
    # 运转
    dev = remote.cl()
    a = tvm.nd.array(np.random.uniform(size=1024).astype(A.dtype), dev)
    b = tvm.nd.array(np.zeros(1024, dtype=A.dtype), dev)
    func(a, b)
    np.testing.assert_equal(b.numpy(), a.numpy() + 1)
    print("OpenCL test passed!")

总结

本教程介绍了 TVM 中的穿插编译和 RPC 功用。

  • 在长途设备上设置 RPC 服务器。
  • 设置方针设备装备，使得可在本地机器上穿插编译内核。
  • 经过 RPC API 长途上传和运转内核。

下载 Python 源代码
下载 Jupyter Notebook

以上就是该文档的全部内容，检查更多 TVM 中文文档，请拜访→TVM 中文站