Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[Fix] Fix some errors in unittests #12170

Merged
merged 9 commits into from
Jul 26, 2022
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions tests/python/unittest/test_arith_domain_touched.py
Original file line number Diff line number Diff line change
Expand Up @@ -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]
Expand All @@ -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):
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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])
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why the order matters?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

T.allocate_const will create a concise scoping here. Then the following T.match_buffer's are within the scope of T.allocate_const, which is not right before first block?
In fact, if you print that script, the T.allocate_const will be printed after all T.match_buffer, as the changes in this pr.

# body
PaddedInput_25 = T.allocate([131072], "int16", "global")
for i1_35, i2_46, i3_47 in T.grid(16, 16, 512):
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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]

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


Expand Down
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
Expand Up @@ -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])
Expand All @@ -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]
Expand Down Expand Up @@ -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():
Expand Down Expand Up @@ -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
Expand All @@ -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
Expand Down Expand Up @@ -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")

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


Expand All @@ -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)


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