☰
TVM算子开发:IR、类型、dtype与infer shape排查
2026/10/26 15:00:59 网站建设 项目流程

从算子定义写完到第一次跑通编译,中间最容易卡住的地方,往往不是调度写不出来,而是类型和 shape 对不上。TVM 的 IR 定义、计算类型、数据类型和 infer shape 这四件事,本质上是一条链上的四个环节:IR 描述"算什么",类型系统描述"数据长什么样",dtype 描述"每个数字占几个字节、怎么解释",shape 推导则负责把前者逐层传递下去,直到后端能算出每一块内存的大小。任何一环断了,报错信息通常都不会直接指向真正的元凶。这篇文章面向已经能跑通几个 Demo、但一改算子或一换输入就报错的人,我会把这四层拆开讲清楚它们各自的职责边界,再给出可以直接抄的排查路径和一段完整的手写 shape 推导示例。

1. 把四个概念摆到同一张桌子上

1.1 一个算子从定义到落地的四层信息

我习惯把 TVM 里一个算子的完整信息拆成四层来看,这个视角帮我在遇到报错时能快速判断该往哪个方向查。最上层是IR 定义,也就是这个算子在图里以什么形式存在:它是一个 Relay 的 Call 节点,还是一个 TIR 里的 BufferStore 序列;它接收几个参数,有没有属性 attrs,能不能被融合。这一层决定了后续所有 pass 能对这个节点做什么。

第二层是计算类型,我这里指的是两件事:一是它属于哪种计算模式,逐元素、广播、规约还是不可融合的黑盒;二是它的 Type 是什么,TensorType、TupleType 还是更基础的基础类型。计算模式决定了 shape 怎么传播,Type 决定了类型检查能不能通过,这两者经常被混在一起谈,但排查时得分清楚。

第三层是数据类型,也就是 dtype。很多人觉得 dtype 就是个字符串,写"float32"还是"float16"而已。实际上它在 TVM 内部是一个三元组结构,code 表示整数、无符号整数、浮点、布尔还是句柄,bits 表示位宽,lanes 表示向量通道数。这三者任意一个错,后端生成的代码就会在寄存器分配或者 intrinsic 选择上出问题,而且报错往往出现在很后面的代码生成阶段,离真正的原因很远。

第四层是infer shape。它不是一个函数,而是一条贯穿编译期的推导链:Relay 层通过类型关系把每个节点的输出 shape 算出来,TIR 层则通过 Buffer 的 shape 字段和符号变量把循环边界确定下来。这条链上的任何一环给出Any或者推导失败,后面所有的内存规划、调度决策都会失去依据。

1.2 为什么这四个东西必须一起理解

单独看任何一层都不难,难在它们互相咬合。举个我自己踩过的例子:我写了一个自定义算子,Relay 层面 shape 推导一切正常,打印出来是TensorType([1, 3, 224, 224], "float32"),看着特别对。但一降到 TIR,循环边界就变成了零,生成的 kernel 什么都不算。查了两个小时才发现,我在注册类型关系时返回的 shape 用的是 Python 的 int,而下游某个 pass 期望的是 PrimExpr,中间某次转换把1当成了条件判断的布尔值处理掉了。这一类问题,只有当你清楚"IR 里的 shape 字段到底是什么类型"才能反应过来。

再比如 dtype 和 shape 的联动。做量化的时候,权重是int8,激活是uint8,累加器必须是int32,输出又要转回int8。如果你在自定义算子的类型推导里把输出 dtype 硬编码成了"float32",编译能过,类型检查也能过,但数值全错,而且错得很隐蔽,因为 shape 完全正确。

我写这一节的目的就是先建立一个心智模型:IR 是骨架,Type 是标签,dtype 是刻度,shape 是尺寸。骨架决定结构,标签决定能不能挂上去,刻度决定怎么读数,尺寸决定要开多大的内存。四者缺一,整条链就断。

1.3 哪些人该重点看哪几节

