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/11/26 20:56:12 UTC
[GitHub] [tvm] manupa-arm opened a new pull request #9597: [microNPU] Move the compilation to use Target Hooks.
manupa-arm opened a new pull request #9597:
URL: https://github.com/apache/tvm/pull/9597
This commits moves the current compilation flow
to use target hooks, so that the generated TIR
is provided to unified module to for unified
optimizations.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759925705
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
Review comment:
Now pushed the follow up here as this PR got conflicted.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-983976948
I've created another follow up PR to modify the CMake to keep utils.cc that carries a Object definition that needs to be compiled always.
https://github.com/apache/tvm/pull/9630
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759905816
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
Review comment:
We generally want to stick the case of the original function definition, therefore this one is in python.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-983771282
Thanks! @leandron and Im happy to take further comments on this PR which I will include in #9605
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759906982
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
+
Parameters
----------
- ext_func : tvm.relay.function.Function
- The partitioned relay function
+ ext_func : relay.Function
+ This is the partitioned relay function
+
Returns
-------
- cs : str
- An hex string of the bytes of command stream
- encoded_constants : str
- An hex string of the bytes that includes concat'd
- encoded weights, encoded biases and scales.
- scratch_size : int
- The size of the scratch buffer needed.
+ primfunc : tir.PrimFunc
+ This returns the scheduled PrimFunc
"""
+ assert len(ext_func.params) == 1
Review comment:
The reasoning as said before.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-981420754
@grant-arm @Mousius
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759000377
##########
File path: src/relay/backend/contrib/ethosu/source_module.cc
##########
@@ -41,34 +41,37 @@
#include <vector>
#include "../../../../runtime/file_utils.h"
+#include "utils.h"
namespace tvm {
namespace runtime {
+using CompilationArtifact = relay::contrib::ethosu::CompilationArtifact;
+
// The runtime.Module that contains the host-side c code
// required for invoking the NPU with the command stream
class EthosUModuleNode : public ModuleNode {
public:
/*!
- * \brief The ethos runtime module.
+ * \brief The microNPU runtime module.
*
- * \param func_name_ name of the should be codegen'd function
- * \param cmms_hex_ command stream for the NPU in hex
- * \param weights_bias_hex_ the encoded biases and weights for the NPU in hex
- * \param scratch_size_ the size of the scratch memory required for command stream
- * \param input_size_ the size (in bytes) for the input tensor
- * \param output_size_ the size (in bytes) for the output tensor
+ * \param compilation_artifacts
+ * This is an array of CompilationArtifacts that is produced via
+ * lowering each PrimFunc to command stream. Here, those artifacts
+ * will be used to create the c-source.
*/
- explicit EthosUModuleNode(const String& func_name_, const String& cmms_hex_,
- const String& weights_bias_hex_, const Integer& scratch_size_,
- const Integer& input_size_, const Integer& output_size_) {
- func_name = func_name_;
- cmms_hex = std::move(cmms_hex_);
- weights_bias_hex = std::move(weights_bias_hex_);
- scratch_size = scratch_size_->value;
- input_size = input_size_->value;
- output_size = output_size_->value;
- c_source = GenerateSource();
+ explicit EthosUModuleNode(Array<CompilationArtifact> compilation_artifacts)
+ : compilation_artifacts_(compilation_artifacts) {
+ c_source += "#include <stdio.h>\n";
+ c_source += "#include <stdlib.h>\n";
+ c_source += "#include <tvm/runtime/crt/module.h>\n";
+ c_source += "#include <tvm_ethosu_runtime.h>\n";
+ c_source += "\n";
Review comment:
Done in #9605
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759925787
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
Review comment:
Done
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759925444
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
Review comment:
Changed params to const_dict :)
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759908850
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
+
+ Returns
+ -------
+ CompilationArtifact
+ This is a structure that holds the binary artifacts
+ for the microNPU
+ """
+ symbol = str(primfunc.attrs["global_symbol"])
+ params = primfunc.attrs["ethos-u.constants"]
+ input_size = primfunc.attrs["ethos-u.input_size"]
+ output_size = primfunc.attrs["ethos-u.output_size"]
+ tir_mod = tvm.IRModule()
+ tir_mod[symbol] = primfunc
+
+ params_with_int_keys = dict()
+ for idx in params.keys():
+ params_with_int_keys[int(idx)] = params[idx].numpy()
Review comment:
No it is not possible -- this is internally generated to always be an int.
--
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] grant-arm commented on a change in pull request #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
grant-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r758278217
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
Review comment:
```suggestion
TIR PrimFunc that has undergone unified optimization
```
##########
File path: src/relay/backend/contrib/ethosu/source_module.cc
##########
@@ -79,7 +82,6 @@ class EthosUModuleNode : public ModuleNode {
*/
void SaveToFile(const std::string& file_name, const std::string& format) final {
std::string fmt = GetFileFormat(file_name, format);
- LOG(INFO) << "format=" << fmt << ";;\n";
Review comment:
Thanks Manupa, I've been meaning to remove this for a while.
##########
File path: src/relay/backend/contrib/ethosu/source_module.cc
##########
@@ -41,34 +41,37 @@
#include <vector>
#include "../../../../runtime/file_utils.h"
+#include "utils.h"
namespace tvm {
namespace runtime {
+using CompilationArtifact = relay::contrib::ethosu::CompilationArtifact;
+
// The runtime.Module that contains the host-side c code
// required for invoking the NPU with the command stream
class EthosUModuleNode : public ModuleNode {
public:
/*!
- * \brief The ethos runtime module.
+ * \brief The microNPU runtime module.
*
- * \param func_name_ name of the should be codegen'd function
- * \param cmms_hex_ command stream for the NPU in hex
- * \param weights_bias_hex_ the encoded biases and weights for the NPU in hex
- * \param scratch_size_ the size of the scratch memory required for command stream
- * \param input_size_ the size (in bytes) for the input tensor
- * \param output_size_ the size (in bytes) for the output tensor
+ * \param compilation_artifacts
+ * This is an array of CompilationArtifacts that is produced via
+ * lowering each PrimFunc to command stream. Here, those artifacts
+ * will be used to create the c-source.
*/
- explicit EthosUModuleNode(const String& func_name_, const String& cmms_hex_,
- const String& weights_bias_hex_, const Integer& scratch_size_,
- const Integer& input_size_, const Integer& output_size_) {
- func_name = func_name_;
- cmms_hex = std::move(cmms_hex_);
- weights_bias_hex = std::move(weights_bias_hex_);
- scratch_size = scratch_size_->value;
- input_size = input_size_->value;
- output_size = output_size_->value;
- c_source = GenerateSource();
+ explicit EthosUModuleNode(Array<CompilationArtifact> compilation_artifacts)
+ : compilation_artifacts_(compilation_artifacts) {
+ c_source += "#include <stdio.h>\n";
+ c_source += "#include <stdlib.h>\n";
+ c_source += "#include <tvm/runtime/crt/module.h>\n";
+ c_source += "#include <tvm_ethosu_runtime.h>\n";
+ c_source += "\n";
Review comment:
```suggestion
c_source += "#include <tvm_ethosu_runtime.h>\n\n";
```
##########
File path: src/relay/backend/contrib/ethosu/source_module.cc
##########
@@ -41,34 +41,37 @@
#include <vector>
#include "../../../../runtime/file_utils.h"
+#include "utils.h"
namespace tvm {
namespace runtime {
+using CompilationArtifact = relay::contrib::ethosu::CompilationArtifact;
+
// The runtime.Module that contains the host-side c code
// required for invoking the NPU with the command stream
class EthosUModuleNode : public ModuleNode {
public:
/*!
- * \brief The ethos runtime module.
+ * \brief The microNPU runtime module.
*
- * \param func_name_ name of the should be codegen'd function
- * \param cmms_hex_ command stream for the NPU in hex
- * \param weights_bias_hex_ the encoded biases and weights for the NPU in hex
- * \param scratch_size_ the size of the scratch memory required for command stream
- * \param input_size_ the size (in bytes) for the input tensor
- * \param output_size_ the size (in bytes) for the output tensor
+ * \param compilation_artifacts
+ * This is an array of CompilationArtifacts that is produced via
+ * lowering each PrimFunc to command stream. Here, those artifacts
+ * will be used to create the c-source.
*/
- explicit EthosUModuleNode(const String& func_name_, const String& cmms_hex_,
- const String& weights_bias_hex_, const Integer& scratch_size_,
- const Integer& input_size_, const Integer& output_size_) {
- func_name = func_name_;
- cmms_hex = std::move(cmms_hex_);
- weights_bias_hex = std::move(weights_bias_hex_);
- scratch_size = scratch_size_->value;
- input_size = input_size_->value;
- output_size = output_size_->value;
- c_source = GenerateSource();
+ explicit EthosUModuleNode(Array<CompilationArtifact> compilation_artifacts)
+ : compilation_artifacts_(compilation_artifacts) {
+ c_source += "#include <stdio.h>\n";
+ c_source += "#include <stdlib.h>\n";
+ c_source += "#include <tvm/runtime/crt/module.h>\n";
+ c_source += "#include <tvm_ethosu_runtime.h>\n";
+ c_source += "\n";
+ for (const CompilationArtifact& compilation_artifact : compilation_artifacts) {
+ c_source += GenerateSource(compilation_artifact);
+ c_source += "\n";
+ c_source += "\n";
Review comment:
```suggestion
c_source += "\n\n";
```
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-981724791
Thanks @ekalda @grant-arm for speedy reviews :).
I've created the follow up to anwser the comments : https://github.com/apache/tvm/pull/9605
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759908850
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
+
+ Returns
+ -------
+ CompilationArtifact
+ This is a structure that holds the binary artifacts
+ for the microNPU
+ """
+ symbol = str(primfunc.attrs["global_symbol"])
+ params = primfunc.attrs["ethos-u.constants"]
+ input_size = primfunc.attrs["ethos-u.input_size"]
+ output_size = primfunc.attrs["ethos-u.output_size"]
+ tir_mod = tvm.IRModule()
+ tir_mod[symbol] = primfunc
+
+ params_with_int_keys = dict()
+ for idx in params.keys():
+ params_with_int_keys[int(idx)] = params[idx].numpy()
Review comment:
No it is not possible -- this internally generated to always be an int.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759906329
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
Review comment:
I dont think "a hook" is correct here. This is the only hook to do this.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759907446
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
Review comment:
Same reasoning as before.
--
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] leandron merged pull request #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
leandron merged pull request #9597:
URL: https://github.com/apache/tvm/pull/9597
--
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] huajsj commented on a change in pull request #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
huajsj commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759767264
##########
File path: include/tvm/tir/transform.h
##########
@@ -287,6 +287,11 @@ TVM_DLL Pass LowerThreadAllreduce();
*/
TVM_DLL Pass InferFragment();
+/*!
+ * \brief This annotation for nodes to be disabled for builtin lowering
Review comment:
"* \brief This annotation is for nodes where you want to disable the built-in lowering."
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
Review comment:
This name is dazzling, how about use "camel-case" like what you did for relay.ext.ethos-u.CompilationArtifact?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
Review comment:
The TIR PrimFunc which has undergone the unified optimization.
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
+
+ Returns
+ -------
+ CompilationArtifact
+ This is a structure that holds the binary artifacts
+ for the microNPU
+ """
+ symbol = str(primfunc.attrs["global_symbol"])
+ params = primfunc.attrs["ethos-u.constants"]
+ input_size = primfunc.attrs["ethos-u.input_size"]
+ output_size = primfunc.attrs["ethos-u.output_size"]
+ tir_mod = tvm.IRModule()
+ tir_mod[symbol] = primfunc
+
+ params_with_int_keys = dict()
+ for idx in params.keys():
+ params_with_int_keys[int(idx)] = params[idx].numpy()
Review comment:
is it possible that the key "idx" is a "str"?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
Review comment:
ethos-u.constants is little confused because params is a dict instead of constant, should it be "params"?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
Review comment:
"This is a hook for the python-based lowering of relay function that gets offloaded to the microNPU."
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
+
Parameters
----------
- ext_func : tvm.relay.function.Function
- The partitioned relay function
+ ext_func : relay.Function
+ This is the partitioned relay function
+
Returns
-------
- cs : str
- An hex string of the bytes of command stream
- encoded_constants : str
- An hex string of the bytes that includes concat'd
- encoded weights, encoded biases and scales.
- scratch_size : int
- The size of the scratch buffer needed.
+ primfunc : tir.PrimFunc
+ This returns the scheduled PrimFunc
"""
+ assert len(ext_func.params) == 1
Review comment:
why len(ext_func.params) only can be 1, is this a limitation?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
Review comment:
"camel-case" for primfunc_to_artifact?
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
Review comment:
```
for idx, val in params.items():
params[idx] = tvm.nd.array(val)
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
+
Parameters
----------
- ext_func : tvm.relay.function.Function
- The partitioned relay function
+ ext_func : relay.Function
+ This is the partitioned relay function
+
Returns
-------
- cs : str
- An hex string of the bytes of command stream
- encoded_constants : str
- An hex string of the bytes that includes concat'd
- encoded weights, encoded biases and scales.
- scratch_size : int
- The size of the scratch buffer needed.
+ primfunc : tir.PrimFunc
+ This returns the scheduled PrimFunc
"""
+ assert len(ext_func.params) == 1
Review comment:
better to use RuntimeError with a error report message.
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
Review comment:
"This is a hook for the python-based lowering of the TIR PrimFunc which has undergone the unified optimization for the microNPU"
not sure if the said sentence match with original meaning, but at least need to add "a", "the" and lowercase the first letter of "Compilation" and "Artifact".
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
+
+ Returns
+ -------
+ CompilationArtifact
+ This is a structure that holds the binary artifacts
+ for the microNPU
+ """
+ symbol = str(primfunc.attrs["global_symbol"])
+ params = primfunc.attrs["ethos-u.constants"]
+ input_size = primfunc.attrs["ethos-u.input_size"]
+ output_size = primfunc.attrs["ethos-u.output_size"]
+ tir_mod = tvm.IRModule()
+ tir_mod[symbol] = primfunc
+
+ params_with_int_keys = dict()
+ for idx in params.keys():
+ params_with_int_keys[int(idx)] = params[idx].numpy()
Review comment:
```
for idx, val in params.items():
params_with_int_keys[int(idx)] = val.numpy()
```
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-981419708
@mbaret @leandron @ekalda PTAL when you have some time.
We need this in a bit urgently as it is a refactor of the codegen and it has the tendency to get conflicted.
If you can provide feedback, I ll create a followup ASAP with addressing them while we could take this in as it is green now.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759000219
##########
File path: src/relay/backend/contrib/ethosu/source_module.cc
##########
@@ -41,34 +41,37 @@
#include <vector>
#include "../../../../runtime/file_utils.h"
+#include "utils.h"
namespace tvm {
namespace runtime {
+using CompilationArtifact = relay::contrib::ethosu::CompilationArtifact;
+
// The runtime.Module that contains the host-side c code
// required for invoking the NPU with the command stream
class EthosUModuleNode : public ModuleNode {
public:
/*!
- * \brief The ethos runtime module.
+ * \brief The microNPU runtime module.
*
- * \param func_name_ name of the should be codegen'd function
- * \param cmms_hex_ command stream for the NPU in hex
- * \param weights_bias_hex_ the encoded biases and weights for the NPU in hex
- * \param scratch_size_ the size of the scratch memory required for command stream
- * \param input_size_ the size (in bytes) for the input tensor
- * \param output_size_ the size (in bytes) for the output tensor
+ * \param compilation_artifacts
+ * This is an array of CompilationArtifacts that is produced via
+ * lowering each PrimFunc to command stream. Here, those artifacts
+ * will be used to create the c-source.
*/
- explicit EthosUModuleNode(const String& func_name_, const String& cmms_hex_,
- const String& weights_bias_hex_, const Integer& scratch_size_,
- const Integer& input_size_, const Integer& output_size_) {
- func_name = func_name_;
- cmms_hex = std::move(cmms_hex_);
- weights_bias_hex = std::move(weights_bias_hex_);
- scratch_size = scratch_size_->value;
- input_size = input_size_->value;
- output_size = output_size_->value;
- c_source = GenerateSource();
+ explicit EthosUModuleNode(Array<CompilationArtifact> compilation_artifacts)
+ : compilation_artifacts_(compilation_artifacts) {
+ c_source += "#include <stdio.h>\n";
+ c_source += "#include <stdlib.h>\n";
+ c_source += "#include <tvm/runtime/crt/module.h>\n";
+ c_source += "#include <tvm_ethosu_runtime.h>\n";
+ c_source += "\n";
+ for (const CompilationArtifact& compilation_artifact : compilation_artifacts) {
+ c_source += GenerateSource(compilation_artifact);
+ c_source += "\n";
+ c_source += "\n";
Review comment:
Done in #9605
--
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 edited a comment on pull request #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm edited a comment on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-982358197
@leandron , do you think whether we can get this version in? Ive opened a follow up : #9605 to capture and address feedback. This should avoid grief from conflicts.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-982358197
@leandron , do you think whether we can get this version in? Ive opened a follow up : #9605 to capture and address feedback. This show avoid grief from conflicts.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#issuecomment-983374035
Thanks @huajsj . I ve modified and responded to your queries.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759926116
##########
File path: include/tvm/tir/transform.h
##########
@@ -287,6 +287,11 @@ TVM_DLL Pass LowerThreadAllreduce();
*/
TVM_DLL Pass InferFragment();
+/*!
+ * \brief This annotation for nodes to be disabled for builtin lowering
Review comment:
Done
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759000103
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
Review comment:
Done in #9605
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759906816
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -52,25 +35,25 @@ def constant_updater(expr, symbol): # pylint: disable=unused-argument
return dict()
-def _compile(ext_func):
+@tvm._ffi.register_func("relay.ext.ethos-u.relay_to_tir_func")
+def relay_to_tir_func(ext_func: relay.Function) -> tvm.tir.PrimFunc:
"""
- This is the main wrapper that accepts an external
- relay function and runs all the passes to lower it down
- to command stream
+ This is hook for python-based lowering of relay function
+ that gets offloaded to the microNPU.
+
Parameters
----------
- ext_func : tvm.relay.function.Function
- The partitioned relay function
+ ext_func : relay.Function
+ This is the partitioned relay function
+
Returns
-------
- cs : str
- An hex string of the bytes of command stream
- encoded_constants : str
- An hex string of the bytes that includes concat'd
- encoded weights, encoded biases and scales.
- scratch_size : int
- The size of the scratch buffer needed.
+ primfunc : tir.PrimFunc
+ This returns the scheduled PrimFunc
"""
+ assert len(ext_func.params) == 1
Review comment:
This is not a runtime error -- we have a pass in the pipeline to concat all inputs to one. So it is a internal developer facing assertion. User will never cause this error.
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759907993
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
+
+ primfunc = tir_mod["main"]
+ primfunc = primfunc.with_attr("global_symbol", ext_func.attrs["global_symbol"])
+ primfunc = primfunc.with_attr("ethos-u.constants", params)
+ primfunc = primfunc.with_attr("ethos-u.input_size", input_size)
+ primfunc = primfunc.with_attr("ethos-u.output_size", output_size)
+ return primfunc
+
+
+@tvm._ffi.register_func("relay.ext.ethos-u.primfunc_to_artifact")
+def primfunc_to_artifact(primfunc: tvm.tir.PrimFunc) -> util.CompilationArtifact:
+ """
+ This is hook for python-based lowering of TIR PrimFunc
+ that has undergone unified optimization to Compilation
+ Artifact destined for the microNPU.
+
+ Parameters
+ ----------
+ primfunc : tir.PrimFunc
+ TIR PrimFuncthat has undergone unified optimization
Review comment:
This is aleady corrected in the followup : #9605
--
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 #9597: [microNPU] Move the compilation to use Target Hooks.
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #9597:
URL: https://github.com/apache/tvm/pull/9597#discussion_r759907365
##########
File path: python/tvm/relay/backend/contrib/ethosu/codegen.py
##########
@@ -80,5 +63,50 @@ def _compile(ext_func):
# that can perform scheduling based on user inputs such as
# scratch memory size.
tir_mod, params = lower_to_tir(mod["main"], copy_constants())
- cmms, encoded_constants, scratch_size = tir_to_cs_translator.translate(tir_mod, params)
- return cmms, encoded_constants, scratch_size
+
+ for idx in params.keys():
+ params[idx] = tvm.nd.array(params[idx])
Review comment:
Seems like a stylistic choice, I would prefer to keep the current implementation as it does not have a meaningful difference.
--
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