[meta_schedule]Try to use cuda-tensorcore in meta schedule, but get not registered error

here is my code

When I try to create ‘space’ , I get the error “ValueError: TensorIntrin ‘wmma_fill_16x16x16_f32’ is not registered”.

According to my test, this bug only appears when I set sch_rules to “cuda-tensorcore”. The gpu I am using is 3090.

The PrimFunc I want to optimize is my custom calculation. I don’t know if this will affect.

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.

We will need to import this module to register intrinsics:

import tvm.tir.tensor_intrin  # pylint: disable=unused-import

It’s usually lazily imported when creating MetaSchedule tuning contexts, but for some reason it’s not triggered in your case

1 Like

I debugged my program. When the space variable is created, it goes to the “PostOrderApply” class and run ‘init’. When running to line 53:

“sch_rules, postprocs, mutator_probs = _normalize_rules(sch_rules, postprocs, mutator_probs)”

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.

Does it solve your problem if you added a line in your script importing tvm.tir.tensor_intrin?

Yes, ‘import tvm.tir.tensor_intrin’ solves my problem. :wink:

1 Like

Is anyone successuly tune onnx model using meta scheduler?

@junrushao

when importing the line, i got

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'

Hi can you please share a script to reproduce the error?