如果你只是调用现成的算子搭网络,最该看的是第 5 章和第 6 章,因为你的问题九成出在 shape 推导和 dtype 不匹配上。如果你在写自定义算子或者做算子融合,第 2 章和第 3 章更重要,因为你必须知道 IR 节点长什么样、类型关系怎么注册。如果你在做量化或者低精度推理,第 4 章要逐字看,dtype 的组合规则比想象中严格得多。

至于完全没接触过 TVM 的读者,我建议先跳到第 2 章看 IR 的两套语法,再回来读第 1 章,顺序反过来会有点抽象。

2. IR定义:TVM的两套语法骨架

2.1 Relay的Expr体系与Function/Module的层级关系

Relay 是 TVM 的高层图 IR,它的设计目标是表达"算子之间的数据流",而不是"怎么算"。理解 Relay 的关键是搞清楚它的节点类型,因为你在做图变换、写 pass、注册算子的时候,操作的永远是这些节点。

最基础的是Expr,所有节点都继承它。常用的几个:Var表示变量,可以是函数参数也可以是中间变量;Constant表示常量张量,内部包着一个 NDArray;Call表示一次算子调用,带一个Op或者函数和一个参数列表;Tuple和TupleGetItem处理多返回值;Let做局部绑定;If、Match处理控制流。还有一组RefCreate、RefRead、RefWrite,用于表达可变状态,一般做图优化的人才会碰。

往上走一层是Function,它包含参数列表params、函数体body、返回值类型ret_type和属性attrs。attrs里经常藏着关键信息,比如Primitive标记表示这个函数是一个已经融合好的子图,num_inputs记录了外部输入个数。再往上是IRModule,它是一个GlobalVar到Function的映射,可以理解为"一个编译单元"。实际编译时,你传给relay.build的就是一个 IRModule。

这个层级关系很重要,因为很多报错信息里会提到 "global function" 或 "primitive function",分不清Function和IRModule就容易找错地方。我一般的做法是先把 module 打印出来,看目标是顶层函数还是某个被融合的子函数。

import tvm from tvm import relay x = relay.var("x", shape=(1, 3, 32, 32), dtype="float32") w = relay.var("w", shape=(8, 3, 3, 3), dtype="float32") y = relay.nn.conv2d(x, w, padding=(1, 1), channels=8, kernel_size=(3, 3)) f = relay.Function([x, w], y) mod = tvm.IRModule.from_expr(f) print(mod["main"].ret_type)

这段跑完打印出来是TensorType([1, 8, 32, 32], float32)。注意这里 shape 推导已经发生了——nn.conv2d注册了类型关系,构造 Call 的时候并不会立刻算,但from_expr之后如果你手动跑一次InferType,结果就填进去了。这个细节我后面第 5 章会展开。

2.2 TIR的PrimExpr与Stmt:从表达式树到可执行的循环

TIR 是 TVM 的低层 IR,它表达的是"具体怎么算"。跟 Relay 最大的区别是,TIR 里所有东西都带上了显式的内存视角:有 Buffer、有循环、有索引计算、有线程绑定。搞不定调度的人,多半是在 TIR 这一层缺少手感。

TIR 的节点分成两大类。一类是PrimExpr,包括算术表达式Add、Sub、Mul、Div、Mod,比较表达式LT、LE、EQ,逻辑表达式And、Or、Not,按位操作BitwiseAnd等,还有IntImm、FloatImm、StringImm三种立即数,以及Var和Cast。注意这里的Var是tir.Var,跟 Relay 的relay.Var不是同一个东西,类型也不同——tir.Var带的是dtype和 name,不带 shape。

另一类是Stmt,包括LetStmt、For、IfThenElse、SeqStmt、Evaluate、BufferStore、BufferLoad(在表达式那一侧)、Allocate、AssertStmt、AttrStmt,以及 TensorIR 引入的Block和BlockRealize。真正写成 TIR 的时候,你不会手写这些节点,而是用 TVMScript 或者 TE 来生成,但排查问题时需要能读懂它们。

最典型的一段 TIR 长这样:

