【TVM 教程】TVM 中的 Schedule 原语

TVM(端到端深度学习编译器)简介 算子编译器前后端引入 引入 2021.2.8 和振伟聊了关于HW OD的工作内容,深度上很好的方向, 目前纯OP的high level的设计,业界已经内卷了即难以进步。 torch已经很好用了社区也好,想从纯软件上进步很难。 而硬骨头OP在硬件上的优化比如GPU有前途但研发周期长,人才紧缺,还在缓慢发展中。 TVM是业界有名气的平台,HW的TBE对标TVM 要做一个接口在前端(各种框架)和后端(异构)。 有了前后端的经验,那么整个行业的Edge端都可以用到,市场很大。 ... 阅读详情

Apache TVM 是一个端到端的深度学习编译框架,适用于 CPU、GPU 和各种机器学习加速芯片。更多 TVM 中文文档可访问 → https://tvm.hyper.ai/

作者:Ziheng Jiang

TVM 是一种用于高效构建内核的领域特定语言。

本教程展示了如何通过 TVM 提供的各种原语来调度计算。

from __future__ import absolute_import, print_function

import tvm
from tvm import te
import numpy as np

计算相同结果的方法众多,然而,不同的方法会导致局部性和性能各异,因此 TVM 要求用户借助 Schedule 执行计算。

Schedule 是一组计算转换,可用于转换程序中的循环计算。

# 声明变量,供之后使用
n = te.var("n")
m = te.var("m")

Schedule 可由算子列表创建,它默认以行优先的方式串行计算张量。

# 声明一个矩阵元素乘法
A = te.placeholder((m, n), name="A")
B = te.placeholder((m, n), name="B")
C = te.compute((m, n), lambda i, j: A[i, j] * B[i, j], name="C")

s = te.create_schedule([C.op])
# lower 会将计算从定义转换为实际可调用的函数。
# 使用参数 `simple_mode=True` 会返回一个可读的类 C 的语句,这里用它来打印 schedule 结果。
print(tvm.lower(s, [A, B, C], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"),
             C: Buffer(C_2: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B, C_1: C}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m, n: int32], [stride, stride_3: int32], type="auto"), B_1: B_3: Buffer(B_2, float32, [m, n], [stride_1, stride_4: int32], type="auto"), C_1: C_3: Buffer(C_2, float32, [m, n], [stride_2, stride_5: int32], type="auto")} {
 for (i: int32, 0, m) {
 for (j: int32, 0, n) {
      C[((i*stride_2) + (j*stride_5))] = (A[((i*stride) + (j*stride_3))]*B[((i*stride_1) + (j*stride_4))])
 }
 }
}

一个 Schedule 由多个 Stage 组成,一个 Stage 代表一个操作的 schedule。每个 stage 的调度都有多种方法:

split​

split 可根据 factor 将指定 axis 拆分为两个 axis。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i] * 2, name="B")

s = te.create_schedule(B.op)
xo, xi = s[B].split(B.op.axis[0], factor=32)
print(tvm.lower(s, [A, B], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto")} {
 for (i.outer: int32, 0, floordiv((m + 31), 32)) {
 for (i.inner: int32, 0, 32) {
 if @tir.likely((((i.outer*32) + i.inner) < m), dtype=bool) {
 let cse_var_1: int32 = ((i.outer*32) + i.inner)
        B[(cse_var_1*stride_1)] = (A[(cse_var_1*stride)]*2f32)
 }
 }
 }
}

也可用 nparts 来拆分 axis,它拆分 axis 的方式与 factor 相反。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i], name="B")

s = te.create_schedule(B.op)
bx, tx = s[B].split(B.op.axis[0], nparts=32)
print(tvm.lower(s, [A, B], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto")} {
 for (i.outer: int32, 0, 32) {
 for (i.inner: int32, 0, floordiv((m + 31), 32)) {
 if @tir.likely(((i.inner + (i.outer*floordiv((m + 31), 32))) < m), dtype=bool) {
        B[((i.inner + (i.outer*floordiv((m + 31), 32)))*stride_1)] = A[((i.inner + (i.outer*floordiv((m + 31), 32)))*stride)]
 }
 }
 }
}

tile​

tile 可在两个 axis 上逐块执行计算。

A = te.placeholder((m, n), name="A")
B = te.compute((m, n), lambda i, j: A[i, j], name="B")

s = te.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))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m, n: int32], [stride, stride_2: int32], type="auto"), B_1: B_3: Buffer(B_2, float32, [m, n], [stride_1, stride_3: int32], type="auto")} {
 for (i.outer: int32, 0, floordiv((m + 9), 10)) {
 for (j.outer: int32, 0, floordiv((n + 4), 5)) {
 for (i.inner: int32, 0, 10) {
 if @tir.likely((((i.outer*10) + i.inner) < m), dtype=bool) {
          for (j.inner: int32, 0, 5) {
            if @tir.likely((((j.outer*5) + j.inner) < n), dtype=bool) {
              let cse_var_2: int32 = ((j.outer*5) + j.inner)
              let cse_var_1: int32 = ((i.outer*10) + i.inner)
              B[((cse_var_1*stride_1) + (cse_var_2*stride_3))] = A[((cse_var_1*stride) + (cse_var_2*stride_2))]
 }
 }
 }
 }
 }
 }
}

