交叉编译和RPC

本文翻译自Cross Compilation and RPC — tvm 0.9.dev0 documentation

本教程介绍在TVM中交叉编译和使用RPC进行远程设备执行。

使用交叉编译和RPC，您可以在本地机器上编译程序，然后在远程设备上运行它。当远程设备资源有限时，它是有用的，如树莓派和移动平台。在本教程中，我们将使用树莓派的CPU示例和萤火虫- rk3399的OpenCL示例。

在设备上编译TVM运行时

第一步是在远程设备上构建TVM运行时。

注意：本节和下一节中的所有指令都应该在目标设备上执行，例如树莓派。我们假设目标是运行Linux系统。

由于我们在本地机器上进行编译，因此远程设备只用于运行生成的代码。我们只需要在远程设备上构建TVM运行时。

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

成功构建运行时之后，我们需要在~/.bashrc文件中设置环境变量。我们可以使用vi ~/.bashrc命令编辑~/.bashrc，并添加下面的行(假设您的TVM目录在~/tvm中):

export PYTHONPATH=$PYTHONPATH:~/tvm/python

执行source ~/.bashrc更新环境变量

在设备上启动RPC Server

 在远程设备上运行以下命令(本例中为Raspberry Pi)启动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

 在本地机器上声明和交叉编译内核

注意：现在我们回到本地机器，它已经安装了一个完整的TVM(包括LLVM)。

这里我们将在本地机器上声明一个简单的内核:

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的目标应该是' 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")
# save the lib at a local temp folder
temp = utils.tempdir()
path = temp.relpath("lib.tar")
func.export_library(path)

注意：

如果要在实际的远程设备运行本教程，请将local_demo更改为False，并在构建中使用适合设备的目标三元组替换target。不同的设备的目标三元组可能是不一样的。例如，对于树莓派3B来说是'llvm -mtriple=armv7l-linux-gnueabihf'，对于RK3399来说是'llvm -mtriple=aarch64-linux-gnu'。

一般情况下，您可以通过在设备上运行gcc -v来查询目标，并查找以Target:开头的行(尽管它可能仍然是一个松散的配置)。

除了-mtriple，您还可以设置其他编译选项，比如:

 -mcpu=<cpuname>

在当前架构中指定要为其生成代码的特定芯片。默认情况下，这是从目标三元组推断出来的，并根据当前体系结构自动检测到。

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

覆盖或控制目标的特定属性，例如是否启用SIMD操作。默认属性是由当前CPU设置的。要获得可用属性列表，你可以这样做:

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

这些选项和llc是一致的。建议设置目标三元组和特征集来包含特定的可用特性，这样我们就可以充分利用单板的特性。你可以在LLVM交叉编译指南中找到更多关于交叉编译属性的细节。

使用RPC运行远程CPU内核

 我们将展示如何在远程设备上运行生成的CPU内核。首先，我们从远程设备获取RPC会话。 

if local_demo:
    remote = rpc.LocalSession()
else:
    # The following is my environment, change this to the IP address of your target device
    host = "10.77.1.162"
    port = 9090
    remote = rpc.connect(host, port)

将库上传到远程设备，然后调用设备本地编译器来重新链接它们。现在func是一个远程模块对象

remote.upload(path)
func = remote.load_module("lib.tar")

# create arrays on the remote device
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)
# the function will run on the remote device
func(a, b)
np.testing.assert_equal(b.numpy(), a.numpy() + 1)

 当您想要评估远程设备上内核的性能时，必须注意网络开销。time_evaluator将返回一个会运行待测试函数多次的远程函数，测量每次在远程设备上运行的开销，并返回测量的成本。这其中不包括网络开销。

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

输出

1.299e-07 secs/op

通过RPC远程运行OpenCL内核

远程OpenCL设备的工作流程几乎与上述相同。您可以定义内核、上传文件和通过RPC运行。

注意：

树莓派不支持OpenCL，以下代码在Firefly-RK3399上测试。您可以按照本教程设置RK3399的操作系统和OpenCL驱动程序。

此外，我们还需要在rk3399板上启用OpenCL来构建运行时。在TVM根目录下执行 

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

The following function shows how we run an OpenCL kernel remotely

def run_opencl():
    # NOTE: This is the setting for my rk3399 board. You need to modify
    # them according to your environment.
    opencl_device_host = "10.77.1.145"
    opencl_device_port = 9090
    target = tvm.target.Target("opencl", host="llvm -mtriple=aarch64-linux-gnu")

    # create schedule for the above "add one" compute declaration
    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)

    # export and upload
    path = temp.relpath("lib_cl.tar")
    func.export_library(path)
    remote.upload(path)
    func = remote.load_module("lib_cl.tar")

    # run
    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远程上传和运行内核。