You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2021/07/14 16:32:32 UTC
[GitHub] [tvm] manupa-arm opened a new pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
manupa-arm opened a new pull request #8469:
URL: https://github.com/apache/tvm/pull/8469
This is an optional pass to convert all parallel for loops in TIR to serial ones for different reasons such as executor does not support parallel launch of for loops (e.g., AoT) or allocating space for parallel for loops might not be desired.
* Additionally adding FFI scaffolding for USMP
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] areusch merged pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
areusch merged pull request #8469:
URL: https://github.com/apache/tvm/pull/8469
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#issuecomment-880132196
given the particular pass can be applied to any TIR Module, we might be able to just put it under tir namespace withour introducing the usmp namespace.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on a change in pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#discussion_r724034826
##########
File path: tests/python/unittest/test_tir_transform_convert_for_loops_serial.py
##########
@@ -0,0 +1,62 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+import tvm
+from tvm import tir, script
+from tvm.script import ty
+from tvm.tir import stmt_functor
+
+# fmt: off
+@tvm.script.tir
+def fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2(placeholder_30: ty.handle, placeholder_31: ty.handle, placeholder_32: ty.handle, T_cast_8: ty.handle) -> None:
+ # function attr dict
+ tir.func_attr({"global_symbol": "fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2", "tir.noalias": True})
+ placeholder_33 = tir.match_buffer(placeholder_30, [1, 28, 28, 192], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_34 = tir.match_buffer(placeholder_31, [1, 1, 192, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_35 = tir.match_buffer(placeholder_32, [1, 1, 1, 16], dtype="int32", elem_offset=0, align=128, offset_factor=1)
+ T_cast_9 = tir.match_buffer(T_cast_8, [1, 28, 28, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ # body
+ PaddedInput_3 = tir.allocate([1, 28, 28, 192], "int16", "global")
+ for i0_i1_fused_3 in tir.parallel(0, 28):
+ for i2_3, i3_3 in tir.grid(28, 192):
+ tir.store(PaddedInput_3, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3), tir.load("int16", placeholder_33.data, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3)), True)
+ for ax0_ax1_fused_ax2_fused_3 in tir.parallel(0, 784):
+ for ax3_2 in tir.serial(0, 16):
+ Conv2dOutput_3 = tir.allocate([1, 1, 1, 1], "int32", "global")
+ tir.store(Conv2dOutput_3, 0, 0, True)
+ for rc_3 in tir.serial(0, 192):
+ tir.store(Conv2dOutput_3, 0, (tir.load("int32", Conv2dOutput_3, 0) + (tir.cast(tir.load("int16", PaddedInput_3, ((ax0_ax1_fused_ax2_fused_3*192) + rc_3)), "int32")*tir.cast(tir.load("int16", placeholder_34.data, ((rc_3*16) + ax3_2)), "int32"))), True)
+ tir.store(T_cast_9.data, ((ax0_ax1_fused_ax2_fused_3*16) + ax3_2), tir.cast(tir.cast(tir.max(tir.min(tir.q_multiply_shift((tir.load("int32", Conv2dOutput_3, 0) + tir.load("int32", placeholder_35.data, ax3_2)), 1764006585, 31, -7, dtype="int32"), 255), 0), "uint8"), "int16"), True)
+# fmt: on
+
+
+def test_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2():
+ primfunc = fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2
+ mod = tvm.IRModule.from_expr(primfunc)
Review comment:
Do we need to ? I mean its written in the test.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] areusch commented on a change in pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
areusch commented on a change in pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#discussion_r726672928
##########
File path: tests/python/unittest/test_tir_transform_convert_for_loops_serial.py
##########
@@ -0,0 +1,62 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+import tvm
+from tvm import tir, script
+from tvm.script import ty
+from tvm.tir import stmt_functor
+
+# fmt: off
+@tvm.script.tir
+def fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2(placeholder_30: ty.handle, placeholder_31: ty.handle, placeholder_32: ty.handle, T_cast_8: ty.handle) -> None:
+ # function attr dict
+ tir.func_attr({"global_symbol": "fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2", "tir.noalias": True})
+ placeholder_33 = tir.match_buffer(placeholder_30, [1, 28, 28, 192], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_34 = tir.match_buffer(placeholder_31, [1, 1, 192, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_35 = tir.match_buffer(placeholder_32, [1, 1, 1, 16], dtype="int32", elem_offset=0, align=128, offset_factor=1)
+ T_cast_9 = tir.match_buffer(T_cast_8, [1, 28, 28, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ # body
+ PaddedInput_3 = tir.allocate([1, 28, 28, 192], "int16", "global")
+ for i0_i1_fused_3 in tir.parallel(0, 28):
+ for i2_3, i3_3 in tir.grid(28, 192):
+ tir.store(PaddedInput_3, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3), tir.load("int16", placeholder_33.data, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3)), True)
+ for ax0_ax1_fused_ax2_fused_3 in tir.parallel(0, 784):
+ for ax3_2 in tir.serial(0, 16):
+ Conv2dOutput_3 = tir.allocate([1, 1, 1, 1], "int32", "global")
+ tir.store(Conv2dOutput_3, 0, 0, True)
+ for rc_3 in tir.serial(0, 192):
+ tir.store(Conv2dOutput_3, 0, (tir.load("int32", Conv2dOutput_3, 0) + (tir.cast(tir.load("int16", PaddedInput_3, ((ax0_ax1_fused_ax2_fused_3*192) + rc_3)), "int32")*tir.cast(tir.load("int16", placeholder_34.data, ((rc_3*16) + ax3_2)), "int32"))), True)
+ tir.store(T_cast_9.data, ((ax0_ax1_fused_ax2_fused_3*16) + ax3_2), tir.cast(tir.cast(tir.max(tir.min(tir.q_multiply_shift((tir.load("int32", Conv2dOutput_3, 0) + tir.load("int32", placeholder_35.data, ax3_2)), 1764006585, 31, -7, dtype="int32"), 255), 0), "uint8"), "int16"), True)
+# fmt: on
+
+
+def test_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2():
+ primfunc = fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2
+ mod = tvm.IRModule.from_expr(primfunc)
Review comment:
let's consider this blocked on testing infrastructure. a common pattern in tests is for the data to be re-used in multiple tests and then lose the "why" behind the test. that's where my request is coming from.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#issuecomment-882623193
Hi @tqchen ,
Yes make sense, will do.
Hi @junrushao1994 ,
I think TOPI and other places define that as a property of the op implementation/scheduling. However, this is used as post-scheduling transformation to force the parallel loops to be serial to reduce memory and/or explicitly depict what the intended executor / backend supports.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#issuecomment-882623193
Hi @tqchen ,
Yes make sense, will do.
Hi @junrushao1994 ,
I think TOPI and other places define that as a property of the op implementation/scheduling. However, this is used as post-scheduling transformation to force the parallel loops to be serial to reduce memory and/or explicitly depict what the intended executor / backend supports.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#issuecomment-882623193
Hi @tqchen ,
Yes make sense, will do.
Hi @junrushao1994 ,
I think TOPI and other places define that as a property of the op implementation/scheduling. However, this is used as post-scheduling transformation to force the parallel loops to be serial to reduce memory and/or explicitly depict what the intended executor / backend supports.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] areusch commented on a change in pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
areusch commented on a change in pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#discussion_r707535161
##########
File path: tests/python/unittest/test_tir_transform_convert_for_loops_serial.py
##########
@@ -0,0 +1,62 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+import tvm
+from tvm import tir, script
+from tvm.script import ty
+from tvm.tir import stmt_functor
+
+# fmt: off
+@tvm.script.tir
+def fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2(placeholder_30: ty.handle, placeholder_31: ty.handle, placeholder_32: ty.handle, T_cast_8: ty.handle) -> None:
+ # function attr dict
+ tir.func_attr({"global_symbol": "fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2", "tir.noalias": True})
+ placeholder_33 = tir.match_buffer(placeholder_30, [1, 28, 28, 192], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_34 = tir.match_buffer(placeholder_31, [1, 1, 192, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_35 = tir.match_buffer(placeholder_32, [1, 1, 1, 16], dtype="int32", elem_offset=0, align=128, offset_factor=1)
+ T_cast_9 = tir.match_buffer(T_cast_8, [1, 28, 28, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ # body
+ PaddedInput_3 = tir.allocate([1, 28, 28, 192], "int16", "global")
+ for i0_i1_fused_3 in tir.parallel(0, 28):
+ for i2_3, i3_3 in tir.grid(28, 192):
+ tir.store(PaddedInput_3, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3), tir.load("int16", placeholder_33.data, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3)), True)
+ for ax0_ax1_fused_ax2_fused_3 in tir.parallel(0, 784):
+ for ax3_2 in tir.serial(0, 16):
+ Conv2dOutput_3 = tir.allocate([1, 1, 1, 1], "int32", "global")
+ tir.store(Conv2dOutput_3, 0, 0, True)
+ for rc_3 in tir.serial(0, 192):
+ tir.store(Conv2dOutput_3, 0, (tir.load("int32", Conv2dOutput_3, 0) + (tir.cast(tir.load("int16", PaddedInput_3, ((ax0_ax1_fused_ax2_fused_3*192) + rc_3)), "int32")*tir.cast(tir.load("int16", placeholder_34.data, ((rc_3*16) + ax3_2)), "int32"))), True)
+ tir.store(T_cast_9.data, ((ax0_ax1_fused_ax2_fused_3*16) + ax3_2), tir.cast(tir.cast(tir.max(tir.min(tir.q_multiply_shift((tir.load("int32", Conv2dOutput_3, 0) + tir.load("int32", placeholder_35.data, ax3_2)), 1764006585, 31, -7, dtype="int32"), 255), 0), "uint8"), "int16"), True)
+# fmt: on
+
+
+def test_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2():
+ primfunc = fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2
+ mod = tvm.IRModule.from_expr(primfunc)
Review comment:
want to assert that you find at least one kParallel for loop in here?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#issuecomment-881766585
Thanks for the contribution! Just curious, why we do such a pass instead of removing `sch.parallel` in scheduling?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on a change in pull request #8469: [TIR][USMP] Add a parallel to serial for loop converter pass
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #8469:
URL: https://github.com/apache/tvm/pull/8469#discussion_r724034826
##########
File path: tests/python/unittest/test_tir_transform_convert_for_loops_serial.py
##########
@@ -0,0 +1,62 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+import tvm
+from tvm import tir, script
+from tvm.script import ty
+from tvm.tir import stmt_functor
+
+# fmt: off
+@tvm.script.tir
+def fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2(placeholder_30: ty.handle, placeholder_31: ty.handle, placeholder_32: ty.handle, T_cast_8: ty.handle) -> None:
+ # function attr dict
+ tir.func_attr({"global_symbol": "fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2", "tir.noalias": True})
+ placeholder_33 = tir.match_buffer(placeholder_30, [1, 28, 28, 192], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_34 = tir.match_buffer(placeholder_31, [1, 1, 192, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ placeholder_35 = tir.match_buffer(placeholder_32, [1, 1, 1, 16], dtype="int32", elem_offset=0, align=128, offset_factor=1)
+ T_cast_9 = tir.match_buffer(T_cast_8, [1, 28, 28, 16], dtype="int16", elem_offset=0, align=128, offset_factor=1)
+ # body
+ PaddedInput_3 = tir.allocate([1, 28, 28, 192], "int16", "global")
+ for i0_i1_fused_3 in tir.parallel(0, 28):
+ for i2_3, i3_3 in tir.grid(28, 192):
+ tir.store(PaddedInput_3, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3), tir.load("int16", placeholder_33.data, (((i0_i1_fused_3*5376) + (i2_3*192)) + i3_3)), True)
+ for ax0_ax1_fused_ax2_fused_3 in tir.parallel(0, 784):
+ for ax3_2 in tir.serial(0, 16):
+ Conv2dOutput_3 = tir.allocate([1, 1, 1, 1], "int32", "global")
+ tir.store(Conv2dOutput_3, 0, 0, True)
+ for rc_3 in tir.serial(0, 192):
+ tir.store(Conv2dOutput_3, 0, (tir.load("int32", Conv2dOutput_3, 0) + (tir.cast(tir.load("int16", PaddedInput_3, ((ax0_ax1_fused_ax2_fused_3*192) + rc_3)), "int32")*tir.cast(tir.load("int16", placeholder_34.data, ((rc_3*16) + ax3_2)), "int32"))), True)
+ tir.store(T_cast_9.data, ((ax0_ax1_fused_ax2_fused_3*16) + ax3_2), tir.cast(tir.cast(tir.max(tir.min(tir.q_multiply_shift((tir.load("int32", Conv2dOutput_3, 0) + tir.load("int32", placeholder_35.data, ax3_2)), 1764006585, 31, -7, dtype="int32"), 255), 0), "uint8"), "int16"), True)
+# fmt: on
+
+
+def test_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2():
+ primfunc = fused_nn_conv2d_add_fixed_point_multiply_clip_cast_cast_2
+ mod = tvm.IRModule.from_expr(primfunc)
Review comment:
Do we need to ? I mean its written in the test.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org