From ddeb8685fb86b4358e3dd979a43040fab1a261b7 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Thu, 14 Jul 2022 23:32:27 -0700 Subject: [PATCH 1/9] unittests fix 0 --- .../unittest/test_arith_domain_touched.py | 8 ++++---- ...est_tir_transform_compact_buffer_region.py | 20 +++++++++---------- .../test_tir_transform_storage_flatten.py | 2 +- ...st_tir_usmp_analysis_extract_bufferinfo.py | 2 +- tests/python/unittest/test_tir_usmp_utils.py | 2 +- .../unittest/test_tvmscript_roundtrip.py | 18 ++++++++--------- 6 files changed, 26 insertions(+), 26 deletions(-) diff --git a/tests/python/unittest/test_arith_domain_touched.py b/tests/python/unittest/test_arith_domain_touched.py index 8b982e65a055..4e508b283df1 100644 --- a/tests/python/unittest/test_arith_domain_touched.py +++ b/tests/python/unittest/test_arith_domain_touched.py @@ -23,8 +23,8 @@ def scalar_func(a: T.handle, b: T.handle): m = T.var("int32") n = T.int32(100) - A = T.match_buffer(a, (n, m), name="A") - B = T.match_buffer(b, (n, m), name="B") + A = T.match_buffer(a, (n, m)) + B = T.match_buffer(b, (n, m)) for i, j in T.grid(n, m): A[i, j] = B[i - 1, j + 1] + A[i - 1, j - 1] @@ -34,8 +34,8 @@ def scalar_func(a: T.handle, b: T.handle): def vector_func(a: T.handle, b: T.handle): n = T.var("int32") m = T.int32(128) - A = T.match_buffer(a, (n, m), name="A") - B = T.match_buffer(b, (n, m), name="B") + A = T.match_buffer(a, (n, m)) + B = T.match_buffer(b, (n, m)) for i in T.serial(n): for j in T.vectorized(m): diff --git a/tests/python/unittest/test_tir_transform_compact_buffer_region.py b/tests/python/unittest/test_tir_transform_compact_buffer_region.py index 5d8b99e7d055..cbc3e7fba3d1 100644 --- a/tests/python/unittest/test_tir_transform_compact_buffer_region.py +++ b/tests/python/unittest/test_tir_transform_compact_buffer_region.py @@ -372,7 +372,7 @@ def compacted_storage_align_func(a: T.handle, c: T.handle) -> None: with T.block(): T.reads(A[i, 0:16]) T.writes(C[i, 0:16]) - B = T.alloc_buffer((1, 16), strides=(31, 1), dtypes="float32") + B = T.alloc_buffer((1, 16), strides=(31, 1), dtype="float32") for j in range(0, 16): with T.block(): T.reads(A[i, j]) @@ -391,7 +391,7 @@ def padding_pattern_func(a: T.handle, c: T.handle) -> None: A = T.match_buffer(a, (16, 16), "float32") C = T.match_buffer(c, (20, 20), "float32") with T.block(): - B = T.alloc_buffer((20, 20), dtypes="float32") + B = T.alloc_buffer((20, 20), dtype="float32") for i, j in T.grid(16, 16): with T.block(): B[i, j] = A[i, j] @@ -473,10 +473,10 @@ def compacted_padding_pattern_inlined( def mem_access_in_branch_func(a: T.handle) -> None: A = T.match_buffer(a, (224, 224), "float32") with T.block(): - B1 = T.alloc_buffer((224, 224), dtypes="float32") - B2 = T.alloc_buffer((224, 224), dtypes="float32") - B3 = T.alloc_buffer((224, 224), dtypes="float32") - B4 = T.alloc_buffer((224, 224), dtypes="float32") + B1 = T.alloc_buffer((224, 224), dtype="float32") + B2 = T.alloc_buffer((224, 224), dtype="float32") + B3 = T.alloc_buffer((224, 224), dtype="float32") + B4 = T.alloc_buffer((224, 224), dtype="float32") for i in range(0, 224): for j in range(0, 224): with T.block(): @@ -519,8 +519,8 @@ def compacted_mem_access_in_branch_func(a: T.handle) -> None: def opaque_access_annotated_func(a: T.handle) -> None: A = T.match_buffer(a, (1024,), "float32") with T.block(): - B = T.alloc_buffer((1024,), dtypes="float32") - C = T.alloc_buffer((1024,), dtypes="float32") + B = T.alloc_buffer((1024,), dtype="float32") + C = T.alloc_buffer((1024,), dtype="float32") for i in range(0, 512): with T.block(): # no annotation, opaque access will cover full region @@ -541,8 +541,8 @@ def opaque_access_annotated_func(a: T.handle) -> None: def compacted_opaque_access_annotated_func(a: T.handle) -> None: A = T.match_buffer(a, (1024,), "float32") with T.block(): - B = T.alloc_buffer((1024,), dtypes="float32") - C = T.alloc_buffer((520,), dtypes="float32") + B = T.alloc_buffer((1024,), dtype="float32") + C = T.alloc_buffer((520,), dtype="float32") for i in range(0, 512): with T.block(): # no annotation, opaque access will cover full region diff --git a/tests/python/unittest/test_tir_transform_storage_flatten.py b/tests/python/unittest/test_tir_transform_storage_flatten.py index b84b3479fe9a..dbf6f18980af 100644 --- a/tests/python/unittest/test_tir_transform_storage_flatten.py +++ b/tests/python/unittest/test_tir_transform_storage_flatten.py @@ -142,7 +142,7 @@ def main(): A_data: T.Ptr[T.int32] = T.call_extern("dummy_extern_function", dtype="handle") # and a buffer is backed by that pointer, - A: T.Buffer = T.buffer_decl([1], dtype="float32", data=A_data) + A = T.buffer_decl([1], dtype="float32", data=A_data) T.evaluate(A[0]) # then the call to StorageFlatten would result in an exception diff --git a/tests/python/unittest/test_tir_usmp_analysis_extract_bufferinfo.py b/tests/python/unittest/test_tir_usmp_analysis_extract_bufferinfo.py index 301dc16d2127..d4e62362495c 100644 --- a/tests/python/unittest/test_tir_usmp_analysis_extract_bufferinfo.py +++ b/tests/python/unittest/test_tir_usmp_analysis_extract_bufferinfo.py @@ -111,7 +111,7 @@ class LinearStructure: def tvmgen_default_fused_cast_subtract(placeholder_2: T.handle, placeholder_3: T.handle, T_subtract: T.handle) -> None: # function attr dict T.func_attr({"global_symbol": "tvmgen_default_fused_cast_subtract", "tir.noalias": True}) - placeholder_4 = T.match_buffer(placeholder_2, [150528], dTpe="uint8", elem_offset=0, align=128, offset_factor=1) + placeholder_4 = T.match_buffer(placeholder_2, [150528], dtype="uint8", elem_offset=0, align=128, offset_factor=1) placeholder_5 = T.match_buffer(placeholder_3, [1], dtype="int16", elem_offset=0, align=128, offset_factor=1) T_subtract_1 = T.match_buffer(T_subtract, [452], dtype="int16", elem_offset=0, align=128, offset_factor=1) # body diff --git a/tests/python/unittest/test_tir_usmp_utils.py b/tests/python/unittest/test_tir_usmp_utils.py index 2034b072838d..6e53bcb5e597 100644 --- a/tests/python/unittest/test_tir_usmp_utils.py +++ b/tests/python/unittest/test_tir_usmp_utils.py @@ -31,7 +31,7 @@ class LinearStructure: def tvmgen_default_fused_cast_subtract(placeholder_2: T.handle, placeholder_3: T.handle, T_subtract: T.handle) -> None: # function attr dict T.func_attr({"global_symbol": "tvmgen_default_fused_cast_subtract", "tir.noalias": True}) - placeholder_4 = T.match_buffer(placeholder_2, [150528], dTpe="uint8", elem_offset=0, align=128, offset_factor=1) + placeholder_4 = T.match_buffer(placeholder_2, [150528], dtype="uint8", elem_offset=0, align=128, offset_factor=1) placeholder_5 = T.match_buffer(placeholder_3, [1], dtype="int16", elem_offset=0, align=128, offset_factor=1) T_subtract_1 = T.match_buffer(T_subtract, [150528], dtype="int16", elem_offset=0, align=128, offset_factor=1) # body diff --git a/tests/python/unittest/test_tvmscript_roundtrip.py b/tests/python/unittest/test_tvmscript_roundtrip.py index 306f60f1b1ba..3b8a6e3e3dd7 100644 --- a/tests/python/unittest/test_tvmscript_roundtrip.py +++ b/tests/python/unittest/test_tvmscript_roundtrip.py @@ -206,29 +206,29 @@ def mmult( A_data: T.Ptr[T.int32] = T.tvm_struct_get(arg0, 0, 1, dtype="handle") T.attr(A_data, "storage_alignment", 128) - A: T.Buffer = T.buffer_decl([1024 * 1024], dtype="int32", data=A_data) + A = T.buffer_decl([1024 * 1024], dtype="int32", data=A_data) buf0_shape_data: T.Ptr[T.int32] = T.tvm_struct_get(arg0, 0, 2, dtype="handle") - buf0_shape: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf0_shape_data) + buf0_shape = T.buffer_decl([2], dtype="int32", data=buf0_shape_data) buf0_strides_data: T.Ptr[T.int32] = T.tvm_struct_get(arg0, 0, 3, dtype="handle") - buf0_strides: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf0_strides_data) + buf0_strides = T.buffer_decl([2], dtype="int32", data=buf0_strides_data) dev_id: T.int32 = T.tvm_struct_get(arg0, 0, 9, dtype="int32") B_data: T.Ptr[T.int32] = T.tvm_struct_get(arg1, 0, 1, dtype="handle") T.attr(B_data, "storage_alignment", 128) - B: T.Buffer = T.buffer_decl([1024 * 1024], dtype="int32", data=B_data) + B = T.buffer_decl([1024 * 1024], dtype="int32", data=B_data) buf1_shape_data: T.Ptr[T.int32] = T.tvm_struct_get(arg1, 0, 2, dtype="handle") - buf1_shape: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf1_shape_data) + buf1_shape = T.buffer_decl([2], dtype="int32", data=buf1_shape_data) buf1_strides_data: T.Ptr[T.int32] = T.tvm_struct_get(arg1, 0, 3, dtype="handle") - buf1_strides: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf1_strides_data) + buf1_strides = T.buffer_decl([2], dtype="int32", data=buf1_strides_data) C_data: T.Ptr[T.int32] = T.tvm_struct_get(arg2, 0, 1, dtype="handle") T.attr(C_data, "storage_alignment", 128) - C: T.Buffer = T.buffer_decl([1024 * 1024], dtype="int32", data=C_data) + C = T.buffer_decl([1024 * 1024], dtype="int32", data=C_data) buf2_shape_data: T.Ptr[T.int32] = T.tvm_struct_get(arg2, 0, 2, dtype="handle") - buf2_shape: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf2_shape_data) + buf2_shape = T.buffer_decl([2], dtype="int32", data=buf2_shape_data) buf2_strides_data: T.Ptr[T.int32] = T.tvm_struct_get(arg2, 0, 3, dtype="handle") - buf2_strides: T.Buffer = T.buffer_decl([2], dtype="int32", data=buf2_strides_data) + buf2_strides = T.buffer_decl([2], dtype="int32", data=buf2_strides_data) assert (((arg0_code == 3) or (arg0_code == 13)) or (arg0_code == 7)) or ( arg0_code == 4 From e37b0fda71ce7edebf067d77eb63c87d2d912cc5 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Fri, 15 Jul 2022 14:48:34 -0700 Subject: [PATCH 2/9] fix unittests --- .../unittest/test_tir_transform_renormalize_split_pattern.py | 4 ++-- tests/python/unittest/test_tvmscript_roundtrip.py | 1 - 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/tests/python/unittest/test_tir_transform_renormalize_split_pattern.py b/tests/python/unittest/test_tir_transform_renormalize_split_pattern.py index 872afeeba5c7..fd08f7e2249a 100644 --- a/tests/python/unittest/test_tir_transform_renormalize_split_pattern.py +++ b/tests/python/unittest/test_tir_transform_renormalize_split_pattern.py @@ -132,7 +132,7 @@ def impossible_equality(n: T.int32): if 2 == 0: # Then this expression evaluates n/2, using the min/max values # of "2", which is caught as a divide by zero error. - if n / 2 >= 16: + if n // 2 >= 16: T.evaluate(0) @@ -141,7 +141,7 @@ def impossible_inequality(n: T.int32): # Prior to bugfix, this conditional set up a range of possible # values for the expression "-2" as [0, kPosInf]. if -1 < -2: - if n / (-2) >= 16: + if n // (-2) >= 16: T.evaluate(0) diff --git a/tests/python/unittest/test_tvmscript_roundtrip.py b/tests/python/unittest/test_tvmscript_roundtrip.py index 3b8a6e3e3dd7..3806f312ae2b 100644 --- a/tests/python/unittest/test_tvmscript_roundtrip.py +++ b/tests/python/unittest/test_tvmscript_roundtrip.py @@ -3036,7 +3036,6 @@ def func_div_mod(): b = T.var("int32") T.evaluate(a // b) T.evaluate(a % b) - T.evaluate(a / b) T.evaluate(T.truncmod(a, b)) return func_div_mod From d3889e0f502105dbdf863576ae6063ce8779d8f9 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Fri, 15 Jul 2022 14:55:39 -0700 Subject: [PATCH 3/9] fix unittests --- .../python/unittest/test_tir_transform_compact_buffer_region.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/unittest/test_tir_transform_compact_buffer_region.py b/tests/python/unittest/test_tir_transform_compact_buffer_region.py index cbc3e7fba3d1..31bb9b8b7cdb 100644 --- a/tests/python/unittest/test_tir_transform_compact_buffer_region.py +++ b/tests/python/unittest/test_tir_transform_compact_buffer_region.py @@ -739,6 +739,7 @@ def func_with_let_binding(): @T.prim_func def func_with_non_index_let_binding(): + A = T.alloc_buffer((64), "float32") x1 = T.call_extern("get", dtype="float16") x2 = T.call_extern("get", dtype="float32") x3 = T.call_extern("get", dtype="float64") @@ -746,7 +747,6 @@ def func_with_non_index_let_binding(): x5 = T.call_extern("get", dtype="int32x16") x6 = T.call_extern("get", dtype="handle") x7 = T.call_extern("get", dtype="") - A = T.alloc_buffer((64), "float32") for rk in range(64): A[rk] = T.call_extern("load_ptr", x1, x2, x3, x4, x5, x6, x7, dtype="float32") From a1b2038ee07bd4dded0c184827d43e64001e78be Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Fri, 15 Jul 2022 17:11:35 -0700 Subject: [PATCH 4/9] fix unittest --- tests/python/unittest/test_tir_schedule_transform_layout.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/unittest/test_tir_schedule_transform_layout.py b/tests/python/unittest/test_tir_schedule_transform_layout.py index 205bd5091268..0332df7fd312 100644 --- a/tests/python/unittest/test_tir_schedule_transform_layout.py +++ b/tests/python/unittest/test_tir_schedule_transform_layout.py @@ -103,7 +103,7 @@ def elementwise(A: T.Buffer[(128, 128), "float32"], B: T.Buffer[(128, 128), "flo def elementwise_transformed(A: T.Buffer[(128, 128), "float32"], B: T.Buffer[(128, 128), "float32"]) -> None: for i in range(16384): with T.block("B"): - vi, = T.axis.remap("S", [i]) + vi = T.axis.remap("S", [i]) B[vi // 128, vi % 128] = A[vi // 128, vi % 128] * 2.0 From 09b6b410bc51e91ca256e888d380e5648739d257 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Fri, 15 Jul 2022 17:14:02 -0700 Subject: [PATCH 5/9] fix unittest --- tests/python/unittest/test_tir_ptx_mma.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/python/unittest/test_tir_ptx_mma.py b/tests/python/unittest/test_tir_ptx_mma.py index 23405fdee98a..8f653c614d42 100644 --- a/tests/python/unittest/test_tir_ptx_mma.py +++ b/tests/python/unittest/test_tir_ptx_mma.py @@ -1311,7 +1311,6 @@ def gemm_mma_m16n8k256_row_col_b1b1s32(a: T.handle, b: T.handle, c: T.handle): Accum.data, 0, False, - "xor", dtype="int32", ) ) From da0594a6b9ef7138453adf5fa9cddcee58c42fa7 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Fri, 15 Jul 2022 22:09:58 -0700 Subject: [PATCH 6/9] fix unittest --- tests/python/unittest/test_tvmscript_syntax_sugar.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/unittest/test_tvmscript_syntax_sugar.py b/tests/python/unittest/test_tvmscript_syntax_sugar.py index aebc606528ba..7248a3a5f47a 100644 --- a/tests/python/unittest/test_tvmscript_syntax_sugar.py +++ b/tests/python/unittest/test_tvmscript_syntax_sugar.py @@ -200,7 +200,7 @@ def test_dynamic_shape_gemm(): @T.prim_func def preflattened_buffer_map(A: T.handle, B: T.handle): A_1 = T.match_buffer(A, [1]) - T.preflattened_buffer(A_1, [1], align=T.int32(1), offset_factor=T.int64(2)) + T.preflattened_buffer(A_1, [1], align=1, offset_factor=2) B_1 = T.match_buffer(B, [1]) T.preflattened_buffer(B_1, [1]) B_1[0] = A_1[0] From 1a2de6a1ca5d018b38efa81b06d936f105876886 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Sat, 16 Jul 2022 03:51:19 -0700 Subject: [PATCH 7/9] Revert "fix unittest" This reverts commit 09b6b410bc51e91ca256e888d380e5648739d257. --- tests/python/unittest/test_tir_ptx_mma.py | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/python/unittest/test_tir_ptx_mma.py b/tests/python/unittest/test_tir_ptx_mma.py index 8f653c614d42..23405fdee98a 100644 --- a/tests/python/unittest/test_tir_ptx_mma.py +++ b/tests/python/unittest/test_tir_ptx_mma.py @@ -1311,6 +1311,7 @@ def gemm_mma_m16n8k256_row_col_b1b1s32(a: T.handle, b: T.handle, c: T.handle): Accum.data, 0, False, + "xor", dtype="int32", ) ) From a24dcfb249912ca458f6b006c3aa2fe73f038b11 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Sun, 24 Jul 2022 16:11:58 -0700 Subject: [PATCH 8/9] fix unittests --- tests/python/unittest/test_tir_analysis_calculate_workspace.py | 2 +- .../unittest/test_tir_analysis_get_block_access_region.py | 2 -- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/tests/python/unittest/test_tir_analysis_calculate_workspace.py b/tests/python/unittest/test_tir_analysis_calculate_workspace.py index a5408ef069e1..8d3163c111c8 100644 --- a/tests/python/unittest/test_tir_analysis_calculate_workspace.py +++ b/tests/python/unittest/test_tir_analysis_calculate_workspace.py @@ -57,11 +57,11 @@ def primfunc_global_allocates(placeholder_144: T.handle, placeholder_145: T.hand def primfunc_local_allocates(placeholder_162: T.handle, placeholder_163: T.handle, placeholder_164: T.handle, T_cast_76: T.handle) -> None: # function attr dict T.func_attr({"global_symbol": "fused_nn_conv2d_add_cast_fixed_point_multiply_clip_cast_cast_9", "tir.noalias": True}) - sid_21 = T.allocate_const([0,1,2,3,4,5,6,7], "int8", [8]) placeholder_165 = T.match_buffer(placeholder_162, [100352], dtype="int16", elem_offset=0, align=128, offset_factor=1) placeholder_166 = T.match_buffer(placeholder_163, [4608], dtype="int16", elem_offset=0, align=128, offset_factor=1) placeholder_167 = T.match_buffer(placeholder_164, [512], dtype="int32", elem_offset=0, align=128, offset_factor=1) T_cast_77 = T.match_buffer(T_cast_76, [100352], dtype="int16", elem_offset=0, align=128, offset_factor=1) + sid_21 = T.allocate_const([0,1,2,3,4,5,6,7], "int8", [8]) # body PaddedInput_25 = T.allocate([131072], "int16", "global") for i1_35, i2_46, i3_47 in T.grid(16, 16, 512): diff --git a/tests/python/unittest/test_tir_analysis_get_block_access_region.py b/tests/python/unittest/test_tir_analysis_get_block_access_region.py index 8a10cbd072f8..21d832848e83 100644 --- a/tests/python/unittest/test_tir_analysis_get_block_access_region.py +++ b/tests/python/unittest/test_tir_analysis_get_block_access_region.py @@ -156,8 +156,6 @@ def gemm() -> None: T.reads(A[vi, vk], B[vj, vk]) T.writes(C[vi, vj]) with T.init(): - T.reads([]) - T.writes(C[vi, vj]) C[vi, vj] = 0 C[vi, vj] += A[vi, vk] * B[vj, vk] From 997e165be9b889483a5d37ed7c4aa09b479f2316 Mon Sep 17 00:00:00 2001 From: Yaxing Cai Date: Mon, 25 Jul 2022 16:34:36 -0700 Subject: [PATCH 9/9] fix --- tests/python/unittest/test_tvmscript_roundtrip.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/python/unittest/test_tvmscript_roundtrip.py b/tests/python/unittest/test_tvmscript_roundtrip.py index 3806f312ae2b..94f29ef50e2a 100644 --- a/tests/python/unittest/test_tvmscript_roundtrip.py +++ b/tests/python/unittest/test_tvmscript_roundtrip.py @@ -3048,8 +3048,7 @@ def test_div_mod(): assert isinstance(func.body[0].value, tvm.tir.FloorDiv) assert isinstance(func.body[1].value, tvm.tir.FloorMod) - assert isinstance(func.body[2].value, tvm.tir.Div) - assert isinstance(func.body[3].value, tvm.tir.Mod) + assert isinstance(func.body[2].value, tvm.tir.Mod) def loop_extent_dependent():