fuse​

fuse 可将一个计算的两个连续 axis 融合。

A = te.placeholder((m, n), name="A")
B = te.compute((m, n), lambda i, j: A[i, j], name="B")

s = te.create_schedule(B.op)
# 首先调用 tile 平铺到四个 axis: (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, yi)
print(tvm.lower(s, [A, B], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m, n: int32], [stride, stride_2: int32], type="auto"), B_1: B_3: Buffer(B_2, float32, [m, n], [stride_1, stride_3: int32], type="auto")} {
 for (i.outer: int32, 0, floordiv((m + 9), 10)) {
 for (j.outer: int32, 0, floordiv((n + 4), 5)) {
 for (i.inner.j.inner.fused: int32, 0, 50) {
 if @tir.likely((((i.outer*10) + floordiv(i.inner.j.inner.fused, 5)) < m), dtype=bool) {
 if @tir.likely((((j.outer*5) + floormod(i.inner.j.inner.fused, 5)) < n), dtype=bool) {
 let cse_var_2: int32 = ((j.outer*5) + floormod(i.inner.j.inner.fused, 5))
 let cse_var_1: int32 = ((i.outer*10) + floordiv(i.inner.j.inner.fused, 5))
            B[((cse_var_1*stride_1) + (cse_var_2*stride_3))] = A[((cse_var_1*stride) + (cse_var_2*stride_2))]
 }
 }
 }
 }
 }
}

reorder​

reorder 可按指定的顺序对 axis 重新排序。

A = te.placeholder((m, n), name="A")
B = te.compute((m, n), lambda i, j: A[i, j], name="B")

s = te.create_schedule(B.op)
# 首先调用 tile 平铺到四个轴: (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)
# 然后将 axis 重新排序:(i.inner,j.outer,i.outer,j.inner)
s[B].reorder(xi, yo, xo, yi)
print(tvm.lower(s, [A, B], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m, n: int32], [stride, stride_2: int32], type="auto"), B_1: B_3: Buffer(B_2, float32, [m, n], [stride_1, stride_3: int32], type="auto")} {
 for (i.inner: int32, 0, 10) {
 for (j.outer: int32, 0, floordiv((n + 4), 5)) {
 for (i.outer: int32, 0, floordiv((m + 9), 10)) {
 if @tir.likely((((i.outer*10) + i.inner) < m), dtype=bool) {
          for (j.inner: int32, 0, 5) {
            if @tir.likely((((j.outer*5) + j.inner) < n), dtype=bool) {
              let cse_var_2: int32 = ((j.outer*5) + j.inner)
              let cse_var_1: int32 = ((i.outer*10) + i.inner)
              B[((cse_var_1*stride_1) + (cse_var_2*stride_3))] = A[((cse_var_1*stride) + (cse_var_2*stride_2))]
 }
 }
 }
 }
 }
 }
}

bind​

bind 可将指定 axis 与线程 axis 绑定,常用于 GPU 编程。

A = te.placeholder((n,), name="A")
B = te.compute(A.shape, lambda i: A[i] * 2, name="B")

s = te.create_schedule(B.op)
bx, tx = s[B].split(B.op.axis[0], factor=64)
s[B].bind(bx, te.thread_axis("blockIdx.x"))
s[B].bind(tx, te.thread_axis("threadIdx.x"))
print(tvm.lower(s, [A, B], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*n: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*n)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [n], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [n], [stride_1], type="auto")} {
  attr [IterVar(blockIdx.x: int32, (nullptr), "ThreadIndex", "blockIdx.x")] "thread_extent" = floordiv((n + 63), 64);
  attr [IterVar(threadIdx.x: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 64;
 if @tir.likely((((blockIdx.x*64) + threadIdx.x) < n), dtype=bool) {
    B[(((blockIdx.x*64) + threadIdx.x)*stride_1)] = (A[(((blockIdx.x*64) + threadIdx.x)*stride)]*2f32)
 }
}

compute_at​

对于包含多个算子的 schedule,TVM 默认会分别计算 root 处的张量。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i] + 1, name="B")
C = te.compute((m,), lambda i: B[i] * 2, name="C")

