从算子定义写完到第一次跑通编译,中间最容易卡住的地方,往往不是调度写不出来,而是类型和 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 类 | 核心字段 | 典型出现位置 | 排查时的关注点 |
|---|---|---|---|
| PrimType | dtype | TIR 表达式、循环变量 | 是否被误当成张量传入 |
| TensorType | shape, dtype | Relay 节点输出、算子签名 | shape 是否含 Any,dtype 是否一致 |
| TupleType | fields | 多返回值算子 | 索引是否对得上 |
| FuncType | arg_types, ret_type, type_params | 高阶函数、算子注册 | 参数个数与类型是否匹配 |
| IncompleteType | 无 | 未注解的 Var | 是否一路传播到后端 |
| TypeVar | kind | 多态算子 | 是否缺少约束 |
这张表我贴在显示器边上贴了很久,遇到"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 | 注意事项 |
|---|---|---|---|---|
| 对称量化卷积 | int8 | int32 | int8 | 缩放因子需在类型推导里保留 |
| 非对称量化卷积 | uint8 | int32 | uint8 | 零点偏移要参与计算 |
| 混合精度注意力 | float16 | float32 | float16 | 累加溢出风险高 |
| bfloat16 训练 | bfloat16 | float32 | bfloat16 | 后端支持度差异大 |
| 整数矩阵乘 | int8 | int32 | int32 | 输出常需要再饱和 |
在这张表之外,还有一条我自己的经验:凡是做跨 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 注册链路是通的。这些看起来是额外工作,但比事后在图里大海捞针要省时间得多。