Skip to content

Commit

Permalink
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Update EthosU unit tests to expect DeclBuffer nodes
Browse files Browse the repository at this point in the history
Lunderberg committed Dec 19, 2023

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature.
1 parent 270ffb5 commit 9875bd3
Showing 8 changed files with 189 additions and 133 deletions.
15 changes: 12 additions & 3 deletions tests/python/contrib/test_ethosu/cascader/test_integration.py
Original file line number Diff line number Diff line change
@@ -109,7 +109,10 @@ def test_single_conv_compute_cycles_hint():
for single convolution.
"""
primfunc = _compile_model(_create_single_conv2d())
ops = primfunc.body.body.seq
body = primfunc
while not isinstance(body, tvm.tir.SeqStmt):
body = body.body
ops = body.seq
compute_cycles_hints = [2944, 320]
for op, compute_cycle_hint in zip(ops, compute_cycles_hints):
assert op.attr_key == "pragma_compute_cycles_hint"
@@ -122,7 +125,10 @@ def test_double_conv_compute_cycles_hint():
for double convolution.
"""
primfunc = _compile_model(_create_double_conv2d())
ops = primfunc.body.body.body.body.seq
body = primfunc
while not isinstance(body, tvm.tir.SeqStmt):
body = body.body
ops = body.seq
compute_cycles_hints = [2944, 1408, 320, 240]
for op, compute_cycle_hint in zip(ops, compute_cycles_hints):
assert op.attr_key == "pragma_compute_cycles_hint"
@@ -135,7 +141,10 @@ def test_scalar_add_compute_cycles_hint():
for add with scalar values.
"""
primfunc = _compile_model(_create_scalar_add())
ops = primfunc.body.body.seq
body = primfunc
while not isinstance(body, tvm.tir.SeqStmt):
body = body.body
ops = body.seq

compute_cycles_hints = [16, 24]
for op, compute_cycle_hint in zip(ops, compute_cycles_hints):
170 changes: 101 additions & 69 deletions tests/python/contrib/test_ethosu/test_encode_constants.py

Large diffs are not rendered by default.

5 changes: 4 additions & 1 deletion tests/python/contrib/test_ethosu/test_identity_optimizer.py
Original file line number Diff line number Diff line change
@@ -311,7 +311,10 @@ def get_graph():

# Check for hints in the TIR prim func that the identity optimization pass
# has ran. There should not be an identity in the prim func.
assert prim_func.body.value.args[0] == "ethosu_pooling"
body = prim_func
while not isinstance(body, tvm.tir.Evaluate):
body = body.body
assert body.value.args[0] == "ethosu_pooling"


def test_same_output():
2 changes: 1 addition & 1 deletion tests/python/contrib/test_ethosu/test_layout_optimizer.py
Original file line number Diff line number Diff line change
@@ -794,7 +794,7 @@ def get_graph():
prim_func = mod[external_gv_name]

# Check for hints in the TIR prim func that the layout optimization pass has ran
ops = prim_func.body.body.seq
ops = prim_func.body.body.body.body.body.seq
max_pool1, max_pool2 = ops

assert str(max_pool1.value.args[31]) == '"NHCWB16"'
8 changes: 4 additions & 4 deletions tests/python/contrib/test_ethosu/test_remove_concatenates.py
Original file line number Diff line number Diff line change
@@ -35,9 +35,9 @@ def main(input_placeholder: T.Buffer((1,8,12,16), "int8"), input_placeholder_1:
# function attr dict
T.func_attr({"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True})

placeholder = T.Buffer(1536, dtype="int8", data=input_placeholder.data)
placeholder_1 = T.Buffer(1280, dtype="int8", data=input_placeholder_1.data)
T_concat = T.Buffer(4096, dtype="int8", data=input_T_concat.data)
placeholder = T.decl_buffer(1536, dtype="int8", data=input_placeholder.data)
placeholder_1 = T.decl_buffer(1280, dtype="int8", data=input_placeholder_1.data)
T_concat = T.decl_buffer(4096, dtype="int8", data=input_T_concat.data)

buffer = T.Buffer([2992], "uint8")
buffer_1 = T.Buffer([160], "uint8")
@@ -49,7 +49,7 @@ def main(input_placeholder: T.Buffer((1,8,12,16), "int8"), input_placeholder_1:
buffer_7 = T.Buffer([160], "uint8")
# body
T_concat_1_data = T.allocate([2816], "int8", "global", annotations={"disable_lower_builtin":True})
T_concat_1 = T.Buffer([2816], "int8", data=T_concat_1_data)
T_concat_1 = T.decl_buffer([2816], "int8", data=T_concat_1_data)
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 8, 10, 16, 8, 0, 10, placeholder_1[0], 0, 0, 0, T.float32(0.5), 10, "NHWC", 160, 16, 1, "int8", 8, 10, 16, 8, 0, 10, T_concat_1[192], 0, 0, 0, T.float32(0.25), 14, "NHWC", 352, 16, 1, 3, 3, 1, 1, 1, 1, buffer[0], 2992, T.int8(-1), T.int8(-1), 12, buffer_1[0], 160, T.int8(-1), T.int8(-1), 1, 1, 1, 1, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 8, 10, 16, 8, 0, 10, T_concat_1[192], 0, 0, 0, T.float32(0.5), 10, "NHWC", 352, 16, 1, "int8", 8, 10, 16, 8, 0, 10, T_concat[352], 0, 0, 0, T.float32(0.25), 14, "NHWC", 512, 16, 1, 3, 3, 1, 1, 1, 1, buffer_2[0], 2992, T.int8(-1), T.int8(-1), 12, buffer_3[0], 160, T.int8(-1), T.int8(-1), 1, 1, 1, 1, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 8, 12, 16, 8, 0, 12, placeholder[0], 0, 0, 0, T.float32(0.5), 10, "NHWC", 192, 16, 1, "int8", 8, 12, 16, 8, 0, 12, T_concat_1[0], 0, 0, 0, T.float32(0.25), 14, "NHWC", 352, 16, 1, 3, 3, 1, 1, 1, 1, buffer_4[0], 2992, T.int8(-1), T.int8(-1), 12, buffer_5[0], 160, T.int8(-1), T.int8(-1), 1, 1, 1, 1, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))
70 changes: 37 additions & 33 deletions tests/python/contrib/test_ethosu/test_replace_conv2d.py

Large diffs are not rendered by default.

32 changes: 18 additions & 14 deletions tests/python/contrib/test_ethosu/test_replace_copy.py
Original file line number Diff line number Diff line change
@@ -35,14 +35,16 @@
class ReferenceModule:
@T.prim_func
def main(input_placeholder_3: T.Buffer((1, 16, 16, 32), "int8"), input_ethosu_write_1: T.Buffer((1, 16, 16, 8), "int8")) -> None:
# function attr dict
T.func_attr({"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True})
buffer_1 = T.Buffer([384], "uint8")
placeholder_3 = T.Buffer([8192], dtype="int8", data=input_placeholder_3.data)
ethosu_write_1 = T.Buffer([2048], dtype="int8", data=input_ethosu_write_1.data)
# body

placeholder_3 = T.decl_buffer([8192], dtype="int8", data=input_placeholder_3.data)
ethosu_write_1 = T.decl_buffer([2048], dtype="int8", data=input_ethosu_write_1.data)

placeholder_global_data = T.allocate([384], "uint8", "global", annotations={"disable_lower_builtin": True})
placeholder_global = T.Buffer([384], "uint8", data=placeholder_global_data)
placeholder_global = T.decl_buffer([384], "uint8", data=placeholder_global_data)

buffer_1 = T.Buffer([384], "uint8")

T.evaluate(T.call_extern("ethosu_copy", buffer_1[0], 384, placeholder_global[0], dtype="handle"))
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 16, 16, 32, 16, 0, 16, placeholder_3[0], 0, 0, 0, T.float32(0.5), 10, "NHWC", 512, 32, 1, "int8", 16, 16, 8, 16, 0, 16, ethosu_write_1[0], 0, 0, 0, T.float32(0.25), 14, "NHWC", 128, 8, 1, 1, 1, 1, 1, 1, 1, placeholder_global[0], 304, T.int8(-1), T.int8(-1), 12, placeholder_global[304], 80, T.int8(-1), T.int8(-1), 0, 0, 0, 0, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))
__tvm_meta__ = None
@@ -79,17 +81,19 @@ def _get_func():
class WeightStream:
@T.prim_func
def main(input_placeholder_5: T.Buffer((1, 16, 16, 32), "int8"), input_ethosu_write_1: T.Buffer((1, 16, 16, 16), "int8")) -> None:
# function attr dict
T.func_attr({"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True})
buffer = T.Buffer([528], "uint8")
buffer_2 = T.Buffer([336], "uint8")
placeholder_5 = T.Buffer([8192], dtype="int8", data=input_placeholder_5.data)
ethosu_write_1 = T.Buffer([4096], dtype="int8", data=input_ethosu_write_1.data)
# body

placeholder_5 = T.decl_buffer([8192], dtype="int8", data=input_placeholder_5.data)
ethosu_write_1 = T.decl_buffer([4096], dtype="int8", data=input_ethosu_write_1.data)

placeholder_d_global_data = T.allocate([528], "uint8", "global", annotations={"disable_lower_builtin": True})
placeholder_d_global = T.Buffer([528], "uint8", data=placeholder_d_global_data)
placeholder_d_global = T.decl_buffer([528], "uint8", data=placeholder_d_global_data)
placeholder_d_global_1_data = T.allocate([336], "uint8", "global", annotations={"disable_lower_builtin": True})
placeholder_d_global_1 = T.Buffer([336], "uint8", data=placeholder_d_global_1_data)
placeholder_d_global_1 = T.decl_buffer([336], "uint8", data=placeholder_d_global_1_data)

buffer = T.Buffer([528], "uint8")
buffer_2 = T.Buffer([336], "uint8")

T.evaluate(T.call_extern("ethosu_copy", buffer[0], 528, placeholder_d_global[0], dtype="handle"))
T.evaluate(T.call_extern("ethosu_copy", buffer_2[0], 336, placeholder_d_global_1[0], dtype="handle"))
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 16, 16, 32, 16, 0, 16, placeholder_5[0], 0, 0, 0, T.float32(0.5), 10, "NHWC", 512, 32, 1, "int8", 16, 16, 10, 16, 0, 16, ethosu_write_1[0], 0, 0, 0, T.float32(0.25), 14, "NHWC", 256, 16, 1, 1, 1, 1, 1, 1, 1, placeholder_d_global[0], 416, T.int8(-1), T.int8(-1), 12, placeholder_d_global[416], 112, T.int8(-1), T.int8(-1), 0, 0, 0, 0, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))
20 changes: 12 additions & 8 deletions tests/python/contrib/test_ethosu/test_scheduler.py
Original file line number Diff line number Diff line change
@@ -198,18 +198,22 @@ class DiamondGraphTir:
@T.prim_func
def main(input_placeholder: T.Buffer((1, 56, 56, 96), "int8"), input_ethosu_write: T.Buffer((1, 56, 56, 24), "int8")) -> None:
T.func_attr({"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True})
placeholder = T.Buffer([301056], dtype='int8', data=input_placeholder.data)
ethosu_write = T.Buffer([75264], dtype='int8', data=input_ethosu_write.data)
buffer1 = T.Buffer([2848], "uint8")
buffer3 = T.Buffer([976], "uint8")

placeholder = T.decl_buffer([301056], dtype='int8', data=input_placeholder.data)
ethosu_write = T.decl_buffer([75264], dtype='int8', data=input_ethosu_write.data)

p1_data = T.allocate([2848], "uint8", "global", annotations={"disable_lower_builtin":True})
p1 = T.Buffer([2848], "uint8", data=p1_data)
p1 = T.decl_buffer([2848], "uint8", data=p1_data)
p2_data = T.allocate([976], "uint8", "global", annotations={"disable_lower_builtin":True})
p2 = T.Buffer([976], "uint8", data=p2_data)
p2 = T.decl_buffer([976], "uint8", data=p2_data)
p5_data = T.allocate([75264], "int8", "global", annotations={"disable_lower_builtin":True})
p5 = T.Buffer([75264], "int8", data=p5_data)
p5 = T.decl_buffer([75264], "int8", data=p5_data)
p6_data = T.allocate([75264], "int8", "global", annotations={"disable_lower_builtin":True})
p6 = T.Buffer([75264], "int8", data=p6_data)
p6 = T.decl_buffer([75264], "int8", data=p6_data)

buffer1 = T.Buffer([2848], "uint8")
buffer3 = T.Buffer([976], "uint8")

T.evaluate(T.call_extern("ethosu_copy", buffer1[0], 2848, p1[0], dtype="handle"))
T.evaluate(T.call_extern("ethosu_copy", buffer3[0], 976, p2[0], dtype="handle"))
T.evaluate(T.call_extern("ethosu_conv2d", "int8", 56, 56, 96, 56, 0, 56, placeholder[0], 0, 0, 0, T.float32(0.5), 10, "NHWC", 5376, 96, 1, "int8", 56, 56, 24, 56, 0, 56, p5[0], 0, 0, 0, T.float32(0.25), 14, "NHWC", 1344, 24, 1, 1, 1, 1, 1, 1, 1, p1[0], 2608, T.int8(-1), T.int8(-1), 12, p1[2608], 240, T.int8(-1), T.int8(-1), 0, 0, 0, 0, "NONE", 0, 0, "TFL", "NONE", 0, 0, 0, dtype="handle"))

0 comments on commit 9875bd3

Please sign in to comment.