areusch commented on a change in pull request #7785:
URL: https://github.com/apache/tvm/pull/7785#discussion_r618503865
##########
File path: src/relay/backend/build_module.cc
##########
@@ -598,6 +587,8 @@ class RelayBuildModule : public runtime::ModuleNode {
std::unordered_map<std::string, runtime::NDArray> params_;
/*! \brief building output */
BuildOutput ret_;
+ /*! \brief Executor used to execute the graph */
Review comment:
specify the possible values
##########
File path: src/target/metadata_module.h
##########
@@ -33,12 +33,15 @@
#include <string>
#include <unordered_map>
+#include "../runtime/meta_data.h"
+
namespace tvm {
namespace codegen {
runtime::Module CreateMetadataModule(
const std::unordered_map<std::string, runtime::NDArray>& params,
- tvm::runtime::Module target_module, const Array<runtime::Module>&
ext_modules, Target target);
+ tvm::runtime::Module target_module, const Array<runtime::Module>&
ext_modules, Target target,
+ int num_inputs = 1, int num_outputs = 1);
Review comment:
I think this disagrees w/ the function impl
##########
File path: src/runtime/meta_data.h
##########
@@ -37,6 +37,12 @@
#include "runtime_base.h"
+/*! \brief Value used to indicate the graph executor. */
+static constexpr const char* kTvmExecutorGraph = "graph";
Review comment:
let's place these inside a top-level namespace e.g. tvm::runtime
##########
File path: src/target/source/codegen_c_host.cc
##########
@@ -258,24 +304,13 @@ void CodeGenCHost::VisitExpr_(const CallNode* op,
std::ostream& os) { // NOLINT
this->stream << "TVMValue " << stack_name << "[" << size << "];\n";
os << stack_name;
} else if (op->op.same_as(builtin::tvm_call_packed_lowered())) {
- const StringImmNode* s = op->args[0].as<StringImmNode>();
- ICHECK(s != nullptr) << "tvm_call_packed_lowered expects first argument as
function name";
- int64_t begin = op->args[3].as<IntImmNode>()->value;
- int64_t end = op->args[4].as<IntImmNode>()->value;
- int64_t num_args = end - begin;
- ICHECK_GE(num_args, 0);
- std::string func_name = s->value;
- // NOTE: cannot rely on GetUnique for global decl_stream declarations
- // because it is reset between AddFunction().
- std::string packed_func_name = func_name + "_packed";
- if (declared_globals_.insert(packed_func_name).second) {
- // Still reserve the name among unique names.
- ICHECK(GetUniqueName(packed_func_name) == packed_func_name)
- << "Expected name " << packed_func_name << " to not be taken";
- decl_stream << "static void* " << packed_func_name << " = NULL;\n";
- }
- this->PrintGetFuncFromBackend(func_name, packed_func_name);
- this->PrintFuncCall(packed_func_name, num_args);
+ auto function_info = GetFunctionInfo(op);
+ this->PrintGetFuncFromBackend(function_info.func_name,
function_info.func_name_packed);
+ this->PrintFuncCall(function_info.func_name, function_info.num_args);
Review comment:
I think `function_info.func_name_packed`
##########
File path: src/runtime/meta_data.h
##########
@@ -32,13 +32,54 @@
#include <string>
#include <unordered_map>
+#include <utility>
#include <vector>
#include "runtime_base.h"
+/*! \brief Value used to indicate the graph executor. */
+static constexpr const char* kTvmExecutorGraph = "graph";
+
+/*! \brief Value used to indicate the aot executor. */
+static constexpr const char* kTvmExecutorAot = "aot";
+
namespace tvm {
namespace runtime {
+/*!
+ * \brief Structure that can be optionally used by the executor codegen
+ */
+class MetadataNode : public Object {
+ public:
+ /*! \brief number of inputs of the main function */
+ int num_inputs = 1;
Review comment:
would prefer to place the defaults on a constructor somewhere
##########
File path: src/target/source/codegen_c_host.cc
##########
@@ -47,6 +47,7 @@ void CodeGenCHost::Init(bool output_ssa, bool emit_asserts,
std::string target_s
decl_stream << "#define TVM_EXPORTS\n";
decl_stream << "#include \"tvm/runtime/c_runtime_api.h\"\n";
decl_stream << "#include \"tvm/runtime/c_backend_api.h\"\n";
+
Review comment:
needed?
##########
File path: src/runtime/crt/memory/stack_allocator.c
##########
@@ -0,0 +1,48 @@
+/*
+ * 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.
+ */
+
+// LINT_C_FILE
+
+#include <tvm/runtime/crt/stack_allocator.h>
+
+void* StackMemoryManager_Allocate(tvm_workspace_t* tvm_runtime_workspace,
int32_t nbytes) {
+ uint32_t offset_bytes = (~nbytes + 1) & (TVM_RUNTIME_ALLOC_ALIGNMENT_BYTES -
1);
Review comment:
don't think we followed-up on these?
##########
File path: python/tvm/relay/backend/executor_factory.py
##########
@@ -60,37 +147,17 @@ def __init__(self, ir_mod, target, graph_json_str, libmod,
libmod_name, params):
def export_library(self, file_name, fcompile=None, addons=None, **kwargs):
return self.module.export_library(file_name, fcompile, addons,
**kwargs)
- # Sometimes we want to get params explicitly.
- # For example, we want to save its params value to
- # an independent file.
+ def save_executor_config(self):
Review comment:
seems like `get_executor_config` would be a better name since we are not
actually writing to file here. alternatively, I think we want to eventually
move export_model_library_format here, so could just continue using
get_graph_json() conditionally there rather than introducing a new generic
function here
##########
File path: src/tir/transforms/lower_tvm_builtin.cc
##########
@@ -295,7 +297,11 @@ class BuiltinLower : public StmtExprMutator {
Array<PrimExpr> packed_args = {op->args[0], scope.stack_value,
scope.stack_tcode,
ConstInt32(arg_stack_begin),
ConstInt32(arg_stack_begin +
op->args.size() - 1)};
- return Call(DataType::Int(32), builtin::tvm_call_packed_lowered(),
packed_args);
+ if (use_string_lookup) {
Review comment:
can probably even further reduce to:
```
auto builtin_call = use_string_lookup ? builtin::tvm_call_packed_lowered() :
builtin::tvm_call_cpacked_lowered();
return Call(DataType::Int(32), builtin_call, packed_args);
```
##########
File path: tests/cpp/relay_build_module_test.cc
##########
@@ -119,7 +119,7 @@ TEST(Relay, BuildModule) {
targets.Set(0, llvm_tgt);
auto relay_mod = tvm::IRModule::FromExpr(func);
ICHECK(relay_mod.defined()) << "Module must be defined";
- build_f(relay_mod, targets, llvm_tgt);
+ build_f(relay_mod, targets, llvm_tgt, "graph");
Review comment:
use constant?
##########
File path: src/relay/backend/utils.h
##########
@@ -37,12 +37,24 @@
#include <typeinfo>
#include <unordered_map>
#include <unordered_set>
+#include <utility>
#include <vector>
+#include "../../runtime/meta_data.h"
+
namespace tvm {
namespace relay {
namespace backend {
+/*! \brief Lowered outputs */
Review comment:
can you expand the comment now that this is moved up to a .h file?
##########
File path: src/runtime/crt/memory/stack_allocator.c
##########
@@ -0,0 +1,48 @@
+/*
+ * 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.
+ */
+
+// LINT_C_FILE
+
+#include <tvm/runtime/crt/stack_allocator.h>
+
+void* StackMemoryManager_Allocate(tvm_workspace_t* tvm_runtime_workspace,
int32_t nbytes) {
+ uint32_t offset_bytes = (~nbytes + 1) & (TVM_RUNTIME_ALLOC_ALIGNMENT_BYTES -
1);
+ uint8_t* current_alloc = tvm_runtime_workspace->next_alloc;
+ uint8_t* next_alloc = tvm_runtime_workspace->next_alloc + nbytes +
offset_bytes;
+ uint8_t* workspace_end = tvm_runtime_workspace->workspace +
tvm_runtime_workspace->workspace_size;
+
+ if (next_alloc > workspace_end) {
+ return NULL;
+ }
+
+ tvm_runtime_workspace->next_alloc = next_alloc;
+ return current_alloc;
+}
+
+tvm_crt_error_t StackMemoryManager_Free(tvm_workspace_t*
tvm_runtime_workspace, void* ptr) {
+ tvm_runtime_workspace->next_alloc = ptr;
Review comment:
don't think we followed-up on these?
##########
File path: src/relay/backend/graph_plan_memory.cc
##########
@@ -209,15 +209,18 @@ class StorageAllocator : public StorageAllocaBaseVisitor {
for (const auto& kv : token_map_) {
std::vector<Integer> storage_ids;
std::vector<Integer> device_types;
+ std::vector<Integer> sid_sizes_byte;
Review comment:
i'm not sure this is exercised anywhere in unit test. can you add one?
right now we just assert on the len() but not the content
##########
File path: src/tir/transforms/lower_tvm_builtin.cc
##########
@@ -184,7 +184,9 @@ class BuiltinLower : public StmtExprMutator {
}
PrimExpr VisitExpr_(const CallNode* op) final {
if (op->op.same_as(builtin::tvm_call_packed())) {
- return MakeCallPacked(op);
+ return MakeCallPacked(op, true);
Review comment:
can you add /* use_string_lookup */ near true?
##########
File path: python/tvm/relay/build_module.py
##########
@@ -243,10 +248,18 @@ def build(ir_mod, target=None, target_host=None,
params=None, mod_name="default"
mod_name: Optional[str]
The module name we will build
+ executor: Optional[str]
+ The type of executor to be used in order to run the model:
+ - If "graph" is specified, then the graph_executor will be used
+ - If "aot" is specified, then the aot_executor will be used
+
Returns
-------
- graph_json : str
- The json string that can be accepted by graph executor.
+ internal_repr : str or tir.PrimFunc
Review comment:
maybe better termed `executor_config` or something?
##########
File path: tests/crt/aot_memory_test.cc
##########
@@ -0,0 +1,90 @@
+/*
+ * 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.
+ */
+
+#include <gtest/gtest.h>
+#include <tvm/runtime/crt/stack_allocator.h>
+
+/*
Review comment:
it would be great to add a test here for out-of-order freeing
--
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.
For queries about this service, please contact Infrastructure at:
[email protected]
