TVM:交叉編譯和RPC

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 遠程上傳并運行核

本文來自互聯網用戶投稿,該文觀點僅代表作者本人,不代表本站立場。本站僅提供信息存儲空間服務,不擁有所有權,不承擔相關法律責任。
如若轉載,請注明出處:http://www.pswp.cn/news/532676.shtml
繁體地址,請注明出處:http://hk.pswp.cn/news/532676.shtml
英文地址,請注明出處:http://en.pswp.cn/news/532676.shtml

如若內容造成侵權/違法違規/事實不符,請聯系多彩編程網進行投訴反饋email:809451989@qq.com,一經查實,立即刪除!

相關文章

2.3單鏈表的基本使用及其cpp示例

2.3線性表的鏈式表現與實現 2.3.1.1單鏈表 【特點&#xff1a; *用一組任意的存儲單元存儲線性表的數據元素 *利用指針實現用不同相鄰的存儲單元存放邏輯上相鄰的元素 *每個元素ai&#xff0c;除存儲本身信息外&#xff0c;還存儲其直接后繼的元素&#xff08;后一個元素的地址…

TVM:簡介

TVM&#xff1a;簡介概述 Apache TVM 是一個用于 CPU、GPU 和機器學習加速器的開源機器學習編譯器框架。它旨在使機器學習工程師能夠在任何硬件后端上高效地優化和運行計算。本教程的目的是通過定義和演示關鍵概念&#xff0c;引導您了解 TVM 的所有主要功能。新用戶應該能夠從…

2.3.3單鏈表的雙向鏈表

2.3.3雙向鏈表 插入、刪除 指在前驅和后驅方向都能游歷&#xff08;遍歷&#xff09;的線性鏈表 雙向鏈表的每個結點有兩個指針域 【結構】&#xff1a;prior data next 雙鏈表通常采用帶頭結點的循環鏈表形式 可理解為首位相接的數據“圈”&#xff0c;每個結點都可以向前…

nvidia-smi 命令詳解

nvidia-smi 命令詳解 簡介 nvidia-smi - NVIDIA System Management Interface program nvidia smi&#xff08;也稱為NVSMI&#xff09;為來自 Fermi 和更高體系結構系列的 nvidia Tesla、Quadro、GRID 和 GeForce 設備提供監控和管理功能。GeForce Titan系列設備支持大多數…

2.4一元多項式的表示及相加,含cpp算法

2.4一元多項式的表示及相加 n階多項式的表示&#xff1a; n階多項式有n1項 指數按升冪排序 【 優點&#xff1a; 多項式的項數可以動態增長&#xff0c;不存在存儲溢出的問題插入&#xff0c;刪除方便&#xff0c;不移動元素 【表示&#xff1a; 有兩個數據域&#xff0c;一…

TVM:使用Tensor Expression (TE)來處理算子

TVM&#xff1a;使用Tensor Expression (TE)來處理算子 在本教程中&#xff0c;我們將聚焦于在 TVM 中使用張量表達式&#xff08;TE&#xff09;來定義張量計算和實現循環優化。TE用純函數語言描述張量計算&#xff08;即每個表達式都沒有副作用&#xff09;。當在 TVM 的整體…

4-數據結構-串的學習

4.1串類型的定義 1.串&#xff1a;&#xff08;或字符串&#xff09; 串是由多個字符組成的有限序列&#xff0c;記作&#xff1a;S‘c1c2c3…cn’ (n>0) 其中S是串的名字&#xff0c;‘c1c2c3…cn’ 是串值 ci是串中字符 n是串的長度&#xff0c;表示字符的數目 空串&a…

Linux下rm誤刪恢復 extundelete

Linux下rm誤刪恢復 extundelete 誤刪之后要第一時間卸載&#xff08;umount&#xff09;該分區&#xff0c;或者以只讀的方式來掛載&#xff08;mount&#xff09;該分區&#xff0c;否則覆寫了誰也沒辦法恢復。如果誤刪除的是根分區&#xff0c;最好直接斷電&#xff0c;進入…

5-數據結構-數組的學習

5.1數組的定義 定義&#xff1a; 由一組類型相同的數據元素構成的有序集合&#xff0c;每個數據元素稱為一個數據元素&#xff08;簡稱元素&#xff09;&#xff0c;每個元素受n&#xff08;n>1&#xff09;個線性關系的約束&#xff0c;每個元素在n個線性關系中的序號i1、…

timm 視覺庫中的 create_model 函數詳解