@T.prim_func def matmul(A: T.Buffer((128, 64), "float32"), B: T.Buffer((64, 32), "float32"), C: T.Buffer((128, 32), "float32")): for i, j, k in T.grid(128, 32, 64): with T.block("C"): vi, vj, vk = T.axis.remap("SSR", [i, j, k]) with T.init(): C[vi, vj] = T.float32(0) C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vk, vj]

这里的T.Buffer((128, 64), "float32")就是 shape 的落脚点。TIR 里 shape 不是单独存在的,它挂在 Buffer 上,通过buffer.shape访问,是一个Array[PrimExpr]。这一点跟 Relay 的TensorType有本质区别,记住这个区别,后面对排查特别有用。

2.3 IR定义里那些"看起来无所谓其实很致命"的细节

第一个细节:Relay 的Var可以不带 shape 注解。relay.var("x")得到的type_annotation是IncompleteType,而不是报错。它会一直保持未完成状态,直到InferType通过上下文把它推出来。如果推不出来,就会一路传成Any,最后在后端报一个跟源头毫无关系的错。我现在的习惯是,凡是手写 Relay 图,变量一律显式写 shape,哪怕写得不准,也比留空强,因为不准会立刻报冲突,留空会沉默地传播。

第二个细节:relay.Constant的 shape 来自它包着的 NDArray,是只读的。有些人在图变换时想改常量的 shape,直接改Constant.data的 shape,结果下游读到的还是旧值。正确做法是新建一个 NDArray 再包成新的 Constant。

第三个细节:TIR 里Buffer的strides和shape是两个独立字段。多数情况下 strides 是空的,表示紧凑布局;一旦你设了非紧凑的 strides,shape 的语义就不完全是"每个维度多长"了,偏移计算会按 strides 走。做 layout 变换或者做切片传参的时候,这个坑非常常见。

第四个细节:PrimFunc的buffer_map把函数参数(tir.Var)映射到Buffer,这才是 shape 的真正来源。如果你看到一个 PrimFunc 的参数是handle类型的 Var,但buffer_map是空的,那说明这个参数没有 shape 信息,后端无法为它分配内存。这种情况在手工构造 PrimFunc 时不罕见。

3. 计算类型:Type系统到底在约束什么

3.1 从PrimType到TensorType,一张表说清职责

TVM 的 Type 体系是分层的,最底下是PrimType,它表示一个标量值的类型,只有int32、float32、bool、handle、void这几种。注意这里的int32是编译器内部的整形类型,不是张量。handle是一个不透明的指针,很多运行时参数用它表示。

往上走是PointerType和TupleType。前者包装一个 element type,表示指向它的指针;后者是多个 Type 的元组,多返回值算子返回的就是它。再往上就是最常用的TensorType,它有两个核心字段:shape和dtype,外加一个可选的span用于记录源码位置。

还有几个不常见的:FuncType描述函数签名,包含参数类型列表、返回类型和类型参数;TypeVar表示一个类型变量,用于多态;TypeData用于自定义类型;IncompleteType就是前面说的"还没推出来"。

我用一张表把这几个类型的分工整理一下:

Type 类核心字段典型出现位置排查时的关注点
PrimTypedtypeTIR 表达式、循环变量是否被误当成张量传入
TensorTypeshape, dtypeRelay 节点输出、算子签名shape 是否含 Any,dtype 是否一致
TupleTypefields多返回值算子索引是否对得上
FuncTypearg_types, ret_type, type_params高阶函数、算子注册参数个数与类型是否匹配
IncompleteType无未注解的 Var是否一路传播到后端
TypeVarkind多态算子是否缺少约束

这张表我贴在显示器边上贴了很久,遇到"expected TensorType but got ..."这类报错,先查是哪个类型混进来了,基本能定位。

3.2 别把"计算类型"和"计算模式"混为一谈

中文里说"计算类型",很容易产生歧义。一种理解是上面讲的 Type 系统,另一种理解是算子的计算模式,也就是 Relay 里那个TOpPattern。这两个东西完全是两回事,但在 shape 推导这件事上,后者其实影响更大,所以必须一起讲。

