张量程序抽象
在我们深入 TensorIR 的细节之前,首先介绍什么是原始张量函数。原始张量函数是对应于计算操作的单个“单元”的函数。例如,卷积操作可以是一个原始张量函数,融合的卷积 + relu 操作也可以是一个原始张量函数。通常,原始张量函数实现的典型抽象包含以下要素:多维缓冲区、驱动张量计算的循环嵌套,以及最终的计算语句本身。
from tvm.script import tir as T
@T.prim_func
def main(
A: T.Buffer((128,), "float32"),
B: T.Buffer((128,), "float32"),
C: T.Buffer((128,), "float32"),
) -> None:
for i in range(128):
with T.block("C"):
vi = T.axis.spatial(128, i)
C[vi] = A[vi] + B[vi]
张量程序的主要元素
演示的原始张量函数计算两个向量的元素wise和。该函数
接受三个多维缓冲区作为参数,并生成一个多维缓冲区作为输出。
包含一个单独的循环嵌套
i
,用于促进计算。具有一个单一的计算语句,用于计算两个向量的元素wise和。
TensorIR 中的额外结构
关键的是,我们无法对程序执行任意转换,因为某些计算依赖于循环的顺序。幸运的是,我们关注的大多数原始张量函数都具有有利的属性,例如循环迭代之间的独立性。例如,上述程序包括块和迭代注释
块注释
with T.block("C")
表示该块是指定用于调度的基本计算单元。一个块可能包含单个计算语句、带有循环的多个计算语句或不透明的内部函数(如 Tensor Core 指令)。迭代注释
T.axis.spatial
,指示变量vi
映射到i
,并且所有迭代都是独立的。
虽然此信息对于执行特定程序不是至关重要的,但在转换程序时证明很有用。因此,只要我们遍历从 0 到 128 的所有索引元素,我们就可以自信地并行化或重新排序与 vi
关联的循环。