update

Author: Siyuan Feng

docs/deep_dive/tensor_ir/tutorials/tir_transformation.py

@@ -78,7 +78,7 @@ def main(
 
 
 def evaluate(mod: tvm.IRModule):
-    lib = tvm.compile(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)

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.compile(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

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/codegen/test_target_codegen_cuda_fp4.py

@@ -129,7 +129,7 @@ def add(
         sch.bind(tx, "threadIdx.x")
         sch.vectorize(vec)
 
-        fadd = tvm.build(sch.mod, target=target)
+        fadd = tvm.compile(sch.mod, target=target)
 
         numpytype = "float4_e2m1fn"
         promoted_base_dtype = promoted_dtype
@@ -189,7 +189,7 @@ def reinterpret(
     # Part 1. reinterpret float4_e2m1fn to uint8
     for vector_length in [1, 2, 4]:
         mod = get_reinterpret_mod("float4_e2m1fn", "uint8", vector_length)
-        f = tvm.build(mod, target=target)
+        f = tvm.compile(mod, target=target)
         a_np = np.random.uniform(low=-6, high=6, size=(n,)).astype("float4_e2m1fn")
         a = tvm.nd.empty(shape=(n,), dtype="float4_e2m1fn", device=dev)
         a.copyfrom(a_np)
@@ -200,7 +200,7 @@ def reinterpret(
     # Part 2. reinterpret uint8 to float4_e2m1fn
     for vector_length in [1, 2, 4]:
         mod = get_reinterpret_mod("uint8", "float4_e2m1fn", vector_length)
-        f = tvm.build(mod, target=target)
+        f = tvm.compile(mod, target=target)
         a_np = np.random.uniform(low=-6, high=6, size=(n,)).astype("uint8")
         a = tvm.nd.empty(shape=(n,), dtype="uint8", device=dev)
         a.copyfrom(a_np)

tests/python/codegen/test_target_codegen_opencl.py

@@ -38,7 +38,7 @@ def check_if_then_else(dev, n, dtype):
         sch = tvm.tir.Schedule(func)
         (x,) = sch.get_loops(sch.get_block("C"))
         sch.bind(x, "threadIdx.x")
-        fun = tvm.compile(sch.mod, target=target)
+        fun = tvm.tir.build(sch.mod, target=target)
         a = tvm.nd.empty((n,), A.dtype, dev)
         c = tvm.nd.empty((n,), A.dtype, dev)
         # Only need to test compiling here
@@ -55,7 +55,7 @@ def check_select(dev, n, dtype):
         sch = tvm.tir.Schedule(func)
         (x,) = sch.get_loops(sch.get_block("C"))
         sch.bind(x, "threadIdx.x")
-        fun = tvm.compile(sch.mod, target=target)
+        fun = tvm.tir.build(sch.mod, target=target)
 
         a = tvm.nd.empty((n,), A.dtype, dev)
         c = tvm.nd.empty((n,), A.dtype, dev)
@@ -85,7 +85,7 @@ def check_inf_nan(dev, n, value, dtype):
         sch = tvm.tir.Schedule(func)
         (x,) = sch.get_loops(sch.get_block("C"))
         sch.bind(x, "threadIdx.x")
-        fun = tvm.compile(sch.mod, target=target)
+        fun = tvm.tir.build(sch.mod, target=target)
         a = tvm.nd.empty((n,), A.dtype, dev)
         c = tvm.nd.empty((n,), A.dtype, dev)
         # Only need to test compiling here
@@ -113,7 +113,7 @@ def check_max(dev, n, dtype):
         sch = tvm.tir.Schedule(func)
         (x,) = sch.get_loops(sch.get_block("C"))
         sch.bind(x, "threadIdx.x")
-        fun = tvm.compile(sch.mod, target=target)
+        fun = tvm.tir.build(sch.mod, target=target)
 
         a = tvm.nd.empty((n,), A.dtype, dev)
         c = tvm.nd.empty((n,), A.dtype, dev)
@@ -178,7 +178,7 @@ def check_type_casting(ctx, n, dtype):
         sch.bind(tx, "threadIdx.x")
         sch.vectorize(vx)
 
-        fun = tvm.compile(sch.mod, target=target)
+        fun = tvm.tir.build(sch.mod, target=target)
         c = tvm.nd.empty((n,), dtype, ctx)
         assembly = fun.imported_modules[0].get_source()
         lcond = "convert_int4(((convert_uint4(((uint4)(((convert_int(get_local_id(0))) == 3), ((convert_int(get_local_id(0))) == 3), ((convert_int(get_local_id(0))) == 3), ((convert_int(get_local_id(0))) == 3)))))"
@@ -210,7 +210,7 @@ def _check(target, n, dtype):
             (x,) = sch.get_loops(sch.get_block("C"))
             sch.bind(x, "threadIdx.x")
 
-            fun = tvm.compile(sch.mod, target=target)
+            fun = tvm.tir.build(sch.mod, target=target)
             assembly = fun.imported_modules[0].get_source()
             if "adreno" in target:
                 pattern = "convert_float"
@@ -225,7 +225,7 @@ def _get_maximum_kernel_args(source):
     def get_kernel_args(source):
         import re
 
-        p = re.compile(r"__kernel void .+\((.*)\)")
+        p = re.tir.build(r"__kernel void .+\((.*)\)")
         args = p.findall(source)
         return args
 

tests/python/codegen/test_target_codegen_vulkan.py

@@ -82,7 +82,7 @@ def test_vector_comparison(target, dev, dtype):
     sch.vectorize(vx)
 
     # Build
-    f = tvm.compile(sch.mod, target=target)
+    f = tvm.tir.build(sch.mod, target=target)
 
     # Verify we generate the boolx4 type declaration and the OpSelect
     # v4{float,half,int} instruction

tests/python/contrib/test_cblas.py

@@ -230,7 +230,9 @@ def verify(target="llvm"):
             return
         dev = tvm.cpu(0)
         name = "test_batch_matmul"
-        f = tvm.compile(te.create_prim_func([input1_data, input2_data, final_result]), target=target)
+        f = tvm.compile(
+            te.create_prim_func([input1_data, input2_data, final_result]), target=target
+        )
         if target == "c":
             f = compiling(f, name)
         matrix_input1 = tvm.nd.array(np.random.uniform(size=ashape).astype(input1_data.dtype), dev)

tests/python/relax/test_relax_operators.py

@@ -47,7 +47,7 @@ def run_cpu(mod, func_name, *args, exec_mode):
         mod = tvm.IRModule.from_expr(func)
 
     target = tvm.target.Target("llvm")
-    ex = tvm.compile(mod, target, exec_mode=exec_mode)
+    ex = relax.build(mod, target, exec_mode=exec_mode)
     vm = relax.VirtualMachine(ex, tvm.cpu())
 
     return vm[func_name](*args)

tests/python/relax/test_vm_alloc_storage_with_scope.py

@@ -62,7 +62,7 @@ def test_alloc_storage_with_scope_global():
     mod = Module
     target = "llvm"
     with tvm.transform.PassContext(opt_level=3):
-        lib = tvm.compile(mod, target, exec_mode="compiled")
+        lib = tvm.relax.build(mod, target=target, exec_mode="compiled")
 
     dev = tvm.cpu()
     # This is the important line which tests nd allocator
@@ -72,3 +72,7 @@ def test_alloc_storage_with_scope_global():
     vm_rt.invoke_stateful("main")
     output = vm_rt.get_outputs("main").numpy()
     tvm.testing.assert_allclose(output_ref, output)
+
+
+if __name__ == "__main__":
+    tvm.testing.main()