TOpPattern是算子注册时的一个属性,常见取值有kElemWise(逐元素)、kBroadcast(广播)、kInjective(单射)、kCommReduce(规约)、kOutEWiseFusable(输出可逐元素融合)、kOpaque(黑盒)。它决定了算子融合 pass 会不会把这个算子跟邻居融到一起。

关键点在于,不同的 pattern 对应完全不同的 shape 传播规则。逐元素算子的输出 shape 直接等于输入 shape,不需要计算;广播算子要按 NumPy 的广播规则逐维对齐,长度不等的维度里面必须有一个是 1 或者Any;规约算子要在某个轴上把维度消掉或者压成 1。如果你注册自定义算子时 pattern 写错了,会有两种严重后果:写得太"激进"(比如把规约算子标成kElemWise),shape 推导会按错的规则走,得到错误的输出尺寸;写得太"保守"(把逐元素算子标成kOpaque),融合完全失效,性能掉一大截。

我自己写过的一个教训是,把一个带 broadcast 语义的加法标成了kElemWise,输入是(1, 128)和(64, 128),结果输出 shape 被推成了(1, 128),后面接的算子全部跟着错,一直错到内存分配才崩。从那以后我养成了一个习惯:凡是输入 shape 可能不等长的算子,一律先按kBroadcast写,跑通之后再考虑要不要收紧。

3.3 TypeRelation与多态:类型不是静态贴上去的标签

这一点很多人理解不到位。TVM 的 Type 不是你在构造节点时"贴上"就固定了的,它是一个可以被反复推导、被约束、被求解的变量。Relay 的类型系统支持多态,FuncType里可以有type_params,一个算子可以对多种输入类型都有定义,具体用哪个要看调用点的实际参数。

这个机制通过TypeRelation实现。注册类型关系的时候,你拿到的不是具体类型,而是一组arg_types和一组type_vars,你要做的是"给定这些输入类型约束,给出输出类型约束,或者反过来"。比如nn.conv2d的类型关系会检查输入是不是 4 维、通道数是否匹配、dtype 是否一致,然后算出输出的 4 维 shape。

这就带来一个反直觉的现象:类型推导是双向的。正常情况下是从输入推输出,但如果你给函数的返回类型写了注解,推导也能从输出反推输入的一部分信息。这就是为什么有时候你给一个变量不写 shape,它也能自己推出来——上下文里有信息。

InferTypepass 的工作方式是迭代到不动点:遍历所有节点,对每个 Call 调用对应算子的类型关系,更新输出类型;如果某个节点的输入类型变了,就把它加入待处理队列重新处理。所以一个函数的推导可能要跑好几轮才稳定。理解了这一点,你就能明白为什么单独调一次某个算子的推导函数,结果跟在完整图里跑是不一样的——图里还有别的约束在起作用。

4. 数据类型:DLDataType的三个数字

4.1 code、bits、lanes拆开看

TVM 的 dtype 在底层是一个叫DLDataType的结构体,只有三个字段:code、bits、lanes。你在 Python 里看到的"float32"字符串,解析之后就是这个三元组。

code表示数据类别,用整数编码。整型、无符号整型、浮点、不透明句柄、bfloat、复数、布尔,各自对应一个固定值。你在写自定义算子时如果要手工构造DataType,最稳妥的方式是传字符串让它自己解析,而不是记那些数字,因为不同版本的编码可能微调。

bits表示每个元素的位宽,常见的是 8、16、32、64,另外还有一个特殊的 1,用于位掩码或者布尔。做量化的时候会用到非常规位宽,比如int4、uint4、int1,这些在 TVM 里是支持的,但后端不一定都实现了对应的硬件指令,需要额外注册。

lanes表示向量通道数,默认是 1。当 lanes 大于 1 时,表示这是一个 SIMD 向量类型,比如float32x4。这一块跟调度里的T.vectorized直接相关,向量化之后循环体内的表达式就用带 lanes 的类型来表达。