s = te.create_schedule(C.op)
print(tvm.lower(s, [A, B, C], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"),
             C: Buffer(C_2: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B, C_1: C}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto"), C_1: C_3: Buffer(C_2, float32, [m], [stride_2], type="auto")} {
 for (i: int32, 0, m) {
    B[(i*stride_1)] = (A[(i*stride)] + 1f32)
 }
 for (i_1: int32, 0, m) {
    C[(i_1*stride_2)] = (B[(i_1*stride_1)]*2f32)
 }
}

compute_at 可将 B 的计算移动到 C 计算的首个 axis 中。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i] + 1, name="B")
C = te.compute((m,), lambda i: B[i] * 2, name="C")

s = te.create_schedule(C.op)
s[B].compute_at(s[C], C.op.axis[0])
print(tvm.lower(s, [A, B, C], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"),
             C: Buffer(C_2: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B, C_1: C}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto"), C_1: C_3: Buffer(C_2, float32, [m], [stride_2], type="auto")} {
 for (i: int32, 0, m) {
    B[(i*stride_1)] = (A[(i*stride)] + 1f32)
    C[(i*stride_2)] = (B[(i*stride_1)]*2f32)
 }
}

compute_inline​

compute_inline 可将 stage 标记为 inline,然后扩展计算体,并将其插入到需要张量的地址。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i] + 1, name="B")
C = te.compute((m,), lambda i: B[i] * 2, name="C")

