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.mzph.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;删去…