这三者的组合空间比你想象中大。同样叫"32 位",int32、uint32、float32是完全不同的三样东西,在两两混用时会触发类型提升规则。而float32和float32x4在类型系统里是两个不同 dtype,不能直接相加,需要显式处理。

4.2 dtype字符串解析与几个高频误用

写"float32"是对的,写"float"在有些地方会被宽松解析成 32 位浮点,但不要依赖这种宽松。我见过因为写了"float"导致某个版本下被解析成 64 位的情况,整个 kernel 的内存翻倍,性能直接掉一半,而且不报错。

另一个高频误用是"int"和"int32"混用。在 TIR 里,索引计算默认用的是 32 位整型,如果你显式写"int64"去做索引,某些 pass 会插入大量的 cast,生成的代码里全是类型转换指令。除非你的张量元素总数确实超过 32 位能表达的范围,否则索引一律用 32 位。

还有一个坑是布尔类型。Relay 里的比较运算返回bool,但bool在内存里的布局不是标准的一字节。做位运算(bitwise_and之类)时如果传入 bool,某些后端会报错或者产生意外的结果。需要位运算的地方,先 cast 成int8或uint8再做。

关于handle类型,它一般只出现在函数参数、外部调用和设备指针上,不要用它来传递张量。如果你在算子的输出类型里看到了handle,基本可以确定是注册时签名写错了。

4.3 量化与混合精度场景下的dtype组合策略

量化是 dtype 出问题最密集的场景。我总结了一个组合表,能覆盖大部分情况:

场景输入 dtype累加 dtype输出 dtype注意事项
对称量化卷积int8int32int8缩放因子需在类型推导里保留
非对称量化卷积uint8int32uint8零点偏移要参与计算
混合精度注意力float16float32float16累加溢出风险高
bfloat16 训练bfloat16float32bfloat16后端支持度差异大
整数矩阵乘int8int32int32输出常需要再饱和

在这张表之外,还有一条我自己的经验:凡是做跨 dtype 的算子,一定要在类型关系里显式做类型检查,而不是靠下游报错。比如你写了一个卷积,输入是int8、权重是uint8,累加器必须是int32,输出你想给int8。如果只在计算里做转换而类型关系返回的是float32,编译能过,但 Schedule 选型会按浮点算子的逻辑走,出来的 kernel 性能极差而且数值不可信。

另外,量化的 dtype 会直接影响 shape 推导吗?一般不直接影响,但有一个例外:打包量化(packed int4 之类)会把多个元素塞进一个存储单元,这时候逻辑 shape 和物理 shape 不一样,Buffer 的 shape 要按物理布局写,而 TensorType 的 shape 要按逻辑尺寸写。这两者的转换需要在算子级别显式处理,是最容易算错内存大小的地方。

5. infer shape:整条链路里最容易翻车的一环

5.1 Relay的推导流程与pass顺序

Relay 的类型推导由InferType这个 pass 承担。它做的事情前面提过,是迭代到不动点。但真正容易出问题的是它在 pass 序列里的位置。

举个典型情况:如果你先跑FoldConstant(常量折叠),再跑InferType,那次序是对的,因为常量折叠会产生新的 Constant 节点,它们的类型需要被重新推。如果你反过来,先推类型再折叠,折叠产生的新节点可能带着不完整的类型信息进入后续 pass。所以标准的优化序列里,InferType会被多次调用,插入在关键步骤之后。

from tvm import relay mod = relay.transform.InferType()(mod) print(mod["main"].ret_type) seq = tvm.transform.Sequential([ relay.transform.InferType(), relay.transform.SimplifyInference(), relay.transform.FoldConstant(), relay.transform.InferType(), ]) mod = seq(mod)

注意SimplifyInference会做一些算子重写,比如把 BatchNorm 拆开,这必然会改动 shape 的传播路径,所以它后面要再跟一次InferType。这个模式我在自己的 pass 序列里一直保留,省了很多莫名其妙的错。

