|
28 | 28 | from tvm import te |
29 | 29 | from tvm.contrib import clang, utils |
30 | 30 | from tvm.relay.backend import Runtime |
31 | | -from tvm.script import tir as T |
| 31 | +from tvm.script import tir as T, ir as I |
32 | 32 | from tvm.target.codegen import llvm_get_intrinsic_name, llvm_lookup_intrinsic_id |
33 | 33 |
|
34 | 34 |
|
@@ -1007,19 +1007,25 @@ def test_debug_symbol_for_float64(): |
1007 | 1007 | """Check that LLVM can define DWARF debug type for float64 |
1008 | 1008 |
|
1009 | 1009 | In previous versions, only specific data types could exist in the |
1010 | | - function signature. In this test, the "calling_conv" attribute |
1011 | | - prevents lowering to the PackedFunc API. |
| 1010 | + function signature. In this test, the subroutine is a context |
| 1011 | + that isn't lowered to the PackedFunc API. |
1012 | 1012 | """ |
1013 | 1013 |
|
1014 | | - @T.prim_func |
1015 | | - def func(a: T.handle("float64"), b: T.handle("float64"), n: T.int64): |
1016 | | - T.func_attr({"calling_conv": 2}) |
1017 | | - A = T.Buffer(16, "float64", data=a) |
1018 | | - B = T.Buffer(16, "float64", data=b) |
1019 | | - for i in range(n): |
1020 | | - B[i] = A[i] |
1021 | | - |
1022 | | - tvm.build(func, target="llvm") |
| 1014 | + @I.ir_module |
| 1015 | + class mod: |
| 1016 | + @T.prim_func |
| 1017 | + def main(A: T.Buffer(16, "float64"), B: T.Buffer(16, "float64")): |
| 1018 | + T.func_attr({"global_symbol": "main"}) |
| 1019 | + mod.subroutine(A.data, B.data, T.int64(16)) |
| 1020 | + |
| 1021 | + @T.prim_func |
| 1022 | + def subroutine(a: T.handle("float64"), b: T.handle("float64"), n: T.int64): |
| 1023 | + A = T.Buffer(16, "float64", data=a) |
| 1024 | + B = T.Buffer(16, "float64", data=b) |
| 1025 | + for i in range(n): |
| 1026 | + B[i] = A[i] |
| 1027 | + |
| 1028 | + tvm.build(mod, target="llvm") |
1023 | 1029 |
|
1024 | 1030 |
|
1025 | 1031 | if __name__ == "__main__": |
|
0 commit comments