I have fixed this issue, I don’t know if this counts as a bug.
I found that the register of the operator is in the file: tvm/tir/tensor_intrin/cuda.py. But it doesn’t run, so I have to “import tvm.tir.tensor_intrin.cuda” at the beginning of the program. Then it works.
Thanks for reporting the fix! I think this is due to the fact that we have tesnor intrinsics written in TensorIR so we have to import certain files before we run tuning.
It enters the ‘_normalize_rules’ function, run ‘ScheduleRule.create(sch_rules)’ on line 171 of this file
Then, in ‘create’, the program correctly matches “cuda-tensorcore” and “_ffi_api. ScheduleRuleDefaultCUDATensorCore”, and ran the latter.
The program then jumps to the “PackedFuncBase” class in “tvm._ffi._ctypes.packed_func.py” and runs “call”. Here, the error is thrown.
It doesn’t create ‘MetaSchedule tuning contexts’ during working. In fact, I also had this bug on another computer (using tvm.012, but with code from a month and a half ago). Not sure if there is a problem with such a workflow.
error: unsupported operand type(s) for +: 'NoneType' and 'int'
--> Traceback (most recent call last):
File
"/home/acer/tvm/python/tvm/script/parser/core/parser.py",
line 204, in _wrapper
return func(self, node)
File "/home/acer/tvm/python/tvm/script/parser/tir/parser.py",
line 254, in visit_assign
rhs = self.eval_expr(node.value)
File
"/home/acer/tvm/python/tvm/script/parser/core/parser.py",
line 329, in eval_expr
return eval_expr(self, node, var_values)
File
"/home/acer/tvm/python/tvm/script/parser/core/evaluator.py",
line 412, in eval_expr
return ExprEvaluator.eval(parser, value_table, node)
File
"/home/acer/tvm/python/tvm/script/parser/core/evaluator.py",
line 107, in eval
result = self._visit(node) # pylint: disable=protected-access
File
"/home/acer/tvm/python/tvm/script/parser/core/evaluator.py",
line 195, in _visit
s.upper.end_col_offset + 1,
TypeError: unsupported operand type(s) for +: 'NoneType'
and 'int'