[Refactor] Migrate build API to tvm.compile

apps/android_rpc/tests/android_rpc_test.py

@@ -64,7 +64,7 @@ def test_rpc_module():
     sch.bind(xi, "threadIdx.x")
 
     if test_opencl:
-        f = tvm.build(sch.mod, target=tvm.target.Target("opencl", host=target))
+        f = tvm.compile(sch.mod, target=tvm.target.Target("opencl", host=target))
         path_dso_cl = temp.relpath("dev_lib_cl.so")
         f.export_library(path_dso_cl, fcompile=ndk.create_shared)
 

apps/ios_rpc/tests/ios_rpc_test.py

@@ -59,7 +59,7 @@ def test_rpc_module(host, port, key, mode):
 
     # Build the dynamic lib.
     # If we don't want to do metal and only use cpu, just set target to be target
-    f = tvm.build(sch.mod, target=tvm.target.Target("metal", host=target))
+    f = tvm.compile(sch.mod, target=tvm.target.Target("metal", host=target))
     path_dso1 = temp.relpath("dev_lib.dylib")
     f.export_library(path_dso1, fcompile=xcode.create_dylib, arch=arch, sdk=sdk)
 

docs/deep_dive/tensor_ir/tutorials/tir_creation.py

@@ -212,7 +212,7 @@ def evaluate_dynamic_shape(lib: tvm.runtime.Module, m: int, n: int, k: int):
 
 
 # Compile lib only once
-dyn_shape_lib = tvm.build(DynamicShapeModule, target="llvm")
+dyn_shape_lib = tvm.compile(DynamicShapeModule, target="llvm")
 # Able to handle different shapes
 print(evaluate_dynamic_shape(dyn_shape_lib, m=4, n=4, k=4))
 print(evaluate_dynamic_shape(dyn_shape_lib, m=64, n=64, k=128))

docs/deep_dive/tensor_ir/tutorials/tir_transformation.py

