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 2022/03/14 11:05:41 UTC
[GitHub] [tvm] lhutton1 opened a new pull request #10599: [microNPU] Refactor Relay to TIR hook
lhutton1 opened a new pull request #10599:
URL: https://github.com/apache/tvm/pull/10599
Refactors the Relay to TIR python hook for the NPU so that optimizations can be applied across the whole module and not just functions that will be offloaded to the NPU. A pass `OutlineCompilerFunctions` is introduced to outline NPU functions, which now happens before optimization passes are run (this previously happened after the prim_func had been created).
In addition, optimization passes that should only run on NPU functions are now limited to running on outlined functions for the NPU (by checking the "Compiler" attribute). To help avoid code duplication, a helpful decorator `npu_pass` has been created for python passes that should only run on NPU functions.
This refactor helps move a number of passes in the microNPU codegen to use an IRModule -> IRModule philosophy.
cc @manupa-arm @ekalda @NicolaLancellotti @dchauhan-arm
--
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 #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#issuecomment-1072497931
Nice! I ll merge when the CI is green
--
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 #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r825841161
##########
File path: tests/python/contrib/test_ethosu/test_outline_compiler_functions.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.
+
+"""
+Test the outline compiler functions pass.
+"""
+
+import tvm
+from tvm import relay
+from tvm.relay.backend.contrib.ethosu.codegen import OutlineCompilerFunctions
+
+
+def test_outline_compiler_functions():
Review comment:
We need to check whether other "compiler" functions are not affected.
--
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 #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r828318700
##########
File path: python/tvm/relay/backend/contrib/ethosu/util.py
##########
@@ -313,3 +318,37 @@ def __init__(
encoded_constants,
base_addresses,
)
+
+
+def npu_pass(opt_level: int, name: str = ""):
Review comment:
I think that is the definition of the function pass.
See here : https://github.com/apache/tvm/blob/d4739eb80662a71781c46095dc2f3e1903b94f8c/src/relay/ir/transform.cc#L104-L105
So all passes are IRModule --> IRModule transformations (at least we should try to make them be :) )
Function pass is the one that mutates at function level but return the whole IRModule.
--
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 #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r827855234
##########
File path: python/tvm/relay/backend/contrib/ethosu/util.py
##########
@@ -313,3 +318,37 @@ def __init__(
encoded_constants,
base_addresses,
)
+
+
+def npu_pass(opt_level: int, name: str = ""):
Review comment:
I would suggest to use an explicit name such as "create_npu_function_pass", because it has semantics equivalent to a function pass but only works with npu annotated functions.
##########
File path: src/relay/backend/contrib/ethosu/codegen.cc
##########
@@ -176,6 +183,20 @@ tvm::transform::Pass IdentityOptimizer() {
TVM_REGISTER_GLOBAL("relay.ext.ethos-u.IdentityOptimizer").set_body_typed(IdentityOptimizer);
+/*!
+ * \brief This pass will lower NPU functions in a Relay module to scheduled TIR prim functions.
+ */
+tvm::transform::Pass RelayToTIR() {
+ runtime::TypedPackedFunc<IRModule(IRModule, transform::PassContext)> pass_func =
+ [=](IRModule ir_module, transform::PassContext pass_context) {
+ auto relay_to_tir_pf = tvm::runtime::Registry::Get("relay.ext.ethos-u.relay_to_tir");
Review comment:
Following up from the above comment, do we need two passes RelayToTIR (one here and one in Python) ?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -322,43 +354,36 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
-def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir")
+def relay_to_tir(mod: tvm.ir.IRModule) -> tvm.ir.IRModule:
"""
- This is the hook for python-based lowering of relay function
- that gets offloaded to the microNPU.
+ This is the hook for python-based lowering of a Relay module which lowers NPU
+ external functions to TIR.
Parameters
----------
- ext_func : relay.Function
- This is the partitioned relay function
+ mod : tvm.ir.IRModule
+ This is the Relay module.
Returns
-------
- primfunc : tir.PrimFunc
- This returns the scheduled PrimFunc
+ mod : tvm.ir.TRModule
+ The Relay module with scheduled NPU external functions.
"""
- assert len(ext_func.params) == 1
- mod = tvm.IRModule()
- mod["main"] = ext_func
+ mod = OutlineCompilerFunctions("ethos-u")(mod)
mod = LegalizeEthosU()(mod)
mod = LUTsOptimizer()(mod)
mod = IdentityOptimizer()(mod)
mod = LayoutOptimizer()(mod)
mod = relay.transform.InferType()(mod)
- # We are currently using copy_constants scheduler In the long run,
- # this should be a single intelligent and a composite scheduler
- # that can perform scheduling based on user inputs such as
- # scratch memory size.
- tir_mod, const_dict = lower_to_tir(mod["main"], copy_constants())
-
- for param in const_dict.keys():
- const_dict[param] = tvm.nd.array(const_dict[param])
-
- primfunc = tir_mod["main"]
- primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
- primfunc = primfunc.with_attr("ethos-u.constants", const_dict)
- return primfunc
+
+ device_contexts = {
+ gv: "ethos-u" for gv, _ in filter(lambda x: util.is_npu_func(x[1]), mod.functions.items())
+ }
+ mod = mod.with_attr("device_contexts", device_contexts)
+ mod = RelayToTIR()(mod)
Review comment:
This is a bit confusing to have inside the function named as relay_to_tir. Why cant we make everything in the body of this function be inside the pass RelayToTIR() ?
--
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] lhutton1 commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r828313033
##########
File path: python/tvm/relay/backend/contrib/ethosu/util.py
##########
@@ -313,3 +318,37 @@ def __init__(
encoded_constants,
base_addresses,
)
+
+
+def npu_pass(opt_level: int, name: str = ""):
Review comment:
Sounds good, although I'm a bit hesitant to call it a function pass as the pass accepts an IRModule and returns an IRModule. Perhaps we should stick with `create_npu_pass`?
--
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] ekalda commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r829890044
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -322,43 +328,41 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
-def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir")
+def relay_to_tir(mod: tvm.ir.IRModule) -> tvm.ir.IRModule:
"""
- This is the hook for python-based lowering of relay function
- that gets offloaded to the microNPU.
+ This is the hook for python-based lowering of a Relay module which lowers NPU
+ external functions to TIR.
Parameters
----------
- ext_func : relay.Function
- This is the partitioned relay function
+ mod : tvm.ir.IRModule
+ This is the Relay module.
Returns
-------
- primfunc : tir.PrimFunc
- This returns the scheduled PrimFunc
+ mod : tvm.ir.TRModule
Review comment:
nit:
```suggestion
mod : tvm.ir.IRModule
```
--
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] lhutton1 commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r828388268
##########
File path: python/tvm/relay/backend/contrib/ethosu/util.py
##########
@@ -313,3 +318,37 @@ def __init__(
encoded_constants,
base_addresses,
)
+
+
+def npu_pass(opt_level: int, name: str = ""):
Review comment:
In that case I'm okay with it :)
--
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] ekalda commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r829890044
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -322,43 +328,41 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
-def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir")
+def relay_to_tir(mod: tvm.ir.IRModule) -> tvm.ir.IRModule:
"""
- This is the hook for python-based lowering of relay function
- that gets offloaded to the microNPU.
+ This is the hook for python-based lowering of a Relay module which lowers NPU
+ external functions to TIR.
Parameters
----------
- ext_func : relay.Function
- This is the partitioned relay function
+ mod : tvm.ir.IRModule
+ This is the Relay module.
Returns
-------
- primfunc : tir.PrimFunc
- This returns the scheduled PrimFunc
+ mod : tvm.ir.TRModule
Review comment:
nit:
```suggestion
mod : tvm.ir.IRModule
```
--
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] lhutton1 commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r830086770
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -322,43 +328,41 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
-def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir")
+def relay_to_tir(mod: tvm.ir.IRModule) -> tvm.ir.IRModule:
"""
- This is the hook for python-based lowering of relay function
- that gets offloaded to the microNPU.
+ This is the hook for python-based lowering of a Relay module which lowers NPU
+ external functions to TIR.
Parameters
----------
- ext_func : relay.Function
- This is the partitioned relay function
+ mod : tvm.ir.IRModule
+ This is the Relay module.
Returns
-------
- primfunc : tir.PrimFunc
- This returns the scheduled PrimFunc
+ mod : tvm.ir.TRModule
Review comment:
Fixed!
--
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] lhutton1 commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r826088227
##########
File path: tests/python/contrib/test_ethosu/test_outline_compiler_functions.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.
+
+"""
+Test the outline compiler functions pass.
+"""
+
+import tvm
+from tvm import relay
+from tvm.relay.backend.contrib.ethosu.codegen import OutlineCompilerFunctions
+
+
+def test_outline_compiler_functions():
Review comment:
Done, thanks!
--
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] lhutton1 commented on a change in pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#discussion_r828315695
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -322,43 +354,36 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
-def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir")
+def relay_to_tir(mod: tvm.ir.IRModule) -> tvm.ir.IRModule:
"""
- This is the hook for python-based lowering of relay function
- that gets offloaded to the microNPU.
+ This is the hook for python-based lowering of a Relay module which lowers NPU
+ external functions to TIR.
Parameters
----------
- ext_func : relay.Function
- This is the partitioned relay function
+ mod : tvm.ir.IRModule
+ This is the Relay module.
Returns
-------
- primfunc : tir.PrimFunc
- This returns the scheduled PrimFunc
+ mod : tvm.ir.TRModule
+ The Relay module with scheduled NPU external functions.
"""
- assert len(ext_func.params) == 1
- mod = tvm.IRModule()
- mod["main"] = ext_func
+ mod = OutlineCompilerFunctions("ethos-u")(mod)
mod = LegalizeEthosU()(mod)
mod = LUTsOptimizer()(mod)
mod = IdentityOptimizer()(mod)
mod = LayoutOptimizer()(mod)
mod = relay.transform.InferType()(mod)
- # We are currently using copy_constants scheduler In the long run,
- # this should be a single intelligent and a composite scheduler
- # that can perform scheduling based on user inputs such as
- # scratch memory size.
- tir_mod, const_dict = lower_to_tir(mod["main"], copy_constants())
-
- for param in const_dict.keys():
- const_dict[param] = tvm.nd.array(const_dict[param])
-
- primfunc = tir_mod["main"]
- primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
- primfunc = primfunc.with_attr("ethos-u.constants", const_dict)
- return primfunc
+
+ device_contexts = {
+ gv: "ethos-u" for gv, _ in filter(lambda x: util.is_npu_func(x[1]), mod.functions.items())
+ }
+ mod = mod.with_attr("device_contexts", device_contexts)
+ mod = RelayToTIR()(mod)
Review comment:
Agreed, I wasn't too happy with this either. I'll revise this so that more of this functionality is rewritten in C++, although its a bit tricky to define a clear boundary between Python and C++.. see what you think to the update
--
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 #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10599:
URL: https://github.com/apache/tvm/pull/10599#issuecomment-1074183891
Thanks @lhutton1 @NicolaLancellotti @ekalda
--
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 merged pull request #10599: [microNPU] Refactor Relay to TIR hook
Posted by GitBox <gi...@apache.org>.
manupa-arm merged pull request #10599:
URL: https://github.com/apache/tvm/pull/10599
--
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