还有一个值得单独说的 attribute:TShapeDataDependent。有些算子的输出 shape 依赖输入的数据内容,比如nonzero、unique、argwhere。这类算子没法在编译期算出 shape,TVM 的处理方式是给它们的输出标上Any,同时在算子属性里打标,告诉后续 pass 不要尝试常量折叠这些节点。如果你的自定义算子的输出尺寸跟数据值有关,一定要记得打这个标,否则FoldConstant可能真的去执行它并得到错误的结果。

5.2 TIR里的shape从哪来,符号变量怎么处理

到了 TIR 这一层,shape 的载体变成了 Buffer。buffer.shape是一个Array[PrimExpr],每个元素要么是IntImm(编译期常量),要么是tir.Var(符号变量),要么是它们的组合表达式。

符号变量的引入是支持动态 shape 的基础。你写n = tir.Var("n", "int32"),然后构造 shape 为(n, 128)的 Buffer,这个n会作为 PrimFunc 的一个隐式参数或者从外部传入。这时候调度和代码生成都必须处理"长度未知"的情况,比如循环上界是n而不是常数。

from tvm import te n = te.var("n") A = te.placeholder((n, 128), name="A", dtype="float32") B = te.compute((n, 128), lambda i, j: A[i, j] * 2.0, name="B") func = te.create_prim_func([A, B]) print(func.script())

打印出来你会看到A: T.Buffer((n, 128), "float32")这样的签名,其中n是一个符号。这里有个细节:te.var返回的是tir.Var,但它的 dtype 是int32,可以直接用在 shape 表达式里;如果你混用 Python 的 int 和它,构造 shape 时会被自动提升,但混合类型在后续的算术简化里可能得不到最好的优化,我一般会显式tvm.tir.IntImm包装常量,让整个表达式类型统一。

还有一个非常实用但少人提的技巧:tvm.tir.analysis里有工具可以分析 Buffer 的访问范围,用来验证你推导出的 shape 是否覆盖了所有实际访问。如果你怀疑某个 Buffer 的 shape 太小或者 strides 算错,用它扫一遍比逐个循环看索引快得多。

5.3 手把手写一个shape推导:自定义算子的完整流程

光讲原理不够,我给一个能跑的最小例子。假设我要实现一个"带缩放的残差相加",输入两个张量,支持广播,输出是a + b * scale。

第一步是定义算子,注册参数和 pattern:

from tvm import relay, te from tvm.relay.op import op as _op def _scaled_add_shape(attrs, inputs, tinfos): # inputs 是一组 PrimExpr 形式的 shape a_shape, b_shape = inputs out_shape = [] for a, b in zip(reversed(a_shape), reversed(b_shape)): out_shape.append(te.if_then_else(a == b, a, te.max(a, b))) for extra in a_shape[:len(a_shape) - len(b_shape)]: out_shape.append(extra) for extra in b_shape[:len(b_shape) - len(a_shape)]: out_shape.append(extra) return out_shape[::-1]

这段_scaled_add_shape是给动态 shape 场景用的 shape function,注意它接收的是PrimExpr,返回的也是PrimExpr列表,不能混入 Python int。写的时候我踩过一次,te.max(a, b)里如果a是IntImm而b是Var,类型推导能过,但简化阶段会留下冗余的if_then_else。

第二步是注册类型关系,管编译期的静态推导:

@_op.register_infer_type("scaled_add") def _scaled_add_infer(attrs, arg_types, type_vars, base_type): a_ty, b_ty = arg_types assert isinstance(a_ty, relay.TensorType), "输入必须是张量" out_dtype = a_ty.dtype if a_ty.shape is None or b_ty.shape is None: return relay.TensorType(relay.Any(), out_dtype) a_shape = list(a_ty.shape) b_shape = list(b_ty.shape) out_shape = [] for a, b in zip(reversed(a_shape), reversed(b_shape)): if isinstance(a, int) and isinstance(b, int) and a != b and a != 1 and b != 1: raise ValueError(f"维度无法广播: {a} vs {b}") out_shape.append(max(a, b) if not isinstance(a, relay.Any) else b) out_shape += list(a_shape[:len(a_shape) - len(b_shape)]) out_shape += list(b_shape[:len(b_shape) - len(a_shape)]) return relay.TensorType(out_shape[::-1], out_dtype)

