diff --git a/apps/bundle_deploy/Makefile b/apps/bundle_deploy/Makefile index 8a5f1cf95571..38d9d3456d55 100644 --- a/apps/bundle_deploy/Makefile +++ b/apps/bundle_deploy/Makefile @@ -62,6 +62,9 @@ $(endif) CRT_SRCS = $(shell find $(CRT_ROOT)) +MODEL_OBJ = $(build_dir)/model_c/devc.o $(build_dir)/model_c/lib0.o $(build_dir)/model_c/lib1.o +TEST_MODEL_OBJ = $(build_dir)/test_model_c/devc.o $(build_dir)/test_model_c/lib0.o $(build_dir)/test_model_c/lib1.o + demo_dynamic: $(build_dir)/demo_dynamic $(build_dir)/bundle.so $(build_dir)/bundle_c.so $(build_dir)/bundle.so $(build_dir)/graph_cpp.json $(build_dir)/graph_c.json $(build_dir)/params_cpp.bin $(build_dir)/params_c.bin $(build_dir)/cat.bin $(QUIET)TVM_NUM_THREADS=1 $(build_dir)/demo_dynamic $(build_dir)/bundle.so $(build_dir)/graph_cpp.json $(build_dir)/params_cpp.bin $(build_dir)/cat.bin $(QUIET)TVM_NUM_THREADS=1 $(build_dir)/demo_dynamic $(build_dir)/bundle_c.so $(build_dir)/graph_c.json $(build_dir)/params_c.bin $(build_dir)/cat.bin @@ -93,11 +96,11 @@ $(build_dir)/test_dynamic: test.cc ${build_dir}/test_graph_c.json ${build_dir}/t $(QUIET)mkdir -p $(@D) $(QUIET)g++ $(PKG_CXXFLAGS) -o $@ test.cc $(BACKTRACE_OBJS) $(BACKTRACE_LDFLAGS) -$(build_dir)/demo_static: demo_static.c ${build_dir}/bundle_static.o ${build_dir}/model_c.o ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a ${build_dir}/graph_c.json.c ${build_dir}/params_c.bin.c $(BACKTRACE_OBJS) +$(build_dir)/demo_static: demo_static.c ${build_dir}/bundle_static.o $(MODEL_OBJ) ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a ${build_dir}/graph_c.json.c ${build_dir}/params_c.bin.c $(BACKTRACE_OBJS) $(QUIET)mkdir -p $(@D) $(QUIET)gcc $(PKG_CFLAGS) -o $@ $^ $(PKG_LDFLAGS) $(BACKTRACE_LDFLAGS) $(BACKTRACE_CFLAGS) -$(build_dir)/test_static: test_static.c ${build_dir}/bundle_static.o ${build_dir}/test_model_c.o ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) +$(build_dir)/test_static: test_static.c ${build_dir}/bundle_static.o $(TEST_MODEL_OBJ) ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) $(QUIET)mkdir -p $(@D) $(QUIET)gcc $(PKG_CFLAGS) -o $@ $^ $(BACKTRACE_LDFLAGS) @@ -119,11 +122,15 @@ $(build_dir)/params_c.bin.c: $(build_dir)/params_c.bin $(build_dir)/params_cpp.bin.c: $(build_dir)/params_cpp.bin $(QUIET)xxd -i $^ > $@ -$(build_dir)/model_c.o $(build_dir)/graph_c.json $(build_dir)/model_cpp.o $(build_dir)/graph_cpp.json $(build_dir)/params.bin $(build_dir)/cat.bin: build_model.py +$(MODEL_OBJ) $(build_dir)/graph_c.json $(build_dir)/model_cpp.o $(build_dir)/graph_cpp.json $(build_dir)/params.bin $(build_dir)/cat.bin: build_model.py $(QUIET)python3 $< -o $(build_dir) + $(QUIET)mkdir -p build/model_c + $(QUIET)tar -C build/model_c -xvf build/model_c.tar -$(build_dir)/test_model_c.o $(build_dir)/test_graph_c.json $(build_dir)/test_params_c.bin $(build_dir)/test_data_c.bin $(build_dir)/test_output_c.bin $(build_dir)/test_model_cpp.o $(build_dir)/test_graph_cpp.json $(build_dir)/test_params_cpp.bin $(build_dir)/test_data_cpp.bin $(build_dir)/test_output_cpp.bin: build_model.py +$(TEST_MODEL_OBJ) $(build_dir)/test_graph_c.json $(build_dir)/test_params_c.bin $(build_dir)/test_data_c.bin $(build_dir)/test_output_c.bin $(build_dir)/test_model_cpp.o $(build_dir)/test_graph_cpp.json $(build_dir)/test_params_cpp.bin $(build_dir)/test_data_cpp.bin $(build_dir)/test_output_cpp.bin: build_model.py $(QUIET)python3 $< -o $(build_dir) --test + $(QUIET)mkdir -p build/test_model_c + $(QUIET)tar -C build/test_model_c -xvf build/test_model_c.tar # Build our bundle against the serialized bundle.c API, the runtime.cc API, and # the serialized graph.json and params.bin @@ -131,7 +138,7 @@ $(build_dir)/bundle.so: bundle.cc runtime.cc $(build_dir)/model_cpp.o $(QUIET)mkdir -p $(@D) $(QUIET)g++ -shared $(PKG_CXXFLAGS) -fvisibility=hidden -o $@ $^ $(PKG_LDFLAGS) -$(build_dir)/bundle_c.so: bundle.c $(build_dir)/model_c.o ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) +$(build_dir)/bundle_c.so: bundle.c $(MODEL_OBJ) ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) $(QUIET)mkdir -p $(@D) $(QUIET)gcc -shared $(PKG_CFLAGS) -fvisibility=hidden -o $@ $^ $(PKG_LDFLAGS) $(BACKTRACE_LDFLAGS) $(BACKTRACE_CFLAGS) @@ -139,7 +146,7 @@ $(build_dir)/test_bundle.so: bundle.cc runtime.cc $(build_dir)/test_model_cpp.o $(QUIET)mkdir -p $(@D) $(QUIET)g++ -shared $(PKG_CXXFLAGS) -fvisibility=hidden -o $@ $^ $(PKG_LDFLAGS) -$(build_dir)/test_bundle_c.so: bundle.c $(build_dir)/test_model_c.o ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) +$(build_dir)/test_bundle_c.so: bundle.c $(TEST_MODEL_OBJ) ${build_dir}/crt/libmemory.a ${build_dir}/crt/libgraph_runtime.a ${build_dir}/crt/libcommon.a $(BACKTRACE_OBJS) $(QUIET)mkdir -p $(@D) $(QUIET)gcc -shared $(PKG_CFLAGS) -fvisibility=hidden -o $@ $^ $(PKG_LDFLAGS) $(BACKTRACE_LDFLAGS) $(BACKTRACE_CFLAGS) diff --git a/apps/bundle_deploy/build_model.py b/apps/bundle_deploy/build_model.py index a2513c8a46d0..0991ac9ad94b 100644 --- a/apps/bundle_deploy/build_model.py +++ b/apps/bundle_deploy/build_model.py @@ -23,6 +23,7 @@ from tvm import te import logging import json +from tvm.contrib import cc as _cc RUNTIMES = { "c": "{name}_c.{ext}", @@ -51,7 +52,17 @@ def build_module(opts): build_dir = os.path.abspath(opts.out_dir) if not os.path.isdir(build_dir): os.makedirs(build_dir) - lib.save(os.path.join(build_dir, file_format_str.format(name="model", ext="o"))) + ext = "tar" if runtime_name == "c" else "o" + lib_file_name = os.path.join(build_dir, file_format_str.format(name="model", ext=ext)) + if runtime_name == "c": + lib.export_library(lib_file_name) + else: + # NOTE: at present, export_libarary will always create _another_ shared object, and you + # can't stably combine two shared objects together (in this case, init_array is not + # populated correctly when you do that). So for now, must continue to use save() with the + # C++ library. + # TODO(areusch): Obliterate runtime.cc and replace with libtvm_runtime.so. + lib.save(lib_file_name) with open( os.path.join(build_dir, file_format_str.format(name="graph", ext="json")), "w" ) as f_graph_json: @@ -84,7 +95,17 @@ def build_test_module(opts): build_dir = os.path.abspath(opts.out_dir) if not os.path.isdir(build_dir): os.makedirs(build_dir) - lib.save(os.path.join(build_dir, file_format_str.format(name="test_model", ext="o"))) + ext = "tar" if runtime_name == "c" else "o" + lib_file_name = os.path.join(build_dir, file_format_str.format(name="test_model", ext=ext)) + if runtime_name == "c": + lib.export_library(lib_file_name) + else: + # NOTE: at present, export_libarary will always create _another_ shared object, and you + # can't stably combine two shared objects together (in this case, init_array is not + # populated correctly when you do that). So for now, must continue to use save() with the + # C++ library. + # TODO(areusch): Obliterate runtime.cc and replace with libtvm_runtime.so. + lib.save(lib_file_name) with open( os.path.join(build_dir, file_format_str.format(name="test_graph", ext="json")), "w" ) as f_graph_json: diff --git a/python/tvm/driver/build_module.py b/python/tvm/driver/build_module.py index 7ad48e19a1db..5eaecb422163 100644 --- a/python/tvm/driver/build_module.py +++ b/python/tvm/driver/build_module.py @@ -428,12 +428,19 @@ def build(inputs, args=None, target=None, target_host=None, name="default_functi if not isinstance(target_host, Target): target_host = Target(target_host) if ( - "system-lib" in target_host.attrs - and target_host.attrs["system-lib"].value == 1 - and target_host.kind.name == "c" + target_host.attrs.get("runtime", tvm.runtime.String("c++")) == "c" + and target_host.attrs.get("system-lib", 0).value == 1 ): - create_csource_metadata_module = tvm._ffi.get_global_func( - "runtime.CreateCSourceMetadataModule" - ) - return create_csource_metadata_module([rt_mod_host], target_host) + if target_host.kind.name == "c": + create_csource_crt_metadata_module = tvm._ffi.get_global_func( + "runtime.CreateCSourceCrtMetadataModule" + ) + return create_csource_crt_metadata_module([rt_mod_host], target_host) + + if target_host.kind.name == "llvm": + create_llvm_crt_metadata_module = tvm._ffi.get_global_func( + "runtime.CreateLLVMCrtMetadataModule" + ) + return create_llvm_crt_metadata_module([rt_mod_host], target_host) + return rt_mod_host diff --git a/python/tvm/micro/compiler.py b/python/tvm/micro/compiler.py index d0431f42b01d..5bc5aba8a1be 100644 --- a/python/tvm/micro/compiler.py +++ b/python/tvm/micro/compiler.py @@ -81,6 +81,9 @@ def _target_from_sources(cls, sources): target_strs = set() for obj in sources: + if os.path.splitext(obj)[1] not in (".cc", ".c"): + continue + with open(obj) as obj_f: for line in obj_f: m = cls.TVM_TARGET_RE.match(line) @@ -246,7 +249,8 @@ def library(self, output, sources, options=None): ) prefix = self._autodetect_toolchain_prefix(target) - outputs = [] + outputs = [s for s in sources if os.path.splitext(s)[1] == ".o"] + sources = [s for s in sources if s not in outputs] for src in sources: src_base, src_ext = os.path.splitext(os.path.basename(src)) diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc index e2a8553199f0..b37cd73ece04 100644 --- a/src/target/llvm/codegen_cpu.cc +++ b/src/target/llvm/codegen_cpu.cc @@ -123,12 +123,6 @@ void CodeGenCPU::AddFunction(const PrimFunc& f) { << "CodeGenLLVM: Expect PrimFunc to have the global_symbol attribute"; export_system_symbols_.emplace_back( std::make_pair(global_symbol.value().operator std::string(), function_)); - } else if (target_c_runtime_) { - auto global_symbol = f->GetAttr(tvm::attr::kGlobalSymbol); - ICHECK(global_symbol.defined()) - << "CodeGenLLVM: Expect PrimFunc to have the global_symbol attribute"; - registry_functions_.emplace_back( - std::make_pair(global_symbol.value().operator std::string(), function_)); } AddDebugInformation(function_); } @@ -791,47 +785,50 @@ llvm::Value* CodeGenCPU::RuntimeTVMParallelBarrier() { return GetContextPtr(gv_tvm_parallel_barrier_); } -void CodeGenCPU::AddStartupFunction() { - if (registry_functions_.size() != 0) { - ICHECK(is_system_lib_) << "Loading of --system-lib modules is yet to be defined for C runtime"; - Array symbols; - std::vector funcs; - for (auto sym : registry_functions_) { - symbols.push_back(sym.first); - funcs.emplace_back(llvm::ConstantExpr::getBitCast( - sym.second, ftype_tvm_backend_packed_c_func_->getPointerTo())); - } - llvm::DataLayout layout(module_.get()); - llvm::ArrayType* t_tvm_crt_func_ptrs = - llvm::ArrayType::get(ftype_tvm_backend_packed_c_func_->getPointerTo(), funcs.size()); - llvm::GlobalVariable* func_registry_ptrs = new llvm::GlobalVariable( - *module_, t_tvm_crt_func_ptrs, true, llvm::GlobalValue::InternalLinkage, - llvm::ConstantArray::get(t_tvm_crt_func_ptrs, funcs), "_tvm_func_registry_ptrs"); - uint64_t align = layout.getTypeAllocSize(ftype_tvm_backend_packed_c_func_->getPointerTo()); +void CodeGenCPU::DefineFunctionRegistry(Array func_names) { + ICHECK(is_system_lib_) << "Loading of --system-lib modules is yet to be defined for C runtime"; + Array symbols; + std::vector funcs; + for (auto sym : func_names) { + symbols.push_back(sym); + llvm::GlobalVariable* sym_func = new llvm::GlobalVariable( + *module_, ftype_tvm_backend_packed_c_func_, true, llvm::GlobalValue::ExternalLinkage, + nullptr, sym.operator std::string()); + funcs.emplace_back(sym_func); + } + llvm::DataLayout layout(module_.get()); + llvm::ArrayType* t_tvm_crt_func_ptrs = + llvm::ArrayType::get(ftype_tvm_backend_packed_c_func_->getPointerTo(), funcs.size()); + llvm::GlobalVariable* func_registry_ptrs = new llvm::GlobalVariable( + *module_, t_tvm_crt_func_ptrs, true, llvm::GlobalValue::InternalLinkage, + llvm::ConstantArray::get(t_tvm_crt_func_ptrs, funcs), "_tvm_func_registry_ptrs"); + uint64_t align = layout.getTypeAllocSize(ftype_tvm_backend_packed_c_func_->getPointerTo()); #if TVM_LLVM_VERSION >= 100 - func_registry_ptrs->setAlignment(llvm::Align(align)); + func_registry_ptrs->setAlignment(llvm::Align(align)); #else - func_registry_ptrs->setAlignment(align); + func_registry_ptrs->setAlignment(align); #endif - llvm::GlobalVariable* func_registry = new llvm::GlobalVariable( - *module_, t_tvm_crt_func_registry_, true, llvm::GlobalVariable::InternalLinkage, - llvm::ConstantStruct::get( - t_tvm_crt_func_registry_, - {GetConstString(::tvm::target::GenerateFuncRegistryNames(symbols)), - func_registry_ptrs}), - "_tvm_crt_func_registry"); - llvm::GlobalVariable* module = new llvm::GlobalVariable( - *module_, t_tvm_crt_module_, true, llvm::GlobalValue::InternalLinkage, - llvm::ConstantStruct::get(t_tvm_crt_module_, {func_registry}), "_tvm_crt_module"); - - // Now build TVMSystemLibEntryPoint. - llvm::FunctionType* ftype = llvm::FunctionType::get(t_void_p_, {}, false); - function_ = llvm::Function::Create(ftype, llvm::Function::ExternalLinkage, - "TVMSystemLibEntryPoint", module_.get()); - llvm::BasicBlock* entry_point_entry = llvm::BasicBlock::Create(*ctx_, "entry", function_); - builder_->SetInsertPoint(entry_point_entry); - builder_->CreateRet(builder_->CreateBitCast(module, t_void_p_)); - } else { + llvm::GlobalVariable* func_registry = new llvm::GlobalVariable( + *module_, t_tvm_crt_func_registry_, true, llvm::GlobalVariable::InternalLinkage, + llvm::ConstantStruct::get( + t_tvm_crt_func_registry_, + {GetConstString(::tvm::target::GenerateFuncRegistryNames(symbols)), func_registry_ptrs}), + "_tvm_crt_func_registry"); + llvm::GlobalVariable* module = new llvm::GlobalVariable( + *module_, t_tvm_crt_module_, true, llvm::GlobalValue::InternalLinkage, + llvm::ConstantStruct::get(t_tvm_crt_module_, {func_registry}), "_tvm_crt_module"); + + // Now build TVMSystemLibEntryPoint. + llvm::FunctionType* ftype = llvm::FunctionType::get(t_void_p_, {}, false); + function_ = llvm::Function::Create(ftype, llvm::Function::ExternalLinkage, + "TVMSystemLibEntryPoint", module_.get()); + llvm::BasicBlock* entry_point_entry = llvm::BasicBlock::Create(*ctx_, "entry", function_); + builder_->SetInsertPoint(entry_point_entry); + builder_->CreateRet(builder_->CreateBitCast(module, t_void_p_)); +} + +void CodeGenCPU::AddStartupFunction() { + if (!target_c_runtime_) { llvm::FunctionType* ftype = llvm::FunctionType::get(t_void_, {}, false); function_ = llvm::Function::Create(ftype, llvm::Function::InternalLinkage, "__tvm_module_startup", module_.get()); diff --git a/src/target/llvm/codegen_cpu.h b/src/target/llvm/codegen_cpu.h index fc46dc53ce15..d08bd639e131 100644 --- a/src/target/llvm/codegen_cpu.h +++ b/src/target/llvm/codegen_cpu.h @@ -50,6 +50,12 @@ class CodeGenCPU : public CodeGenLLVM { llvm::Value* CreateCallExtern(Type ret_type, String global_symbol, const Array& args, bool skip_first_arg) override; + /*! + * \brief A CPU-specific function to create the FuncRegistry. + * \param func_names List of functions to be included, in order. + */ + void DefineFunctionRegistry(Array func_names); + protected: void AddStartupFunction() final; // meta data diff --git a/src/target/llvm/llvm_module.cc b/src/target/llvm/llvm_module.cc index 43d20971404e..24fb3dc95819 100644 --- a/src/target/llvm/llvm_module.cc +++ b/src/target/llvm/llvm_module.cc @@ -34,6 +34,7 @@ #include "../../runtime/library_module.h" #include "../func_registry_generator.h" #include "codegen_blob.h" +#include "codegen_cpu.h" #include "codegen_llvm.h" #include "llvm_common.h" @@ -445,6 +446,58 @@ TVM_REGISTER_GLOBAL("codegen.codegen_blob") return runtime::Module(n); }); +runtime::Module CreateLLVMCrtMetadataModule(const Array& modules, Target target) { + Array func_names; + for (runtime::Module mod : modules) { + auto pf_funcs = mod.GetFunction("get_func_names"); + if (pf_funcs != nullptr) { + Array func_names_ = pf_funcs(); + for (const auto& fname : func_names_) { + func_names.push_back(fname); + } + } + } + + InitializeLLVM(); + auto tm = GetLLVMTargetMachine(target); + bool system_lib = target->GetAttr("system-lib").value_or(Bool(false)); + bool target_c_runtime = (target->GetAttr("runtime").value_or("") == kTvmRuntimeCrt); + ICHECK(system_lib && target_c_runtime) + << "For LLVM C-runtime metadata module, must include --system-lib and --runtime=c; " + << "got target: " << target->str(); + auto ctx = std::make_shared(); + std::unique_ptr cg{new CodeGenCPU()}; + cg->Init("TVMMetadataMod", tm.get(), ctx.get(), system_lib, system_lib, target_c_runtime); + + cg->DefineFunctionRegistry(func_names); + auto mod = cg->Finish(); + mod->addModuleFlag(llvm::Module::Warning, "tvm_target", + llvm::MDString::get(*ctx, LLVMTargetToString(target))); + mod->addModuleFlag(llvm::Module::Override, "Debug Info Version", llvm::DEBUG_METADATA_VERSION); + + if (tm->getTargetTriple().isOSDarwin()) { + mod->addModuleFlag(llvm::Module::Override, "Dwarf Version", 2); + } + + std::string verify_errors_storage; + llvm::raw_string_ostream verify_errors(verify_errors_storage); + LOG_IF(FATAL, llvm::verifyModule(*mod, &verify_errors)) + << "LLVM module verification failed with the following errors: \n" + << verify_errors.str(); + + auto n = make_object(); + n->Init(std::move(mod), ctx); + for (auto m : modules) { + n->Import(m); + } + return runtime::Module(n); +} + +TVM_REGISTER_GLOBAL("runtime.CreateLLVMCrtMetadataModule") + .set_body_typed([](const Array& modules, Target target) { + return CreateLLVMCrtMetadataModule(modules, target); + }); + } // namespace codegen } // namespace tvm #endif // TVM_LLVM_VERSION diff --git a/src/target/llvm/llvm_module.h b/src/target/llvm/llvm_module.h new file mode 100644 index 000000000000..3eab00c643e5 --- /dev/null +++ b/src/target/llvm/llvm_module.h @@ -0,0 +1,44 @@ +/* + * 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. + */ + +/*! + * \file llvm_module.h + * \brief Declares top-level shared functions related to the LLVM codegen. + */ + +#ifndef TVM_TARGET_LLVM_LLVM_MODULE_H_ +#define TVM_TARGET_LLVM_LLVM_MODULE_H_ + +#include +#include +#include + +#ifdef TVM_LLVM_VERSION + +namespace tvm { +namespace codegen { + +runtime::Module CreateLLVMCrtMetadataModule(const Array& modules, Target target); + +} // namespace codegen +} // namespace tvm + +#endif // TVM_LLVM_VERSION + +#endif // TVM_TARGET_LLVM_LLVM_MODULE_H_ diff --git a/src/target/metadata_module.cc b/src/target/metadata_module.cc new file mode 100644 index 000000000000..e2575c34d8f2 --- /dev/null +++ b/src/target/metadata_module.cc @@ -0,0 +1,136 @@ +/* + * 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. + */ + +/*! + * \file metadata_module.cc + * \brief Defines functions that build MetadataModules for C++ and C runtimes. + */ + +#include "metadata_module.h" + +#include + +#include "../runtime/meta_data.h" +#include "llvm/llvm_module.h" +#include "source/source_module.h" + +namespace tvm { +namespace codegen { + +/*! + * \brief Create a metadata module wrapper. The helper is used by different + * codegens, such as graph runtime codegen and the vm compiler. + * + * \param params The metadata for initialization of all modules. + * \param target_module the internal module that is compiled by tvm. + * \param ext_modules The external modules that needs to be imported inside the metadata + * module(s). + * \param target The target that all the modules are compiled for + * \return The created metadata module that manages initialization of metadata. + */ +runtime::Module CreateMetadataModule( + const std::unordered_map& params, + tvm::runtime::Module target_module, const Array& ext_modules, Target target) { + // Here we split modules into two groups: + // 1. Those modules which can be exported to C-runtime. These are DSO-exportable + // (i.e. llvm or c) modules which return nothing from get_const_vars(). + // 2. Other modules. + Array crt_exportable_modules; + Array non_crt_exportable_modules; + + auto DSOExportable = [](tvm::runtime::Module& mod) { + return !std::strcmp(mod->type_key(), "llvm") || !std::strcmp(mod->type_key(), "c"); + }; + + bool is_targeting_crt = + target.defined() && target->GetAttr("runtime").value_or(String("")) == kTvmRuntimeCrt; + + // Wrap all submodules in the initialization wrapper. + std::unordered_map> sym_metadata; + for (tvm::runtime::Module mod : ext_modules) { + auto pf_sym = mod.GetFunction("get_symbol"); + auto pf_var = mod.GetFunction("get_const_vars"); + std::vector arrays; + if (pf_sym != nullptr && pf_var != nullptr) { + String symbol = pf_sym(); + Array variables = pf_var(); + for (size_t i = 0; i < variables.size(); i++) { + arrays.push_back(variables[i].operator std::string()); + } + ICHECK_EQ(sym_metadata.count(symbol), 0U) << "Found duplicated symbol: " << symbol; + sym_metadata[symbol] = arrays; + } + // We only need loading of serialized constant data + // if there are constants present and required by the + // runtime module to be initialized by the binary + // metadata module. If not rest of the modules are + // wrapped in c-source metadata module. + + // TODO(@manupa-arm) : we should be able to use csource_metadata + // if the variables are empty when all the runtime modules implement get_func_names + if (arrays.empty() && is_targeting_crt && DSOExportable(mod) && + (target->kind->name == "c" || target->kind->name == "llvm")) { + crt_exportable_modules.push_back(mod); + } else { + non_crt_exportable_modules.push_back(mod); + } + } + + if (is_targeting_crt) { + if (!non_crt_exportable_modules.empty()) { + std::string non_exportable_modules; + for (unsigned int i = 0; i < non_crt_exportable_modules.size(); i++) { + if (i > 0) { + non_exportable_modules += ", "; + } + auto mod = non_crt_exportable_modules[i]; + auto pf_sym = mod.GetFunction("get_symbol"); + if (pf_sym != nullptr) { + non_exportable_modules += pf_sym().operator std::string(); + } else { + non_exportable_modules += + std::string{"(module type_key="} + mod->type_key() + std::string{")"}; + } + } + CHECK(false) << "These " << non_crt_exportable_modules.size() + << " modules are not exportable to C-runtime: " << non_exportable_modules; + } + + if (target->kind->name == "c") { + crt_exportable_modules.push_back(target_module); + target_module = CreateCSourceCrtMetadataModule(crt_exportable_modules, target); + } else if (target->kind->name == "llvm") { + crt_exportable_modules.push_back(target_module); + target_module = CreateLLVMCrtMetadataModule(crt_exportable_modules, target); + } + } else { + if (!non_crt_exportable_modules.empty()) { + runtime::Module binary_meta_mod = runtime::MetadataModuleCreate(params, sym_metadata); + binary_meta_mod.Import(target_module); + for (const auto& it : non_crt_exportable_modules) { + binary_meta_mod.Import(it); + } + return binary_meta_mod; + } + } + return target_module; +} + +} // namespace codegen +} // namespace tvm diff --git a/src/target/metadata_module.h b/src/target/metadata_module.h new file mode 100644 index 000000000000..83cb29dd5a46 --- /dev/null +++ b/src/target/metadata_module.h @@ -0,0 +1,46 @@ +/* + * 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. + */ + +/*! + * \file metadata_module.h + * \brief Declares functions that build MetadataModules for C++ and C runtimes. + */ + +#ifndef TVM_TARGET_METADATA_MODULE_H_ +#define TVM_TARGET_METADATA_MODULE_H_ + +#include +#include +#include +#include + +#include +#include + +namespace tvm { +namespace codegen { + +runtime::Module CreateMetadataModule( + const std::unordered_map& params, + tvm::runtime::Module target_module, const Array& ext_modules, Target target); + +} // namespace codegen +} // namespace tvm + +#endif // TVM_TARGET_METADATA_MODULE_H_ diff --git a/src/target/source/codegen_source_base.h b/src/target/source/codegen_source_base.h index ed838f825812..3baa44eb639f 100644 --- a/src/target/source/codegen_source_base.h +++ b/src/target/source/codegen_source_base.h @@ -170,12 +170,13 @@ runtime::Module DeviceSourceModuleCreate( std::string type_key, std::function fget_source = nullptr); /*! - * \brief Wrap the submodules that are to be wrapped in a c-source metadata module. + * \brief Wrap the submodules that are to be wrapped in a c-source metadata module for C runtime. * \param modules The modules to be wrapped. * \param target the target the modules are compiled for. * \return The wrapped module. */ -runtime::Module CreateCSourceMetadataModule(const Array& modules, Target target); +runtime::Module CreateCSourceCrtMetadataModule(const Array& modules, + Target target); } // namespace codegen } // namespace tvm diff --git a/src/target/source/source_module.cc b/src/target/source/source_module.cc index 4b4770a79816..a7732719a699 100644 --- a/src/target/source/source_module.cc +++ b/src/target/source/source_module.cc @@ -21,12 +21,17 @@ * \file source_module.cc * \brief Source code module, only for viewing */ +#include "source_module.h" + #include #include #include +#include +#include +#include + #include "../../runtime/file_utils.h" -#include "../../runtime/meta_data.h" #include "../../support/str_escape.h" #include "../func_registry_generator.h" #include "codegen_source_base.h" @@ -43,73 +48,6 @@ using runtime::GetFileFormat; using runtime::GetMetaFilePath; using runtime::SaveBinaryToFile; -/*! - * \brief Create a metadata module wrapper. The helper is used by different - * codegens, such as graph runtime codegen and the vm compiler. - * - * \param params The metadata for initialization of all modules. - * \param target_module the internal module that is compiled by tvm. - * \param ext_modules The external modules that needs to be imported inside the metadata - * module(s). - * \param target The target that all the modules are compiled for - * \return The created metadata module that manages initialization of metadata. - */ -runtime::Module CreateMetadataModule( - const std::unordered_map& params, - tvm::runtime::Module target_module, const Array& ext_modules, Target target) { - Array csource_modules; - Array binary_modules; - - auto DSOExportable = [](tvm::runtime::Module& mod) { - return !std::strcmp(mod->type_key(), "llvm") || !std::strcmp(mod->type_key(), "c"); - }; - - // Wrap all submodules in the initialization wrapper. - std::unordered_map> sym_metadata; - for (tvm::runtime::Module mod : ext_modules) { - auto pf_sym = mod.GetFunction("get_symbol"); - auto pf_var = mod.GetFunction("get_const_vars"); - std::vector arrays; - if (pf_sym != nullptr && pf_var != nullptr) { - String symbol = pf_sym(); - Array variables = pf_var(); - for (size_t i = 0; i < variables.size(); i++) { - arrays.push_back(variables[i].operator std::string()); - } - ICHECK_EQ(sym_metadata.count(symbol), 0U) << "Found duplicated symbol: " << symbol; - sym_metadata[symbol] = arrays; - } - // We only need loading of serialized constant data - // if there are constants present and required by the - // runtime module to be initialized by the binary - // metadata module. If not rest of the modules are - // wrapped in c-source metadata module. - - // TODO(@manupa-arm) : we should be able to use csource_metadata - // if the variables are empty when all the runtime modules implement get_func_names - if (arrays.empty() && DSOExportable(mod) && target->kind->name == "c") { - csource_modules.push_back(mod); - } else { - binary_modules.push_back(mod); - } - } - - if (target.defined() && target->kind->name == "c") { - csource_modules.push_back(target_module); - target_module = CreateCSourceMetadataModule(csource_modules, target); - } - - if (!binary_modules.empty()) { - runtime::Module binary_meta_mod = runtime::MetadataModuleCreate(params, sym_metadata); - binary_meta_mod.Import(target_module); - for (const auto& it : binary_modules) { - binary_meta_mod.Import(it); - } - return binary_meta_mod; - } - return target_module; -} - // Simulator function class SourceModuleNode : public runtime::ModuleNode { public: @@ -189,9 +127,10 @@ runtime::Module CSourceModuleCreate(const String& code, const String& fmt, return runtime::Module(n); } -class CSourceMetadataModuleNode : public runtime::ModuleNode { +class CSourceCrtMetadataModuleNode : public runtime::ModuleNode { public: - CSourceMetadataModuleNode(const Array& func_names, const std::string& fmt, Target target) + CSourceCrtMetadataModuleNode(const Array& func_names, const std::string& fmt, + Target target) : fmt_(fmt), func_names_(func_names), target_(target) { CreateSource(); } @@ -261,7 +200,8 @@ class CSourceMetadataModuleNode : public runtime::ModuleNode { } }; -runtime::Module CreateCSourceMetadataModule(const Array& modules, Target target) { +runtime::Module CreateCSourceCrtMetadataModule(const Array& modules, + Target target) { Array func_names; for (runtime::Module mod : modules) { auto pf_funcs = mod.GetFunction("get_func_names"); @@ -272,7 +212,7 @@ runtime::Module CreateCSourceMetadataModule(const Array& module } } } - auto n = make_object(func_names, "cc", target); + auto n = make_object(func_names, "cc", target); auto csrc_metadata_module = runtime::Module(n); for (const auto& mod : modules) { csrc_metadata_module.Import(mod); @@ -341,9 +281,9 @@ TVM_REGISTER_GLOBAL("runtime.CSourceModuleCreate") return CSourceModuleCreate(code, fmt, func_names, const_vars); }); -TVM_REGISTER_GLOBAL("runtime.CreateCSourceMetadataModule") +TVM_REGISTER_GLOBAL("runtime.CreateCSourceCrtMetadataModule") .set_body_typed([](const Array& modules, Target target) { - return CreateCSourceMetadataModule(modules, target); + return CreateCSourceCrtMetadataModule(modules, target); }); } // namespace codegen diff --git a/src/target/source/source_module.h b/src/target/source/source_module.h new file mode 100644 index 000000000000..45858b9f4ef2 --- /dev/null +++ b/src/target/source/source_module.h @@ -0,0 +1,46 @@ +/* + * 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. + */ + +/*! + * \file source_module.h + * \brief Source code module + */ + +#ifndef TVM_TARGET_SOURCE_SOURCE_MODULE_H_ +#define TVM_TARGET_SOURCE_SOURCE_MODULE_H_ + +#include +#include +#include + +namespace tvm { +namespace codegen { + +/*! + * \brief Create C-runtime targeted metadata module for "c" backend. + * \param modules Array of modules included in the compilation output. + * \param target TVM target. + */ +runtime::Module CreateCSourceCrtMetadataModule(const Array& modules, + tvm::Target target); + +} // namespace codegen +} // namespace tvm + +#endif // TVM_TARGET_SOURCE_SOURCE_MODULE_H_ diff --git a/tests/python/unittest/test_crt.py b/tests/python/unittest/test_crt.py index 3c68b4090309..4b744b8ee10a 100644 --- a/tests/python/unittest/test_crt.py +++ b/tests/python/unittest/test_crt.py @@ -28,7 +28,6 @@ import pytest import tvm -import tvm.testing import tvm.relay import tvm.testing @@ -103,6 +102,23 @@ def test_compile_runtime(): assert (C_data.asnumpy() == np.array([6, 7])).all() +@tvm.testing.requires_micro +def test_compile_runtime_llvm(): + """Test targeting the on-device runtime with the llvm backend.""" + global TARGET + old_target = TARGET + try: + # NOTE: test_compile_runtime uses the "c" backend--re run it using the llvm backend. + target_str = str(TARGET) + assert target_str.startswith("c ") + TARGET = tvm.target.Target("llvm " + str(TARGET)[len("c ") :]) + + test_compile_runtime() + + finally: + TARGET = old_target + + @tvm.testing.requires_micro def test_reset(): """Test when the remote end resets during a session.""" @@ -124,7 +140,7 @@ def test_graph_runtime(): """Test use of the graph runtime with microTVM.""" import tvm.micro - workspace = tvm.micro.Workspace() + workspace = tvm.micro.Workspace(debug=True) relay_mod = tvm.parser.fromtext( """ #[version = "0.0.5"] @@ -157,6 +173,19 @@ def test_std_math_functions(): """Verify that standard math functions can be used.""" import tvm.micro + workspace = tvm.micro.Workspace() + + with _make_add_sess(workspace) as sess: + A_data = tvm.nd.array(np.array([2, 3], dtype="int8"), ctx=sess.context) + assert (A_data.asnumpy() == np.array([2, 3])).all() + B_data = tvm.nd.array(np.array([4], dtype="int8"), ctx=sess.context) + assert (B_data.asnumpy() == np.array([4])).all() + C_data = tvm.nd.array(np.array([0, 0], dtype="int8"), ctx=sess.context) + assert (C_data.asnumpy() == np.array([0, 0])).all() + + system_lib = sess.get_system_lib() + system_lib.get_function("add")(A_data, B_data, C_data) + workspace = tvm.micro.Workspace() A = tvm.te.placeholder((2,), dtype="float32", name="A") B = tvm.te.compute(A.shape, lambda i: tvm.te.exp(A[i]), name="B") diff --git a/tests/python/unittest/test_link_params.py b/tests/python/unittest/test_link_params.py index 80ea11f6d9aa..ffe859927ad7 100644 --- a/tests/python/unittest/test_link_params.py +++ b/tests/python/unittest/test_link_params.py @@ -21,6 +21,7 @@ import re import struct import sys +import tempfile import numpy as np import pytest @@ -182,31 +183,38 @@ def _add_decl(name, dtype): @tvm.testing.requires_llvm def test_llvm_link_params(): for dtype in LINKABLE_DTYPES: - mod, param_init = _make_mod_and_params(dtype) + ir_mod, param_init = _make_mod_and_params(dtype) rand_input = _make_random_tensor(dtype, INPUT_SHAPE) - main_func = mod["main"] + main_func = ir_mod["main"] target = "llvm --runtime=c --system-lib --link-params" with tvm.transform.PassContext(opt_level=3): - lib = tvm.relay.build(mod, target, params=param_init) + lib = tvm.relay.build(ir_mod, target, params=param_init) + + # NOTE: Need to export_library() and load_library() to link all the Module(llvm, ...) + # against one another. + temp_dir = tempfile.mkdtemp() + export_file = os.path.join(temp_dir, "lib.so") + lib.lib.export_library(export_file) + mod = tvm.runtime.load_module(export_file) assert set(lib.params.keys()) == {"p0", "p1"} # NOTE: op folded + assert mod.get_function("TVMSystemLibEntryPoint") != None - print("graph", lib.graph_json) graph = json.loads(lib.graph_json) for p in lib.params: - _verify_linked_param(dtype, lib, lib.lib, graph, p) or found_one + _verify_linked_param(dtype, lib, mod, graph, p) or found_one # Wrap in function to explicitly deallocate the runtime. - def _run_linked(lib): - graph_json, mod, _ = lib + def _run_linked(lib, mod): + graph_json, _, _ = lib graph_rt = tvm.contrib.graph_runtime.create(graph_json, mod, tvm.cpu(0)) graph_rt.set_input("rand_input", rand_input) # NOTE: params not required. graph_rt.run() return graph_rt.get_output(0) - linked_output = _run_linked(lib) + linked_output = _run_linked(lib, mod) with tvm.transform.PassContext(opt_level=3): - lib = tvm.relay.build(mod, "llvm --system-lib", params=param_init) + lib = tvm.relay.build(ir_mod, "llvm --system-lib", params=param_init) def _run_unlinked(lib): graph_json, mod, lowered_params = lib @@ -266,8 +274,8 @@ def test_c_link_params(): lib = tvm.relay.build(mod, target, params=param_init) assert set(lib.params.keys()) == {"p0", "p1"} # NOTE: op folded - src = lib.lib.imported_modules[0].get_source() - lib.lib.save("test.c", "cc") + src = lib.lib.get_source() + lib.lib.save("test.c", "c") c_dtype = _get_c_datatype(dtype) src_lines = src.split("\n") param = lib.params["p0"].asnumpy().reshape(np.prod(KERNEL_SHAPE)) diff --git a/tests/python/unittest/test_target_codegen_llvm.py b/tests/python/unittest/test_target_codegen_llvm.py index 67c1f6bff429..ec7c5aea333f 100644 --- a/tests/python/unittest/test_target_codegen_llvm.py +++ b/tests/python/unittest/test_target_codegen_llvm.py @@ -17,6 +17,8 @@ import collections import ctypes import json +import sys + import tvm import tvm.testing from tvm import te @@ -26,6 +28,7 @@ import ctypes import math import re +import pytest @tvm.testing.requires_llvm @@ -816,27 +819,4 @@ def do_atomic_add(A): if __name__ == "__main__": - test_multiple_func() - test_llvm_large_uintimm() - test_llvm_import() - test_alignment() - test_rank_zero() - test_rank_zero_bound_checkers() - test_llvm_bool() - test_llvm_persist_parallel() - test_llvm_condition() - test_llvm_vadd_pipeline() - test_llvm_add_pipeline() - test_llvm_intrin() - test_llvm_overloaded_intrin() - test_llvm_flip_pipeline() - test_llvm_madd_pipeline() - test_llvm_temp_space() - test_llvm_lookup_intrin() - test_llvm_div() - test_llvm_fp_math() - test_dwarf_debug_information() - test_llvm_shuffle() - test_llvm_bf16() - test_llvm_crt_static_lib() - test_llvm_gpu_lower_atomic() + sys.exit(pytest.main([__file__] + sys.argv[1:]))