技术标签: split 调度Schedule tile TVM中文教程 TVM深度学习编译器
TVM是用于高效内核代码构建的版本领域专用语言(Domain-Specialed-Language,DSL) 。
这篇教程,我们将展示通过TVM提供的各种原语怎么去调度计算。
from __future__ import absolute_import, print_function
import tvm
import numpy as np
通常存在几种计算相同结果的方法,但是,不同的方法将导致不同的局部性和性能,所以TVM要求用户提供怎么去调用Schedule描述计算是如何执行的。
Schedule是计算的变换的集合,它通过变换程序中的计算循环Loop,实现不同性能。
#定义一些变量
n = tvm.var('n')
m = tvm.var('m')
调度(Schedule)能通过一系列计算ops来定义,默认情况下,调度计算张量以航为顺序。
#定义矩阵元素element-wise乘法
A = tvm.placeholder((m,n), name='A')
B = tvm.placeholder((m,n), name='B')
C = tvm.compute((m,n),lambda i,j: A[i,j] * B[i,j], name ='C')
#创建调度
s = tvm.create_schedule([C.op])
#lower会将计算从定义转换为真正的可调用函数。 使用参数`simple_mode = True`,它将返回一个可读的C伪代码,我们在这里使用它来打印计划结果。
print(tvm.lower(s, [A, B, C], simple_mode=True))
输出:
produce C {
for (i, 0, m) {
for (j, 0, n) {
C[((i*n) + j)] = (A[((i*n) + j)]*B[((i*n) + j)])
}
}
}
一个调度过程由多个阶段组成,一个阶段表示操作的一个调度。我们提供各种方法来调度每个阶段。
split
通过factor
分裂指定轴为两个轴。
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i]*2, name='B')
s = tvm.create_schedule(B.op)
#分裂0轴为两个轴,先计算内循环再计算外循环,xo为外循环,xi为内循环
xo, xi = s[B].split(B.op.axis[0], factor=32)
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
for (i.outer, 0, ((m + 31)/32)) {
for (i.inner, 0, 32) {
if (likely(((i.outer*32) < (m - i.inner)))) {
B[((i.outer*32) + i.inner)] = (A[((i.outer*32) + i.inner)]*2.000000f)
}
}
}
}
使用nparts
与factor
作用相反,nparts指定外循环次数,factor指定内循环次数。
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i], name='B')
s = tvm.create_schedule(B.op)
bx, tx = s[B].split(B.op.axis[0], nparts=32)
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
for (i.outer, 0, 32) {
for (i.inner, 0, ((m + 31)/32)) {
if (likely((i.inner < (m - (i.outer*((m + 31)/32)))))) {
if (likely(((0 - (i.outer*((m + 31)/32))) <= i.inner))) {
B[(i.inner + (i.outer*((m + 31)/32)))] = A[(i.inner + (i.outer*((m + 31)/32)))]
}
}
}
}
}
tile
通过平铺两个轴执行计算图块
A = tvm.placeholder((m, n), name='A')
B = tvm.compute((m, n), lambda i, j: A[i, j], name='B')
s = tvm.create_schedule(B.op)
xo, yo, xi, yi = s[B].tile(B.op.axis[0], B.op.axis[1], x_factor=10, y_factor=5)
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
for (i.outer, 0, ((m + 9)/10)) {
for (j.outer, 0, ((n + 4)/5)) {
//先执行10x5的图块,滑动下一个图块
for (i.inner, 0, 10) {
for (j.inner, 0, 5) {
if (likely(((i.outer*10) < (m - i.inner)))) {
if (likely(((j.outer*5) < (n - j.inner)))) {
B[(((j.outer*5) + (((i.outer*10) + i.inner)*n)) + j.inner)] = A[(((j.outer*5) + (((i.outer*10) + i.inner)*n)) + j.inner)]
}
}
}
}
}
}
}
fuse
能融合一个计算的两个轴
A = tvm.placeholder((m,n),name='A')
B = tvm.compute((m,n), lambda i,j: A[i,j], name='B')
s = tvm.create_schedule(B.op)
#首先平铺成4轴(i.outer,j.outer,i.inner,j.inner)
xo,yo,xi,yi = s[B].tile(B.op.axis[0],B.op.axis[1], x_factor=10, y_factor=5)
#然后融合(i.inner,j.inner)进一个轴:(i.inner.j.inner.fused)
fused = s[B].fuse(xi,yj)
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
for (i.outer, 0, ((m + 9)/10)) {
for (j.outer, 0, ((n + 4)/5)) {
for (i.inner.j.inner.fused, 0, 50) {
if (likely(((i.outer*10) < (m - (i.inner.j.inner.fused/5))))) {
if (likely(((j.outer*5) < (n - (i.inner.j.inner.fused % 5))))) {
B[(((j.outer*5) + (i.inner.j.inner.fused % 5)) + (((i.outer*10) + (i.inner.j.inner.fused/5))*n))] = A[(((j.outer*5) + (i.inner.j.inner.fused % 5)) + (((i.outer*10) + (i.inner.j.inner.fused/5))*n))]
}
}
}
}
}
}
reorder
能按照指定顺序重新排列轴(类似于permute)。
A = tvm.placeholder((m, n), name='A')
B = tvm.compute((m, n), lambda i, j: A[i, j], name='B')
s = tvm.create_schedule(B.op)
#首先平铺成4轴(i.outer,j.outer,i.inner,j.inner)
xo, yo, xi, yi = s[B].tile(B.op.axis[0], B.op.axis[1], x_factor=10, y_factor=5)
s[B].reorder(xi,yo,xo,yi)
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
for (i.inner, 0, 10) {
for (j.outer, 0, ((n + 4)/5)) {
for (i.outer, 0, ((m + 9)/10)) {
for (j.inner, 0, 5) {
if (likely(((i.outer*10) < (m - i.inner)))) {
if (likely(((j.outer*5) < (n - j.inner)))) {
B[(((j.outer*5) + (((i.outer*10) + i.inner)*n)) + j.inner)] = A[(((j.outer*5) + (((i.outer*10) + i.inner)*n)) + j.inner)]
}
}
}
}
}
}
}
bind
可以使用线程轴绑定指定的轴,通常在GPU编程中使用。
A = tvm.placeholder((n,), name='A')
B = tvm.compute(A.shape, lambda i: A[i] * 2, name='B')
s = tvm.create_schedule(B.op)
bx, tx = s[B].split(B.op.axis[0], factor=64)
s[B].bind(bx, tvm.thread_axis("blockIdx.x"))
s[B].bind(tx, tvm.thread_axis("threadIdx.x"))
print(tvm.lower(s, [A, B], simple_mode=True))
输出:
produce B {
// attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = ((n + 63)/64)
// attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = 64
if (likely(((blockIdx.x*64) < (n - threadIdx.x)))) {
B[((blockIdx.x*64) + threadIdx.x)] = (A[((blockIdx.x*64) + threadIdx.x)]*2.000000f)
}
}
对于包含多个算子的调度,TVM默认从root开始遍历计算张量。
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i]+1, name='B')
C = tvm.compute((m,), lambda i: B[i]*2, name='C')
s = tvm.create_schedule(C.op)
print(tvm.lower(s, [A, B, C], simple_mode=True))
输出:
produce B {
for (i, 0, m) {
B[i] = (A[i] + 1.000000f)
}
}
produce C {
for (i, 0, m) {
C[i] = (B[i]*2.000000f)
}
}
compute_at
可以将B
的计算移动到C
的第一个计算轴。
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i]+1, name='B')
C = tvm.compute((m,), lambda i: B[i]*2, name='C')
s = tvm.create_schedule(C.op)
# 移动B循环到C循环的第一个轴
s[B].compute_at(S[C], C.op.axis[0])
print(tvm.lower(s, [A, B, C], simple_mode=True))
输出:
produce C {
for (i, 0, m) {
produce B {
B[i] = (A[i] + 1.000000f)
}
C[i] = (B[i]*2.000000f)
}
}
compute_inline
可以将一个计算阶段标记为内联,然后将计算体扩展并插入需要张量的地址处。(和C中的内联函数一个意思)
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i]+1, name='B')
C = tvm.compute((m,), lambda i: B[i]*2, name='C')
s = tvm.create_schedule(C.op)
s[B].compute_inline()
print(tvm.lower(s, [A, B, C], simple_mode=True))
输出:
produce C {
for (i, 0, m) {
//类似内联函数,直接合成一个循环
C[i] = ((A[i]*2.000000f) + 2.000000f)
}
}
compute_root
可以将一个计算阶段的计算移动到root。(compute_at
的逆过程)
A = tvm.placeholder((m,), name='A')
B = tvm.compute((m,), lambda i: A[i]+1, name='B')
C = tvm.compute((m,), lambda i: B[i]*2, name='C')
s = tvm.create_schedule(C.op)
# B移动到C的0轴
s[B].compute_at(s[C], C.op.axis[0])
# B重新移动回root
s[B].compute_root()
print(tvm.lower(s, [A, B, C], simple_mode=True))
输出:
produce B {
for (i, 0, m) {
B[i] = (A[i] + 1.000000f)
}
}
produce C {
for (i, 0, m) {
C[i] = (B[i]*2.000000f)
}
}
本教程介绍了tvm中的调度原语,它允许用户轻松灵活地调度计算。
为了获得良好性能的内核实现,一般工作流程通常是:
文章浏览阅读119次。该楼层疑似违规已被系统折叠隐藏此楼查看此楼/***Getaparametervalue**@paramkeyString*@paramdefString*@returnString*/publicStringgetParameter(Stringkey,Stringdef){returnisStandalone?System.getProperty(ke..._java http隧道
文章浏览阅读913次。IP主机名备注192.168.117.14keepalived-master主节点192.168.117.15keepalived-slaver备节点192.168.117.100VIP1.主备节点均安装keepalived# yum install -y keepalived httpd2.主备节点均修改keepalived日志存放路径..._keepalived sendmail
文章浏览阅读469次。--==========================================--SPFILE错误导致数据库无法启动(ORA-01565)--========================================== SPFILE错误导致数据库无法启动 SQL> startup ORA-01078: failurein proce_ora01565 ora27046
文章浏览阅读6.1k次,点赞2次,收藏54次。功能测试基础知识总结_功能测试
文章浏览阅读3.2k次,点赞3次,收藏2次。pg 中文首字母排序_pg中文排序
文章浏览阅读3.1w次,点赞23次,收藏109次。本文主要讲解CONVERT函数_mysql convert
文章浏览阅读8.6k次,点赞2次,收藏2次。HTML5 的视频播放事件想必大家已经期待很久了吧,在HTML4.1、4.0之前我们如果在网页上播放视频无外乎两种方法: 第一种:安装FLASH插件或者微软发布的插件 第二种:在本地安装播放器,在线播放组件之类的 因为并不是所有的浏览器都安装了FLASH插件,就算安装也不一定所有的都能安装成功。像苹果系统就是默认禁用FLASH的,安卓虽然一开始的时候支持FLASH,但是在安卓4.0以后也开始不_微信开发者工具视频快进
文章浏览阅读5.4k次,点赞3次,收藏4次。在使用redis的过程常见错误总结1.JedisConnectionException Connection Reset参考这边文章:Connection reset原因分析和解决方案https://blog.csdn.net/cwclw/article/details/527971311.1问题描述Exception in thread "main" redis.clients...._jedisconnectionexception: java.net.socketexception: connection reset
文章浏览阅读8.3k次,点赞8次,收藏42次。目录1.Lua垃圾回收算法原理简述2.Lua垃圾回收中的三种颜色3.Lua垃圾回收详细过程4.步骤源码详解4.1新建对象阶段4.2触发条件4.3 GC函数状态机4.4标记阶段4.5清除阶段5.总结参考资料lua垃圾回收(Garbage Collect)是lua中一个比较重要的部分。由于lua源码版本变迁,目前大多数有关这个方面的文章都还是基于lua5.1版本,有一定的滞后性。因此本文通过参考当前..._lua5.3 gc
文章浏览阅读511次。最近家中的潮人,老妈闲着没事干,开始学玩电脑,引起他的各种好奇心。如看看新闻,上上微信或做做其他的事情。但意料之中的是电脑上会莫名出现各种问题?不翼而飞的图标?照片又不见了?文件被删了,卡机或者黑屏,无声音了,等等问题。常常让她束手无策,求助于我,可惜在电话中说不清,往往只能苦等我回家后才能解决,那种开心乐趣一下子消失了。想想,这样也不是办法啊, 于是,我潜心寻找了两款优秀的远程控制软件。两款软件...
文章浏览阅读1.8k次。二.初始化工作空间三.设置下载地址四.下载功能包此处可能会报错,请看:rosdep update遇到ERROR: error loading sources list: The read operation timed out问题_DD᭄ꦿng的博客-程序员宅基地接下来一次安装所有功能包,注意对应ROS版本 五.编译功能包isolated:单独编译各个功能包,每个功能包之间不产生依赖。编译过程时间比较长,可能需要几分钟时间。此处可能会报错:缺少absl依赖包_ros18.04 安装ca
文章浏览阅读4.1k次,点赞3次,收藏7次。Haobor2.2.1配置(trivy扫描器、镜像签名)docker-compose下载https://github.com/docker/compose/releases安装cp docker-compose /usr/local/binchmod +x /usr/local/bin/docker-composeharbor下载https://github.com/goharbor/harbor/releases解压tar xf xxx.tgx配置harbor根下建立:mkd_init error: db error: failed to download vulnerability db: database download