timm 視覺庫中的 create_model 函數詳解 最近一年 Vision Transformer 及其相關改進的工作層出不窮&#xff0c;在他們開源的代碼中&#xff0c;大部分都用到了這樣一個庫&#xff1a;timm。各位煉丹師應該已經想必已經對其無比熟悉了&#xff0c;本文將介紹其中最關鍵的函數之…

C--數據結構--樹的學習

6.2.1二叉樹的性質 1.二叉樹 性質&#xff1a; 1.若二叉樹的層次從1開始&#xff0c;則在二叉樹的第i層最多有2^(i-1)個結點 2.深度為k的二叉樹最多有2^k -1個結點 &#xff08;k>1&#xff09; 3.對任何一顆二叉樹&#xff0c;如果其葉結點個數為n0,度為2的非葉結點個數…

TVM:使用 Schedule 模板和 AutoTVM 來優化算子

TVM&#xff1a;使用 Schedule 模板和 AutoTVM 來優化算子 在本文中&#xff0c;我們將介紹如何使用 TVM 張量表達式&#xff08;Tensor Expression&#xff0c;TE&#xff09;語言編寫 Schedule 模板&#xff0c;AutoTVM 可以搜索通過這些模板找到最佳 Schedule。這個過程稱為…

TVM:使用 Auto-scheduling 來優化算子

TVM&#xff1a;使用 Auto-scheduling 來優化算子 在本教程中&#xff0c;我們將展示 TVM 的 Auto-scheduling 功能如何在無需編寫自定義模板的情況下找到最佳 schedule。 與基于模板的 AutoTVM 依賴手動模板定義搜索空間不同&#xff0c;auto-scheduler 不需要任何模板。 用…

C語言—sort函數比較大小的快捷使用--algorithm頭文件下

sort函數 一般情況下要將一組數從的大到小排序或從小到大排序&#xff0c;要定義一個新的函數排序。 而我們也可以直接使用在函數下的sort函數&#xff0c;只需加上頭文件&#xff1a; #include<algorithm> using namespace std;sort格式&#xff1a;sort(首元素地址&…

散列的使用

散列 散列簡單來說&#xff1a;給N個正整數和M個負整數&#xff0c;問這M個數中的每個數是否在N中出現過。 比如&#xff1a;N&#xff1a;{1,2,3,4}&#xff0c;M{2,5,7}&#xff0c;其中M的2在N中出現過 對這個問題最直觀的思路是&#xff1a;對M中每個欲查的值x&#xff0…

關于C++中的unordered_map和unordered_set不能直接以pair作為鍵名的問題

關于C中的unordered_map和unordered_set不能直接以pair作為鍵名的問題 在 C STL 中&#xff0c;不同于有序的 std::map 和 std::set 是基于紅黑樹實現的&#xff0c;std::unordered_map 和 std::unordered_set 是基于哈希實現的&#xff0c;在不要求容器內的鍵有序&#xff0c…

AI編譯器與傳統編譯器的聯系與區別

AI編譯器與傳統編譯器的區別與聯系 總結整理自知乎問題 針對神經網絡的編譯器和傳統編譯器的區別和聯系是什么&#xff1f;。 文中提到的答主的知乎主頁&#xff1a;金雪鋒、楊軍、藍色、SunnyCase、貝殼與知了、工藤福爾摩 筆者本人理解 為了不用直接手寫機器碼&#xff0…

python學習1:注釋\變量類型\轉換函數\轉義字符\運算符

python基礎學習 與大多數語言不同&#xff0c;python最具特色的就是使用縮進來表示代碼塊&#xff0c;不需要使用大括號 {} 。縮進的空格數是可變的&#xff0c;但是同一個代碼塊的語句必須包含相同的縮進空格數。 &#xff08;一個tab4個空格&#xff09; Python語言中常見的…

Python、C++ lambda 表達式

Python、C lambda 表達式 lambda函數簡介 匿名函數lambda&#xff1a;是指一類無需定義標識符&#xff08;函數名&#xff09;的函數或子程序。所謂匿名函數&#xff0c;通俗地說就是沒有名字的函數&#xff0c;lambda函數沒有名字&#xff0c;是一種簡單的、在同一行中定義函…

python 學習2 /輸入/ 輸出 /列表 /字典

python基礎學習第二天 輸入輸出 xinput("輸入內容") print(x)input輸出&#xff1a; eval :去掉字符串外圍的引號&#xff0c;按照python的語法執行內容 aeval(12) print(a)eval輸出樣式&#xff1a; 列表 建立&#xff0c;添加&#xff0c;插入&#xff0c;刪去…