@@ -78,7 +78,7 @@ def main(
 
 
 def evaluate(mod: tvm.IRModule):
-    lib = tvm.build(mod, target="llvm")
+    lib = tvm.tir.build(mod, target="llvm")
     # check correctness
     lib(a_nd, b_nd, c_nd)
     np.testing.assert_allclose(c_nd.numpy(), c_np, rtol=1e-5)

docs/get_started/tutorials/ir_module.py

@@ -232,7 +232,7 @@ def main(
 # ~~~~~~~~~~~~~
 # We can deploy the IRModule on CPU by specifying the target as ``llvm``.
 
-exec = relax.build(mod, target="llvm")
+exec = tvm.compile(mod, target="llvm")
 dev = tvm.cpu()
 vm = relax.VirtualMachine(exec, dev)
 
@@ -263,7 +263,7 @@ def main(
 ######################################################################
 # Now we can compile the IRModule on GPU, the similar way as we did on CPU.
 
-exec = relax.build(gpu_mod, target="cuda")
+exec = tvm.compile(gpu_mod, target="cuda")
 dev = tvm.device("cuda", 0)
 vm = relax.VirtualMachine(exec, dev)
 # Need to allocate data and params on GPU device

docs/get_started/tutorials/quick_start.py

@@ -137,7 +137,7 @@ def forward(self, x):
 import numpy as np
 
 target = tvm.target.Target("llvm")
-ex = relax.build(mod, target)
+ex = tvm.compile(mod, target)
 device = tvm.cpu()
 vm = relax.VirtualMachine(ex, device)
 data = np.random.rand(1, 784).astype("float32")

docs/how_to/tutorials/cross_compilation_and_rpc.py

@@ -119,7 +119,7 @@
 else:
     target = "llvm -mtriple=armv7l-linux-gnueabihf"
 
-func = tvm.build(mod, target=target)
+func = tvm.compile(mod, target=target)
 # save the lib at a local temp folder
 temp = utils.tempdir()
 path = temp.relpath("lib.tar")
@@ -237,7 +237,7 @@ def run_opencl():
     xo, xi = sch.split(i, [None, 32])
     sch.bind(x, "blockIdx.x")
     sch.bind(x, "threadIdx.x")
-    func = tvm.build(sch.mod, target=target)
+    func = tvm.compile(sch.mod, target=target)
 
     remote = rpc.connect(opencl_device_host, opencl_device_port)
 

docs/how_to/tutorials/customize_opt.py

@@ -205,7 +205,7 @@ def transform_module(self, mod: IRModule, _ctx: tvm.transform.PassContext) -> IR
 # --------------------------
 # We can build and deploy the optimized model to the TVM runtime.
 
-ex = relax.build(mod, target="cuda")
+ex = tvm.compile(mod, target="cuda")
 dev = tvm.device("cuda", 0)
 vm = relax.VirtualMachine(ex, dev)
 # Need to allocate data and params on GPU device

docs/how_to/tutorials/e2e_opt_model.py

@@ -113,7 +113,7 @@
 # We skip this step in the CI environment.
 
 if not IS_IN_CI:
-    ex = relax.build(mod, target="cuda")
+    ex = tvm.compile(mod, target="cuda")
     dev = tvm.device("cuda", 0)
     vm = relax.VirtualMachine(ex, dev)
     # Need to allocate data and params on GPU device

docs/reference/api/python/driver.rst

@@ -19,4 +19,4 @@ tvm.driver
 ----------
 .. automodule:: tvm.driver
 
-.. autofunction:: tvm.build
+.. autofunction:: tvm.compile

include/tvm/runtime/profiling.h

@@ -504,7 +504,7 @@ String ShapeString(NDArray shape, DLDataType dtype);
 String ShapeString(const std::vector<int64_t>& shape, DLDataType dtype);
 
 /*! \brief Collect performance information of a function execution. Usually
- * used with a compiled PrimFunc (via tvm.build).
+ * used with a compiled PrimFunc (via tvm.compile).
  *
  * This information can include performance counters like cache hits and FLOPs
  * that are useful in debugging performance issues of individual PrimFuncs.

python/tvm/contrib/msc/core/runtime/runner.py

@@ -1453,14 +1453,14 @@ def _build_runnable(self, model: Any) -> Any:
         if self._device == "cpu":
             target = tvm.target.Target("llvm")
             with tvm.transform.PassContext(opt_level=3):
-                self._executable = tvm.relax.build(model, target)
+                self._executable = tvm.compile(model, target)
                 runnable = tvm.relax.VirtualMachine(self._executable, tvm.cpu())
         elif self._device.startswith("cuda"):
             target = tvm.target.Target("cuda")
             with target:
                 model = tvm.tir.transform.DefaultGPUSchedule()(model)
             with tvm.transform.PassContext(opt_level=3):
-                self._executable = tvm.relax.build(model, target)
+                self._executable = tvm.compile(model, target)
                 runnable = tvm.relax.VirtualMachine(self._executable, tvm.cuda())
         else:
             raise NotImplementedError("Unsupported device " + str(self._device))

python/tvm/contrib/msc/framework/tvm/runtime/runner.py

@@ -101,14 +101,14 @@ def _build_runnable(self, model: Any) -> Any:
             if self._device.startswith("cpu"):
                 target = tvm.target.Target("llvm")
                 with tvm.transform.PassContext(opt_level=3):
-                    self._executable = tvm.relax.build(model, target)
+                    self._executable = tvm.compile(model, target)
                     runnable = tvm.relax.VirtualMachine(self._executable, tvm.cpu())
             elif self._device.startswith("cuda"):
                 target = tvm.target.Target("cuda")
                 with target:
                     model = tvm.tir.transform.DefaultGPUSchedule()(model)
                 with tvm.transform.PassContext(opt_level=3):
-                    self._executable = tvm.relax.build(model, target)
+                    self._executable = tvm.compile(model, target)
                     runnable = tvm.relax.VirtualMachine(self._executable, tvm.cuda())
             else:
                 raise NotImplementedError("Unsupported device " + str(self._device))
@@ -248,13 +248,13 @@ def run_native(
             with target:
                 model = tvm.tir.transform.DefaultGPUSchedule()(model)
             with tvm.transform.PassContext(opt_level=3):
-                relax_exec = tvm.relax.build(model, target)
+                relax_exec = tvm.compile(model, target)
                 runnable = tvm.relax.VirtualMachine(relax_exec, tvm.cuda())
             tvm_inputs = [tvm.nd.array(inputs[i], device=tvm.cuda()) for i in input_names]
         else:
             target = tvm.target.Target("llvm")
             with tvm.transform.PassContext(opt_level=3):
-                relax_exec = tvm.relax.build(model, target)
+                relax_exec = tvm.compile(model, target)
                 runnable = tvm.relax.VirtualMachine(relax_exec, tvm.cpu())
             tvm_inputs = [tvm.nd.array(inputs[i]) for i in input_names]
 

python/tvm/dlight/benchmark/bench.py

@@ -121,7 +121,7 @@ def benchmark(
     # append scalar input tensors for rotary embedding
     input_tensors.extend(scalar_input_tensors)
     # build locally
-    rt_mod = tvm.build(mod, target=target)
+    rt_mod = tvm.tir.build(mod, target=target)
     # set up evaluator config
     evaluator_config = EvaluatorConfig._normalized(  # pylint: disable=protected-access
         evaluator_config

python/tvm/exec/gpu_memory_bandwidth.py

@@ -220,7 +220,7 @@ def main():  # pylint: disable=too-many-locals
 
         if rpcConfig is None:
             _, profile_result = local_run(
-                tvm.build(sch.mod, target=target),
+                tvm.compile(sch.mod, target=target),
                 target.kind.name,
                 [a, b],
                 evaluator_config=EvaluatorConfig(
@@ -232,7 +232,7 @@ def main():  # pylint: disable=too-many-locals
             )
         else:
             _, profile_result = rpc_run(
-                tvm.build(sch.mod, target=target),
+                tvm.compile(sch.mod, target=target),
                 target.kind.name,
                 [a, b],
                 evaluator_config=EvaluatorConfig(

python/tvm/ir/instrument.py

@@ -213,7 +213,7 @@ def should_run(self, mod, pass_info)
 
         skip_annotate = SkipPass("AnnotateSpans")
         with tvm.transform.PassContext(instruments=[skip_annotate]):
-            tvm.build(mod, "llvm")
+            tvm.compile(mod, "llvm")
     """
 
     def create_pass_instrument(pi_cls):

python/tvm/meta_schedule/testing/validate_database.py

@@ -554,7 +554,7 @@ def local_build_and_run(
         The running time of running the module
     """
     # potential memory leak https://github.com/apache/tvm/issues/11096
-    lib = tvm.build(mod, target=target)
+    lib = tvm.compile(mod, target=target)
     tvm_inputs = [tvm.nd.array(inp, device=device) for inp in inputs]
     device.sync()
     func = lib.time_evaluator(lib.entry_name, dev=device, number=ARGS.number, repeat=ARGS.repeat)

python/tvm/relax/frontend/nn/extern.py

@@ -156,7 +156,7 @@ def shape_dtype_inference(a, b):
     **A compiler pass `AttachExternModules`.** It is introduced to attach a list of
     `nn.ExternModule`s into an IRModule at any stage of the compilation pipeline,
     and attach the compiled external modules as `runtime.Module`s into IRModule's `external_mods`
-    attribute. It is required by linking in `relax.build`, but with the existence of this pass,
+    attribute. It is required by linking in `tvm.compile`, but with the existence of this pass,
     source compilation can be deferred to arbitrary stage of TVM compilation.
 
     **Caveats.** It is required to call `nn.add_extern` to register external modules exactly once

python/tvm/relax/pipeline.py

@@ -77,7 +77,7 @@ def f_zero_pipeline(mod: tvm.ir.IRModule, ctx: tvm.transform.PassContext) -> tvm
 
 
 def default_build_pipeline():
-    """The default compilation pipeline used in relax.build"""
+    """The default compilation pipeline used in tvm.compile"""
 
     @tvm.transform.module_pass(opt_level=0)
     def _pipeline(mod: tvm.ir.IRModule, _ctx: tvm.transform.PassContext) -> tvm.ir.IRModule:
@@ -144,7 +144,7 @@ def static_shape_tuning_pipeline(
             cpu_weight_prepack=True,
         )(mod)
 
-        ex = relax.build(mod, target=target)
+        ex = tvm.compile(mod, target=target)
         vm = relax.VirtualMachine(ex, device=tvm.cpu())
 
         # Transform the params using the vm function

python/tvm/relax/training/trainer.py

@@ -57,7 +57,7 @@ class Trainer:
             [pred_sinfo, target_sinfo],
         )
         train_mod = setup_trainer(Backbone)
-        ex = relax.build(train_mod, target)
+        ex = tvm.compile(train_mod, target)
         vm = relax.VirtualMachine(ex, dev)
 
         trainer = training.Trainer(train_mod, vm, dev, False)

python/tvm/relax/transform/tuning_api/default_functions.py

@@ -176,7 +176,7 @@ def relax_build(
         ):
             if params:
                 mod = tvm.relax.transform.BindParams("main", params)(mod)
-            relax_exec = tvm.relax.build(mod, target)
+            relax_exec = tvm.compile(mod, target)
             return relax_exec.mod
 
         builder = LocalBuilder(f_build=relax_build)

python/tvm/relax/vm_build.py

@@ -229,7 +229,7 @@ def foo(x: Tensor((3, 4), "float32"), y: Tensor((3, 4), "float32")):
 
         mod = InputModule
         target = tvm.target.Target("llvm", host="llvm")
-        ex = relax.build(mod, target)
+        ex = tvm.compile(mod, target)
     """
 
     def _extract_attrs(mod: tvm.IRModule):

python/tvm/runtime/profiling/__init__.py

@@ -231,7 +231,7 @@ def profile_function(mod, dev, collectors, func_name=None, warmup_iters=10):
 
     .. code-block: python
 
-        f = tvm.build(my_func, target="llvm", name="my_func")
+        f = tvm.compile(my_func, target="llvm", name="my_func")
         prof = tvm.runtime.profiling.profile_function(
             f,
             tvm.cpu(),

python/tvm/runtime/relax_vm.py

@@ -431,7 +431,7 @@ def time_evaluator(
         .. code-block:: python
 
             target = tvm.target.Target("llvm", host="llvm")
-            ex = relax.build(TestTimeEvaluator, target)
+            ex = tvm.compile(TestTimeEvaluator, target)
             vm = relax.VirtualMachine(mod, tvm.cpu())
             timing_res = vm.time_evaluator("func_name", tvm.cpu())(arg0, arg1, ..., argn)
 
@@ -440,7 +440,7 @@ def time_evaluator(
         .. code-block:: python
 
             target = tvm.target.Target("llvm", host="llvm")
-            ex = relax.build(TestTimeEvaluator, target)
+            ex = tvm.compile(TestTimeEvaluator, target)
             vm = relax.VirtualMachine(mod, tvm.cpu())
             vm.set_input("func_name", arg0, arg1, ..., argn)
             timing_res = vm.time_evaluator("invoke_stateful", tvm.cpu())("func_name")
@@ -451,7 +451,7 @@ def time_evaluator(
         .. code-block:: python
 
             target = tvm.target.Target("llvm", host="llvm")
-            ex = relax.build(TestTimeEvaluator, target)
+            ex = tvm.compile(TestTimeEvaluator, target)
             vm = relax.VirtualMachine(mod, tvm.cpu())
             vm.save_function("func_name", "func_name_saved", arg0, arg1, ..., argn)
             timing_res = vm.time_evaluator("func_name_saved", tvm.cpu())()

python/tvm/testing/utils.py

@@ -327,7 +327,7 @@ def _compute_body(*us):
 
         A = tvm.te.compute([r.extent.value for v, r in vranges.items()], _compute_body)
         args = [tvm.nd.empty(A.shape, A.dtype)]
-        mod = tvm.build(tvm.IRModule.from_expr(tvm.te.create_prim_func([A])))
+        mod = tvm.compile(tvm.IRModule.from_expr(tvm.te.create_prim_func([A])))
         mod(*args)
         return args[0].numpy()
 

python/tvm/topi/sort.py

@@ -103,7 +103,7 @@ def argsort(data, valid_count=None, axis=-1, is_ascend=1, dtype="float32"):
         out = argsort(data, axis=axis, is_ascend=is_ascend)
         np_data = np.random.uniform(dshape)
         s = topi.generic.schedule_argsort(out)
-        f = tvm.build(s, [data, out], "llvm")
+        f = tvm.compile(s, [data, out], "llvm")
         dev = tvm.cpu()
         tvm_data = tvm.nd.array(np_data, dev)
         tvm_out = tvm.nd.array(np.zeros(dshape, dtype=data.dtype), dev)

src/target/source/codegen_c_host.cc

@@ -31,10 +31,6 @@
 #include <utility>
 #include <vector>
 
-#include "../../support/str_escape.h"
-#include "../build_common.h"
-#include "codegen_params.h"
-
 namespace tvm {
 namespace codegen {
 
@@ -51,6 +47,7 @@ void CodeGenCHost::Init(bool output_ssa, bool emit_asserts, bool emit_fwd_func_d
   decl_stream << "#include \"tvm/runtime/c_backend_api.h\"\n";
   decl_stream << "#include <math.h>\n";
   decl_stream << "#include <stdbool.h>\n";
+  CodeGenCHost::InitGlobalContext();
   CodeGenC::Init(output_ssa);
 }
 
@@ -92,7 +89,6 @@ void CodeGenCHost::AddFunction(const GlobalVar& gvar, const PrimFunc& func,
 }
 
 void CodeGenCHost::GenerateForwardFunctionDeclarations(String global_symbol,
-
                                                        const Array<Type>& arg_types,
                                                        const Type& ret_type) {
   if (!emit_fwd_func_decl_) {
@@ -443,9 +439,6 @@ runtime::Module BuildCHost(IRModule mod, Target target) {
     return sort_key(kv_a) < sort_key(kv_b);
   });
 
-  // Declare all functions first.  This ensures that all functions,
-  // including the __tvm_main__ used in AOT, have access to forward
-  // declarations of other functions in the IRModule.
   for (const auto& [gvar, prim_func] : funcs) {
     cg.DeclareFunction(gvar, prim_func);
   }
@@ -457,11 +450,6 @@ runtime::Module BuildCHost(IRModule mod, Target target) {
     cg.AddFunction(gvar, prim_func, emit_fwd_func_decl);
   }
 
-  if (target->GetAttr<Bool>("system-lib").value_or(Bool(false))) {
-    ICHECK_EQ(target->GetAttr<String>("runtime").value_or(""), "c")
-        << "c target only supports generating C runtime SystemLibs";
-  }
-
   std::string code = cg.Finish();
   return CSourceModuleCreate(code, "c", cg.GetFunctionNames());
 }

tests/python/all-platform-minimal-test/test_minimal_target_codegen_llvm.py

@@ -46,7 +46,7 @@ def test_llvm_add_pipeline():
 
     def check_llvm():
         # BUILD and invoke the kernel.
-        f = tvm.build(sch.mod, target="llvm")
+        f = tvm.compile(sch.mod, target="llvm")
         dev = tvm.cpu(0)
         # launch the kernel.
         n = nn