From 78ae3f9faee111ef594200a9e404e5078d4a9c54 Mon Sep 17 00:00:00 2001 From: tqchen Date: Sat, 4 Apr 2020 09:34:30 -0700 Subject: [PATCH] Remove LoweredFunc. --- apps/lldb/tvm.py | 1 - docs/dev/codebase_walkthrough.rst | 9 -- include/tvm/driver/driver_api.h | 3 +- include/tvm/target/codegen.h | 12 -- include/tvm/tir/ir_pass.h | 1 - include/tvm/tir/lowered_func.h | 149 ------------------ python/tvm/driver/build_module.py | 9 +- python/tvm/relay/backend/_backend.py | 12 +- .../relay/backend/graph_runtime_codegen.py | 6 +- python/tvm/target/build_config.py | 10 +- python/tvm/tir/__init__.py | 2 +- python/tvm/tir/stmt.py | 8 - src/contrib/hybrid/codegen_hybrid.h | 1 - src/driver/driver_api.cc | 1 - src/relay/backend/build_module.cc | 10 +- src/relay/backend/compile_engine.h | 1 - src/relay/backend/graph_runtime_codegen.cc | 2 +- src/relay/backend/vm/compiler.cc | 2 - src/relay/transforms/gradient.cc | 1 - src/target/build_common.h | 17 -- src/target/codegen.cc | 47 ------ src/target/llvm/codegen_llvm.cc | 2 +- src/target/source/codegen_c.h | 1 - src/target/spirv/codegen_spirv.h | 1 - src/target/stackvm/codegen_stackvm.h | 1 - src/tir/ir/lowered_func.cc | 37 ----- src/tir/pass/ffi_api.cc | 7 - src/tir/pass/storage_rewrite.cc | 23 --- src/tir/transforms/split_host_device.cc | 1 - tutorials/dev/low_level_custom_pass.py | 2 +- 30 files changed, 23 insertions(+), 356 deletions(-) delete mode 100644 include/tvm/tir/lowered_func.h delete mode 100644 src/tir/ir/lowered_func.cc diff --git a/apps/lldb/tvm.py b/apps/lldb/tvm.py index 811d32db6c75b..135aeff5258aa 100644 --- a/apps/lldb/tvm.py +++ b/apps/lldb/tvm.py @@ -46,7 +46,6 @@ def __lldb_init_module(debugger, _): "tvm::IterVarAttr", "tvm::IterVarRelation", "tvm::Layout", - "tir::LoweredFunc", "tvm::Map", "tvm::Map", "tvm::MemoryInfo", diff --git a/docs/dev/codebase_walkthrough.rst b/docs/dev/codebase_walkthrough.rst index b7eb06b32df64..a66328fef7c92 100644 --- a/docs/dev/codebase_walkthrough.rst +++ b/docs/dev/codebase_walkthrough.rst @@ -145,15 +145,6 @@ After lowering is done, ``build()`` function generates target machine code from Code generation is done by ``build_module()`` function, defined in ``python/tvm/target/codegen.py``. On the C++ side, code generation is implemented in ``src/target/codegen`` subdirectory. ``build_module()`` Python function will reach ``Build()`` function below in ``src/target/codegen/codegen.cc``: -:: - - runtime::Module Build(const Array& funcs, - const std::string& target) { - std::string build_f_name = "codegen.build_" + target; - const PackedFunc* bf = runtime::Registry::Get(build_f_name); - runtime::Module m = (*bf)(funcs, target); - return m; - } The ``Build()`` function looks up the code generator for the given target in the ``PackedFunc`` registry, and invokes the function found. For example, ``codegen.build_cuda`` function is registered in ``src/codegen/build_cuda_on.cc``, like this: diff --git a/include/tvm/driver/driver_api.h b/include/tvm/driver/driver_api.h index cec1b2a049464..e6d4427544465 100644 --- a/include/tvm/driver/driver_api.h +++ b/include/tvm/driver/driver_api.h @@ -34,7 +34,6 @@ #include #include #include -#include #include #include @@ -44,7 +43,7 @@ namespace tvm { /*! -* \brief Build a LoweredFunc given a schedule, args and binds +* \brief Build an IRModule given a schedule, args and binds * \param sch The schedule to lower. * \param args The arguments to the function. * \param name The name of the lowered function. diff --git a/include/tvm/target/codegen.h b/include/tvm/target/codegen.h index c604eb5c93de4..4b7ea56e705de 100644 --- a/include/tvm/target/codegen.h +++ b/include/tvm/target/codegen.h @@ -27,7 +27,6 @@ #include #include #include -#include #include #include @@ -41,17 +40,6 @@ using runtime::PackedFunc; using runtime::TVMArgs; using runtime::TVMRetValue; -/*! - * \brief Temporary backward compatible function to convert a list - * of LoweredFunc to a IRModule of PrimfFuncs - * \param funcs The input lowered function. - * \return The IRModule. - * - * \note This function is only used for code refactor and will be - * removed once the refactor completes. - */ -IRModule ToIRModule(const Array& funcs); - /*! * \brief Build a module from array of lowered function. * \param mod The Module to be built diff --git a/include/tvm/tir/ir_pass.h b/include/tvm/tir/ir_pass.h index 0a89c04d1b7cc..e228ce32adab1 100644 --- a/include/tvm/tir/ir_pass.h +++ b/include/tvm/tir/ir_pass.h @@ -31,7 +31,6 @@ #include #include #include -#include #include #include diff --git a/include/tvm/tir/lowered_func.h b/include/tvm/tir/lowered_func.h deleted file mode 100644 index 2d01c8958aefd..0000000000000 --- a/include/tvm/tir/lowered_func.h +++ /dev/null @@ -1,149 +0,0 @@ -/* - * 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 tvm/tir/lowered_func.h - * \brief Information about a lowered TVM function. - * This data structure is final step toward codegen. - */ -#ifndef TVM_TIR_LOWERED_FUNC_H_ -#define TVM_TIR_LOWERED_FUNC_H_ - -#include -#include -#include -#include - -namespace tvm { -namespace tir { - -// Internal node container of lowered function. -class LoweredFuncNode; - -/*! - * \brief LoweredFunc represents function after lowering. - * This is the final IR representation before codegen. - */ -class LoweredFunc : public FunctionRef { - public: - LoweredFunc() {} - explicit LoweredFunc(ObjectPtr n) : FunctionRef(n) {} - /*! - * \brief access the internal node container - * \return the pointer to the internal node container - */ - inline const LoweredFuncNode* operator->() const; - /*! \brief specify container node */ - using ContainerType = LoweredFuncNode; -}; - -/*! \brief specific type of lowered function */ -enum LoweredFuncType : int { - /*! \brief Function that can mix device and host calls */ - kMixedFunc = 0, - /*! \brief Only contains host code */ - kHostFunc = 1, - /*! \brief Only contains device code */ - kDeviceFunc = 2 -}; - -/*! \brief Node container of LoweredFunc */ -class LoweredFuncNode : public tir::FunctionBaseNode { - public: - /*! \brief The name of the function */ - std::string name; - /*! - * \brief The arguments of the function - * This function can only take pod type(int, float) and void* as arguments. - */ - Array args; - /*! - * \brief The IterVar axis of threads - * Each axis need host function to specify a size. - * \note Calling convention into LoweredFunc - * - * Assume we have a LoweredFunc f, a call into f - * Call(f, arg1, arg2, ..., arg_n, - * size_axis_1, size_axis_2, ... size_axis_m) - * - * Here n = len(args), m = len(thread_axis) - * - * The CodeGen should take this and translate this call - * to corresponding API specific kernel launchs or function calls. - */ - Array thread_axis; - /*! - * \brief The hint data type of Var handles defined in LetStmt - * Can be used as hint when generating type signiture. - * The creation rule is given by - * handle_data_type[var_handle] = make_const(the_type, 0); - * - * \note Expr is used instead Type, because Type cannot be hold by Map. - * constant Expr of given type is used. - */ - Map handle_data_type; - /*! \brief The type of the function */ - LoweredFuncType func_type{kMixedFunc}; - /*! \brief Whether this function is packed function */ - bool is_packed_func{true}; - /*! - * \brief Whether function ensures that argument pointers do not alias. - * This corresponds to restrict keyword in C. - */ - bool is_restricted{true}; - /*! \brief The body statment of the function */ - Stmt body; - /*! \return name of the operation */ - const std::string& func_name() const final { - return name; - } - // there is no return value, but return 1 - // to enable Call into this function. - int num_outputs() const final { - return 1; - } - void VisitAttrs(AttrVisitor* v) { - v->Visit("name", &name); - v->Visit("args", &args); - v->Visit("thread_axis", &thread_axis); - v->Visit("handle_data_type", &handle_data_type); - v->Visit("func_type", &func_type); - v->Visit("is_packed_func", &is_packed_func); - v->Visit("is_restricted", &is_restricted); - v->Visit("body", &body); - } - - static constexpr const char* _type_key = "LoweredFunc"; - TVM_DECLARE_FINAL_OBJECT_INFO(LoweredFuncNode, Object); -}; - -// Implementations of inline functions -inline const LoweredFuncNode* LoweredFunc::operator->() const { - return static_cast(get()); -} -} // namespace tir -} // namespace tvm - -namespace std { -template <> -struct hash<::tvm::tir::LoweredFunc> : public tvm::ObjectHash { -}; -} - -#endif // TVM_TIR_LOWERED_FUNC_H_ diff --git a/python/tvm/driver/build_module.py b/python/tvm/driver/build_module.py index a30bec29e152c..0dd01e1860344 100644 --- a/python/tvm/driver/build_module.py +++ b/python/tvm/driver/build_module.py @@ -371,9 +371,7 @@ def build(inputs, elif isinstance(inputs, tvm.IRModule): input_mod = inputs elif not isinstance(inputs, (dict, container.Map)): - raise ValueError("inputs must be Schedule, LoweredFunc, list of " - "LoweredFunc, or dict of target to list of " - "LoweredFunc.") + raise ValueError("inputs must be Schedule, IRModule or dict of target to IRModule") if not isinstance(inputs, (dict, container.Map)): target = _target.Target.current() if target is None else target @@ -387,9 +385,8 @@ def build(inputs, raise ValueError("The key of inputs must be str or " "_target.Target when inputs is dict.") if not isinstance(mod, tvm.IRModule): - raise ValueError("inputs must be Schedule, LoweredFunc, IRModule," - "or dict of str to list of " - "LoweredFunc.") + raise ValueError("inputs must be Schedule, IRModule," + "or dict of str to IRModule.") if not target_host: for tar, _ in target_input_mod.items(): diff --git a/python/tvm/relay/backend/_backend.py b/python/tvm/relay/backend/_backend.py index 70a89c2518559..641ff04adab01 100644 --- a/python/tvm/relay/backend/_backend.py +++ b/python/tvm/relay/backend/_backend.py @@ -39,7 +39,7 @@ def lower(sch, inputs, func_name, source_func): Returns ------- - lowered_funcs : List[tvm.LoweredFunc] + mod : tvm.IRModule The result of lowering. """ # pylint: disable=broad-except, import-outside-toplevel @@ -59,15 +59,13 @@ def lower(sch, inputs, func_name, source_func): @tvm._ffi.register_func("relay.backend.build") -def build(funcs, target, target_host=None): +def build(mod, target, target_host=None): """Backend build function. Parameters ---------- - funcs : List[tvm.LoweredFunc] or Dict[str, List[tvm.LoweredFunc]] - A list of lowered functions or dictionary mapping from targets to - lowered functions. - + mod : tvm.IRModule or Dict[str, tvm.IRModule] + Input module target : tvm.Target The target to run the code on. @@ -82,7 +80,7 @@ def build(funcs, target, target_host=None): """ if target_host == "": target_host = None - return tvm.driver.build(funcs, target=target, target_host=target_host) + return tvm.driver.build(mod, target=target, target_host=target_host) @tvm._ffi.register_func("relay._tensor_value_repr") diff --git a/python/tvm/relay/backend/graph_runtime_codegen.py b/python/tvm/relay/backend/graph_runtime_codegen.py index 762210dbe4288..3e5f0157b32f5 100644 --- a/python/tvm/relay/backend/graph_runtime_codegen.py +++ b/python/tvm/relay/backend/graph_runtime_codegen.py @@ -48,7 +48,7 @@ def __init__(self, mod, target): self._get_graph_json = self._mod["get_graph_json"] self._list_params_name = self._mod["list_params_name"] self._get_param_by_name = self._mod["get_param_by_name"] - self._get_lowered_funcs = self._mod["get_lowered_funcs"] + self._get_irmodule = self._mod["get_irmodule"] self._setup(mod, target) def _setup(self, mod, target): @@ -74,14 +74,14 @@ def codegen(self, func): ------- graph_json : str The graph json that can be consumed by runtime. - lowered_funcs : List[tvm.LoweredFunc] or Dict[str, List[tvm.LoweredFunc]] + mod : IRModule or Dict[str, IRModule] The lowered functions. params : Dict[str, tvm.nd.NDArray] Additional constant parameters. """ self._codegen(func) graph_json = self._get_graph_json() - lowered_func = self._get_lowered_funcs() + lowered_func = self._get_irmodule() param_names = self._list_params_name() params = {} for name in param_names: diff --git a/python/tvm/target/build_config.py b/python/tvm/target/build_config.py index c105175d3e261..6a0dcf743a0df 100644 --- a/python/tvm/target/build_config.py +++ b/python/tvm/target/build_config.py @@ -20,9 +20,7 @@ import tvm.ir from tvm.runtime import Object -from tvm.ir import container from tvm.tir import Stmt -from tvm.tir.stmt import LoweredFunc from . import _ffi_api @@ -48,17 +46,13 @@ def decorate(self, func): def dump(*args, **kwargs): """dump function""" retv = func(*args, **kwargs) - if not isinstance(retv, (Stmt, LoweredFunc, container.Array)): + if not isinstance(retv, (Stmt,)): return retv fname = func.func_name if hasattr(func, 'func_name') else func.__name__ pname = str(self._pass_id) + "_" + fname + "_ir.cc" with open(pname, "a") as f: - out = retv.body if isinstance(retv, LoweredFunc) else retv + out = retv f.write(str(out)) - if isinstance(retv, container.Array): - for x in retv: - out = x.body if isinstance(x, LoweredFunc) else x - f.write("---------%s\n%s\n-----------\n"%(x.name, str(out))) self._pass_id += 1 return retv return dump diff --git a/python/tvm/tir/__init__.py b/python/tvm/tir/__init__.py index bd8e33fe4c3b2..b5d9fb147722e 100644 --- a/python/tvm/tir/__init__.py +++ b/python/tvm/tir/__init__.py @@ -29,7 +29,7 @@ from .stmt import Stmt, LetStmt, AssertStmt, ProducerConsumer, For from .stmt import BufferStore, Store, Provide, Allocate, AttrStmt, Free, Realize, SeqStmt -from .stmt import IfThenElse, Evaluate, Prefetch, LoweredFunc, stmt_seq, stmt_list +from .stmt import IfThenElse, Evaluate, Prefetch, stmt_seq, stmt_list from .function import PrimFunc diff --git a/python/tvm/tir/stmt.py b/python/tvm/tir/stmt.py index 0badad3c092f1..4531cdfc35ac3 100644 --- a/python/tvm/tir/stmt.py +++ b/python/tvm/tir/stmt.py @@ -385,14 +385,6 @@ def __init__(self, func, value_index, dtype, bounds): _ffi_api.Prefetch, func, value_index, dtype, bounds) -@tvm._ffi.register_object -class LoweredFunc(Object): - """Represent a LoweredFunc in TVM.""" - MixedFunc = 0 - HostFunc = 1 - DeviceFunc = 2 - - def stmt_seq(*args): """Make sequence of statements diff --git a/src/contrib/hybrid/codegen_hybrid.h b/src/contrib/hybrid/codegen_hybrid.h index 6491491ec2b3b..9784defcba884 100644 --- a/src/contrib/hybrid/codegen_hybrid.h +++ b/src/contrib/hybrid/codegen_hybrid.h @@ -27,7 +27,6 @@ #include #include #include -#include #include #include #include diff --git a/src/driver/driver_api.cc b/src/driver/driver_api.cc index 53c08ad59cee4..ae1d5393ad341 100644 --- a/src/driver/driver_api.cc +++ b/src/driver/driver_api.cc @@ -41,7 +41,6 @@ namespace tvm { using runtime::TVMArgs; using runtime::TVMRetValue; using runtime::PackedFunc; -using tir::LoweredFunc; bool LLVMEnabled() { const runtime::PackedFunc* pf = runtime::Registry::Get("target.build.llvm"); diff --git a/src/relay/backend/build_module.cc b/src/relay/backend/build_module.cc index 495922a4bd391..eaf78bc1b0f78 100644 --- a/src/relay/backend/build_module.cc +++ b/src/relay/backend/build_module.cc @@ -80,8 +80,8 @@ struct GraphCodegen { return CallFunc>("get_external_modules", nullptr); } - Map GetLoweredFunc() { - return CallFunc>("get_lowered_funcs", nullptr); + Map GetIRModule() { + return CallFunc>("get_irmodule", nullptr); } std::unordered_map GetParams() { @@ -151,9 +151,9 @@ class RelayBuildModule : public runtime::ModuleNode { this->SetParam(kv.first, kv.second->data); } }); - } else if (name == "get_lowered_funcs") { + } else if (name == "get_irmodule") { return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { - *rv = this->graph_codegen_->GetLoweredFunc(); + *rv = this->graph_codegen_->GetIRModule(); }); } else if (name == "get_external_modules") { return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { @@ -451,7 +451,7 @@ class RelayBuildModule : public runtime::ModuleNode { ret_.graph_json = graph_codegen_->GetJSON(); ret_.params = graph_codegen_->GetParams(); - auto lowered_funcs = graph_codegen_->GetLoweredFunc(); + auto lowered_funcs = graph_codegen_->GetIRModule(); // When there is no lowered_funcs due to reasons such as optimization. if (lowered_funcs.size() == 0) { diff --git a/src/relay/backend/compile_engine.h b/src/relay/backend/compile_engine.h index bb2d5ca4f96a7..4a3a04d02dcd1 100644 --- a/src/relay/backend/compile_engine.h +++ b/src/relay/backend/compile_engine.h @@ -27,7 +27,6 @@ #include #include -#include #include #include #include diff --git a/src/relay/backend/graph_runtime_codegen.cc b/src/relay/backend/graph_runtime_codegen.cc index 2f35669c898ae..c7f1be82c3710 100644 --- a/src/relay/backend/graph_runtime_codegen.cc +++ b/src/relay/backend/graph_runtime_codegen.cc @@ -646,7 +646,7 @@ class GraphRuntimeCodegenModule : public runtime::ModuleNode { CHECK_GT(this->output_.params.count(key), 0); *rv = this->output_.params[key]; }); - } else if (name == "get_lowered_funcs") { + } else if (name == "get_irmodule") { return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->output_.lowered_funcs; }); diff --git a/src/relay/backend/vm/compiler.cc b/src/relay/backend/vm/compiler.cc index 6fec6f2da3a9a..78ebb0fc5383c 100644 --- a/src/relay/backend/vm/compiler.cc +++ b/src/relay/backend/vm/compiler.cc @@ -962,8 +962,6 @@ void VMCompiler::PopulateGlobalMap() { } void VMCompiler::Codegen() { - using tir::LoweredFunc; - if (!context_.module.defined()) { LOG(WARNING) << "Did you forget to call VMCompiler::Lower?"; return; diff --git a/src/relay/transforms/gradient.cc b/src/relay/transforms/gradient.cc index a3728e9059224..d0ff169445fb2 100644 --- a/src/relay/transforms/gradient.cc +++ b/src/relay/transforms/gradient.cc @@ -22,7 +22,6 @@ * \brief API for Automatic Differentiation for the Relay IR. */ #include -#include #include #include #include diff --git a/src/target/build_common.h b/src/target/build_common.h index 47ec8f032c404..fc45cef3a874e 100644 --- a/src/target/build_common.h +++ b/src/target/build_common.h @@ -31,29 +31,12 @@ #include #include #include -#include #include #include #include "../runtime/meta_data.h" namespace tvm { namespace codegen { -// Extract function information from device function. -inline std::unordered_map -ExtractFuncInfo(const Array& funcs) { - std::unordered_map fmap; - for (tir::LoweredFunc f : funcs) { - runtime::FunctionInfo info; - for (size_t i = 0; i < f->args.size(); ++i) { - info.arg_types.push_back(f->args[i].dtype()); - } - for (size_t i = 0; i < f->thread_axis.size(); ++i) { - info.thread_axis_tags.push_back(f->thread_axis[i]->thread_tag); - } - fmap[f->name] = info; - } - return fmap; -} inline std::unordered_map ExtractFuncInfo(const IRModule& mod) { diff --git a/src/target/codegen.cc b/src/target/codegen.cc index 703328f8761f0..0eceea81da178 100644 --- a/src/target/codegen.cc +++ b/src/target/codegen.cc @@ -43,50 +43,6 @@ namespace tvm { namespace codegen { -// convert legacy LoweredFunc to PrimFunc. -tir::PrimFunc ToPrimFunc(tir::LoweredFunc from) { - // remap args to attach type annotations. - Array args; - Map remap_vars; - - for (auto var : from->args) { - auto it = from->handle_data_type.find(var); - if (it != from->handle_data_type.end()) { - tir::Var new_var(var->name_hint, - PointerType(PrimType((*it).second->dtype))); - args.push_back(new_var); - remap_vars.Set(var, new_var); - } else { - args.push_back(var); - } - } - tir::PrimFunc func(args, Substitute(from->body, remap_vars)); - - func = WithAttr(std::move(func), attr::kGlobalSymbol, runtime::String(from->name)); - func = WithAttr(std::move(func), tir::attr::kDeviceThreadAxis, from->thread_axis); - if (from->func_type == tir::LoweredFuncType::kDeviceFunc) { - func = WithAttr(std::move(func), - attr::kCallingConv, Integer(CallingConv::kDeviceKernelLaunch)); - } - if (from->is_restricted) { - func = WithAttr(std::move(func), tir::attr::kNoAlias, Integer(1)); - } - return func; -} - -IRModule ToIRModule(const Array& funcs) { - Map functions; - for (size_t i = 0; i < funcs.size(); ++i) { - auto f = funcs[i]; - tir::PrimFunc pf = ToPrimFunc(f); - if (i == 0) { - pf = WithAttr(std::move(pf), tir::attr::kIsEntryFunc, Integer(1)); - } - functions.Set(GlobalVar(f->name), pf); - } - return IRModule(functions); -} - runtime::Module Build(IRModule mod, const Target& target) { if (BuildConfig::Current()->disable_assert) { mod = tir::transform::SkipAssert()(mod); @@ -284,9 +240,6 @@ runtime::Module PackImportsToLLVM(const runtime::Module& mod, TVM_REGISTER_GLOBAL("target.Build") .set_body_typed(Build); -TVM_REGISTER_GLOBAL("testing.LoweredFuncsToIRModule") -.set_body_typed(ToIRModule); - // Export two auxiliary function to the runtime namespace. TVM_REGISTER_GLOBAL("runtime.ModulePackImportsToC") .set_body_typed(PackImportsToC); diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc index 31465cd56bcbe..450ebbcd02b84 100644 --- a/src/target/llvm/codegen_llvm.cc +++ b/src/target/llvm/codegen_llvm.cc @@ -448,7 +448,7 @@ CodeGenLLVM::CreateDebugInfo(llvm::Module* module) { auto debug_info = llvm::make_unique(); debug_info->di_builder_ = llvm::make_unique(*module); #endif - // TODO(tulloch): pass this information through relay::Span classes to the LoweredFunc instance? + // TODO(tulloch): pass this information through relay::Span classes to the IRModule instance? debug_info->file_ = debug_info->di_builder_->createFile("model.tvm", "/tmp/"); debug_info->compilation_unit_ = debug_info->di_builder_->createCompileUnit( llvm::dwarf::DW_LANG_C, debug_info->file_, "TVM", 0, "", 0, "", diff --git a/src/target/source/codegen_c.h b/src/target/source/codegen_c.h index c1894a379ddb6..30ad890c923d8 100644 --- a/src/target/source/codegen_c.h +++ b/src/target/source/codegen_c.h @@ -29,7 +29,6 @@ #include #include #include -#include #include #include #include diff --git a/src/target/spirv/codegen_spirv.h b/src/target/spirv/codegen_spirv.h index a5ccd549633db..edcee20f173f4 100644 --- a/src/target/spirv/codegen_spirv.h +++ b/src/target/spirv/codegen_spirv.h @@ -27,7 +27,6 @@ #include #include #include -#include #include #include diff --git a/src/target/stackvm/codegen_stackvm.h b/src/target/stackvm/codegen_stackvm.h index 041c7a7225cf6..fd370d285ea81 100644 --- a/src/target/stackvm/codegen_stackvm.h +++ b/src/target/stackvm/codegen_stackvm.h @@ -26,7 +26,6 @@ #include #include -#include #include #include #include diff --git a/src/tir/ir/lowered_func.cc b/src/tir/ir/lowered_func.cc deleted file mode 100644 index 8790f2b12e396..0000000000000 --- a/src/tir/ir/lowered_func.cc +++ /dev/null @@ -1,37 +0,0 @@ -/* - * 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 lowered_func.cc - */ -#include - -namespace tvm { -namespace tir { -TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable) -.set_dispatch([](const ObjectRef& node, ReprPrinter* p) { - auto* op = static_cast(node.get()); - p->stream << "LoweredFunc(" << op->name << ", " << op << ")"; -}); - -TVM_REGISTER_NODE_TYPE(LoweredFuncNode); - - -} // namespace tir -} // namespace tvm diff --git a/src/tir/pass/ffi_api.cc b/src/tir/pass/ffi_api.cc index 99a0a03099594..3083b68796352 100644 --- a/src/tir/pass/ffi_api.cc +++ b/src/tir/pass/ffi_api.cc @@ -105,13 +105,6 @@ TVM_REGISTER_GLOBAL("ir_pass.PostOrderVisit") }); }); -TVM_REGISTER_GLOBAL("ir_pass.LowerStorageAccess") -.set_body([](TVMArgs args, TVMRetValue *ret) { - LoweredFunc f = args[0]; - auto n = make_object(*f.operator->()); - n->body = LowerStorageAccessInfo(f->body); - *ret = LoweredFunc(n); -}); // make from two arguments #define REGISTER_PASS(PassName) \ diff --git a/src/tir/pass/storage_rewrite.cc b/src/tir/pass/storage_rewrite.cc index b4e6061a35d02..f3604b6403498 100644 --- a/src/tir/pass/storage_rewrite.cc +++ b/src/tir/pass/storage_rewrite.cc @@ -994,29 +994,6 @@ class VectorAllocRewriter : public StmtExprMutator { }; -LoweredFunc PointerValueTypeRewrite(LoweredFunc f) { - auto n = make_object(*f.operator->()); - VectorAllocRewriter rewriter; - n->body = rewriter(n->body); - for (Var arg : f->args) { - if (arg.dtype().is_handle()) { - const auto& tvec = rewriter.acc_map_[arg.get()]; - if (tvec.size() == 1) { - PrimExpr dtype = make_const(tvec[0], 0); - n->handle_data_type.Set(arg, dtype); - } else { - // always set data type to be non vectorized so - // load/store can still work via scalarization - if (tvec.size() != 0 && !n->handle_data_type.count(arg)) { - PrimExpr dtype = make_const(tvec[0].with_lanes(1), 0); - n->handle_data_type.Set(arg, dtype); - } - } - } - } - return LoweredFunc(n); -} - PrimFunc PointerValueTypeRewrite(PrimFunc f) { auto* n = f.CopyOnWrite(); VectorAllocRewriter rewriter; diff --git a/src/tir/transforms/split_host_device.cc b/src/tir/transforms/split_host_device.cc index 838ad82d974f9..ae32bdcbadeac 100644 --- a/src/tir/transforms/split_host_device.cc +++ b/src/tir/transforms/split_host_device.cc @@ -264,7 +264,6 @@ class HostDeviceSplitter : public StmtMutator { std::string name_prefix_; // Number of device functions. int device_func_counter_{0}; - std::vector device_funcs_; std::unordered_map handle_data_type_; }; diff --git a/tutorials/dev/low_level_custom_pass.py b/tutorials/dev/low_level_custom_pass.py index 298b24f6d046e..25ca279bf3399 100644 --- a/tutorials/dev/low_level_custom_pass.py +++ b/tutorials/dev/low_level_custom_pass.py @@ -36,7 +36,7 @@ - Visitor design pattern. Otherwise, check the `Python AST module `_ to see how an AST visitor is implemented. -- How a HalideIR/Schedule is lowered to either a LoweredFunc class or a LLVM module. Otherwise, +- How a Schedule is lowered to either an IRModule class or a LLVM module. Otherwise, take a look at ``python/tvm/build_module.py`` to get some basics. """