这里有几处是刻意写详细的,都是实战中会踩的点。一是shape可能为None,必须先判空;二是静态推导阶段拿到的维度既可能是 Python int 也可能是Any,做比较前要分别处理;三是广播规则要按从后往前的顺序对齐,长度不等时把长的那段直接接到前面。至于注册接口的模块路径,不同版本会有点差异,如果tvm.relay.op.op里找不到,试试tvm.relay.op,这两个位置我在不同版本上都见过。

第三步是验证。构造图,跑InferType,然后对比手算结果:

a = relay.var("a", shape=(1, 128), dtype="float32") b = relay.var("b", shape=(64, 128), dtype="float32") out = _op.get("scaled_add")(a, b, scale=0.5) f = relay.Function([a, b], out) mod = tvm.IRModule.from_expr(f) mod = relay.transform.InferType()(mod) print(mod["main"].ret_type) # 期望 (64, 128)

如果这里打印出来是(1, 128),说明你的广播规则写反了,max取成了第一项。如果打印出来带Any,说明某条分支没走到,检查是不是shape is None的判断把它短路了。如果直接抛了Cannot infer type,多半是_op.get时属性没注册全,比如漏了TOpPattern,Relay 在构造 Call 时找不到合法的推断路径。

第四步是降层验证。类型推对了还不够,得确认降到 TIR 之后 Buffer 的 shape 也是对的:

from tvm.relay.backend import te_compiler funcs = te_compiler.get_te_compiler().lower_call(...)

这一步我一般会直接打一个简单的relay.build并 dump 出 TIR,看prim_func里 Buffer 的 shape 是不是(64, 128)。这一步的价值在于,它能暴露 Relay 层看不出来的问题,比如 dtype 在降层时被意外提升,或者某个符号变量被错误地常量折叠掉了。

6. 常见问题与排查实录

6.1 报错信息与根因的对照

下面这张表是我这几年攒下来的,左边是报错原文(或它的关键片段),右边是我实际查到的根因。有些报错信息很误导人,光看字面会查错方向。

报错或现象常见根因首选排查动作
Cannot infer type of ...变量没有 shape 注解,上下文也推不出打印 module,看IncompleteType出现在哪个 Var
The type relation of op xxx is not implemented自定义算子没注册FInferType检查注册装饰器是否生效,版本路径是否正确
shape 推出来是Any输入含Any,或算子被标为 data dependent逐层打印中间节点的ret_type
编译通过但结果全零TIR 循环边界为 0,shape 表达式被误判为布尔dump TIR,看T.grid的参数
expected shape ... but got ...上下游 shape 不一致,通常是 broadcast 规则写错对比两侧 shape,检查 pattern 属性
dtype 不匹配相关报错量化链路中累加类型写错逐节点打印 dtype,检查类型关系里的输出 dtype
内存占用异常大索引用了 64 位,或 shape 被误推成超大值检查 dtype 字符串是否被宽松解析
融合后性能反而下降自定义算子 pattern 标成了kOpaque改成kElemWise或kBroadcast重试

这张表最有用的其实是第一行和第三行。因为Any是会传播的,一旦某个节点推不出来,它的所有下游都会变成Any,最后报错的位置往往是最后一个算子,而不是最开始的源头。我的做法是从图的开头按拓扑序扫一遍,找到第一个ret_type里出现Any或IncompleteType的节点,那才是真正的问题点。

6.2 一套可以照抄的排查顺序

我自己总结了一个五步排查法,顺序不能乱,乱了容易做无用功。

第一步,打印 IR。print(mod)或者mod.show(),把整个图结构看清楚。这一步看似废话,但我见过太多人直接跳进去改代码,结果改了半天发现报错的函数根本不是他以为的那个。图里如果有primitive标记的函数,要特别注意,那是融合后的子图,shape 推导是在融合前做的。

第二步,跑一次InferType并逐节点检查。用post_order_visit遍历所有 Call 节点,打印它们的checked_type。找到第一个类型不对的节点,问题基本就在它或者它的直接上游。