s = te.create_schedule(C.op)
s[B].compute_inline()
print(tvm.lower(s, [A, B, C], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"),
             C: Buffer(C_2: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B, C_1: C}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto"), C_1: C_3: Buffer(C_2, float32, [m], [stride_2], type="auto")} {
 for (i: int32, 0, m) {
    C[(i*stride_2)] = ((A[(i*stride)] + 1f32)*2f32)
 }
}

compute_root​

compute_root 可将一个 stage 的计算移动到 root。

A = te.placeholder((m,), name="A")
B = te.compute((m,), lambda i: A[i] + 1, name="B")
C = te.compute((m,), lambda i: B[i] * 2, name="C")

s = te.create_schedule(C.op)
s[B].compute_at(s[C], C.op.axis[0])
s[B].compute_root()
print(tvm.lower(s, [A, B, C], simple_mode=True))

输出结果:

@main = primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"),
             B: Buffer(B_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"),
             C: Buffer(C_2: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")}
  buffer_map = {A_1: A, B_1: B, C_1: C}
  preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [m], [stride], type="auto"), B_1: B_3: Buffer(B_2, float32, [m], [stride_1], type="auto"), C_1: C_3: Buffer(C_2, float32, [m], [stride_2], type="auto")} {
 for (i: int32, 0, m) {
    B[(i*stride_1)] = (A[(i*stride)] + 1f32)
 }
 for (i_1: int32, 0, m) {
    C[(i_1*stride_2)] = (B[(i_1*stride_1)]*2f32)
 }
}

总结​

本教程介绍了 TVM 中的 schedule 原语(使得用户可以轻松、灵活地调度计算)。

为提高内核的性能,一般的工作流程如下:

  • 通过一系列操作描述你的计算。
  • 使用原语来调度计算。
  • 编译并运行,查看性能差异。
  • 根据运行结果来调整 schedule。

下载 Python 源代码:schedule_primitives.py

下载 Jupyter Notebook:schedule_primitives.ipynb

TVM之Schedule TVM之Schedule1. schedule2. stage3. methods for stage3.1 split3.2 tile3.3 fuse3.4 reorder3.5 bind3.6 compute_at3.7 compute_inline3.8 compute_root 最近在阅读TVM官网的文档,在阅读的过程中将keywords以及相应理解记录了下来,有理解不对的地方,欢迎批评指... 阅读详情

相关推荐

TVM基础编程示例分析

TVM基础编程示例分析 一.TVM编程基础示例 前言 继前图灵奖获得者Hennessy和Patterson在ISCA 2018提出“A New Golden Age for Computer Architecture”,编译器界大神Chris Lattner在ASPLOS 2021提出了“The Golden Age of Compiler Design”。另一方面,2020年图灵奖授予了编译器“龙书”作者Jeffrey Ullman和Alfred Aho。编译器技术在新的时代背景下似乎又再次焕发了新的活力

吴建明wujianming_110117 1491

TVM概述

TVM TVM是陈天奇领导的一个DL加速框架项目。它处于DL框架(如tensorflow、pytorch)和硬件后端(如CUDA、OpenCL)之间,兼顾了前者的易用性和后者的执行效率。 官网: https://tvm.apache.org/ 代码: https://github.com/apache/tvm 论文: 《TVM: End-to-End Optimization Stack for Deep Learning》 和同类项目的差异: TFLite和ONNXRuntime只能接收特定格式的模

antkillerfarm的专栏 1万+

TVM自定义算子+测试

TVM自定义算子定义算子编译器后端算子的测试变量的声明调用底层算子部署到特定硬件上运行全部代码 定义算子 TVM分为编译器前端和编译器后端。编译器前端主要解决的事情是将其他框架(pytorch等)转换成自己的symbolic的表示方式,所有的基本操作都会转换成sym.基本算子的形式,然后通过注册的形式,使得sym.基本算子调用后端的算子实现。编译器后端的算子实现分为两部分,第一部分是compute的部分,compute定义了基本的计算规则,schedule定义了基本计算规则在不同的硬件上的实际执行的方式。

free1993的博客 2957

TVM简介

TVM软件栈 工作流程 从已有的深度学习框架中获取一个模型并将此模型转换为计算图表示(深度学习框架的前端主要是计算图表示以及自动梯度); 图中 Section 3 使用一些方法优化当前的计算图得到优化后的计算图(操作融合,可以将多个操作合并为一个kernel,以此节省存储中间结果的空间;数据布局转换,改变计算图中张量的存储方式以适应目标硬件设备并提升执行效率); 图中 Section 4 针对...

MLer 4万+

【TVM系列二】TVM介绍

文章同步更新在公众号 AIPlayer,欢迎扫码关注,共同进步 目录 一、TVM的工作流程 1、整体流程 2、关键数据结构 3、Transformations 4、搜索空间和基于机器学习的转换 5、目标代码转化 二、逻辑架构组件 三、运行TVM实例 1、交叉编译runtime 2、编译模型 3、运行模型 四、总结 一、TVM的工作流程 TVM主要由两个部分组成: (1)TVM编译器:负责编译和优化模型 (2)TVM runtime:提供目标设备上运行模型的API ..

Flying Cat 6695

TVM学习(五)schedule

Schedule是和硬件体系结构相关的一些列优化,Halide在其文章中对其做了以下定义: 1 When and where should be the value at each coordinate in each function be computed? 2 Where should they be stored? 3 How long are values cached and communicated across multiple consumers, and when are they in

qq_33287871的博客 1231

TVM Compiler中文教程:TVM调度原语(Schedule Primitives)

文章目录TVM调度原语(Schedule Primitives)分裂split平铺tile TVM调度原语(Schedule Primitives) TVM是用于高效内核代码构建的版本领域专用语言(Domain-Specialed-Language,DSL) 。 这篇教程,我们将展示通过TVM提供的各种原语怎么去调度计算。 from __future__ import absolute_impor...

Mars_WH的专栏 8660

【CATIA的二次开发38】对象Selection交互选择部分方法介绍01

在CATIA V5的VBA开发中,Selection对象是用户交互的核心组件,用于管理用户在图形区域或特征树中的选择操作。Selection 对象是 CATIA VBA 中的中央交互枢纽,充当用户界面与程序逻辑之间的桥梁。它代表当前在图形区域或特征树中被选中的元素集合,是自动化操作的基础。

学习共勉:不积跬步,无以至千里; 不积小流,无以成江海。 1342

图解TVM中的调度原语(Schedule Primitives)

上面简要介绍了TVM中基本的调度原语(Schedule Primitives),我们可以通过组合这些基本元素,达到一个满意的优化结果。

生活需要深度 1126

树莓派4B没有serial0或ttyS0,serial0没有映射到ttyS0,serial0和serial1调换映射

环境: 树莓派4B,ubuntu mate20 现象 终端输入ls -l /dev,发现只有serial1->ttyAMA0,没有serial0页没有ttyS0,更没有二者的映射。网上很多方法都是树莓派3B和3B+的,我自己试过,不太适用于4B,下面的方法是4B的,我自己用是成功的。 解决方法 cd /boot/firmware sudo gedit config.txt 注意不是/boot下的config,而是/boot/firmware下的config 打开文件后,找到enable_uart=0

qq_41208487的博客 7891

逆向选择启示录:一个失败高频策略如何意外跑赢市场

本文探讨了高频交易中的逆向选择现象,揭示了一个失败的高频策略如何通过简化意外跑赢市场。通过加密货币市场的实证分析和蒙特卡洛模拟,研究发现低预测性的简单策略反而更具鲁棒性,有效规避了逆向选择风险。文章还提供了小币种做市策略的优化框架和实践建议,为量化交易者提供了宝贵洞见。

weixin_29191669的博客 401

TVM:使用 Schedule 模板和 AutoTVM 来优化算子

TVM:使用 Schedule 模板和 AutoTVM 来优化算子 在本文中,我们将介绍如何使用 TVM 张量表达式(Tensor Expression,TE)语言编写 Schedule 模板,AutoTVM 可以搜索通过这些模板找到最佳 Schedule。这个过程称为自动调整(Auto Tuning),它有助于自动优化张量计算的过程。 本教程需基于之前介绍的如何使用 TE 来写一个矩阵乘法的教程。 Auto Tuning 有两步: 第一步是定义一个搜索空间 第二步是运行相应的搜索算法来探索这个空间 本

weixin_44966641的博客 1650

告别手动复制粘贴!用Python脚本+pywinauto实现PC微信消息自动发送(附完整源码)

本文详细介绍了如何利用Python和pywinauto实现PC微信消息的自动化发送,告别繁琐的手动操作。从环境搭建、UI控件定位到构建健壮的发送流程,再到高级功能扩展和项目打包部署,全面解析了自动化办公的实现方案,显著提升工作效率。

weixin_30315723的博客 585

TVM的“hello world“基础流程 I

前言 继前图灵奖获得者Hennessy和Patterson在ISCA 2018提出“A New Golden Age for Computer Architecture”,编译器大神Chris Lattner在ASPLOS 2021提出了“The Golden Agen of Compiler Design”。另一方面,2020年图灵奖授予了编译器“龙书”作者Jeffrey Ullman和Alfred Aho。编译器技术在新的时代背景下似乎又焕发了新的活力,成为了业界的热点。 而作为现在最热门的AI计算场景

世事难料,保持低调 6092

大语言模型在构建UNSPSC 分类数据中的应用

UNSPSC 是联合国标准产品和服务代码。本文研究如何利用大语言模型生成分类系统的分析

姚家湾博客 1865

模拟IC设计笔记:拆解一个经典的时钟占空比校正(DCC)电路模块,从理论到SPICE仿真

本文深入解析了时钟占空比校正(DCC)电路在模拟IC设计中的关键作用,特别是在高速ADC应用中。通过晶体管级拆解,详细介绍了基于电荷泵和积分器的模拟反馈型DCC电路设计,包括占空比检测、误差放大与校正执行单元的实现。文章还提供了SPICE仿真验证方法,确保电路在100MHz时钟下的稳定性和性能。

weixin_33725722的博客 428

【TVM教程】设计与架构

IRModule:包含 relay.Function 和 tir.PrimFuncruntime.Module:包含 runtime.PackedFunc编译基本是在进行关键数据结构之间的转换。relay/transform 和 tir/transform 是确定性的基于规则的转换meta-schedule 则包含基于搜索的转换最后,编译流程示例只是 TVM 堆栈的一个典型用例。将这些关键数据结构和转换提供给 Python 和 C++ API。

HyperAI超神经 1245

WIN下使用visual studio 2022安装GMSSL命令行工具

2.1 使用管理员权限打开visual studio(非常重要,否则最后一步 nmake install会失败)右键vs2022,以管理员身份运行即可。2.4 配置环境变量,gmssl可执行程序默认会被安装到如下目录。将如下路径添加到系统环境变量即可。添加后的环境变量如下。2.编译解压后的gmssl源码,编译过程主要参考install.md中的内容,README.md中的步骤不完整。2.2 打开vs2022的命令终端,右键文件,选择在终端中打开。1.下载gmssl源码,解压到相应目录。

qq_45794382的博客 1269

TVM自动优化:VibeThinker生成Schedule Template

微博开源的小模型VibeThinker-1.5B能生成针对TVM编译器的调度建议,将自然语言转化为硬件感知的优化策略。通过解析模型输出并结合TVM流程,可实现从提示到高效内核的端到端优化,尤其适用于资源受限场景下的算子级性能提升。

weixin_29607511的博客 951

保姆级教程:用Docker Compose一键部署WVP-PRO+ZLM+录像服务(含Nginx反代)

本文提供了一套完整的容器化GB28181视频平台部署方案,通过Docker Compose实现WVP-PRO、ZLMediaKit(ZLM)和录像服务的一键式部署。详细介绍了环境准备、架构设计、容器化部署实战以及高级配置优化,帮助用户快速搭建高性能视频监控管理平台,适用于安防行业和物联网设备管理场景。

weixin_30415113的博客 861

jre9下载 jdk9

(Java SE Runtime Environment 9.0) 9.0u175 官方最新版 【x86|x64】

TVM Ansor与Meta Schedule对比:自动调度算法的演进与实现

本文深入对比分析了深度学习编译器TVM中的两大自动调度系统:Ansor与Meta Schedule。文章系统阐述了两代算法的核心设计哲学、演进路径与实现差异,重点解析了Ansor基于程序生成的梯度优化搜索策略,以及Meta Schedule基于预代价模型与强化学习的进化式成本模型。通过一个自定义算子的完整实战案例,详细展示了从算子定义、搜索空间构建到性能调优的自动调度全流程,为读者理解现代AI编译器的自动代码生成技术提供了全面而深入的视角。

AladdinEdu,你的AI学习实践工作坊。让想法落地,让研究加速。助力高校AI人才成长,点亮创新未来。 874

文档:亲测可用解决vnc灰屏,黑屏,鼠标大黑叉情况

主要解决vnc-viewer灰屏,黑屏,鼠标大黑叉等现象,详细见文档。

【TVM系列五】添加Relay自定义算子

这就需要实现一个函数,传递相应的参数给对应的relay算子,并且返回对应算子的Call Node(这个算子最终在Relay表达式的AST里面)。A:在函数中会创建输入与输出的TypeVarNode,然后创建TypeRelationNode将类型关系函数管理起来,并定义一个FuncTypeNode将这些定义好的对象作为输入,最终赋值给op_type。本文以实现一个axis_abs的自定义算子为例介绍如何在tvm中添加新的relay算子,该算子实现的功能是以输入的3维tensor取某一维度的指定切片取绝对值。

生活需要深度 1195

efm32 入门.pdf

本PDF共计22页,介绍的是操作驱动API,下载与否自己斟酌。

TMS_WebGMaps_1.9.8.0_Full_Source

TMS_WebGMaps_1.9.8.0_Full_Source 全源码 delphi cbuilder

【TVM 教程】规约(reduce)

除了te.sumtvm.te.min和tvm.te.max等内置规约操作外,还可以通过定义交换规约操作。本教程演示了如何规约 schedule。用 reduce_axis 描述规约。如需并行性(parallelism),用 rfactor 来分解轴。通过定义新的规约操作。

HyperAI超神经 1122

TVM基本调度原语

TVM调度原语官方DOC #!/usr/bin/env python # coding: utf-8 # 通常有几种方法来计算相同的结果,但是不同的方法会导致不同的局部性和性能。所以TVM要求用户提供如何执行称为Schedule的计算。 # Schedule是一组转换计算的集合,它转换程序中的计算循环。 # In[1]: #兼容python2.x from __future__ import print_function #在python2.x中也要像python3.x那样使用print() fr

weixin_45021321的博客 475

【TVM 教程】TVM 代码库实例讲解

Apache TVM 是一个深度的深度学习编译框架,适用于 CPU、GPU 和各种机器学习加速芯片。

HyperAI超神经 902

tvm.build过程

代码示例 import tvm from tvm import te n = 2 A = te.placeholder((n,), name='A') B = te.placeholder((n,), name='B') C = te.compute(A.shape, lambda *i: A(*i) + B(*i), name='C') s = tvm.te.create_schedule(C.op) m = tvm.lower(s, [A, B, C], name="test_add") rt_mod

kh815的博客 747
上一篇: Agent心理诊所上线!基于1.3K抑郁症问诊对话,上海交大团队搭建大模型对话Agent,可初诊抑郁症
下一篇: 实现蛋白质动态对接预测!上海交大/星药科技/中山大学等联合推出几何深度生成模型DynamicBind
HyperAI超神经
博客等级 码龄7年 8873粉丝 1753原创
评论
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付:10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值