TVM:交叉編譯和RPC
之前我們介紹了 TVM 的安裝、本機demo和樹莓派遠程demo。本文將介紹了在 TVM 中使用 RPC 進行交叉編譯和遠程設備執行。
通過交叉編譯和 RPC,我們可以在本地機器上編譯程序,然后在遠程設備上運行它。 當遠程設備資源有限時很有用,例如 Raspberry Pi 和移動平臺。 在本文中,我們將使用 Raspberry Pi 作為 CPU 示例,使用 Firefly-RK3399 作為 OpenCL 示例。
在遠程設備上構建 TVM Runtime
首先我們要在遠程設備上編譯安裝 TVM Runtime,注意這里我們對模型的編譯是在本機進行的,而遠程設備只需要運行模型即可,因此只需要構建 TVM Runtime。
注意:本節和下一節中的所有指令都應在目標設備上執行,例如樹莓派。 我們假設它運行著 Linux。
git clone --recursive https://github.com/apache/tvm tvm
cd tvm
make runtime -j2
將 Python 路徑添加到環境變量:
export PYTHONPATH=$PYTHONPATH:/path/to/tvm/python
在遠程設備上設置 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 npimport tvm
from tvm import te
from tvm import rpc
from tvm.contrib import utilsn = 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’
。如果真的有一個遠程設備樹莓派的話可以將下面的 local_demo
改為 False
,否則還是保留為 True
使得本 demo 可以正常運行。
local_demo = Trueif 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 triple)。不同設備的(target triple)可能不同。例如,對于 Raspberry Pi 3B,它是'llvm -mtriple=armv7l-linux-gnueabihf'
,對于 RK3399,它是'llvm -mtriple=aarch64-linux-gnu'
。通常,我們可以通過在您的設備上運行
gcc -v
來查詢目標,并查找以Target:
開頭的行(盡管它可能仍然是一個寬松的配置。)除了
-mtriple
,您還可以設置其他編譯選項,例如:
-mcpu=<cpuname>
指定當前架構中的特定芯片以為其生成代碼。默認情況下,這是從 target triple 推斷出來的,并自動檢測到當前架構。-mattr=a1,+a2,-a3,…
覆蓋或控制目標的特定屬性,例如是否啟用 SIMD 操作。默認屬性集由當前 CPU 設置。要獲取可用屬性列表,我們可以執行以下操作:
llc -mtriple=\<your device target triple\> -mattr=help
這些選項與 llc 一致。建議將 target triple和功能集設置為包含可用的特定功能,以便我們可以充分利用板的功能。可以到 LLVM 交叉編譯指南中找到有關交叉編譯屬性的更多詳細信息。
通過RPC在遠程運行CPU核
接下來是如何將生成的CPU核運行在遠程設備上,首先我們建立與遠程設備的 RPC 會話。
if local_demo:remote = rpc.LocalSession()
else:# 下面的IP是筆者的,請大家換成自己的遠程設備的IPhost = "10.206.105.111"port = 9090remote = 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 將會運行在遠程設備上
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.178e-07 secs/op
通過RPC在遠程設備上運行OpenCL核
對于遠程 OpenCL 設備,整個流程和上面幾乎是一樣的。我們定義自己的和核,上傳文件,并通過 RPC 運行。
注意:樹莓派并不支持 OpenCL,以下代碼是在 Firefly-RK3399 上進行測試的。大家可以通過這個教程來為 RK3399 配置操作系統和 OpenCL。
同樣我們需要再 RK3399 上構建 TVM Runtiime(注意要在 config.cmake
中啟用 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 = 9090target = tvm.target.Target("opencl", host="llvm -mtriple=aarch64-linux-gnu")# 為上述 'add one' 計算聲明創建 schedules = 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 服務器
- 設置設備的 target配置 并在本機上交叉編譯核
- 通過 RPC 遠程上傳并運行核