()
| 435 | |
| 436 | |
| 437 | def test_call_tir(): |
| 438 | # also from test_parser |
| 439 | @tvm.script.ir_module |
| 440 | class TestCallTIR: |
| 441 | @T.prim_func(s_tir=True) |
| 442 | def addone(A_handle: T.handle, B_handle: T.handle) -> None: |
| 443 | m = T.int64() |
| 444 | n = T.int64() |
| 445 | A = T.match_buffer(A_handle, (m, n), "float32") |
| 446 | B = T.match_buffer(B_handle, (m, n), "float32") |
| 447 | T.func_attr({"global_symbol": "addone"}) |
| 448 | for i, j in T.grid(m, n): |
| 449 | with T.sblock("addone"): |
| 450 | vi, vj = T.axis.remap("SS", [i, j]) |
| 451 | B[vi, vj] = A[vi, vj] + T.int32(1) |
| 452 | |
| 453 | @R.function |
| 454 | def foo(x: R.Tensor(("m", "n"), "float32")): |
| 455 | m, n = T.int64(), T.int64() |
| 456 | gv0 = R.call_tir(TestCallTIR.addone, (x,), R.Tensor((m, n), dtype="float32")) |
| 457 | return gv0 |
| 458 | |
| 459 | mod = TestCallTIR |
| 460 | foo = mod["foo"] |
| 461 | |
| 462 | foo_str = strip_whitespace( |
| 463 | dump_ast( |
| 464 | foo, |
| 465 | include_struct_info_annotations=False, |
| 466 | include_call_attrs=False, |
| 467 | ) |
| 468 | ) |
| 469 | assert foo_str.startswith('Function(params=[Var(name_hint="x")]') |
| 470 | |
| 471 | # call_tir is an op in Relax and it takes an extern func as an argument |
| 472 | assert isinstance(foo.body, rx.SeqExpr) |
| 473 | tir_call = foo.body.blocks[0].bindings[0].value |
| 474 | tir_call_text = dump_ast( |
| 475 | tir_call, |
| 476 | include_struct_info_annotations=False, |
| 477 | include_call_attrs=False, |
| 478 | ) |
| 479 | assert_fields( |
| 480 | "Call", |
| 481 | { |
| 482 | "op": 'Op(name="relax.call_tir")', |
| 483 | "args": """[ |
| 484 | GlobalVar(name_hint="addone"), |
| 485 | Tuple(fields=[Var(name_hint="x")]) |
| 486 | ]""", |
| 487 | "sinfo_args": """[ |
| 488 | TensorStructInfo( |
| 489 | dtype=float32, |
| 490 | shape=ShapeExpr( |
| 491 | values=[ |
| 492 | PrimExpr(value=`m`), |
| 493 | PrimExpr(value=`n`) |
| 494 | ] |
nothing calls this directly
no test coverage detected
searching dependent graphs…