Skip to content

Commit

Permalink
[Fix] Fix some errors in unittests (#12170)
Browse files Browse the repository at this point in the history
* unittests fix 0

* fix unittests

* fix unittests

* fix unittest

* fix unittest

* fix unittest

* Revert "fix unittest"

This reverts commit 09b6b41.

* fix unittests

* fix
cyx-6 authored Jul 26, 2022

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature. The key has expired.
1 parent ea6ea42 commit eada707
Showing 11 changed files with 33 additions and 37 deletions.
8 changes: 4 additions & 4 deletions tests/python/unittest/test_arith_domain_touched.py
Original file line number Diff line number Diff line change
@@ -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):
Original file line number Diff line number Diff line change
@@ -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):
Original file line number Diff line number Diff line change
@@ -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]

Original file line number Diff line number Diff line change
@@ -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


22 changes: 11 additions & 11 deletions tests/python/unittest/test_tir_transform_compact_buffer_region.py
Original file line number Diff line number Diff line change
@@ -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
@@ -739,14 +739,14 @@ 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")
x4 = T.call_extern("get", dtype="uint8")
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")

Original file line number Diff line number Diff line change
@@ -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)


Original file line number Diff line number Diff line change
@@ -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
Original file line number Diff line number Diff line change
@@ -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
2 changes: 1 addition & 1 deletion tests/python/unittest/test_tir_usmp_utils.py
Original file line number Diff line number Diff line change
@@ -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
22 changes: 10 additions & 12 deletions tests/python/unittest/test_tvmscript_roundtrip.py
Original file line number Diff line number Diff line change
@@ -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
@@ -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
@@ -3049,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():
2 changes: 1 addition & 1 deletion tests/python/unittest/test_tvmscript_syntax_sugar.py
Original file line number Diff line number Diff line change
@@ -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]

0 comments on commit eada707

Please sign in to comment.