关于算法:交叉编译和-RPC

本篇文章译自英文文档 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 中文站