from tvm.relay.analysis import post_order_visit def collect(node): if isinstance(node, relay.Call): print(node.op, "->", node.checked_type) post_order_visit(mod["main"].body, collect)

第三步,核对 dtype。把每个节点的 dtype 打印出来,跟你的预期逐个比。量化场景下这一步几乎每次都能捞到问题,因为累加类型经常被忽略。我一般会专门写个小函数,把整条链的 dtype 打成一个列表,一眼扫过去就能看出哪里跳变了。

第四步,降层看 TIR。如果 Relay 层全对,问题就在降层。dump 出 TIR 之后,重点看三处:PrimFunc 的buffer_map里每个 Buffer 的 shape、循环的上界、以及索引表达式里的 cast。这三处任意一处不对,都能解释大部分"编译过了但结果不对"的问题。

第五步,用最小用例复现。如果前面四步都定位不到,把图裁到一个算子、两个输入,单独跑。我遇到过好几次问题只在多算子组合时才出现,裁到最小之后反而全对了,那说明是 pass 顺序或者属性传播的问题,而不是单个算子的错。

6.3 几个只在实际项目里才会遇到的坑

第一个坑跟动态 shape 有关。当你在 Relay 层用relay.Any()表示动态维度时,编译期一切正常,但到了运行时,后端的 shape function 可能没被注册。这个错误的报错信息通常很晦涩,提到shape function或vm相关的东西。根因是:编译期类型推导和运行期 shape 计算是两套机制,前者靠FInferType,后者靠FShapeFunc,两个都得注册。我一开始只注册了前者,静态 shape 场景全对,一上动态 shape 就崩,查了很久才想明白这两套机制是分开的。

第二个坑跟Any的传播有关。只要一个Any进了图,它就会沿着数据流一路往下传,中间只要有一个算子能做常量折叠,它就可能被折叠成一个具体的错误值。所以我的原则是:编译期能确定 shape 的地方,绝不引入Any。只有在输入尺寸真的逐次变化时(比如自然语言处理的变长序列),才用动态维度,而且要在图的边界处就把它收住。

第三个坑跟 layout 变换有关。某些算子(尤其是卷积类)在优化时会做 layout 重排,比如从通道优先变成通道在后。这个重排会改变 shape 的维度顺序,而FInferCorrectLayout就是负责这件事的。如果你自定义的算子没有正确实现这个属性,图优化可能会给它插入不必要甚至错误的 layout 变换,shape 跟着变形。排查时如果发现某个节点的 shape 维度顺序跟你写的不一样,先查这个属性。

第四个坑是常量折叠和 shape 推导的顺序冲突。前面提过FoldConstant会产生新的 Constant,但如果这个 Constant 的来源是一个 data dependent 的算子,折叠就会出错。表现是编译期不报错,运行时结果完全不对。防御手段是给这类算子打上TShapeDataDependent标,让折叠 pass 跳过它们。

第五个坑跟 dtype 的隐式提升有关。TVM 在做算术时会按一套规则提升类型,int8 + int8在某些路径下会被提升成int32,float16 * float32会提升成float32。这些提升在数值上是安全的,但会让你的中间张量内存翻倍。做低精度推理时,我一般会在关键节点后显式插入 cast,把类型收回到目标精度,而不是依赖自动提升。

排查这件事说到底就是个耐心活,工具都很简单,难的是知道往哪看。上面这套顺序我用了很久,基本能覆盖八成以上的情况。剩下两成,通常是多个问题叠加,得拆开逐个解决,这时候最小复现用例的价值就体现出来了——裁到最后,往往能发现两个看起来毫无关系的错误其实是同一个根因。

至于日常开发中怎么少踩坑,我的做法是凡是手写的算子,shape 推导和 dtype 一定先单独测通再往图里放,别指望在大图里调试;凡是引入动态维度,先在最小模型上跑一遍端到端,确认 shape function 注册链路是通的。这些看起来是额外工作,但比事后在图里大海捞针要省时间得多。

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询