diff --git a/apps/android_rpc/tests/android_rpc_test.py b/apps/android_rpc/tests/android_rpc_test.py index ba6c0f9c9679..b9c6995729d0 100644 --- a/apps/android_rpc/tests/android_rpc_test.py +++ b/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) diff --git a/apps/ios_rpc/tests/ios_rpc_test.py b/apps/ios_rpc/tests/ios_rpc_test.py index 40149fba1468..0e563ee1b688 100644 --- a/apps/ios_rpc/tests/ios_rpc_test.py +++ b/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) diff --git a/docs/deep_dive/tensor_ir/tutorials/tir_creation.py b/docs/deep_dive/tensor_ir/tutorials/tir_creation.py index 51481fb2e325..3d07f6227b96 100644 --- a/docs/deep_dive/tensor_ir/tutorials/tir_creation.py +++ b/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)) diff --git a/docs/deep_dive/tensor_ir/tutorials/tir_transformation.py b/docs/deep_dive/tensor_ir/tutorials/tir_transformation.py index 1dcf8e7ab5c8..9dae16294426 100644 --- a/docs/deep_dive/tensor_ir/tutorials/tir_transformation.py +++ b/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) diff --git a/docs/get_started/tutorials/ir_module.py b/docs/get_started/tutorials/ir_module.py index 0a825c3da757..c53d0ca5ef74 100644 --- a/docs/get_started/tutorials/ir_module.py +++ b/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 diff --git a/docs/get_started/tutorials/quick_start.py b/docs/get_started/tutorials/quick_start.py index a4edf0b7c4fe..1153108c9632 100644 --- a/docs/get_started/tutorials/quick_start.py +++ b/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") diff --git a/docs/how_to/tutorials/cross_compilation_and_rpc.py b/docs/how_to/tutorials/cross_compilation_and_rpc.py index 94a6f48b4b73..159acc1eb04f 100644 --- a/docs/how_to/tutorials/cross_compilation_and_rpc.py +++ b/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) diff --git a/docs/how_to/tutorials/customize_opt.py b/docs/how_to/tutorials/customize_opt.py index aab1de1b904d..d215654019f0 100644 --- a/docs/how_to/tutorials/customize_opt.py +++ b/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 diff --git a/docs/how_to/tutorials/e2e_opt_model.py b/docs/how_to/tutorials/e2e_opt_model.py index f74b827fe21f..88cc86bfa800 100644 --- a/docs/how_to/tutorials/e2e_opt_model.py +++ b/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 diff --git a/docs/reference/api/python/driver.rst b/docs/reference/api/python/driver.rst index 97c30ec2d25b..0e921eb127a8 100644 --- a/docs/reference/api/python/driver.rst +++ b/docs/reference/api/python/driver.rst @@ -19,4 +19,4 @@ tvm.driver ---------- .. automodule:: tvm.driver -.. autofunction:: tvm.build +.. autofunction:: tvm.compile diff --git a/include/tvm/runtime/profiling.h b/include/tvm/runtime/profiling.h index 2a41aeff6bff..7ec0b0860fac 100644 --- a/include/tvm/runtime/profiling.h +++ b/include/tvm/runtime/profiling.h @@ -504,7 +504,7 @@ String ShapeString(NDArray shape, DLDataType dtype); String ShapeString(const std::vector& 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. diff --git a/python/tvm/contrib/msc/core/runtime/runner.py b/python/tvm/contrib/msc/core/runtime/runner.py index 8b0646b1d927..074c7048c5e9 100644 --- a/python/tvm/contrib/msc/core/runtime/runner.py +++ b/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)) diff --git a/python/tvm/contrib/msc/framework/tvm/runtime/runner.py b/python/tvm/contrib/msc/framework/tvm/runtime/runner.py index 642a88c93386..c6ae512a64e6 100644 --- a/python/tvm/contrib/msc/framework/tvm/runtime/runner.py +++ b/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] diff --git a/python/tvm/dlight/benchmark/bench.py b/python/tvm/dlight/benchmark/bench.py index a5c6f05f74d0..a0ca2dff0bfe 100644 --- a/python/tvm/dlight/benchmark/bench.py +++ b/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 diff --git a/python/tvm/exec/gpu_memory_bandwidth.py b/python/tvm/exec/gpu_memory_bandwidth.py index 4fdc6aef2f64..e7e38638d715 100644 --- a/python/tvm/exec/gpu_memory_bandwidth.py +++ b/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( diff --git a/python/tvm/ir/instrument.py b/python/tvm/ir/instrument.py index 72d4777194e5..63b0e4cde4d5 100644 --- a/python/tvm/ir/instrument.py +++ b/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): diff --git a/python/tvm/meta_schedule/testing/validate_database.py b/python/tvm/meta_schedule/testing/validate_database.py index a790bb49f73e..ec7a4237546b 100644 --- a/python/tvm/meta_schedule/testing/validate_database.py +++ b/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) diff --git a/python/tvm/relax/frontend/nn/extern.py b/python/tvm/relax/frontend/nn/extern.py index 198ef0f23c46..e797e4493b90 100644 --- a/python/tvm/relax/frontend/nn/extern.py +++ b/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 diff --git a/python/tvm/relax/pipeline.py b/python/tvm/relax/pipeline.py index ddf88e8bd08f..37ef6156e4e7 100644 --- a/python/tvm/relax/pipeline.py +++ b/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 diff --git a/python/tvm/relax/training/trainer.py b/python/tvm/relax/training/trainer.py index d2c417a651ae..fbf48fece9f6 100644 --- a/python/tvm/relax/training/trainer.py +++ b/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) diff --git a/python/tvm/relax/transform/tuning_api/default_functions.py b/python/tvm/relax/transform/tuning_api/default_functions.py index efdc9b13e3fe..9cc7aad47125 100644 --- a/python/tvm/relax/transform/tuning_api/default_functions.py +++ b/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) diff --git a/python/tvm/relax/vm_build.py b/python/tvm/relax/vm_build.py index 564c4a747d48..3332df6c8735 100644 --- a/python/tvm/relax/vm_build.py +++ b/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): diff --git a/python/tvm/runtime/profiling/__init__.py b/python/tvm/runtime/profiling/__init__.py index 23ce5476f5b0..903cd443876f 100644 --- a/python/tvm/runtime/profiling/__init__.py +++ b/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(), diff --git a/python/tvm/runtime/relax_vm.py b/python/tvm/runtime/relax_vm.py index 1dc4aef6f4b4..411835cd216f 100644 --- a/python/tvm/runtime/relax_vm.py +++ b/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())() diff --git a/python/tvm/testing/utils.py b/python/tvm/testing/utils.py index 1fc3b0e873d2..d5d25fb6c2ff 100644 --- a/python/tvm/testing/utils.py +++ b/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() diff --git a/python/tvm/topi/sort.py b/python/tvm/topi/sort.py index 5b8e33413d65..f75e5db4b9b1 100644 --- a/python/tvm/topi/sort.py +++ b/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) diff --git a/src/target/source/codegen_c_host.cc b/src/target/source/codegen_c_host.cc index 9b062ef488bb..b00763a8a36e 100644 --- a/src/target/source/codegen_c_host.cc +++ b/src/target/source/codegen_c_host.cc @@ -31,10 +31,6 @@ #include #include -#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 \n"; decl_stream << "#include \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& 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("system-lib").value_or(Bool(false))) { - ICHECK_EQ(target->GetAttr("runtime").value_or(""), "c") - << "c target only supports generating C runtime SystemLibs"; - } - std::string code = cg.Finish(); return CSourceModuleCreate(code, "c", cg.GetFunctionNames()); } diff --git a/tests/python/all-platform-minimal-test/test_minimal_target_codegen_llvm.py b/tests/python/all-platform-minimal-test/test_minimal_target_codegen_llvm.py index 67d5d84a0c1d..4767c24b693a 100644 --- a/tests/python/all-platform-minimal-test/test_minimal_target_codegen_llvm.py +++ b/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 diff --git a/tests/python/codegen/test_gpu_codegen_allreduce.py b/tests/python/codegen/test_gpu_codegen_allreduce.py index fed4e4c04d32..5a24fcb197a5 100644 --- a/tests/python/codegen/test_gpu_codegen_allreduce.py +++ b/tests/python/codegen/test_gpu_codegen_allreduce.py @@ -71,7 +71,7 @@ def test_allreduce_sum(dims, target, dev): sch.bind(j, "threadIdx.z") sch.bind(k, "threadIdx.y") sch.bind(l, "threadIdx.x") - f = tvm.build(sch.mod["main"], target=target) + f = tvm.compile(sch.mod["main"], target=target) # prepare input and output array a_np = np.random.rand(1, d1, d2, d3).astype("float32") @@ -123,7 +123,7 @@ def test_allreduce_sum_compile(optional_metal_compile_callback): sch.bind(j, "threadIdx.z") sch.bind(k, "threadIdx.y") sch.bind(l, "threadIdx.x") - tvm.build(sch.mod["main"], target=target) + tvm.compile(sch.mod["main"], target=target) @tvm.testing.parametrize_targets("cuda", "metal") @@ -138,7 +138,7 @@ def test_allreduce_max(dims, target, dev): sch.bind(j, "threadIdx.z") sch.bind(k, "threadIdx.y") sch.bind(l, "threadIdx.x") - f = tvm.build(sch.mod["main"], target=target) + f = tvm.compile(sch.mod["main"], target=target) # prepare input and output array a_np = -np.random.rand(1, d1, d2, d3).astype("float32") diff --git a/tests/python/codegen/test_inject_ptx_ldg32.py b/tests/python/codegen/test_inject_ptx_ldg32.py index 4a6d4c366a61..a45e8f57f38f 100644 --- a/tests/python/codegen/test_inject_ptx_ldg32.py +++ b/tests/python/codegen/test_inject_ptx_ldg32.py @@ -46,7 +46,7 @@ def test_inject_ptx_intrin(): # Require at least SM80 return with tvm.transform.PassContext(config={"tir.ptx_ldg32": True}): - mod = tvm.build(f, target="cuda") + mod = tvm.compile(f, target="cuda") A_np = np.random.rand(16).astype("float32") B_np = np.zeros((32)).astype("float32") dev = tvm.cuda(0) diff --git a/tests/python/codegen/test_target_codegen.py b/tests/python/codegen/test_target_codegen.py index bae15b5377e3..3332d015a818 100644 --- a/tests/python/codegen/test_target_codegen.py +++ b/tests/python/codegen/test_target_codegen.py @@ -31,7 +31,7 @@ def func(b: T.handle): err_msg = "Predicated buffer store is not supported." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target(target): - tvm.build(func) + tvm.compile(func) @tvm.testing.parametrize_targets("cuda", "opencl", "metal", "rocm", "vulkan -from_device=0") @@ -49,7 +49,7 @@ def func(a: T.handle, b: T.handle): err_msg = "Predicated buffer store is not supported." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target(target): - tvm.build(func) + tvm.compile(func) @tvm.testing.parametrize_targets("c") @@ -67,7 +67,7 @@ def func(a: T.handle, b: T.handle): err_msg = "Predicated buffer load is not supported." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target(target): - tvm.build(func) + tvm.compile(func) @tvm.testing.parametrize_targets("cuda", "opencl", "metal", "rocm", "vulkan -from_device=0") @@ -85,7 +85,7 @@ def func(a: T.handle, b: T.handle): err_msg = "Predicated buffer load is not supported." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target(target): - tvm.build(func) + tvm.compile(func) if __name__ == "__main__": diff --git a/tests/python/codegen/test_target_codegen_aarch64.py b/tests/python/codegen/test_target_codegen_aarch64.py index 8bd0cb17267d..43870044d528 100644 --- a/tests/python/codegen/test_target_codegen_aarch64.py +++ b/tests/python/codegen/test_target_codegen_aarch64.py @@ -43,7 +43,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] * B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and mul instructions using z registers assembly = f.get_source("asm") @@ -73,7 +73,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] + B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and add instructions using z registers assembly = f.get_source("asm") @@ -103,7 +103,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] - B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and sub instructions using z registers assembly = f.get_source("asm") @@ -134,7 +134,7 @@ def check_correct_assembly(type): B = te.placeholder(m, dtype=type, name="B") C = te.placeholder(m, dtype=type, name="C") D = te.compute((m), lambda i: A[i] * B[i] + C[i], name="D") - f = tvm.build(te.create_prim_func([A, B, C, D]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C, D]), target=target) # Verify we see SVE load instructions and either mad or mla instructions using z registers assembly = f.get_source("asm") @@ -164,7 +164,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: tvm.te.max(A[i], B[i])) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and cmgt + sel instructions or a max instruction, all using z registers assembly = f.get_source("asm") @@ -198,7 +198,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: tvm.te.min(A[i], B[i])) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and cmgt + sel instructions or a min instruction, all using z registers assembly = f.get_source("asm") @@ -232,7 +232,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: tvm.te.div(A[i], B[i])) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and div instructions using z registers assembly = f.get_source("asm") @@ -261,7 +261,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: tvm.te.floormod(A[i], B[i]), name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and mls instructions using z registers assembly = f.get_source("asm") @@ -291,7 +291,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] == B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and cmpeq or cmeq instructions using z registers assembly = f.get_source("asm") @@ -321,7 +321,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] != B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and cmpgt, cmgt, cmpne or cmne instructions, all using z registers assembly = f.get_source("asm") @@ -350,7 +350,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] | B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and orr instructions using z registers assembly = f.get_source("asm") @@ -379,7 +379,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype=type, name="B") C = te.compute((m), lambda i: A[i] & B[i], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see SVE load instructions and and instructions using z registers assembly = f.get_source("asm") @@ -407,7 +407,7 @@ def check_correct_assembly(type): m = te.var("m") A = te.placeholder(m, dtype=type, name="A") C = te.compute((m), lambda i: ~A[i], name="C") - f = tvm.build(te.create_prim_func([A, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, C]), target=target) # Verify we see SVE load instructions and eor instructions using z registers assembly = f.get_source("asm") @@ -440,7 +440,7 @@ def check_correct_assembly(type): A = te.placeholder(m, dtype=type, name="A") B = te.placeholder(m, dtype="int32", name="B") C = te.compute((m), lambda i: A[B[i]], name="C") - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, B, C]), target=target) # Verify we see gather instructions in the assembly assembly = f.get_source("asm") @@ -463,7 +463,7 @@ def main(A: T.Buffer((5,), "int32")): for i in range(5): A[i] = 2 * vscale - build_mod = tvm.build(main, target=target) + build_mod = tvm.tir.build(main, target=target) llvm = build_mod.get_source() assert re.findall(r"llvm.vscale.i32", llvm), "No vscale in generated LLVM." @@ -482,7 +482,7 @@ def my_func(a: T.handle, b: T.handle): T.func_attr({"global_symbol": "my_module", "tir.noalias": True}) B[T.ramp(0, 1, 4 * T.vscale())] = A[T.ramp(0, 1, 4 * T.vscale())] - mod = tvm.build(my_func, target=target) + mod = tvm.tir.build(my_func, target=target) llvm = mod.get_source("ll") assert re.findall(r"load ", llvm), "No scalable load in generated LLVM." @@ -501,7 +501,7 @@ def my_func(a: T.handle): T.func_attr({"global_symbol": "my_module", "tir.noalias": True}) A[T.ramp(0, 1, 4 * T.vscale())] = T.broadcast(1, 4 * T.vscale()) - mod = tvm.build(my_func, target=target) + mod = tvm.tir.build(my_func, target=target) llvm = mod.get_source("ll") assert re.findall( @@ -529,7 +529,7 @@ def test_vscale_range_function_attribute(mattr, expect_attr): m = te.var("m") A = te.placeholder(m, dtype="float32", name="A") C = te.compute((m), lambda i: A[i] + 1, name="C") - f = tvm.build(te.create_prim_func([A, C]), target=target) + f = tvm.tir.build(te.create_prim_func([A, C]), target=target) # Check if the vscale_range() attribute exists ll = f.get_source("ll") @@ -558,7 +558,7 @@ def before(a: T.handle): A[i : i + T.vscale() * 4] = T.get_active_lane_mask("uint1xvscalex4", i, 30) with tvm.target.Target(target): - out = tvm.build(before) + out = tvm.tir.build(before) ll = out.get_source("ll") assert "get.active.lane.mask" in ll @@ -581,7 +581,7 @@ def before(a: T.handle, b: T.handle): B[i_0 * 4 * T.vscale() + i_1] = A[i_0 * 4 * T.vscale() + i_1] + 1.0 with tvm.target.Target(target): - out = tvm.build(before) + out = tvm.tir.build(before) ll = out.get_source("ll") assert "get.active.lane.mask" in ll diff --git a/tests/python/codegen/test_target_codegen_arm.py b/tests/python/codegen/test_target_codegen_arm.py index 9357d38e667b..d22e528770b3 100644 --- a/tests/python/codegen/test_target_codegen_arm.py +++ b/tests/python/codegen/test_target_codegen_arm.py @@ -17,8 +17,6 @@ import tvm from tvm import te import re -import os -import ctypes def test_popcount(): @@ -30,7 +28,7 @@ def check_correct_assembly(type, elements, counts): B = te.compute(A.shape, lambda i: tvm.tir.popcount(A[i]), name="B") sch = tvm.tir.Schedule(te.create_prim_func([A, B])) sch.vectorize(sch.get_loops("B")[0]) - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) # Verify we see the correct number of vpaddl and vcnt instructions in the assembly assembly = f.get_source("asm") matches = re.findall("vpaddl", assembly) @@ -60,7 +58,7 @@ def check_correct_assembly(N): ) sch = tvm.tir.Schedule(te.create_prim_func([A, B, C])) sch.vectorize(sch.get_loops("C")[0]) - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) # Verify we see the correct number of vmlal.s16 instructions assembly = f.get_source("asm") @@ -84,7 +82,7 @@ def check_broadcast_correct_assembly(N): ) sch = tvm.tir.Schedule(te.create_prim_func([A, B, C])) sch.vectorize(sch.get_loops("C")[0]) - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) # Verify we see the correct number of vmlal.s16 instructions assembly = f.get_source("asm") diff --git a/tests/python/codegen/test_target_codegen_blob.py b/tests/python/codegen/test_target_codegen_blob.py index a61e6e894c83..39373c4d840c 100644 --- a/tests/python/codegen/test_target_codegen_blob.py +++ b/tests/python/codegen/test_target_codegen_blob.py @@ -57,8 +57,8 @@ def my_inplace_update(x: T.Buffer((12), "float32")) -> None: temp = utils.tempdir() target = tvm.target.Target("cuda", host="llvm") - libA = tvm.build(ModA, target=target) - libB = tvm.build(ModB, target=target) + libA = tvm.compile(ModA, target=target) + libB = tvm.compile(ModB, target=target) pathA = temp.relpath("libA.tar") pathB = temp.relpath("libB.tar") diff --git a/tests/python/codegen/test_target_codegen_bool.py b/tests/python/codegen/test_target_codegen_bool.py index a575c0cec9c9..96bd21329c93 100644 --- a/tests/python/codegen/test_target_codegen_bool.py +++ b/tests/python/codegen/test_target_codegen_bool.py @@ -52,7 +52,7 @@ def get_module(target, compute): @tvm.testing.uses_gpu def test_cmp_load_store(target, dev, arr_size, compute, get_module): A, B, _, D = compute - f = tvm.build(get_module, target=target) + f = tvm.compile(get_module, target=target) a_np = np.random.uniform(size=arr_size).astype(A.dtype) b_np = np.random.uniform(size=arr_size).astype(B.dtype) diff --git a/tests/python/codegen/test_target_codegen_c_host.py b/tests/python/codegen/test_target_codegen_c_host.py index d7a7cbc8a44b..af94cae71f1c 100644 --- a/tests/python/codegen/test_target_codegen_c_host.py +++ b/tests/python/codegen/test_target_codegen_c_host.py @@ -33,7 +33,7 @@ def test_add(): C = te.compute(A.shape, lambda *i: A(*i) + B(*i), name="C") def check_c(): - mhost = tvm.build( + mhost = tvm.compile( tvm.IRModule.from_expr( te.create_prim_func([A, B, C]).with_attr("global_symbol", "test_fadd") ), @@ -65,7 +65,7 @@ def test_reinterpret(): ) def check_c(): - mhost = tvm.build( + mhost = tvm.compile( tvm.IRModule.from_expr( te.create_prim_func([A, B]).with_attr("global_symbol", "test_reinterpret") ), @@ -93,7 +93,7 @@ def test_ceil(): B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.ceil", A(*i)), name="B") def check_c(): - mhost = tvm.build( + mhost = tvm.compile( tvm.IRModule.from_expr( te.create_prim_func([A, B]).with_attr("global_symbol", "test_ceil") ), @@ -121,7 +121,7 @@ def test_floor(): B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.floor", A(*i)), name="B") def check_c(): - mhost = tvm.build( + mhost = tvm.compile( tvm.IRModule.from_expr( te.create_prim_func([A, B]).with_attr("global_symbol", "test_floor") ), @@ -149,7 +149,7 @@ def test_round(): B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.round", A(*i)), name="B") def check_c(): - mhost = tvm.build( + mhost = tvm.compile( tvm.IRModule.from_expr( te.create_prim_func([A, B]).with_attr("global_symbol", "test_round") ), @@ -182,7 +182,7 @@ def subroutine(A_data: T.handle("float32")): A = T.decl_buffer(1, dtype="float32", data=A_data) A[0] = 42.0 - built = tvm.build(mod, target="c") + built = tvm.tir.build(mod, target="c") func_names = list(built["get_func_names"]()) assert ( diff --git a/tests/python/codegen/test_target_codegen_cross_llvm.py b/tests/python/codegen/test_target_codegen_cross_llvm.py index 9dc001e1949a..c126e531090e 100644 --- a/tests/python/codegen/test_target_codegen_cross_llvm.py +++ b/tests/python/codegen/test_target_codegen_cross_llvm.py @@ -49,7 +49,7 @@ def verify_elf(path, e_machine): def build_i386(): temp = utils.tempdir() target = "llvm -mtriple=i386-pc-linux-gnu" - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) path = temp.relpath("myadd.o") f.save(path) verify_elf(path, 0x03) @@ -60,7 +60,7 @@ def build_arm(): print("Skip because %s is not enabled.." % target) return temp = utils.tempdir() - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) path = temp.relpath("myadd.o") f.save(path) verify_elf(path, 0x28) diff --git a/tests/python/codegen/test_target_codegen_cuda.py b/tests/python/codegen/test_target_codegen_cuda.py index b3cad9acd38e..e0e37660fc75 100644 --- a/tests/python/codegen/test_target_codegen_cuda.py +++ b/tests/python/codegen/test_target_codegen_cuda.py @@ -15,19 +15,16 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -import os -import re -import tvm -from tvm import te import numpy as np -from tvm import topi -from tvm.contrib.nvcc import have_fp16, have_int8, have_bf16 -from tvm.contrib import utils -from tvm.script import tir as T -import tvm.testing import pytest +import tvm +import tvm.testing +from tvm import te, topi +from tvm.contrib.nvcc import have_bf16, have_fp16, have_int8 +from tvm.script import tir as T + @tvm.testing.requires_gpu @tvm.testing.requires_cuda @@ -48,7 +45,7 @@ def check_cuda(dtype, n, lanes): xo, xi = sch.split(sch.get_loops("B")[0], factors=[None, num_thread]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") dev = tvm.cuda(0) a = tvm.nd.empty((n,), A.dtype, dev).copyfrom(np.random.uniform(size=(n, lanes))) @@ -103,7 +100,7 @@ def check_cuda(n, lanes): with tvm.transform.PassContext( disabled_pass=["tir.BF16Promote", "tir.BF16CastElimination", "tir.BF16TypeLowering"] ): - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") dev = tvm.cuda(0) np_a = np.random.uniform(size=(n, lanes)).astype("float32") np_a = np_bf162np_float(np_float2np_bf16(np_a)) @@ -138,7 +135,7 @@ def check_cuda(dtype, n, lanes): xo, xi = sch.split(sch.get_loops("D")[0], factors=[None, num_thread]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") np_a = np.random.randint(low=-128, high=127, size=(n, lanes)) np_b = np.random.randint(low=-128, high=127, size=(n, lanes)) @@ -169,7 +166,7 @@ def check_cuda(dtype, n, lanes): xo, xi = sch.split(sch.get_loops("B")[0], factors=[None, num_thread]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") np_a = np.random.randint(low=-128, high=127, size=(n, lanes)) a = tvm.nd.empty((n,), A.dtype, dev).copyfrom(np_a) @@ -196,7 +193,7 @@ def check_cuda(n, value, lanes): y, x = sch.get_loops("A") sch.vectorize(x) sch.bind(y, "blockIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") np_a = np.full((n, lanes), value, dtype=dtype) a = tvm.nd.empty(np_a.shape, dtype, dev) @@ -225,7 +222,7 @@ def check_cuda(n, value, lanes): y, x = sch.get_loops("A") sch.vectorize(x) sch.bind(y, "blockIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") np_a = np.full((n, lanes), value, dtype="int8") a = tvm.nd.empty((n, lanes), dtype, dev) @@ -256,7 +253,7 @@ def check_inf_nan(dev, n, value, dtype): xo, xi = sch.split(sch.get_loops("C")[0], factors=[None, 8]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") a = tvm.nd.empty((n,), A.dtype, dev) c = tvm.nd.empty((n,), A.dtype, dev) @@ -287,7 +284,7 @@ def sched(nthd): ko, _ = sch.split(k, factors=[nthd, None]) sch.bind(ko, "threadIdx.x") sch.bind(x, "blockIdx.x") - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") return fun def verify(nthd): @@ -325,7 +322,7 @@ def sched(nthdx, nthdy): sch.bind(k0o, "threadIdx.x") sch.bind(k1o, "threadIdx.y") sch.bind(x, "blockIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") return func def verify(nthdx, nthdy): @@ -359,7 +356,7 @@ def test_cuda_reduction_binding(): sch.reorder(k, x) mo, _ = sch.split(x, factors=[None, 32]) sch.bind(mo, "blockIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") @tvm.testing.requires_gpu @@ -377,7 +374,7 @@ def test_cuda_const_float_to_half(): xo, xi = sch.split(sch.fuse(*sch.get_loops("C")), factors=[None, 64]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") dev = tvm.cuda(0) a_np = np.random.uniform(size=shape).astype(a.dtype) @@ -404,7 +401,7 @@ def test_cuda_floordiv_with_vectorization(): sch.vectorize(xii) sch.bind(xo, "blockIdx.x") sch.bind(xio, "threadIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") dev = tvm.cuda(0) a_np = np.random.uniform(size=(n,)).astype(A.dtype) @@ -430,7 +427,7 @@ def test_cuda_floormod_with_vectorization(): sch.vectorize(xii) sch.bind(xo, "blockIdx.x") sch.bind(xio, "threadIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") dev = tvm.cuda(0) a_np = np.random.uniform(size=(n,)).astype(A.dtype) @@ -460,7 +457,7 @@ def check(t0, t1, factor): ob, ib = sch.split(sch.get_loops("C")[0], factors=[None, factor]) sch.vectorize(ib) sch.bind(ob, "threadIdx.x") - func = tvm.build(sch.mod, target="cuda") + func = tvm.compile(sch.mod, target="cuda") # correctness dev = tvm.cuda(0) @@ -514,7 +511,7 @@ def sched(A, B): sch.vectorize(iiii) sch.bind(io, "blockIdx.x") sch.bind(iio, "threadIdx.x") - return tvm.build(sch.mod, target="cuda") + return tvm.compile(sch.mod, target="cuda") @tvm.testing.requires_gpu @@ -650,7 +647,7 @@ def check_cuda(dtype, n, l, padding, lanes): sch.bind(block, "blockIdx.x") sch.bind(thread, "threadIdx.x") sch.vectorize(vectorize) - fun = tvm.build(sch.mod, target="cuda") + fun = tvm.compile(sch.mod, target="cuda") np_a = np.random.randint(low=-128, high=127, size=(n, l)).astype(A.dtype) a = tvm.nd.empty((n, l), A.dtype, dev).copyfrom(np_a) @@ -691,7 +688,7 @@ def build(A, C, N, C_N): sch.bind(oi, "threadIdx.x") sch.vectorize(ii) # BUG: misalignment - f = tvm.build(sch.mod, target="cuda") + f = tvm.tir.build(sch.mod, target="cuda") kernel_source = f.imported_modules[0].get_source() dev = tvm.cuda() @@ -757,13 +754,13 @@ def func3(A: T.Buffer((4, 4), "float32")) -> None: mod = tvm.IRModule({"main": func1}) with pytest.raises(tvm.error.InternalError): - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") mod = tvm.IRModule({"main": func2}) - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") mod = tvm.IRModule({"main": func3}) - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") @tvm.testing.requires_cuda @@ -774,7 +771,7 @@ def func(A: T.Buffer((4,), "uint32"), B: T.Buffer((4,), "uint8")) -> None: B[tx] = T.call_intrin("uint8", "tir.reinterpret", A[tx]) with pytest.raises(tvm.error.TVMError): - tvm.build(func, target="cuda") + tvm.compile(func, target="cuda") if __name__ == "__main__": diff --git a/tests/python/codegen/test_target_codegen_cuda_fp4.py b/tests/python/codegen/test_target_codegen_cuda_fp4.py index 46825826a9d1..0a170026c96b 100644 --- a/tests/python/codegen/test_target_codegen_cuda_fp4.py +++ b/tests/python/codegen/test_target_codegen_cuda_fp4.py @@ -63,7 +63,7 @@ def add( sch.bind(tx, "threadIdx.x") target = "cuda" - fadd = tvm.build(sch.mod, target=target) + fadd = tvm.compile(sch.mod, target=target) dev = tvm.device(target, 0) numpytype = "float4_e2m1fn" @@ -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) diff --git a/tests/python/codegen/test_target_codegen_cuda_fp8.py b/tests/python/codegen/test_target_codegen_cuda_fp8.py index b91efd61922f..5e0a4c3000da 100644 --- a/tests/python/codegen/test_target_codegen_cuda_fp8.py +++ b/tests/python/codegen/test_target_codegen_cuda_fp8.py @@ -67,7 +67,7 @@ def add( sch.bind(tx, "threadIdx.x") target = "cuda" - fadd = tvm.build(sch.mod, target=target) + fadd = tvm.compile(sch.mod, target=target) cuda_src = fadd.imported_modules[0].get_source() assert "__nv_fp8_e4m3" in cuda_src, "FP8E4M3 (fp8_e4_t) datatype not found in generated CUDA" @@ -125,7 +125,7 @@ def add( sch.bind(tx, "threadIdx.x") target = "cuda" - f = tvm.build(sch.mod, target=target) + f = tvm.compile(sch.mod, target=target) dev = tvm.device(target, 0) numpytype = "float8_e4m3fn" @@ -179,7 +179,7 @@ def add( sch.bind(tx, "threadIdx.x") target = "cuda" - fadd = tvm.build(sch.mod, target=target) + fadd = tvm.compile(sch.mod, target=target) cuda_src = fadd.imported_modules[0].get_source() dev = tvm.device(target, 0) @@ -230,7 +230,7 @@ def vector_broadcast(a: T.Buffer((), dtype), vec: T.Buffer((bcast_length,), dtyp sch.bind(tx, "threadIdx.x") target = "cuda" - func = tvm.build(sch.mod, target=target) + func = tvm.compile(sch.mod, target=target) dev = tvm.device(target, 0) a_np = np.random.uniform(low=0, high=4, size=()).astype(dtype) @@ -263,7 +263,7 @@ def vector_load( B[i] = A[vec_index] target = "cuda" - f = tvm.build(vector_load, target=target) + f = tvm.compile(vector_load, target=target) dev = tvm.device(target, 0) a_np = np.random.uniform(low=0, high=1, size=(length,)).astype(dtype) @@ -312,7 +312,7 @@ def add( sch.bind(tx, "threadIdx.x") target = "cuda" - fadd = tvm.build(sch.mod, target=target) + fadd = tvm.compile(sch.mod, target=target) dev = tvm.device(target, 0) a_np = np.random.uniform(-1, 1, (length, vector_length)).astype(dtype) @@ -670,7 +670,7 @@ def compile_quant_and_dequant_by_scale( dl.gpu.GeneralReduction(), dl.gpu.Fallback(), )(quant_mod) - ex_1 = relax.build(quant_mod, target=target) + ex_1 = tvm.compile(quant_mod, target=target) vm_1 = relax.VirtualMachine(ex_1, dev) dequant_mod = cls.create_dequantize_func( @@ -694,13 +694,13 @@ def compile_quant_and_dequant_by_scale( )(dequant_mod) dequant_mod.show() - ex_2 = relax.build(dequant_mod, target=target) + ex_2 = tvm.compile(dequant_mod, target=target) vm_2 = relax.VirtualMachine(ex_2, dev) def print_cuda(target, mod, name=None): if name: mod = mod[name] - f = tvm.build(mod, target=target) + f = tvm.compile(mod, target=target) cuda_src = f.imported_modules[0].get_source() print(cuda_src) @@ -817,7 +817,7 @@ def func(A: T.Buffer((4,), dtype)) -> None: A[tx] = A_local[tx] mod = tvm.IRModule({"main": func}) - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") @tvm.testing.requires_cuda_compute_version(8, 9) @@ -846,7 +846,7 @@ def func( B[tx, i] = A[tx, i] mod = tvm.IRModule({"main": func}) - rtmod = tvm.build(mod, target="cuda") + rtmod = tvm.compile(mod, target="cuda") num_experts = 8 @@ -942,7 +942,7 @@ def _pipeline(mod: tvm.ir.IRModule) -> tvm.ir.IRModule: target = tvm.target.Target("cuda") with tvm.transform.PassContext(config={"relax.backend.use_cuda_graph": False}) and target: mod = _pipeline(mod) - rt_mod = tvm.relax.build(mod, target=target) + rt_mod = tvm.compile(mod, target=target) dev = tvm.cuda(0) x_data = np.zeros((1, reduce_size), dtype=np.float16) diff --git a/tests/python/codegen/test_target_codegen_device.py b/tests/python/codegen/test_target_codegen_device.py index ad27356961aa..1adb337de0c4 100644 --- a/tests/python/codegen/test_target_codegen_device.py +++ b/tests/python/codegen/test_target_codegen_device.py @@ -48,7 +48,7 @@ def check_target(device): if not tvm.testing.device_enabled(device): return dev = tvm.device(device, 0) - f = tvm.build(sch.mod, target=device) + f = tvm.compile(sch.mod, target=device) # launch the kernel. a = tvm.nd.empty((n,), dtype=A.dtype, device=dev) f(a) @@ -94,7 +94,7 @@ def check_target(device, host="stackvm"): return dev = tvm.device(device, 0) target = tvm.target.Target(device, host) - mhost = tvm.build(sch.mod, target=target) + mhost = tvm.tir.build(sch.mod, target=target) f = mhost.entry_func # launch the kernel. n = 1027 diff --git a/tests/python/codegen/test_target_codegen_extern.py b/tests/python/codegen/test_target_codegen_extern.py index 378eb427fd54..99069b1bd1ed 100644 --- a/tests/python/codegen/test_target_codegen_extern.py +++ b/tests/python/codegen/test_target_codegen_extern.py @@ -69,7 +69,7 @@ def check_target(target): mod = mod_gpu if target in ["opencl", "cuda"] else mod_cpu C = C_gpu if target in ["opencl", "cuda"] else C_cpu # build and invoke the kernel. - f = tvm.build(mod, target=target) + f = tvm.compile(mod, target=target) dev = tvm.device(target, 0) # launch the kernel. n = nn @@ -105,7 +105,7 @@ def check_target(target): if not tvm.testing.device_enabled(target): return # build and invoke the kernel. - f = tvm.build(mod, target=target) + f = tvm.compile(mod, target=target) dev = tvm.cpu(0) # launch the kernel. n = nn @@ -137,7 +137,7 @@ def check_target(target): if not tvm.testing.device_enabled(target): return # build and invoke the kernel. - f = tvm.build(mod, target=target) + f = tvm.compile(mod, target=target) dev = tvm.cpu(0) # launch the kernel. n = nn diff --git a/tests/python/codegen/test_target_codegen_gpu_common.py b/tests/python/codegen/test_target_codegen_gpu_common.py index 2941f366a43b..08f43a114084 100644 --- a/tests/python/codegen/test_target_codegen_gpu_common.py +++ b/tests/python/codegen/test_target_codegen_gpu_common.py @@ -40,7 +40,7 @@ def run_test(tvm_intrin, np_func, dtype): sch = tvm.tir.Schedule(func) (x,) = sch.get_loops(sch.get_block("B")) sch.bind(x, "threadIdx.x") - f = tvm.build(sch.mod, target=target) + f = tvm.compile(sch.mod, target=target) a = tvm.nd.array(np.random.randint(0, 100000, size=n).astype(A.dtype), dev) b = tvm.nd.array(np.zeros(shape=(n,)).astype(B.dtype), dev) f(a, b) diff --git a/tests/python/codegen/test_target_codegen_hexagon.py b/tests/python/codegen/test_target_codegen_hexagon.py index 37e62e5b34ef..c0665ce316ad 100644 --- a/tests/python/codegen/test_target_codegen_hexagon.py +++ b/tests/python/codegen/test_target_codegen_hexagon.py @@ -45,7 +45,7 @@ def check_add(): B = tvm.te.placeholder((128,), dtype="uint8", name="A") C = tvm.te.compute((128,), lambda i: A[i] + B[i], name="C") mod = tvm.IRModule.from_expr(te.create_prim_func([C, A, B])) - hexm = tvm.build(mod, target=tvm.target.Target(target, target)) + hexm = tvm.compile(mod, target=tvm.target.Target(target, target)) asm = hexm.get_source("s") vadds = re.findall(r"v[0-9]+.b = vadd\(v[0-9]+.b,v[0-9]+.b\)", asm) assert vadds # Check that it's non-empty @@ -60,7 +60,7 @@ def test_llvm_target_features(): A = tvm.te.placeholder((128,), dtype="uint8", name="A") C = tvm.te.compute((128,), lambda i: A[i] + 1, name="C") mod = tvm.IRModule.from_expr(te.create_prim_func([C, A]).with_attr("global_symbol", "add_one")) - m = tvm.build(mod, target=tvm.target.Target(target, target)) + m = tvm.compile(mod, target=tvm.target.Target(target, target)) llvm_ir = m.get_source("ll") # Make sure we find +hvx-length128b in "attributes". fs = re.findall(r"attributes.*\+hvx-length128b", llvm_ir) @@ -74,7 +74,7 @@ def test_llvm_options(): mod = tvm.IRModule.from_expr(te.create_prim_func([Zero])) # Check that BuildHexagon hasn't crashed because of target attribute # type mismatch. - tvm.build(mod, target=tvm.target.Target(target, target)) + tvm.compile(mod, target=tvm.target.Target(target, target)) assert re.search("-hexagon-noopt", str(target)) diff --git a/tests/python/codegen/test_target_codegen_llvm.py b/tests/python/codegen/test_target_codegen_llvm.py index 304c79559cbb..e4a28453184d 100644 --- a/tests/python/codegen/test_target_codegen_llvm.py +++ b/tests/python/codegen/test_target_codegen_llvm.py @@ -14,21 +14,18 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -import collections -import ctypes -import json import math +import re + import numpy as np import pytest -import re -import sys import tvm import tvm.testing -from tvm import te -from tvm import tir +from tvm import te, tir from tvm.contrib import clang, utils -from tvm.script import tir as T, ir as I +from tvm.script import ir as I +from tvm.script import tir as T from tvm.target.codegen import llvm_get_intrinsic_name, llvm_lookup_intrinsic_id @@ -42,7 +39,7 @@ def test_llvm_intrin(): body = ib.get() mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([A], body).with_attr("global_symbol", "prefetch")) - fcode = tvm.build(mod, None) + fcode = tvm.compile(mod) @tvm.testing.requires_llvm @@ -54,7 +51,7 @@ def test_llvm_void_intrin(): ib.emit(x) body = ib.get() mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([A], body).with_attr("global_symbol", "main")) - fcode = tvm.build(mod, None) + fcode = tvm.compile(mod) @tvm.testing.requires_llvm @@ -92,7 +89,7 @@ def use_llvm_intrinsic(A, C): sch = tir.Schedule(mod) # Build from scheduled TIR - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") @tvm.testing.requires_llvm @@ -106,7 +103,7 @@ def test_llvm_lookup_intrin(): ib.emit(x) body = ib.get() mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([A], body).with_attr("global_symbol", "main")) - fcode = tvm.build(mod, None) + fcode = tvm.compile(mod, None) @tvm.testing.requires_llvm @@ -120,7 +117,7 @@ def test_llvm_large_uintimm(): sch = tir.Schedule(mod) def check_llvm(): - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) # launch the kernel. a = tvm.nd.empty((), dtype=A.dtype, device=dev) @@ -162,7 +159,7 @@ def test_llvm_multi_parallel(): 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. a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), dev) @@ -194,7 +191,7 @@ def check_llvm(nn, base): sch.vectorize(xi) # 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 @@ -228,7 +225,7 @@ def test_llvm_vadd_pipeline(): _, inner = sch.split(loop, factors=[None, 4]) sch.vectorize(inner) # Build and verify - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) n = 128 a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), dev) @@ -259,7 +256,7 @@ def check_llvm(nn, base, stride): sch.vectorize(xi) # 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 @@ -289,7 +286,7 @@ def test_llvm_temp_space(): 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 @@ -322,7 +319,7 @@ def test_multiple_func(): ) # Build and verify - f = tvm.build(mod, target="llvm") + f = tvm.compile(mod, target="llvm") dev = tvm.cpu(0) n = 10 a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), dev) @@ -347,7 +344,7 @@ def check_llvm(n, offset): sch = tir.Schedule(mod) # 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. a = tvm.nd.array(np.random.uniform(size=(n,)).astype(A.dtype), dev) @@ -371,7 +368,7 @@ def check_llvm(n): sch = tir.Schedule(mod) # 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. a = tvm.nd.array(np.random.randint(0, 2, size=(n,)).astype(A.dtype), dev) @@ -397,7 +394,7 @@ def check_llvm(n): sch = tir.Schedule(mod) # 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. a = tvm.nd.array(np.random.randint(0, 2, size=(n,)).astype(A.dtype), dev) @@ -425,7 +422,7 @@ def check_llvm(n): sch = tir.Schedule(mod) # 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. a = tvm.nd.array(np.random.randint(0, 2, size=(n,)).astype(A.dtype), dev) @@ -457,7 +454,7 @@ def test_alignment(): sch.vectorize(tx) # Build with name - f = tvm.build(sch.mod, target="llvm") + f = tvm.tir.build(sch.mod, target="llvm") lines = f.get_source().split("\n") @@ -533,7 +530,7 @@ def clipb(x): sch = tir.Schedule(mod) # Build from scheduled TIR - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") # Fill input arrays with values A_arr = tvm.nd.empty((end - start + 1,), dtype) @@ -639,7 +636,7 @@ def check_llvm_reciprocal(n): sch = tir.Schedule(mod) # Build from scheduled TIR - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") a = tvm.nd.array(np.full((n,), 100, "float32")) b = tvm.nd.empty((n,), "float32") @@ -659,7 +656,7 @@ def check_llvm_sigmoid(n): sch = tir.Schedule(mod) # Build from scheduled TIR - f = tvm.build(sch.mod, target="llvm") + f = tvm.compile(sch.mod, target="llvm") a = tvm.nd.array(np.full((n,), -1000, "float32")) b = tvm.nd.empty((n,), "float32") @@ -704,7 +701,7 @@ def check_llvm_object(): "fadd2": sch.mod["main"].with_attr("global_symbol", "fadd2"), } ) - m = tvm.build(mod, target="llvm") + m = tvm.compile(mod, target="llvm") temp = utils.tempdir() o_path = temp.relpath("temp.o") m.save(o_path) @@ -742,7 +739,7 @@ def check_llvm_ir(): "fadd2": sch.mod["main"].with_attr("global_symbol", "fadd2"), } ) - m = tvm.build(mod, target="llvm -mtriple=aarch64-linux-gnu") + m = tvm.tir.build(mod, target="llvm -mtriple=aarch64-linux-gnu") ll = m.get_source("ll") # On non-Darwin OS, don't explicitly specify DWARF version. @@ -752,7 +749,7 @@ def check_llvm_ir(): assert re.search(r"""llvm.dbg.value""", ll) # Try Darwin, require DWARF-2 - m = tvm.build(mod, target="llvm -mtriple=x86_64-apple-darwin-macho") + m = tvm.tir.build(mod, target="llvm -mtriple=x86_64-apple-darwin-macho") ll = m.get_source("ll") assert re.search(r"""i32 4, !"Dwarf Version", i32 2""", ll) assert re.search(r"""llvm.dbg.value""", ll) @@ -808,7 +805,7 @@ def dotest(do_vectorize): if do_vectorize: sch.vectorize(loop) - module = tvm.build(sch.mod, target="llvm") + module = tvm.compile(sch.mod, target="llvm") npa = np.random.rand(32).astype("float32") npb = np.random.rand(32).astype("float32") va = np_bf16_cast_and_cast_back(npa) @@ -830,7 +827,7 @@ def test_llvm_crt_static_lib(): B = te.placeholder((32,), dtype="bfloat16") d = te.compute((32,), lambda x: A[x] + B[x]) mod = tvm.IRModule.from_expr(te.create_prim_func([A, B, d])) - module = tvm.build( + module = tvm.tir.build( mod.with_attr("system_lib_prefix", ""), target=tvm.target.Target("llvm"), ) @@ -860,7 +857,7 @@ def make_call_extern(caller, callee): "Kirby": make_call_extern("Kirby", "Fred"), } mod = tvm.IRModule(functions=functions) - ir_text = tvm.build(mod, target="llvm").get_source("ll") + ir_text = tvm.tir.build(mod, target="llvm").get_source("ll") # Skip functions whose names start with _. matches = re.findall(r"^define[^@]*@([a-zA-Z][a-zA-Z0-9_]*)", ir_text, re.MULTILINE) assert matches == sorted(matches) @@ -896,7 +893,7 @@ def check_llvm(use_file): else: sch.annotate(sch.get_loops("B")[0], "pragma_import_llvm", ll_code) # 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. a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), dev) @@ -921,7 +918,7 @@ def test_llvm_scalar_concat(): # This will crash in LLVM codegen if CodeGenLLVM::CreateVecConcat doesn't convert # scalars to single-lane LLVM vectors. with tvm.transform.PassContext(config={"tir.disable_assert": True}): - m = tvm.build(mod, target="llvm") + m = tvm.compile(mod, target="llvm") @tvm.testing.requires_llvm @@ -936,7 +933,7 @@ def threadpool_nested_parallel_loop( B[i, j] = A[i, j] * 2.0 with pytest.raises(tvm.TVMError) as e: - tvm.build(tvm.IRModule.from_expr(threadpool_nested_parallel_loop), target="llvm") + tvm.compile(tvm.IRModule.from_expr(threadpool_nested_parallel_loop), target="llvm") msg = str(e) assert msg.find("Nested parallel loop is not supported") != -1 @@ -959,7 +956,7 @@ def test_llvm_target_attributes(): target_llvm = "llvm -mtriple=x86_64-linux-gnu -mcpu=skylake -mattr=+avx512f" target = tvm.target.Target(target_llvm, host=target_llvm) - module = tvm.build(sch.mod, target=target) + module = tvm.tir.build(sch.mod, target=target) llvm_ir = module.get_source() llvm_ir_lines = llvm_ir.split("\n") @@ -1010,7 +1007,7 @@ def tir_assume_func(A: T.Buffer((4, 4), "int32"), B: T.Buffer((14,), "int32")): mod = tvm.IRModule.from_expr(tir_assume_func) inp = te.placeholder((4, 4), name="A", dtype="int32") out = te.placeholder((14,), name="B", dtype="int32") - m = tvm.build(mod, target="llvm") + m = tvm.compile(mod, target="llvm") @tvm.testing.requires_llvm @@ -1030,7 +1027,7 @@ def func(a: T.handle("float64"), b: T.handle("float64"), n: T.int64): for i in range(n): B[i] = A[i] - tvm.build(func, target="llvm") + tvm.compile(func, target="llvm") @tvm.testing.requires_llvm @@ -1053,7 +1050,7 @@ def subroutine(A_data: T.handle("float32")): target = "llvm" dev = tvm.cpu() - built = tvm.build(mod) + built = tvm.compile(mod) arr = tvm.nd.array(np.zeros([1], "float32"), device=dev) built["main"](arr) @@ -1086,7 +1083,7 @@ def func(): # Error occurred during build, as part of # CodeGenCPU::MakeCallPackedLowered. - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") @tvm.testing.requires_llvm @@ -1105,7 +1102,7 @@ def func(A: T.Buffer(1, "float32")): T.Call("int32", tvm.ir.Op.get("tir.tvm_call_packed"), [A.data]) with pytest.raises(tvm.TVMError): - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") @tvm.testing.requires_llvm @@ -1117,7 +1114,7 @@ def func(): T.func_attr({"global_symbol": "func"}) T.Call("void", tvm.ir.Op.get("tir.call_extern"), ["dummy_function_name"]) - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") def test_invalid_volatile_masked_buffer_load(): @@ -1132,7 +1129,7 @@ def func(b: T.handle): err_msg = "The masked load intrinsic does not support declaring load as volatile." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target("llvm"): - tvm.build(func) + tvm.compile(func) def test_invalid_volatile_masked_buffer_store(): @@ -1146,7 +1143,7 @@ def func(): err_msg = "The masked store intrinsic does not support declaring store as volatile." with pytest.raises(tvm.TVMError, match=err_msg): with tvm.target.Target("llvm"): - tvm.build(func) + tvm.compile(func) def test_int_parameter(): @@ -1160,7 +1157,7 @@ def func(arg: T.int32) -> T.int32: else: return 20 - built = tvm.build(func) + built = tvm.compile(func) output = built(True) assert output == 10 @@ -1179,7 +1176,7 @@ def func(arg: T.bool) -> T.int32: else: return 20 - built = tvm.build(func) + built = tvm.compile(func) output = built(1) assert output == 10 @@ -1198,7 +1195,7 @@ def func(value: T.int32) -> T.bool: T.func_attr({"target": T.target("llvm")}) return value < 10 - built = tvm.build(func) + built = tvm.compile(func) assert isinstance(built(0), bool) assert built(0) diff --git a/tests/python/codegen/test_target_codegen_metal.py b/tests/python/codegen/test_target_codegen_metal.py index b4e747a0b4d8..2d669081e347 100644 --- a/tests/python/codegen/test_target_codegen_metal.py +++ b/tests/python/codegen/test_target_codegen_metal.py @@ -36,7 +36,7 @@ def check_inf_nan(dev, n, value, dtype): sch = tvm.tir.Schedule(prim_func) (x,) = sch.get_loops(sch.get_block("C")) sch.bind(x, "threadIdx.x") - fun = tvm.build(sch.mod, target=target) + fun = tvm.compile(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 @@ -72,7 +72,7 @@ def main(A: T.Buffer((2, 3), "float32"), B: T.Buffer((6,), "float32")): a = (np.arange(6).reshape(2, 3)).astype("float32") a_nd = tvm.nd.array(a, dev) b_nd = tvm.nd.empty((6,), "float32", dev) - f = tvm.build(IRModule, target=target) + f = tvm.compile(IRModule, target=target) f(a_nd, b_nd) np.testing.assert_allclose(b_nd.numpy(), a.reshape(6), atol=1e-5, rtol=1e-5) @@ -89,7 +89,7 @@ def check_erf(dev, n, dtype): sch = tvm.tir.Schedule(func) (x,) = sch.get_loops(sch.get_block("C")) sch.bind(x, "threadIdx.x") - fun = tvm.build(sch.mod, target=target) + fun = tvm.compile(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 @@ -117,7 +117,7 @@ def main(A: T.Buffer((1, 2), "int32")): r = T.ramp(tx, 3, 2) A[0, T.ramp(0, 1, 2)] = r - f = tvm.build(IRModule, target=target) + f = tvm.compile(IRModule, target=target) dev = tvm.metal() a_nd = tvm.nd.empty((1, 2), "int32", dev) f(a_nd) @@ -143,7 +143,7 @@ def main(A: T.Buffer((6), "float32"), B: T.Buffer((6,), "float32")): a = np.arange(6).astype("float32") a_nd = tvm.nd.array(a, dev) b_nd = tvm.nd.empty((6,), "float32", dev) - f = tvm.build(IRModule, target=target) + f = tvm.compile(IRModule, target=target) f(a_nd, b_nd) a.reshape(3, 2)[:, 1] = 0 np.testing.assert_allclose(b_nd.numpy(), a, atol=1e-5, rtol=1e-5) @@ -164,7 +164,7 @@ def func(A: T.Buffer((16), "uint8"), B: T.Buffer((16), "float32")): a = np.arange(16).astype("uint8") a_nd = tvm.nd.array(a, dev) b_nd = tvm.nd.empty((16,), "float32", dev) - f = tvm.build(func, target="metal") + f = tvm.compile(func, target="metal") f(a_nd, b_nd) np.testing.assert_allclose(b_nd.numpy(), a.astype("float32"), atol=1e-5, rtol=1e-5) @@ -186,7 +186,7 @@ def compile_metal(src, target): mod = tvm.IRModule({"main": func}) - f = tvm.build(mod, target="metal") + f = tvm.compile(mod, target="metal") src: str = f.imported_modules[0].get_source() occurrences = src.count("struct func_kernel_args_t") assert occurrences == 1, occurrences diff --git a/tests/python/codegen/test_target_codegen_opencl.py b/tests/python/codegen/test_target_codegen_opencl.py index 90af959472c5..cbdb60477b06 100644 --- a/tests/python/codegen/test_target_codegen_opencl.py +++ b/tests/python/codegen/test_target_codegen_opencl.py @@ -14,12 +14,11 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. +import re + import tvm -from tvm import te import tvm.testing -import re -import pytest -import numpy as np +from tvm import te target = "opencl" @@ -39,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.build(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 @@ -56,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.build(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) @@ -86,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.build(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 @@ -114,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.build(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) @@ -139,7 +138,7 @@ def check_erf(dev, n, dtype): sch = tvm.tir.Schedule(func) (x,) = sch.get_loops(sch.get_block("C")) sch.bind(x, "threadIdx.x") - fun = tvm.build(sch.mod, target=target) + fun = tvm.tir.build(sch.mod, target=target) source_str = fun.imported_modules[0].get_source() matches = re.findall("erf", source_str) @@ -179,7 +178,7 @@ def check_type_casting(ctx, n, dtype): sch.bind(tx, "threadIdx.x") sch.vectorize(vx) - fun = tvm.build(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)))))" @@ -211,7 +210,7 @@ def _check(target, n, dtype): (x,) = sch.get_loops(sch.get_block("C")) sch.bind(x, "threadIdx.x") - fun = tvm.build(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" @@ -226,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 diff --git a/tests/python/codegen/test_target_codegen_riscv.py b/tests/python/codegen/test_target_codegen_riscv.py index 2f3aae29c343..b06aeb4ced06 100644 --- a/tests/python/codegen/test_target_codegen_riscv.py +++ b/tests/python/codegen/test_target_codegen_riscv.py @@ -15,6 +15,7 @@ # specific language governing permissions and limitations # under the License. import tvm +import tvm.testing from tvm.script import tir as T from tvm.target.codegen import target_has_features @@ -33,7 +34,7 @@ def load_vec(A: T.Buffer((N,), "int8")): for j in T.vectorized(0, extent): A[j] = 1 - f = tvm.build(load_vec, target) + f = tvm.tir.build(load_vec, target) # Check RVV `vsetvli` prensence assembly = f.get_source("asm") if target_has_features("v"): @@ -46,4 +47,4 @@ def load_vec(A: T.Buffer((N,), "int8")): if __name__ == "__main__": - test_rvv() + tvm.testing.main() diff --git a/tests/python/codegen/test_target_codegen_rocm.py b/tests/python/codegen/test_target_codegen_rocm.py index 4c7592034ef0..a89d71f2be48 100644 --- a/tests/python/codegen/test_target_codegen_rocm.py +++ b/tests/python/codegen/test_target_codegen_rocm.py @@ -31,7 +31,7 @@ def check_inf_nan(dev, n, value, dtype): xo, xi = sch.split(sch.get_loops("C")[0], factors=[None, 128]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, "rocm") + fun = tvm.compile(sch.mod, "rocm") a = tvm.nd.empty((n,), A.dtype, dev) c = tvm.nd.empty((n,), A.dtype, dev) # Only need to test compiling here @@ -76,7 +76,7 @@ def check_rocm(dtype, n, lanes): xo, xi = sch.split(sch.get_loops("B")[0], factors=[None, 4]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - fun = tvm.build(sch.mod, target="rocm") + fun = tvm.compile(sch.mod, target="rocm") dev = tvm.rocm(0) a = tvm.nd.empty((n,), A.dtype, dev).copyfrom(np.random.uniform(size=(n, lanes))) @@ -107,7 +107,7 @@ def func( A_local[0] = T.tvm_warp_shuffle(mask[0], A_local[0], 0, 32, 32) A[tx] = A_local[0] - mod = tvm.build(func, target="rocm") + mod = tvm.compile(func, target="rocm") dev = tvm.rocm(0) a = tvm.nd.array(np.random.uniform(size=(32,)).astype("float32"), dev) mod(a) @@ -130,7 +130,7 @@ def func( for i in T.vectorized(0, 4): B[i] = T.exp2(A[i]) - mod = tvm.build(func, target="rocm") + mod = tvm.compile(func, target="rocm") dev = tvm.rocm(0) a = tvm.nd.array(np.ones((4,)).astype("float32"), dev) b = tvm.nd.array(np.zeros((4,)).astype("float32"), dev) diff --git a/tests/python/codegen/test_target_codegen_vulkan.py b/tests/python/codegen/test_target_codegen_vulkan.py index 0e1aa1a0403b..b661ce486981 100644 --- a/tests/python/codegen/test_target_codegen_vulkan.py +++ b/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.build(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 @@ -121,7 +121,7 @@ def test_array_vectorize_add(target, dev, dtype): xo, xi = sch.split(sch.get_loops("B")[0], factors=[None, 4]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - f = tvm.build(sch.mod, target=target) + f = tvm.compile(sch.mod, target=target) a = tvm.nd.empty((arr_size,), A.dtype, dev).copyfrom(np.random.uniform(size=(arr_size, lanes))) c = tvm.nd.empty((arr_size,), B.dtype, dev) @@ -142,7 +142,7 @@ def test_vulkan_bool_load(target, dev): sch.bind(xi, "threadIdx.x") # Build - f = tvm.build(sch.mod, target=target) + f = tvm.compile(sch.mod, target=target) a_np = np.random.uniform(size=arr_size) > 0.5 b_np = np.zeros((arr_size,), dtype="int32") @@ -194,7 +194,7 @@ def test_vulkan_constant_passing(target, dev, vulkan_parameter_impl, vulkan_para xo, xi = sch.split(sch.get_loops("B")[0], factors=[None, 64]) sch.bind(xo, "blockIdx.x") sch.bind(xi, "threadIdx.x") - f_add = tvm.build(sch.mod, target=target) + f_add = tvm.compile(sch.mod, target=target) n = 1024 scalars = np.array([1 for _ in scalars]).astype(dtype) @@ -242,7 +242,7 @@ def do_compute(A, B, n): sch = tir.Schedule(mod) # Build - func = tvm.build(sch.mod, target=target) + func = tvm.compile(sch.mod, target=target) a = tvm.nd.array(np.array([5], dtype=A.dtype), dev) b = tvm.nd.array(np.zeros(n, dtype=A.dtype), dev) @@ -290,7 +290,7 @@ def do_compute(A, B, n): sch = tir.Schedule(mod) # Build - func = tvm.build(sch.mod, target=target) + func = tvm.compile(sch.mod, target=target) n = 32 a_np = np.arange(n).astype(dtype=A.dtype) @@ -383,7 +383,7 @@ def do_compute(ins, outs): return tvm.IRModule.from_expr(te.create_prim_func([A, R, B])) def test_ramp_broadcast_index(self, target, dev, mod, ref_data): - f = tvm.build(mod, target=target) + f = tvm.compile(mod, target=target) a_np, reorder_np, b_np = ref_data a = tvm.nd.array(a_np, dev) @@ -424,7 +424,7 @@ def func(A: T.Buffer((N, 2), "int32")): sch.bind(sch.get_loops("A")[0], "threadIdx.x") func = sch.mod["main"] - built = tvm.build(func, target=target) + built = tvm.compile(func, target=target) a_dev = tvm.nd.empty([N, 2], "int32", dev) built(a_dev) @@ -534,7 +534,7 @@ def tensorize_load(block, dim): tgt_attrs = tvm.target.Target(target).attrs if tgt_attrs.get("supports_cooperative_matrix"): - f = tvm.build(sch.mod, target=target) + f = tvm.compile(sch.mod, target=target) dev = tvm.device(target, 0) diff --git a/tests/python/codegen/test_target_codegen_x86.py b/tests/python/codegen/test_target_codegen_x86.py index f433964f7f5d..51d648f2c4a9 100644 --- a/tests/python/codegen/test_target_codegen_x86.py +++ b/tests/python/codegen/test_target_codegen_x86.py @@ -14,11 +14,10 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -import numpy as np import platform -import pytest import re -import textwrap + +import pytest import tvm from tvm import te @@ -40,7 +39,7 @@ def fp16_to_fp32(target, width, match=None, not_match=None): B = te.compute(A.shape, lambda *i: A(*i).astype("float32"), name="B") sch = tvm.tir.Schedule(te.create_prim_func([A, B])) sch.vectorize(sch.get_loops("B")[1]) - f = tvm.build(sch.mod, target=target) + f = tvm.tir.build(sch.mod, target=target) assembly = f.get_source("asm").splitlines() if match: diff --git a/tests/python/contrib/test_cblas.py b/tests/python/contrib/test_cblas.py index b8851e685b13..c0e1553ea782 100644 --- a/tests/python/contrib/test_cblas.py +++ b/tests/python/contrib/test_cblas.py @@ -63,7 +63,7 @@ def verify(target="llvm"): return dev = tvm.cpu(0) name = "test_matmul_add" - f = tvm.build( + f = tvm.compile( te.create_prim_func([input1_data, input2_data, final_result, bias]).with_attr( "global_symbol", name ), @@ -146,7 +146,7 @@ def verify(target="llvm"): print("skip because extern function is not available") return dev = tvm.cpu(0) - f = tvm.build( + f = tvm.compile( te.create_prim_func([input1_data, input2_data, final_result, bias]), target=target ) matrix_input1 = tvm.nd.array( @@ -230,7 +230,9 @@ def verify(target="llvm"): return dev = tvm.cpu(0) name = "test_batch_matmul" - f = tvm.build(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) diff --git a/tests/python/contrib/test_dlpack.py b/tests/python/contrib/test_dlpack.py index 70277cb0ca0a..421853899979 100644 --- a/tests/python/contrib/test_dlpack.py +++ b/tests/python/contrib/test_dlpack.py @@ -51,7 +51,7 @@ def verify_torch_dlpack(): ZZ = te.compute((n, n), lambda i, j: te.sum(XX[i, k] * YY[k, j], axis=k)) # No need to speficy target_host if it's llvm # Otherwise you will need to specify the target and target_host - f = tvm.build(te.create_prim_func([XX, YY, ZZ])) + f = tvm.compile(te.create_prim_func([XX, YY, ZZ])) f_pytorch = to_pytorch_func(f) zz2 = torch.empty(137, 137) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 0f8ca9789a64..8d3ed034b5ea 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -94,7 +94,7 @@ def build_and_run(inputs, func, target: str, target_host: str, *args, **kwargs): """build and run the function func""" schedule, placeholders, binds = func(*args, **kwargs) - func = tvm.build( + func = tvm.compile( schedule, placeholders, target=tvm.target.Target(target, host=target_host), binds=binds ) dev = tvm.device(target) diff --git a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py index 169d868b5479..1f953aa60125 100644 --- a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py +++ b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py @@ -301,14 +301,14 @@ def test_build(self, schedule_args, target_host, input_layout, working_layout, o if uses_2d_memory and not is_hexagon: stack.enter_context(pytest.raises(tvm.TVMError)) - tvm.build(*schedule_args, target=target_host) + tvm.compile(*schedule_args, target=target_host) @tvm.testing.fixture def runtime_module(self, schedule_args, target_host): if target_host.kind.name != "hexagon": pytest.skip("Only running on hexagon") - return tvm.build(*schedule_args, target=target_host) + return tvm.compile(*schedule_args, target=target_host) @tvm.testing.requires_hexagon def test_execute( diff --git a/tests/python/contrib/test_hexagon/test_async_dma_pipeline.py b/tests/python/contrib/test_hexagon/test_async_dma_pipeline.py index efe4920ec0b2..89a1c606bed5 100644 --- a/tests/python/contrib/test_hexagon/test_async_dma_pipeline.py +++ b/tests/python/contrib/test_hexagon/test_async_dma_pipeline.py @@ -277,7 +277,7 @@ def evaluate( "tir.experimental_dma_bypass_cache": 1, } ): - func_tir = tvm.build( + func_tir = tvm.compile( sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) ) module = hexagon_session.load_module(func_tir) diff --git a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py index f0cefa3fe256..346197443e90 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -204,7 +204,7 @@ def _benchmark_hexagon_elementwise_add_kernel( input2 = tvm.te.placeholder(shape, dtype=dtype) output = tvm.te.placeholder(shape, dtype=dtype) - built_module: tvm.driver.build_module.OperatorModule = tvm.build( + built_module: tvm.driver.build_module.OperatorModule = tvm.compile( ns_tir_module, [ input1, diff --git a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py index 9224568407e1..b6b1f8fa73d6 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py @@ -253,7 +253,7 @@ def test_maxpool2d_nhwc( block="tensor", buffer="placeholder", index_map=_int8_nhwc_8h8w32c_map ) - built_module = tvm.build( + built_module = tvm.compile( sch.mod, target=get_hexagon_target("v69"), ) diff --git a/tests/python/contrib/test_hexagon/test_dma_builtin.py b/tests/python/contrib/test_hexagon/test_dma_builtin.py index 86be640689c0..2dc5345b3e36 100644 --- a/tests/python/contrib/test_hexagon/test_dma_builtin.py +++ b/tests/python/contrib/test_hexagon/test_dma_builtin.py @@ -154,7 +154,7 @@ def test_vtcm_alloc_compute(self, hexagon_launcher, mode, module): target_hexagon = tvm.target.hexagon("v69") target = tvm.target.Target(target_hexagon, host=target_hexagon) with tvm.transform.PassContext(opt_level=3, config=[]): - ex = relax.build(mod=module, target=target, exec_mode=mode) + ex = tvm.compile(mod=module, target=target, exec_mode=mode) with hexagon_launcher.create_session() as session: dev = session.device input_arg0_data = np.random.randint(0, 9, size=(12800,), dtype=data_type) diff --git a/tests/python/contrib/test_hexagon/test_memory_alloc.py b/tests/python/contrib/test_hexagon/test_memory_alloc.py index a0e3255a5428..ae086a2b4a02 100644 --- a/tests/python/contrib/test_hexagon/test_memory_alloc.py +++ b/tests/python/contrib/test_hexagon/test_memory_alloc.py @@ -57,7 +57,7 @@ class TestMemoryAlloc: def test_global_axis_separator(self, hexagon_session, shape, dtype, scope, axis_separators): """Test with global axis separator.""" - mod1 = tvm.build( + mod1 = tvm.compile( generated_func(shape, dtype, axis_separators), target=get_hexagon_target("v69"), ) diff --git a/tests/python/contrib/test_hexagon/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index c0c7355a9afa..9f5ca587f8f1 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -156,7 +156,7 @@ def schedule_dense(sch, block, m_size, do_tune): def verify_dense(sch, target, m_size, n_size, k_size, hexagon_session): """Verify dense operator.""" - f = tvm.build(sch.mod["main"], target=target) + f = tvm.compile(sch.mod["main"], target=target) mod = hexagon_session.load_module(f) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/test_parallel_hvx.py b/tests/python/contrib/test_hexagon/test_parallel_hvx.py index 13ad36278bc6..679321aba7fb 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_hvx.py +++ b/tests/python/contrib/test_hexagon/test_parallel_hvx.py @@ -144,7 +144,7 @@ def evaluate(hexagon_session, shape_dtypes, expected_output_producer, sch): """Evaluate schedule.""" a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes - func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v68")) + func_tir = tvm.compile(sch.mod["main"], target=get_hexagon_target("v68")) module = hexagon_session.load_module(func_tir) a = np.random.randint(0, 16, a_shape, dtype=a_dtype) diff --git a/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py b/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py index ee7f789ed103..3c21e8a9f663 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py +++ b/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py @@ -319,7 +319,7 @@ def evaluate_result(operations, tag, time, result, expected_output): def setup_and_run(hexagon_session, sch, a, b, c, operations, mem_scope="global"): """Setup and run operator.""" - func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) + func_tir = tvm.compile(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device, mem_scope=mem_scope) @@ -340,7 +340,7 @@ def setup_and_run(hexagon_session, sch, a, b, c, operations, mem_scope="global") def setup_and_run_preallocated(hexagon_session, sch, a, b, c, operations): """Setup and run for preallocated.""" - func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) + func_tir = tvm.compile(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) a_vtcm = np.zeros((a.size), dtype="uint8") diff --git a/tests/python/contrib/test_hexagon/test_parallel_scalar.py b/tests/python/contrib/test_hexagon/test_parallel_scalar.py index 0ca8c6ba0c47..d76f2b53a112 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_scalar.py +++ b/tests/python/contrib/test_hexagon/test_parallel_scalar.py @@ -87,7 +87,7 @@ def evaluate(hexagon_session, operations, expected, sch): shape = operations dtype = "float64" - func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v68")) + func_tir = tvm.compile(sch.mod["main"], target=get_hexagon_target("v68")) module = hexagon_session.load_module(func_tir) # np.random.random returns float64 by default, but make the cast explicit diff --git a/tests/python/contrib/test_hexagon/test_relax_2d_buffer_allocation.py b/tests/python/contrib/test_hexagon/test_relax_2d_buffer_allocation.py index 62ece5c6b88c..42038b97f90e 100644 --- a/tests/python/contrib/test_hexagon/test_relax_2d_buffer_allocation.py +++ b/tests/python/contrib/test_hexagon/test_relax_2d_buffer_allocation.py @@ -79,7 +79,7 @@ def test_alloc_storage_with_scope_global(hexagon_launcher): target_hexagon = tvm.target.hexagon("v69", vtcm_capacity=4 * 2**20) target = tvm.target.Target(target_hexagon, host=target_hexagon) with tvm.transform.PassContext(opt_level=3): - lib = relax.build(mod, target, exec_mode="compiled") + lib = tvm.compile(mod, target, exec_mode="compiled") with hexagon_launcher.create_session() as session: dev = session.device diff --git a/tests/python/contrib/test_hexagon/test_relax_integration.py b/tests/python/contrib/test_hexagon/test_relax_integration.py index 7d90adbc959a..8b2960ab9e2a 100644 --- a/tests/python/contrib/test_hexagon/test_relax_integration.py +++ b/tests/python/contrib/test_hexagon/test_relax_integration.py @@ -50,7 +50,7 @@ def test_mobilenet_onnx(hexagon_session: Session): relax_mod = relay_translator.from_relay(relay_mod["main"], target_hexagon) # Compile and run on Hexagon. - exe = relax.build(relax_mod, target) + exe = tvm.compile(relax_mod, target) dev = hexagon_session.device vm_mod = hexagon_session.get_executor_from_factory(exe) @@ -62,7 +62,7 @@ def test_mobilenet_onnx(hexagon_session: Session): # Compile and run on LLVM for comparison. relax_mod = relay_translator.from_relay(relay_mod["main"], "llvm") - exe = relax.build(relax_mod, "llvm") + exe = tvm.compile(relax_mod, "llvm") dev = tvm.cpu() vm_rt = relax.VirtualMachine(exe, dev) data = tvm.nd.array(data_np, dev) @@ -84,7 +84,7 @@ def test_mobilenet(hexagon_session: Session): relax_mod = relay_translator.from_relay(relay_mod["main"], target, params) # Compile and run on Hexagon. - exe = relax.build(relax_mod, target) + exe = tvm.compile(relax_mod, target) dev = hexagon_session.device vm_mod = hexagon_session.get_executor_from_factory(exe) @@ -96,7 +96,7 @@ def test_mobilenet(hexagon_session: Session): # Compile and run on LLVM for comparison. relax_mod = relay_translator.from_relay(relay_mod["main"], "llvm", params) - exe = relax.build(relax_mod, "llvm") + exe = tvm.compile(relax_mod, "llvm") dev = tvm.cpu() vm_rt = relax.VirtualMachine(exe, dev) data = tvm.nd.array(data_np, dev) diff --git a/tests/python/contrib/test_hexagon/test_sigmoid.py b/tests/python/contrib/test_hexagon/test_sigmoid.py index 1247d9075972..c6196ce42517 100644 --- a/tests/python/contrib/test_hexagon/test_sigmoid.py +++ b/tests/python/contrib/test_hexagon/test_sigmoid.py @@ -92,7 +92,7 @@ def test_sigmoid( func_name = "sigmoid" with tvm.transform.PassContext(opt_level=3): - runtime_module = tvm.build(tir_s.mod, target=get_hexagon_target("v69")) + runtime_module = tvm.compile(tir_s.mod, target=get_hexagon_target("v69")) assert "hvx_sigmoid" in runtime_module.get_source("asm") assert "vmin" in runtime_module.get_source("asm") diff --git a/tests/python/contrib/test_hexagon/test_software_pipeline_async.py b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py index d45b35befd11..3be9683a7deb 100644 --- a/tests/python/contrib/test_hexagon/test_software_pipeline_async.py +++ b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py @@ -183,7 +183,7 @@ def test_async_software_pipeline( "tir.experimental_dma_bypass_cache": 1, } ): - func = tvm.build(schedule.mod["main"], target=get_hexagon_target("v68")) + func = tvm.compile(schedule.mod["main"], target=get_hexagon_target("v68")) with hexagon_launcher.create_session() as hexagon_session: dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/test_take.py b/tests/python/contrib/test_hexagon/test_take.py index 80c2b053395f..15058e17af5a 100644 --- a/tests/python/contrib/test_hexagon/test_take.py +++ b/tests/python/contrib/test_hexagon/test_take.py @@ -366,7 +366,7 @@ def test_value(): after = generate_take_op.PassReplaceWithTakeOpPrimFuncs()(before) target = tvm.target.Target("llvm", host="llvm") - ex = relax.build(after, target, exec_mode="compiled") + ex = tvm.compile(after, target, exec_mode="compiled") vm = relax.VirtualMachine(ex, tvm.cpu()) res = vm["main"](inp_quant) diff --git a/tests/python/contrib/test_hexagon/test_thread_pool.py b/tests/python/contrib/test_hexagon/test_thread_pool.py index 2fc82cf49984..2dc426749680 100644 --- a/tests/python/contrib/test_hexagon/test_thread_pool.py +++ b/tests/python/contrib/test_hexagon/test_thread_pool.py @@ -75,7 +75,7 @@ def benchmark_func(mod, name, args, hexagon_session): @tvm.testing.requires_hexagon def test_speedup(hexagon_session: Session, capsys): """Test speedup""" - func = tvm.build( + func = tvm.compile( ElemwiseSumIRModule, target=get_hexagon_target("v68"), ) @@ -91,7 +91,7 @@ def test_speedup(hexagon_session: Session, capsys): @tvm.testing.requires_hexagon def test_elemwise_sum_parallel(hexagon_session: Session): """Test parallel elementwise sum""" - func = tvm.build( + func = tvm.compile( ElemwiseSumIRModule, target=get_hexagon_target("v68"), ) diff --git a/tests/python/contrib/test_hexagon/test_vtcm.py b/tests/python/contrib/test_hexagon/test_vtcm.py index 524bcb859990..ea9feb740319 100644 --- a/tests/python/contrib/test_hexagon/test_vtcm.py +++ b/tests/python/contrib/test_hexagon/test_vtcm.py @@ -49,7 +49,7 @@ def test_vtcm_building(): """Test building with vtcm mem scope""" sch = get_scale_by_two_schedule() target = get_hexagon_target("v68") - built = tvm.build(sch.mod, target=target) + built = tvm.compile(sch.mod, target=target) assert "global.vtcm" in built.get_source("asm") @@ -69,18 +69,18 @@ def _raises_exception(f): target = get_hexagon_target("v68", vtcm_capacity=vtcm_capacity) assert ( - _raises_exception(lambda: tvm.build(sch.mod, target=target)) == limited + _raises_exception(lambda: tvm.compile(sch.mod, target=target)) == limited ), "Case 1 - arg. VTCM memory allocation limiter does not work correctly " with target: assert ( - _raises_exception(lambda: tvm.build(sch.mod)) == limited + _raises_exception(lambda: tvm.compile(sch.mod)) == limited ), "Case 2 - with.VTCM memory allocation limiter does not work correctly " with tvm.transform.PassContext(config={"tir.vtcm_capacity": vtcm_capacity}): assert ( _raises_exception( - lambda: tvm.build(sch.mod, target=get_hexagon_target("v68", vtcm_capacity=0)) + lambda: tvm.compile(sch.mod, target=get_hexagon_target("v68", vtcm_capacity=0)) ) == limited ), "Case 3 - context. VTCM memory allocation limiter does not work correctly " diff --git a/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py b/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py index 1ac01efcdefb..9e0b7920b454 100644 --- a/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py +++ b/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py @@ -95,7 +95,7 @@ def evaluate(hexagon_session, sch, size): """Evaluate schedule.""" a_shape = size - func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) + func_tir = tvm.compile(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) a = np.random.randint(-128, 127, a_shape, dtype="int8") diff --git a/tests/python/contrib/test_hipblas.py b/tests/python/contrib/test_hipblas.py index e5df51e62942..33187fa4efba 100644 --- a/tests/python/contrib/test_hipblas.py +++ b/tests/python/contrib/test_hipblas.py @@ -35,7 +35,7 @@ def verify(target="rocm"): print("skip because extern function is not available") return dev = tvm.rocm(0) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.compile(te.create_prim_func([A, B, C]), target=target) a = tvm.nd.array(np.random.uniform(0, 128, size=(n, l)).astype(A.dtype), dev) b = tvm.nd.array(np.random.uniform(0, 128, size=(l, m)).astype(B.dtype), dev) c = tvm.nd.array(np.zeros((n, m), dtype=C.dtype), dev) @@ -57,7 +57,7 @@ def verify_batch_matmul(Ashape, Bshape, Cshape, in_dtype, out_dtype, rtol=1e-5): C = hipblas.batch_matmul(A, B, dtype=out_dtype) dev = tvm.rocm(0) - f = tvm.build(te.create_prim_func([A, B, C]), target="rocm") + f = tvm.compile(te.create_prim_func([A, B, C]), target="rocm") if "int" in in_dtype: a = tvm.nd.array(np.random.uniform(1, 10, size=Ashape).astype(in_dtype), dev) diff --git a/tests/python/contrib/test_mps.py b/tests/python/contrib/test_mps.py index e876672feaed..41847f3b8fea 100644 --- a/tests/python/contrib/test_mps.py +++ b/tests/python/contrib/test_mps.py @@ -35,7 +35,7 @@ def verify(A, B, C): print("skip because extern function is not available") return dev = tvm.metal(0) - f = tvm.build(te.create_prim_func([A, B, C]), target="metal") + f = tvm.compile(te.create_prim_func([A, B, C]), target="metal") a = tvm.nd.array(np.random.uniform(size=(n, l)).astype(A.dtype), dev) b = tvm.nd.array(np.random.uniform(size=(l, m)).astype(B.dtype), dev) c = tvm.nd.array(np.zeros((n, m), dtype=C.dtype), dev) @@ -64,7 +64,7 @@ def verify(A, B, C, target="llvm"): print("skip because extern function is not available") return dev = tvm.metal(0) - f = tvm.build(te.create_prim_func([A, B, C]), target="metal") + f = tvm.compile(te.create_prim_func([A, B, C]), target="metal") a = tvm.nd.array(np.random.uniform(size=(n, h, w, ci)).astype(A.dtype), dev) b = tvm.nd.array(np.random.uniform(size=(co, kh, kw, ci)).astype(B.dtype), dev) c = tvm.nd.array(np.zeros((n, h // stride, w // stride, co), dtype=C.dtype), dev) diff --git a/tests/python/contrib/test_msc/test_plugin.py b/tests/python/contrib/test_msc/test_plugin.py index 5a033b2f6ecb..f4607a35bfa8 100644 --- a/tests/python/contrib/test_msc/test_plugin.py +++ b/tests/python/contrib/test_msc/test_plugin.py @@ -261,7 +261,7 @@ def _run_relax(relax_mod, target_name, data): else: device = tvm.cpu() with tvm.transform.PassContext(opt_level=3): - relax_exec = tvm.relax.build(relax_mod, target) + relax_exec = tvm.compile(relax_mod, target) runnable = tvm.relax.VirtualMachine(relax_exec, device) data = tvm.nd.array(data, device) return runnable["main"](data).asnumpy() diff --git a/tests/python/contrib/test_msc/test_translate_relax.py b/tests/python/contrib/test_msc/test_translate_relax.py index 2c1af75f9a33..41e8f0e44e64 100644 --- a/tests/python/contrib/test_msc/test_translate_relax.py +++ b/tests/python/contrib/test_msc/test_translate_relax.py @@ -53,7 +53,7 @@ def _tvm_runtime_to_np(obj): def _run_relax(relax_mod): relax_mod = tvm.relax.transform.LegalizeOps()(relax_mod) - relax_exec = tvm.relax.build(relax_mod, target) + relax_exec = tvm.compile(relax_mod, target) vm_runner = tvm.relax.VirtualMachine(relax_exec, dev) res = vm_runner["main"](*args) return _tvm_runtime_to_np(res) diff --git a/tests/python/contrib/test_msc/test_translate_tensorrt.py b/tests/python/contrib/test_msc/test_translate_tensorrt.py index 0e009c542ae1..ab0f8be148a3 100644 --- a/tests/python/contrib/test_msc/test_translate_tensorrt.py +++ b/tests/python/contrib/test_msc/test_translate_tensorrt.py @@ -44,7 +44,7 @@ def build_and_run(mod, inputs): with target: mod = tvm.tir.transform.DefaultGPUSchedule()(mod) with tvm.transform.PassContext(opt_level=3): - rt_mod = tvm.relax.build(mod, target) + rt_mod = tvm.compile(mod, target) runnable = tvm.relax.VirtualMachine(rt_mod, tvm.cuda()) res = runnable["main"](*inputs) if isinstance(res, tvm.runtime.NDArray): diff --git a/tests/python/contrib/test_random.py b/tests/python/contrib/test_random.py index be9fed2c6ee8..c8c8054dfb6b 100644 --- a/tests/python/contrib/test_random.py +++ b/tests/python/contrib/test_random.py @@ -39,7 +39,7 @@ def verify(target="llvm"): print("skip because extern function is not available") return dev = tvm.cpu(0) - f = tvm.build(te.create_prim_func([A]), target=target) + f = tvm.compile(te.create_prim_func([A]), target=target) a = tvm.nd.array(np.zeros((m, n), dtype=A.dtype), dev) f(a) na = a.numpy() @@ -64,7 +64,7 @@ def verify(target="llvm"): print("skip because extern function is not available") return dev = tvm.cpu(0) - f = tvm.build(te.create_prim_func([A]), target=target) + f = tvm.compile(te.create_prim_func([A]), target=target) a = tvm.nd.array(np.zeros((m, n), dtype=A.dtype), dev) f(a) na = a.numpy() @@ -89,7 +89,7 @@ def verify(target="llvm"): print("skip because extern function is not available") return dev = tvm.cpu(0) - f = tvm.build(te.create_prim_func([A]), target=target) + f = tvm.compile(te.create_prim_func([A]), target=target) a = tvm.nd.array(np.zeros((m, n), dtype=A.dtype), dev) f(a) na = a.numpy() diff --git a/tests/python/contrib/test_rocblas.py b/tests/python/contrib/test_rocblas.py index 2c1889a0c43b..a715a5bb4a74 100644 --- a/tests/python/contrib/test_rocblas.py +++ b/tests/python/contrib/test_rocblas.py @@ -39,7 +39,7 @@ def verify(target="rocm"): print("skip because extern function is not available") return dev = tvm.rocm(0) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.compile(te.create_prim_func([A, B, C]), target=target) a = tvm.nd.array(np.random.uniform(size=(n, l)).astype(A.dtype), dev) b = tvm.nd.array(np.random.uniform(size=(l, m)).astype(B.dtype), dev) c = tvm.nd.array(np.zeros((n, m), dtype=C.dtype), dev) @@ -72,7 +72,7 @@ def verify(target="rocm"): print("skip because extern function is not available") return dev = tvm.rocm(0) - f = tvm.build(te.create_prim_func([A, B, C]), target=target) + f = tvm.compile(te.create_prim_func([A, B, C]), target=target) a = tvm.nd.array(np.random.uniform(size=ashape).astype(A.dtype), dev) b = tvm.nd.array(np.random.uniform(size=bshape).astype(B.dtype), dev) c = tvm.nd.array(np.zeros((batch, m, n), dtype=C.dtype), dev) diff --git a/tests/python/contrib/test_sort.py b/tests/python/contrib/test_sort.py index 0e0aa71caf10..a853df569498 100644 --- a/tests/python/contrib/test_sort.py +++ b/tests/python/contrib/test_sort.py @@ -52,7 +52,7 @@ def test_sort(): dev = tvm.cpu(0) target = "llvm" - f = tvm.build(te.create_prim_func([data, sort_num, out]), target=target) + f = tvm.compile(te.create_prim_func([data, sort_num, out]), target=target) a = tvm.nd.array(np.array(input_data).astype(data.dtype), dev) b = tvm.nd.array(np.array(sort_num_input).astype(sort_num.dtype), dev) c = tvm.nd.array(np.zeros(a.shape, dtype=out.dtype), dev) @@ -80,7 +80,7 @@ def test_sort_np(): dev = tvm.cpu(0) target = "llvm" - f = tvm.build(te.create_prim_func([data, sort_num, out]), target=target) + f = tvm.compile(te.create_prim_func([data, sort_num, out]), target=target) np_data = np.random.uniform(size=dshape) np_out = np.argsort(np_data, axis=axis) diff --git a/tests/python/contrib/test_tir_triton_integration.py b/tests/python/contrib/test_tir_triton_integration.py index 522351f3dc55..69f727150544 100644 --- a/tests/python/contrib/test_tir_triton_integration.py +++ b/tests/python/contrib/test_tir_triton_integration.py @@ -114,6 +114,6 @@ def add(x_handle: T.handle, y_handle: T.handle, output_handle: T.handle): output_np = x_nd.numpy() + y_nd.numpy() with tvm.target.Target("cuda"): - lib = relax.build(Module) + lib = tvm.compile(Module) output_nd = tvm.runtime.relax_vm.VirtualMachine(lib, device)["main"](x_nd, y_nd) tvm.testing.assert_allclose(output_nd.numpy(), output_np, rtol=1e-5) diff --git a/tests/python/disco/test_callback.py b/tests/python/disco/test_callback.py index 3f8d5e9e525b..d0defa15b869 100644 --- a/tests/python/disco/test_callback.py +++ b/tests/python/disco/test_callback.py @@ -64,7 +64,7 @@ def transform_params( with tvm.target.Target("cuda"): mod = tvm.IRModule.from_expr(transform_params) mod = pipeline(mod) - built = tvm.relax.build(mod, "cuda") + built = tvm.compile(mod, "cuda") num_shards = 2 diff --git a/tests/python/disco/test_ccl.py b/tests/python/disco/test_ccl.py index c29ece957245..f93030f63923 100644 --- a/tests/python/disco/test_ccl.py +++ b/tests/python/disco/test_ccl.py @@ -484,7 +484,7 @@ def relax_build(mod, target): dl.gpu.GeneralReduction(), dl.gpu.Fallback(), )(mod) - return rx.build(mod, target=target) + return tvm.compile(mod, target=target) # pylint: disable=invalid-name X = np.random.randn(128, 128).astype("float32") @@ -623,7 +623,7 @@ def relax_build(mod, target): dl.gpu.GeneralReduction(), dl.gpu.Fallback(), )(mod) - return rx.build(mod, target=target) + return tvm.compile(mod, target=target) # pylint: disable=invalid-name X = np.random.randn(1, 10, 128).astype("float32") diff --git a/tests/python/disco/test_loader.py b/tests/python/disco/test_loader.py index b4e2440857e6..ba0287afc61a 100644 --- a/tests/python/disco/test_loader.py +++ b/tests/python/disco/test_loader.py @@ -270,7 +270,7 @@ def main( def relax_build(mod, target): with target: mod = rx.get_pipeline("zero")(mod) # pylint: disable=no-value-for-parameter - return rx.build(mod, target="cuda") + return tvm.compile(mod, target="cuda") target = Target( { diff --git a/tests/python/disco/test_session.py b/tests/python/disco/test_session.py index 38aa757bf8f1..83002e971cc2 100644 --- a/tests/python/disco/test_session.py +++ b/tests/python/disco/test_session.py @@ -220,7 +220,7 @@ def main(A: R.Tensor((8, 16), dtype="float32")) -> R.Tensor((16, 8), dtype="floa x_np = np.arange(8 * 16).astype("float32").reshape([8, 16]) y_np = x_np.transpose() - rx.build(TestMod, target="llvm").export_library(path) + tvm.compile(TestMod, target="llvm").export_library(path) mod = sess.load_vm_module(path, device=device) x_disc = _numpy_to_worker_0(sess, x_np, device=device) @@ -285,7 +285,7 @@ def transpose_2( x_np = np.arange(8 * 16).astype("float32").reshape([8, 16]) y_np = x_np.transpose() - rx.build(TestMod, target="llvm").export_library(path) + tvm.compile(TestMod, target="llvm").export_library(path) mod = sess.load_vm_module(path, device=device) x_disc = _numpy_to_worker_0(sess, x_np, device=device) diff --git a/tests/python/ir/test_datatype_nv_fp8.py b/tests/python/ir/test_datatype_nv_fp8.py index b812c70ab687..72cdfb469d43 100644 --- a/tests/python/ir/test_datatype_nv_fp8.py +++ b/tests/python/ir/test_datatype_nv_fp8.py @@ -81,7 +81,7 @@ def test_fp8_unary_op(np_dtype, dtype_str): """Skip test if ml_dtypes is not installed""" return - f = tvm.build(func, target="llvm") + f = tvm.compile(func, target="llvm") a = np.random.randn(128).astype(np_dtype) b = np.random.randn(128).astype(np_dtype) a_add_b = np.zeros(128).astype(np_dtype) diff --git a/tests/python/ir/test_pass_instrument.py b/tests/python/ir/test_pass_instrument.py index 718cf3a663e5..aea8492485ed 100644 --- a/tests/python/ir/test_pass_instrument.py +++ b/tests/python/ir/test_pass_instrument.py @@ -38,7 +38,7 @@ def func(a: T.handle, b: T.handle) -> None: B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * 2.0 with tvm.transform.PassContext(opt_level=3, instruments=[PrintBeforeAll(), PrintAfterAll()]): - tvm.build(func) + tvm.compile(func) all_passes_output = capsys.readouterr().out assert "Before Running Pass:" in all_passes_output assert "After Running Pass:" in all_passes_output diff --git a/tests/python/meta_schedule/test_meta_schedule_tune_tir.py b/tests/python/meta_schedule/test_meta_schedule_tune_tir.py index c8fc4a73f56b..0d349752537a 100644 --- a/tests/python/meta_schedule/test_meta_schedule_tune_tir.py +++ b/tests/python/meta_schedule/test_meta_schedule_tune_tir.py @@ -105,7 +105,7 @@ def test_tune_matmul_cuda(): @pytest.mark.skip("Integration test") def test_tune_run_module_via_rpc(): target = tvm.target.Target("llvm") - rt_mod = tvm.build(matmul, target) + rt_mod = tvm.compile(matmul, target) # construct the input input_data = {} diff --git a/tests/python/nightly/test_nnapi/infrastructure.py b/tests/python/nightly/test_nnapi/infrastructure.py index 9f3813ef6179..07aa0bd752df 100644 --- a/tests/python/nightly/test_nnapi/infrastructure.py +++ b/tests/python/nightly/test_nnapi/infrastructure.py @@ -89,7 +89,7 @@ def _build(mod, enable_nnapi): mod = partition_for_nnapi(mod) mod = tvm.relax.transform.RunCodegen()(mod) - ex = tvm.relax.build(mod, target="llvm -mtriple=aarch64-linux-android") + ex = tvm.compile(mod, target="llvm -mtriple=aarch64-linux-android") return ex diff --git a/tests/python/nightly/test_nnapi/test_network.py b/tests/python/nightly/test_nnapi/test_network.py index 742613c25c75..2f9863eb4ee8 100644 --- a/tests/python/nightly/test_nnapi/test_network.py +++ b/tests/python/nightly/test_nnapi/test_network.py @@ -34,7 +34,7 @@ def _build_and_run_network(remote_obj, tracker, mod, input_data): def execute_on_host(mod, inputs): with tvm.transform.PassContext(opt_level=3): - ex = tvm.relax.build(mod, target="llvm") + ex = tvm.compile(mod, target="llvm") dev = tvm.cpu(0) vm = tvm.relax.VirtualMachine(ex, device=dev) output = vm["main"](*inputs) diff --git a/tests/python/nightly/test_nnapi/test_ops.py b/tests/python/nightly/test_nnapi/test_ops.py index 589ff6ee89e7..31d584db1396 100644 --- a/tests/python/nightly/test_nnapi/test_ops.py +++ b/tests/python/nightly/test_nnapi/test_ops.py @@ -34,7 +34,7 @@ def _build_and_run_network(remote_obj, tracker, mod, input_data): def execute_on_host(mod, inputs): with tvm.transform.PassContext(opt_level=3): - ex = tvm.relax.build(mod, target="llvm") + ex = tvm.compile(mod, target="llvm") dev = tvm.cpu(0) vm = tvm.relax.VirtualMachine(ex, device=dev) output = vm["main"](*inputs) diff --git a/tests/python/relax/backend/clml/utils.py b/tests/python/relax/backend/clml/utils.py index 22b587c964ab..a7aace57ce11 100644 --- a/tests/python/relax/backend/clml/utils.py +++ b/tests/python/relax/backend/clml/utils.py @@ -42,7 +42,7 @@ def build_and_run( pipeline = relax.pipeline.get_default_pipeline(tgt) mod = pipeline(mod) if rpc: - ex = relax.build(mod, tgt) + ex = tvm.compile(mod, tgt) temp = utils.tempdir() path = temp.relpath(load_path) path = "./" + load_path @@ -52,7 +52,7 @@ def build_and_run( dev = rpc.cl(0) vm = relax.VirtualMachine(rexec, dev) else: - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.device(target, 0) vm = relax.VirtualMachine(ex, dev) diff --git a/tests/python/relax/nvshmem/test_runtime_builtin_kv_cache_transfer.py b/tests/python/relax/nvshmem/test_runtime_builtin_kv_cache_transfer.py index 5a3d65dbe0de..3e17f6436600 100644 --- a/tests/python/relax/nvshmem/test_runtime_builtin_kv_cache_transfer.py +++ b/tests/python/relax/nvshmem/test_runtime_builtin_kv_cache_transfer.py @@ -169,7 +169,7 @@ def set_global_func(head_dim, dtype): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.compile(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_backend_dispatch_sort_scan.py b/tests/python/relax/test_backend_dispatch_sort_scan.py index 4fe6de9e09c6..004050aaf892 100644 --- a/tests/python/relax/test_backend_dispatch_sort_scan.py +++ b/tests/python/relax/test_backend_dispatch_sort_scan.py @@ -426,7 +426,7 @@ def main(x: R.Tensor(("m", "n"), "int32")): np_cumsum = np.cumsum(np_data, axis=-1) with tvm.target.Target(target): mod = DispatchSortScan()(Module) - ex = tvm.relax.build(mod, target) + ex = tvm.compile(mod, target) vm = tvm.relax.VirtualMachine(ex, dev) tvm_data = tvm.nd.array(np_data, dev) cumsum = vm["main"](tvm_data) diff --git a/tests/python/relax/test_codegen_coreml.py b/tests/python/relax/test_codegen_coreml.py index f33f165e7d13..7b9c22b8b9d8 100644 --- a/tests/python/relax/test_codegen_coreml.py +++ b/tests/python/relax/test_codegen_coreml.py @@ -53,12 +53,12 @@ def verify(mod, inputs): mod1 = relax.transform.LegalizeOps()(mod1) assert relax.analysis.well_formed(mod1) - ex1 = relax.build(mod1, target=target) + ex1 = tvm.compile(mod1, target=target) vm1 = relax.VirtualMachine(ex1, dev, profile=True) out1 = vm1["main"](*inputs) mod2 = relax.transform.LegalizeOps()(mod) - ex2 = relax.build(mod2, target=target) + ex2 = tvm.compile(mod2, target=target) vm2 = relax.VirtualMachine(ex2, dev, profile=True) out2 = vm2["main"](*inputs) diff --git a/tests/python/relax/test_codegen_cublas.py b/tests/python/relax/test_codegen_cublas.py index c5514e272709..dbcb25b69d52 100644 --- a/tests/python/relax/test_codegen_cublas.py +++ b/tests/python/relax/test_codegen_cublas.py @@ -49,7 +49,7 @@ def build_and_run(mod, inputs_np, target, legalize=False, cuda_graph=False): "relax.transform.apply_legalize_ops": legalize, } ): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) f = vm["main"] inputs = [tvm.nd.array(inp, dev) for inp in inputs_np] diff --git a/tests/python/relax/test_codegen_cudnn.py b/tests/python/relax/test_codegen_cudnn.py index 0f9a0bc262a6..990f21138619 100644 --- a/tests/python/relax/test_codegen_cudnn.py +++ b/tests/python/relax/test_codegen_cudnn.py @@ -110,7 +110,7 @@ def build_and_run(mod, inputs_np, target, legalize=False, cuda_graph=False): "relax.transform.apply_legalize_ops": legalize, } ): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) f = vm["main"] inputs = [tvm.nd.array(inp, dev) for inp in inputs_np] diff --git a/tests/python/relax/test_codegen_cutlass.py b/tests/python/relax/test_codegen_cutlass.py index 9d31ced08c9d..88d3fdc57140 100644 --- a/tests/python/relax/test_codegen_cutlass.py +++ b/tests/python/relax/test_codegen_cutlass.py @@ -89,7 +89,7 @@ def build_and_run(mod, inputs_np, target, legalize=True, cuda_graph=False): "relax.transform.apply_legalize_ops": legalize, } ): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.device(target, 0) vm = relax.VirtualMachine(ex, dev) @@ -1477,7 +1477,7 @@ def main_residual( mod_transform, mod_deploy, transform_func_name = split_transform_deploy_mod(mod) - ex = relax.build(mod_transform, target="llvm") + ex = tvm.compile(mod_transform, target="llvm") vm = relax.vm.VirtualMachine(ex, tvm.cpu(0)) packed_weight, scales, bias_trans = vm[transform_func_name]( @@ -1485,7 +1485,7 @@ def main_residual( ) dev = tvm.device("cuda", 0) - ex = relax.build(mod_deploy, target="cuda") + ex = tvm.compile(mod_deploy, target="cuda") vm = relax.vm.VirtualMachine(ex, dev) x_nd = tvm.nd.array(x, dev) @@ -1630,7 +1630,7 @@ def main( mod_transform, mod_deploy, transform_func_name = split_transform_deploy_mod(mod) - ex = relax.build(mod_transform, target="llvm") + ex = tvm.compile(mod_transform, target="llvm") vm = relax.vm.VirtualMachine(ex, tvm.cpu(0)) packed_weight, scales, bias_trans = vm[transform_func_name]( @@ -1638,7 +1638,7 @@ def main( ) dev = tvm.device("cuda", 0) - ex = relax.build(mod_deploy, target="cuda") + ex = tvm.compile(mod_deploy, target="cuda") vm = relax.vm.VirtualMachine(ex, dev) x_nd = tvm.nd.array(x, dev) @@ -1906,13 +1906,13 @@ def main( mod_transform, mod_deploy, transform_func_name = split_transform_deploy_mod(mod) - ex = relax.build(mod_transform, target="llvm") + ex = tvm.compile(mod_transform, target="llvm") vm = relax.vm.VirtualMachine(ex, tvm.cpu(0)) packed_weight, scales = vm[transform_func_name]((tvm.nd.array(y),)) dev = tvm.device("cuda", 0) - ex = relax.build(mod_deploy, target="cuda") + ex = tvm.compile(mod_deploy, target="cuda") vm = relax.vm.VirtualMachine(ex, dev) x_nd = tvm.nd.array(x, dev) @@ -2061,13 +2061,13 @@ def main( mod_transform, mod_deploy, transform_func_name = split_transform_deploy_mod(mod) - ex = relax.build(mod_transform, target="llvm") + ex = tvm.compile(mod_transform, target="llvm") vm = relax.vm.VirtualMachine(ex, tvm.cpu(0)) packed_weight, scales = vm[transform_func_name]((tvm.nd.array(y),)) dev = tvm.device("cuda", 0) - ex = relax.build(mod_deploy, target="cuda") + ex = tvm.compile(mod_deploy, target="cuda") vm = relax.vm.VirtualMachine(ex, dev) x_nd = tvm.nd.array(x, dev) diff --git a/tests/python/relax/test_codegen_dnnl.py b/tests/python/relax/test_codegen_dnnl.py index fe4590f85a12..370c5f03a486 100644 --- a/tests/python/relax/test_codegen_dnnl.py +++ b/tests/python/relax/test_codegen_dnnl.py @@ -57,7 +57,7 @@ def build_and_run(mod, inputs, legalize=False): inputs = [tvm.nd.array(inp, dev) for inp in inputs] with tvm.transform.PassContext(config={"relax.transform.apply_legalize_ops": legalize}): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) f = vm["main"] return f(*inputs).numpy() diff --git a/tests/python/relax/test_codegen_hipblas.py b/tests/python/relax/test_codegen_hipblas.py index 7edbed61bc96..004e70e4e60e 100644 --- a/tests/python/relax/test_codegen_hipblas.py +++ b/tests/python/relax/test_codegen_hipblas.py @@ -42,7 +42,7 @@ def reset_seed(): def build_and_run(mod, inputs_np, target, legalize=False): dev = tvm.device(target, 0) with tvm.transform.PassContext(config={"relax.transform.apply_legalize_ops": legalize}): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) f = vm["main"] inputs = [tvm.nd.array(inp, dev) for inp in inputs_np] diff --git a/tests/python/relax/test_codegen_tensorrt.py b/tests/python/relax/test_codegen_tensorrt.py index 009bb24c63b8..746f4eba6028 100644 --- a/tests/python/relax/test_codegen_tensorrt.py +++ b/tests/python/relax/test_codegen_tensorrt.py @@ -64,7 +64,7 @@ def main( def build_and_run(mod, inputs_np, target, legalize=False): dev = tvm.device(target, 0) with tvm.transform.PassContext(config={"relax.transform.apply_legalize_ops": legalize}): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) f = vm["main"] inputs = [tvm.nd.array(inp, dev) for inp in inputs_np] diff --git a/tests/python/relax/test_codegen_tir_cutlass.py b/tests/python/relax/test_codegen_tir_cutlass.py index a14ca7ac36d7..9670f1598670 100644 --- a/tests/python/relax/test_codegen_tir_cutlass.py +++ b/tests/python/relax/test_codegen_tir_cutlass.py @@ -54,19 +54,13 @@ def build(mod): mod = relax.transform.FuseTIR()(mod) mod = relax.transform.SplitCallTIRByPattern(get_tir_pattern(), cutlass_fcodegen())(mod) mod = relax.transform.DeadCodeElimination()(mod) - print(mod.script()) - f = tempfile.NamedTemporaryFile(suffix=".so", delete=True) - executable = relax_build(mod, target) - - executable.mod.export_library(f.name, **compile_options(target)) - rt_mod = runtime.load_module(f.name) - f.close() - return rt_mod + executable = tvm.compile(mod, target) + return executable.jit(**compile_options(target)) def build_and_run_reference(mod, inputs_np): dev = tvm.device("llvm", 0) - ex = relax.build(mod, "llvm") + ex = tvm.compile(mod, "llvm") vm = relax.VirtualMachine(ex, dev) f = vm["main"] inputs = [tvm.nd.array(inp, dev) for inp in inputs_np] diff --git a/tests/python/relax/test_contrib_vllm.py b/tests/python/relax/test_contrib_vllm.py index f3c4839133e3..0a8d338a455e 100644 --- a/tests/python/relax/test_contrib_vllm.py +++ b/tests/python/relax/test_contrib_vllm.py @@ -43,7 +43,7 @@ def build_and_run(mod, inputs_np, target, legalize=True): mod = tvm.tir.transform.DefaultGPUSchedule()(mod) with tvm.transform.PassContext(): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.device(target, 0) vm = relax.VirtualMachine(ex, dev) diff --git a/tests/python/relax/test_dataflow_inplace.py b/tests/python/relax/test_dataflow_inplace.py index cd6e285de499..754b7a669feb 100644 --- a/tests/python/relax/test_dataflow_inplace.py +++ b/tests/python/relax/test_dataflow_inplace.py @@ -531,7 +531,7 @@ def main( expected = np.zeros((2, 3), dtype="float32") target = tvm.target.Target("llvm") - ex = relax.build(new_mod, target) + ex = tvm.compile(new_mod, target) vm = relax.VirtualMachine(ex, tvm.cpu()) res = vm["main"](x, y) assert (expected == res.numpy()).all() @@ -614,7 +614,7 @@ def main( expected = np.zeros((2, 3), dtype="float32") target = tvm.target.Target("llvm") - ex = relax.build(new_mod, target) + ex = tvm.compile(new_mod, target) vm = relax.VirtualMachine(ex, tvm.cpu()) res = vm["main"](x, y) assert (expected == res.numpy()).all() diff --git a/tests/python/relax/test_dataflow_pattern.py b/tests/python/relax/test_dataflow_pattern.py index ab6c6df31bdc..9be0d761c11f 100644 --- a/tests/python/relax/test_dataflow_pattern.py +++ b/tests/python/relax/test_dataflow_pattern.py @@ -1240,7 +1240,7 @@ def expected( mod = tvm.IRModule() mod["main"] = rewritten - rx.build(mod, target="llvm") + tvm.compile(mod, target="llvm") def test_combine_transposed_matmul_twice(): @@ -1336,7 +1336,7 @@ def rewriter(matchings, _): mod["main"] = rewritten print(mod) - rx.build(mod, target="llvm") + tvm.compile(mod, target="llvm") def test_commutative_pattern_match(): diff --git a/tests/python/relax/test_e2e_op_dynamic.py b/tests/python/relax/test_e2e_op_dynamic.py index 641469172f97..9179802360b3 100644 --- a/tests/python/relax/test_e2e_op_dynamic.py +++ b/tests/python/relax/test_e2e_op_dynamic.py @@ -28,7 +28,7 @@ def build(mod): - exe = relax.build(mod, target=target) + exe = tvm.compile(mod, target=target) return relax.VirtualMachine(exe, dev) diff --git a/tests/python/relax/test_frontend_nn_extern_module.py b/tests/python/relax/test_frontend_nn_extern_module.py index 6eaf1fbfc805..ef97cfd9056c 100644 --- a/tests/python/relax/test_frontend_nn_extern_module.py +++ b/tests/python/relax/test_frontend_nn_extern_module.py @@ -190,7 +190,7 @@ def test_sym(self, a: nn.Tensor, b: nn.Tensor): # pylint: disable=invalid-name _check_ir_equality(mod) mod = AttachExternModules(ext_mods)(mod) # pylint: disable=not-callable compiled = tvm.runtime.relax_vm.VirtualMachine( - relax.build(mod, target="llvm"), + tvm.compile(mod, target="llvm"), device=tvm.cpu(), ) _test_scalar_add(compiled["scalar_add"]) @@ -239,7 +239,7 @@ def test_sym(self, a: nn.Tensor, b: nn.Tensor): # pylint: disable=invalid-name _check_ir_equality(mod) mod = AttachExternModules(ext_mods)(mod) # pylint: disable=not-callable compiled = tvm.runtime.relax_vm.VirtualMachine( - relax.build(mod, target="llvm"), + tvm.compile(mod, target="llvm"), device=tvm.cpu(), ) _test_scalar_add(compiled["scalar_add"]) diff --git a/tests/python/relax/test_frontend_nn_op.py b/tests/python/relax/test_frontend_nn_op.py index 6a337b34c114..682f805026cb 100644 --- a/tests/python/relax/test_frontend_nn_op.py +++ b/tests/python/relax/test_frontend_nn_op.py @@ -855,7 +855,7 @@ def test(self): return result irmodule, _ = Model().export_tvm(spec={"test": {}}, debug=True) - ex = relax.build(irmodule, "llvm") + ex = tvm.compile(irmodule, "llvm") vm = relax.VirtualMachine(ex, tvm.cpu()) effects = vm["_initialize_effect"]() vm["test"](*effects) @@ -912,7 +912,7 @@ def foo(prob: R.Tensor((3, 5), dtype="float32"), uniform_sample: R.Tensor((6, 1) with target: mod = relax.backend.DispatchSampling()(mod) mod = tir.transform.DefaultGPUSchedule()(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.device(str(target), 0) vm = relax.VirtualMachine(ex, dev) @@ -1044,7 +1044,7 @@ def foo(prob: R.Tensor((2, 3), dtype="float32"), index: R.Tensor((2, 3), dtype=" with target: mod = tir.transform.DefaultGPUSchedule()(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.cuda(0) vm = relax.VirtualMachine(ex, dev) @@ -1160,7 +1160,7 @@ def foo(prob: R.Tensor((2, 3), dtype="float32"), sorted_prob: R.Tensor((2, 3), d mod = relax.transform.LegalizeOps()(mod) mod = tir.transform.DefaultGPUSchedule()(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) dev = tvm.cuda(0) vm = relax.VirtualMachine(ex, dev) diff --git a/tests/python/relax/test_frontend_onnx.py b/tests/python/relax/test_frontend_onnx.py index 16445a7914c9..10c185ae09d6 100644 --- a/tests/python/relax/test_frontend_onnx.py +++ b/tests/python/relax/test_frontend_onnx.py @@ -133,7 +133,7 @@ def check_correctness( tvm_model, params = relax.frontend.detach_params(tvm_model) # Compile the relax graph into a VM then run. with tvm.transform.PassContext(opt_level=3): - ex = relax.build(tvm_model, target="llvm") + ex = tvm.compile(tvm_model, target="llvm") vm = relax.VirtualMachine(ex, tvm.cpu()) # Prepare inputs. input_list = [ diff --git a/tests/python/relax/test_frontend_stablehlo.py b/tests/python/relax/test_frontend_stablehlo.py index 63defaf7d73b..4f049555f148 100644 --- a/tests/python/relax/test_frontend_stablehlo.py +++ b/tests/python/relax/test_frontend_stablehlo.py @@ -119,7 +119,7 @@ def check_correctness( # "llvm" should be good for this check target = tvm.target.Target("llvm", host="llvm") # Compile and run - ex = relax.build(ir_mod, target) + ex = tvm.compile(ir_mod, target) vm = relax.VirtualMachine(ex, tvm.cpu()) vm.set_input("main", *inputs_np) vm.invoke_stateful("main") @@ -156,7 +156,7 @@ def get_vm_res( """ target = tvm.target.Target("llvm", host="llvm") # Compile and run - ex = relax.build(ir_mod, target) + ex = tvm.compile(ir_mod, target) vm = relax.VirtualMachine(ex, tvm.cpu()) vm.set_input("main", *weights) vm.invoke_stateful("main") diff --git a/tests/python/relax/test_op_gradient_numeric.py b/tests/python/relax/test_op_gradient_numeric.py index acf0f615dd94..710c98b46c84 100644 --- a/tests/python/relax/test_op_gradient_numeric.py +++ b/tests/python/relax/test_op_gradient_numeric.py @@ -122,7 +122,7 @@ def _is_call_no_grad(expr): out = forward_bb.emit_output(call) forward_bb.emit_func_output(out) forward_mod = forward_bb.get() - forward_ex = relax.build(forward_mod, target) + forward_ex = tvm.compile(forward_mod, target) forward_vm = relax.VirtualMachine(forward_ex, dev) # Generate weights @@ -186,7 +186,7 @@ def forward(*inputs): grad_bb.emit_func_output(out) grad_mod = grad_bb.get() - grad_ex = relax.build(grad_mod, target) + grad_ex = tvm.compile(grad_mod, target) grad_vm = relax.VirtualMachine(grad_ex, dev) # tvm.runtime.NDArray inputs diff --git a/tests/python/relax/test_op_inspect.py b/tests/python/relax/test_op_inspect.py index ca4b0fc440be..f8326db8ddcf 100644 --- a/tests/python/relax/test_op_inspect.py +++ b/tests/python/relax/test_op_inspect.py @@ -53,7 +53,7 @@ class mod: def main(A: R.Tensor): return A.dtype.type_code - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty([16], dtype) @@ -70,7 +70,7 @@ class mod: def main(A: R.Tensor): return A.dtype.bits - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty([16], dtype) @@ -87,7 +87,7 @@ class mod: def main(A: R.Tensor): return A.dtype.lanes - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty([16], dtype) @@ -104,7 +104,7 @@ class mod: def main(A: R.Tensor): return A.ndim - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty(shape, "int32") @@ -120,7 +120,7 @@ class mod: def main(A: R.Tensor, axis: R.Prim("int64")): return A.shape[axis] - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty(shape, "int32") @@ -146,7 +146,7 @@ class mod: def main(A: R.Tensor, axis: R.Prim("int64")): return A.strides[axis] - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) arg = tvm.nd.empty(shape, "int32") @@ -168,7 +168,7 @@ class mod: def main(A: R.Tensor, axis: R.Prim("int64")): return A.strides[axis] - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) backing_ndarray = tvm.nd.empty(backing_shape, "int32") @@ -198,7 +198,7 @@ class mod: def main(A: R.Tensor): return A.byte_offset - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) backing_ndarray = tvm.nd.empty(backing_shape, "int32") @@ -230,7 +230,7 @@ class mod: def main(A: R.Tensor): return A.elem_offset - built = relax.build(mod) + built = tvm.compile(mod) vm = relax.VirtualMachine(built, tvm.cpu()) backing_ndarray = tvm.nd.empty(backing_shape, dtype) diff --git a/tests/python/relax/test_op_take.py b/tests/python/relax/test_op_take.py index babf91869a41..15ca5f4c8975 100644 --- a/tests/python/relax/test_op_take.py +++ b/tests/python/relax/test_op_take.py @@ -40,7 +40,7 @@ def main(A: R.Tensor([16, 16], "float16")): output = R.take(A, R.const(1), axis=axis) return output - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) np_input = np.random.random(size=[16, 16]).astype("float16") @@ -66,7 +66,7 @@ def main(A: R.Tensor([16, 16], "float16")): output = R.take(A, R.const([1]), axis=axis) return output - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) np_input = np.random.random(size=[16, 16]).astype("float16") @@ -88,7 +88,7 @@ def main(A: R.Tensor([16, 16], "float16")): output = R.take(A, R.const([[1, 3], [5, 7]]), axis=axis) return output - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) np_input = np.random.random(size=[16, 16]).astype("float16") @@ -115,7 +115,7 @@ def main(A: R.Tensor([16, 16], "float16")): output = R.take(A, R.prim_value(1), axis=axis) return output - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) np_input = np.random.random(size=[16, 16]).astype("float16") @@ -143,7 +143,7 @@ def main(A: R.Tensor(["n", "n"], "float16")): output = R.take(A, R.prim_value(n - 1), axis=axis) return output - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) np_input = np.random.random(size=[16, 16]).astype("float16") diff --git a/tests/python/relax/test_op_view.py b/tests/python/relax/test_op_view.py index 0900e1be306b..fc9458827b26 100644 --- a/tests/python/relax/test_op_view.py +++ b/tests/python/relax/test_op_view.py @@ -660,7 +660,7 @@ def main(A: R.Tensor([4096], "float32")): B = R.memory.view(A) return B - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, device=dev) np_input = np.random.random([4096]).astype("float32") @@ -680,7 +680,7 @@ def main(A: R.Tensor([4096], "float32")): B = R.memory.view(A, shape=R.shape([64, 64])) return B - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, device=dev) np_input = np.random.random([4096]).astype("float32") @@ -704,7 +704,7 @@ def main(A: R.Tensor([4096], "float32")): ) return B - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, device=dev) np_input = np.random.random([4096]).astype("float32") @@ -724,7 +724,7 @@ def main(A: R.Tensor([4096], "float32")): B = R.memory.view(A, dtype="uint32") return B - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, device=dev) np_input = np.random.random([4096]).astype("float32") @@ -754,7 +754,7 @@ def main(A: R.Tensor([4096], "uint8")): ) return (B, C) - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, device=dev) np_input = np.random.randint(0, 255, size=[4096]).astype("uint8") diff --git a/tests/python/relax/test_pipeline.py b/tests/python/relax/test_pipeline.py index 6d4d44c9e5a7..34d0ca9e36d2 100644 --- a/tests/python/relax/test_pipeline.py +++ b/tests/python/relax/test_pipeline.py @@ -37,7 +37,7 @@ def main(x: R.Tensor((3, 4), "float32"), y: R.Tensor((3, 4), "float32")): mod = Mod mod = pipeline(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) x_np = np.random.rand(3, 4).astype(np.float32) y_np = np.random.rand(3, 4).astype(np.float32) x = tvm.nd.array(x_np) @@ -95,7 +95,7 @@ def main( mod = Mod mod = pipeline(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) num_steps = 8 cache_np = np.empty((num_steps, 4), dtype="float32") diff --git a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_cpu.py b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_cpu.py index 0d6ee7b54e50..1941edeaa715 100644 --- a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_cpu.py +++ b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_cpu.py @@ -139,7 +139,7 @@ def set_global_func(head_dim, dtype): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.tir.build(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_flashinfer.py b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_flashinfer.py index 589184b51091..41743efeeea2 100644 --- a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_flashinfer.py +++ b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_flashinfer.py @@ -155,7 +155,7 @@ def load_module(name: str, static_modules: List[tvm.runtime.Module]): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.compile(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_flashinfer.py b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_flashinfer.py index 2aeb1b158bf4..84b50125ee15 100644 --- a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_flashinfer.py +++ b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_flashinfer.py @@ -168,7 +168,7 @@ def load_module(name: str, static_modules: List[tvm.runtime.Module]): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.compile(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_tir.py b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_tir.py index bee4cfe1a3cf..b2982abdb0a5 100644 --- a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_tir.py +++ b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_mla_tir.py @@ -133,7 +133,7 @@ def set_global_func(dtype): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.tir.build(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_tir.py b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_tir.py index 70a016697758..8cd3a737402e 100644 --- a/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_tir.py +++ b/tests/python/relax/test_runtime_builtin_paged_attention_kv_cache_tir.py @@ -141,7 +141,7 @@ def set_global_func(head_dim, dtype): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) - f = tvm.build(mod["main"], target=target) + f = tvm.tir.build(mod["main"], target=target) builts.append(f.entry_func) ( diff --git a/tests/python/relax/test_runtime_builtin_rnn_state.py b/tests/python/relax/test_runtime_builtin_rnn_state.py index de35ad5d7793..095aba8b83e5 100644 --- a/tests/python/relax/test_runtime_builtin_rnn_state.py +++ b/tests/python/relax/test_runtime_builtin_rnn_state.py @@ -80,7 +80,7 @@ def _build(tir_func): mod = tvm.IRModule({"main": tir_func}) with target: mod = dl.ApplyDefaultSchedule(dl.gpu.Fallback())(mod) # pylint: disable=not-callable - f = tvm.build(mod["main"], target=target) + f = tvm.tir.build(mod["main"], target=target) return f.entry_func _f_tir_gets, _f_tir_sets = [], [] diff --git a/tests/python/relax/test_tir_call_source_kernel.py b/tests/python/relax/test_tir_call_source_kernel.py index 9a877ad35f8f..9aba489f480b 100644 --- a/tests/python/relax/test_tir_call_source_kernel.py +++ b/tests/python/relax/test_tir_call_source_kernel.py @@ -95,6 +95,6 @@ def add(x_handle: T.handle, y_handle: T.handle, output_handle: T.handle): output_np = x_nd.numpy() + y_nd.numpy() with tvm.target.Target("cuda"): - lib = relax.build(Module) + lib = tvm.compile(Module) output_nd = tvm.runtime.relax_vm.VirtualMachine(lib, device)["main"](x_nd, y_nd) tvm.testing.assert_allclose(output_nd.numpy(), output_np, rtol=1e-5) diff --git a/tests/python/relax/test_training_optimizer_numeric.py b/tests/python/relax/test_training_optimizer_numeric.py index 8acf7ad66b2a..937582adf715 100644 --- a/tests/python/relax/test_training_optimizer_numeric.py +++ b/tests/python/relax/test_training_optimizer_numeric.py @@ -29,7 +29,7 @@ def _legalize_and_build(mod: IRModule, target, dev): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = VirtualMachine(ex, dev) return vm diff --git a/tests/python/relax/test_training_trainer_numeric.py b/tests/python/relax/test_training_trainer_numeric.py index 23a0ee3ab0ff..6737136cbf5c 100644 --- a/tests/python/relax/test_training_trainer_numeric.py +++ b/tests/python/relax/test_training_trainer_numeric.py @@ -64,7 +64,7 @@ def test_execute(target, dev): ) train_mod = setup_trainer(backbone) - ex = relax.build(train_mod, target) + ex = tvm.compile(train_mod, target) vm = relax.VirtualMachine(ex, dev, profile=True) trainer = Trainer(train_mod, vm, dev, False) @@ -89,7 +89,7 @@ def test_execute_numeric(target, dev): ) train_mod = setup_trainer(backbone) - ex = relax.build(train_mod, target) + ex = tvm.compile(train_mod, target) vm = relax.VirtualMachine(ex, dev) trainer = Trainer(train_mod, vm, dev, False) @@ -118,7 +118,7 @@ def test_load_export_params(target, dev): ) train_mod = setup_trainer(backbone) - ex = relax.build(train_mod, target) + ex = tvm.compile(train_mod, target) vm = relax.VirtualMachine(ex, dev) trainer = Trainer(train_mod, vm, dev, False) @@ -153,7 +153,7 @@ def test_setting_error(target, dev): ) train_mod = setup_trainer(backbone) - ex = relax.build(train_mod, target) + ex = tvm.compile(train_mod, target) vm = relax.VirtualMachine(ex, dev) trainer = Trainer(train_mod, vm, dev, False) diff --git a/tests/python/relax/test_transform_bind_params.py b/tests/python/relax/test_transform_bind_params.py index 9e212693f969..2e9845f73f40 100644 --- a/tests/python/relax/test_transform_bind_params.py +++ b/tests/python/relax/test_transform_bind_params.py @@ -60,11 +60,11 @@ def main( assert len(mod["main"].params) == 1 target = tvm.target.Target("llvm") - ex_after = relax.build(mod, target) + ex_after = tvm.compile(mod, target) vm_after = relax.VirtualMachine(ex_after, tvm.cpu()) res_after = vm_after["main"](x_tvm) - ex_before = relax.build(InputModule, target) + ex_before = tvm.compile(InputModule, target) vm_before = relax.VirtualMachine(ex_before, tvm.cpu()) res_before = vm_before["main"](x_tvm, w_tvm) diff --git a/tests/python/relax/test_transform_codegen_pass.py b/tests/python/relax/test_transform_codegen_pass.py index cb827e9734e3..bc0bc2dda154 100644 --- a/tests/python/relax/test_transform_codegen_pass.py +++ b/tests/python/relax/test_transform_codegen_pass.py @@ -78,7 +78,7 @@ def gen_ground_truth(mod, target, dev, inputs): ) new_mod = seq(mod) assert relax.analysis.well_formed(new_mod) - exec = relax.build(new_mod, target, params={}) + exec = tvm.compile(new_mod, target, params={}) vm = relax.VirtualMachine(exec, dev) return vm["main"](*inputs) @@ -145,7 +145,7 @@ def test_tensorrt_only(entry_func_name): ] )(mod) - ex0 = relax.build(new_mod, target, params={}) + ex0 = tvm.compile(new_mod, target, params={}) # Sanity check for the correctness and roundtrip check_roundtrip(ex0, dev, inputs, expected, entry_func_name) @@ -179,7 +179,7 @@ def test_mix_use_tensorrt_and_tvm(): )(mod) assert relax.analysis.well_formed(new_mod) with transform.PassContext(opt_level=0): - ex0 = relax.build(new_mod, target, params={}) + ex0 = tvm.compile(new_mod, target, params={}) # Sanity check for the correctness and roundtrip check_roundtrip(ex0, dev, inputs, expected) diff --git a/tests/python/relax/test_transform_few_shot_tuning.py b/tests/python/relax/test_transform_few_shot_tuning.py index 52870a82b4be..c640deee5496 100644 --- a/tests/python/relax/test_transform_few_shot_tuning.py +++ b/tests/python/relax/test_transform_few_shot_tuning.py @@ -341,7 +341,7 @@ def _expected_results( ) -> np.ndarray: func = _get_single_prim_func(mod) func = func.with_attr("global_symbol", "main") - rt_mod = tvm.build(func, target="llvm") + rt_mod = tvm.compile(func, target="llvm") data = [ tvm.nd.array(x) for x in [ @@ -357,7 +357,7 @@ def _actual_results( actual: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str ): target = _target() - actual_rt_mod = tvm.build(actual, target=target) + actual_rt_mod = tvm.compile(actual, target=target) actual_data = [ tvm.nd.array(x, device=tvm.cuda() if target.kind.name == "cuda" else tvm.cpu()) for x in [ diff --git a/tests/python/relax/test_transform_fold_batch_norm_to_conv2d.py b/tests/python/relax/test_transform_fold_batch_norm_to_conv2d.py index fc68f51b9f6b..4b17829fa0d7 100644 --- a/tests/python/relax/test_transform_fold_batch_norm_to_conv2d.py +++ b/tests/python/relax/test_transform_fold_batch_norm_to_conv2d.py @@ -90,14 +90,14 @@ def test_fold_batchnorm_info_conv2d(): # Normal build mod = tvm.relax.transform.DecomposeOpsForInference()(mod) - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, tvm.cpu()) out = vm["main"](data_in) # Fold BN to Conv2D mod_fold = relax.transform.FoldBatchnormToConv2D()(mod_fold) mod_fold = relax.transform.FoldConstant()(mod_fold) - ex_fold = relax.build(mod_fold, target) + ex_fold = tvm.compile(mod_fold, target) vm_fold = relax.VirtualMachine(ex_fold, tvm.cpu()) out_fold = vm_fold["main"](data_in) diff --git a/tests/python/relax/test_transform_gradient_numeric.py b/tests/python/relax/test_transform_gradient_numeric.py index 22241a2fbd65..70d6da8d7109 100644 --- a/tests/python/relax/test_transform_gradient_numeric.py +++ b/tests/python/relax/test_transform_gradient_numeric.py @@ -28,7 +28,7 @@ def rand(dtype, *shape): def _legalize_and_build(mod, target, dev): - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) vm = relax.VirtualMachine(ex, dev) return vm diff --git a/tests/python/relax/test_transform_lazy_transform_params.py b/tests/python/relax/test_transform_lazy_transform_params.py index 87a5698f1bf8..25d483fc449c 100644 --- a/tests/python/relax/test_transform_lazy_transform_params.py +++ b/tests/python/relax/test_transform_lazy_transform_params.py @@ -651,7 +651,7 @@ def transform_params( mod = TransformModule mod = relax.transform.LazyTransformParams()(mod) mod = relax.transform.LegalizeOps()(mod) - built = relax.build(mod, target=target) + built = tvm.compile(mod, target=target) params = [ np.random.random(size=(3, 64, 3, 3)).astype("float32"), diff --git a/tests/python/relax/test_vm_alloc_storage_with_scope.py b/tests/python/relax/test_vm_alloc_storage_with_scope.py index 17ae449a5d6a..ec6696000429 100644 --- a/tests/python/relax/test_vm_alloc_storage_with_scope.py +++ b/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 = relax.build(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() diff --git a/tests/python/relax/test_vm_builtin.py b/tests/python/relax/test_vm_builtin.py index c3272055fc5f..04e2ae1bf339 100644 --- a/tests/python/relax/test_vm_builtin.py +++ b/tests/python/relax/test_vm_builtin.py @@ -40,7 +40,7 @@ def foo(x: R.Tensor((3, 5), "float32"), y: R.Tensor((3, 1), "float32")): mod = CallSample target = tvm.target.Target("llvm", host="llvm") - ex = relax.build(mod, target) + ex = tvm.compile(mod, target) np_rand = np.random.rand(3, 5).astype(np.float32) # normalize it to get the random prob np_prob = np_rand / np_rand.sum(axis=1, keepdims=True) @@ -73,7 +73,7 @@ def main(): output = R.builtin.alloc_tensor(R.shape([1024, 1024, 1024, 1024, 1024]), "uint8", 0) return output - built = relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = relax.VirtualMachine(built, dev) with pytest.raises(Exception, match="CUDA: out of memory"): diff --git a/tests/python/relax/test_vm_callback_function.py b/tests/python/relax/test_vm_callback_function.py index 29a502ad7f98..73336db559ee 100644 --- a/tests/python/relax/test_vm_callback_function.py +++ b/tests/python/relax/test_vm_callback_function.py @@ -37,7 +37,11 @@ def relax_func( _ = callback(B) return R.tuple() - ex = tvm.relax.build(tvm.IRModule.from_expr(relax_func), target=target, exec_mode=exec_mode) + ex = tvm.relax.build( + tvm.IRModule.from_expr(relax_func), + target=target, + exec_mode=exec_mode, + ) vm = tvm.relax.VirtualMachine(ex, dev) from_callback = None diff --git a/tests/python/relax/test_vm_cuda_graph.py b/tests/python/relax/test_vm_cuda_graph.py index b6c8cdfdeea4..69c04eeca4ef 100644 --- a/tests/python/relax/test_vm_cuda_graph.py +++ b/tests/python/relax/test_vm_cuda_graph.py @@ -168,7 +168,7 @@ def main(A: R.Tensor([16], "float16")): "to have been captured by RewriteCUDAGraph." ) - built = tvm.relax.build(Module, target=target) + built = tvm.compile(Module, target=target) vm = tvm.relax.VirtualMachine(built, dev) arg = tvm.nd.array(np.arange(16).astype("float16"), dev) diff --git a/tests/python/relax/test_vm_execbuilder.py b/tests/python/relax/test_vm_execbuilder.py index b2d9edd34661..861ec9f8b041 100644 --- a/tests/python/relax/test_vm_execbuilder.py +++ b/tests/python/relax/test_vm_execbuilder.py @@ -277,7 +277,7 @@ def main(inp: R.Tensor((10, 10), dtype="float32")) -> R.Tensor((10, 10), dtype=" R.output(gv) return gv - ex = relax.build(Module, "llvm") + ex = tvm.compile(Module, "llvm") vm = relax.VirtualMachine(ex, tvm.cpu()) correct_input = tvm.nd.array(np.random.normal(size=(10, 10)).astype("float32")) diff --git a/tests/python/relax/test_vm_instrument.py b/tests/python/relax/test_vm_instrument.py index 615854264454..c5f293114f3c 100644 --- a/tests/python/relax/test_vm_instrument.py +++ b/tests/python/relax/test_vm_instrument.py @@ -46,7 +46,7 @@ def get_exec(data_shape): mod = relax.transform.BindParams("main", params)(mod) target = "llvm" - return relax.build(mod, target) + return tvm.compile(mod, target) def get_exec_int32(data_shape): @@ -61,7 +61,7 @@ def get_exec_int32(data_shape): mod = builder.get() target = "llvm" - return relax.build(mod, target) + return tvm.compile(mod, target) def test_conv2d_cpu(): diff --git a/tests/python/relax/test_vm_multi_device.py b/tests/python/relax/test_vm_multi_device.py index 73c78d70f042..b6a39437f2fc 100644 --- a/tests/python/relax/test_vm_multi_device.py +++ b/tests/python/relax/test_vm_multi_device.py @@ -37,7 +37,7 @@ def compile( mod = relax.transform.LegalizeOps()(mod) mod = tvm.tir.transform.DefaultGPUSchedule()(mod) # no need to feed target argument for mult-target compilation - ex = relax.build(mod) + ex = tvm.compile(mod) return relax.VirtualMachine(ex, device) diff --git a/tests/python/relax/test_vm_profiler.py b/tests/python/relax/test_vm_profiler.py index 114596741113..eaf914560530 100644 --- a/tests/python/relax/test_vm_profiler.py +++ b/tests/python/relax/test_vm_profiler.py @@ -47,7 +47,7 @@ def get_exec(data_shape): mod = relax.transform.BindParams("main", params)(mod) target = "llvm" - return relax.build(mod, target) + return tvm.compile(mod, target) def test_conv2d_cpu(): @@ -115,7 +115,7 @@ def main( return ((x, (x,)), x) target = "llvm" - ex = relax.build(NestedTuple, target) + ex = tvm.compile(NestedTuple, target) data_np = np.random.randn(16).astype("float32") diff --git a/tests/python/runtime/test_evaluator_with_preproc.py b/tests/python/runtime/test_evaluator_with_preproc.py index d9ef63054030..fd8f8e95b0bf 100644 --- a/tests/python/runtime/test_evaluator_with_preproc.py +++ b/tests/python/runtime/test_evaluator_with_preproc.py @@ -15,13 +15,13 @@ # specific language governing permissions and limitations # under the License. -import tvm -from tvm import te -from tvm.script import tir as T -import tvm.testing import numpy as np import pytest +import tvm +import tvm.testing +from tvm.script import tir as T + @T.prim_func def matmul(a: T.handle, b: T.handle, c: T.handle) -> None: @@ -45,7 +45,7 @@ def test_time_evalutor_with_preproc(f_preproc: str): i, j, k = sch.get_loops(blk) sch.bind(i, "blockIdx.x") sch.bind(j, "threadIdx.x") - f = tvm.build(sch.mod["main"], target="cuda") + f = tvm.tir.build(sch.mod["main"], target="cuda") dev = tvm.cuda(0) evaluator = f.time_evaluator(f.entry_name, dev, repeat=1000, number=1, f_preproc=f_preproc) diff --git a/tests/python/runtime/test_runtime_dlpack.py b/tests/python/runtime/test_runtime_dlpack.py index 60a86f662c6c..201037c6e469 100644 --- a/tests/python/runtime/test_runtime_dlpack.py +++ b/tests/python/runtime/test_runtime_dlpack.py @@ -35,7 +35,7 @@ def test_from_dlpack_shape_one(): B = te.placeholder((rows, 16), name="B") C = te.compute(A.shape, lambda i, j: A[i, j] + B[i, j], name="C") - fadd = tvm.build(te.create_prim_func([A, B, C]), target=tgt) + fadd = tvm.compile(te.create_prim_func([A, B, C]), target=tgt) dev = tvm.device(tgt.kind.name, 0) diff --git a/tests/python/runtime/test_runtime_extension.py b/tests/python/runtime/test_runtime_extension.py index be4ca55ebf8b..9b2f1adf7caa 100644 --- a/tests/python/runtime/test_runtime_extension.py +++ b/tests/python/runtime/test_runtime_extension.py @@ -43,7 +43,7 @@ def test_dltensor_compatible(): stmt = ib.get() mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr("global_symbol", "arange")) - f = tvm.build(mod, target="stackvm") + f = tvm.compile(mod, target="llvm") a = tvm.nd.array(np.zeros(10, dtype=dtype)) aview = MyTensorView(a) f(aview) diff --git a/tests/python/runtime/test_runtime_measure.py b/tests/python/runtime/test_runtime_measure.py index 4b39cef18bc5..ef27feb26398 100644 --- a/tests/python/runtime/test_runtime_measure.py +++ b/tests/python/runtime/test_runtime_measure.py @@ -35,7 +35,7 @@ def my_debug(filename): fout.write("c") X = te.compute((), lambda: tvm.tir.call_packed("my_debug", filename)) - func = tvm.build(te.create_prim_func([X])) + func = tvm.tir.build(te.create_prim_func([X])) x = tvm.nd.empty((), dtype="int32") ftimer = func.time_evaluator(func.entry_name, tvm.cpu(), number=1, repeat=1) diff --git a/tests/python/runtime/test_runtime_module_export.py b/tests/python/runtime/test_runtime_module_export.py index 1dff6c42502e..8897837a26af 100644 --- a/tests/python/runtime/test_runtime_module_export.py +++ b/tests/python/runtime/test_runtime_module_export.py @@ -35,8 +35,8 @@ def test_import_static_library(): te.create_prim_func([A, B]).with_attr("global_symbol", "myadd1") ) - mod0 = tvm.build(irmod0, target="llvm") - mod1 = tvm.build(irmod1, target="llvm") + mod0 = tvm.tir.build(irmod0, target="llvm") + mod1 = tvm.tir.build(irmod1, target="llvm") assert mod0.implements_function("myadd0") assert mod1.implements_function("myadd1") diff --git a/tests/python/runtime/test_runtime_module_load.py b/tests/python/runtime/test_runtime_module_load.py index 71f6a793877d..79b95256f9fa 100644 --- a/tests/python/runtime/test_runtime_module_load.py +++ b/tests/python/runtime/test_runtime_module_load.py @@ -62,7 +62,7 @@ def save_object(names): mod = tvm.IRModule.from_expr( tvm.tir.PrimFunc([Ab], stmt).with_attr("global_symbol", "main") ) - m = tvm.driver.build(mod, target=target) + m = tvm.tir.build(mod, target=target) for name in names: m.save(name) @@ -114,7 +114,7 @@ def check_device(device): print("Skip because %s is not enabled" % device) return temp = utils.tempdir() - f = tvm.build(sch.mod, target=device) + f = tvm.compile(sch.mod, target=device) path_dso = temp.relpath("dev_lib.so") # test cross compiler function @@ -122,7 +122,6 @@ def check_device(device): def popen_check(): import tvm - import sys f1 = tvm.runtime.load_module(path_dso) a = tvm.nd.array(np.random.uniform(size=1024).astype(A.dtype), dev) @@ -135,24 +134,20 @@ def popen_check(): worker.send(popen_check) worker.recv() - def check_stackvm(device): + def check_c(device): dev = tvm.device(device, 0) if not tvm.testing.device_enabled(device): print("Skip because %s is not enabled" % device) return - temp = utils.tempdir() - f = tvm.build(sch.mod, target=tvm.target.Target(device, host="stackvm")) - path_dso = temp.relpath("dev_lib.stackvm") - f.export_library(path_dso) - f1 = tvm.runtime.load_module(path_dso) + f = tvm.compile(sch.mod, target=tvm.target.Target(device, host="c")) a = tvm.nd.array(np.random.uniform(size=1024).astype(A.dtype), dev) b = tvm.nd.array(np.zeros(1024, dtype=A.dtype), dev) - f(a, b) + f["main"](a, b) np.testing.assert_equal(b.numpy(), a.numpy() + 1) for device in ["cuda", "vulkan", "opencl", "metal"]: check_device(device) - check_stackvm(device) + check_c(device) @tvm.testing.requires_llvm @@ -169,8 +164,8 @@ def test_combine_module_llvm(): def check_llvm(): dev = tvm.cpu(0) temp = utils.tempdir() - fadd1 = tvm.build(mod1, "llvm") - fadd2 = tvm.build(mod2, "llvm") + fadd1 = tvm.tir.build(mod1, "llvm") + fadd2 = tvm.tir.build(mod2, "llvm") path1 = temp.relpath("myadd1.o") path2 = temp.relpath("myadd2.o") path_dso = temp.relpath("mylib.so") @@ -195,8 +190,8 @@ def check_system_lib(): return temp = utils.tempdir() print("Running popen check") - fadd1 = tvm.build(mod1.with_attr("system_lib_prefix", ""), "llvm") - fadd2 = tvm.build(mod2.with_attr("system_lib_prefix", ""), "llvm") + fadd1 = tvm.tir.build(mod1.with_attr("system_lib_prefix", ""), "llvm") + fadd2 = tvm.tir.build(mod2.with_attr("system_lib_prefix", ""), "llvm") path1 = temp.relpath("myadd1.o") path2 = temp.relpath("myadd2.o") path_dso = temp.relpath("mylib.so") diff --git a/tests/python/runtime/test_runtime_module_property.py b/tests/python/runtime/test_runtime_module_property.py index 97c51ff93996..83e535e1ac83 100644 --- a/tests/python/runtime/test_runtime_module_property.py +++ b/tests/python/runtime/test_runtime_module_property.py @@ -33,7 +33,7 @@ def create_csource_module(): def create_llvm_module(): A = te.placeholder((1024,), name="A") B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B") - return tvm.build(te.create_prim_func([A, B]), target="llvm") + return tvm.tir.build(te.create_prim_func([A, B]), target="llvm") def test_property(): diff --git a/tests/python/runtime/test_runtime_rpc.py b/tests/python/runtime/test_runtime_rpc.py index 717cc8fffa05..d702b0e068bf 100644 --- a/tests/python/runtime/test_runtime_rpc.py +++ b/tests/python/runtime/test_runtime_rpc.py @@ -73,7 +73,7 @@ def test_bigendian_rpc(): def verify_rpc(remote, target, shape, dtype): A = te.placeholder(shape, dtype=dtype) B = te.compute(A.shape, lambda i: A[i] + tvm.tir.const(1, A.dtype)) - f = tvm.build(te.create_prim_func([A, B]), target=target) + f = tvm.compile(te.create_prim_func([A, B]), target=target) dev = remote.cpu(0) a = tvm.nd.array(np.random.randint(0, 256, size=shape).astype(A.dtype), device=dev) @@ -264,7 +264,7 @@ def test_rpc_remote_module(): def check_remote(remote): temp = utils.tempdir() dev = remote.cpu(0) - f = tvm.build(mod, "llvm") + f = tvm.compile(mod, "llvm") path_dso = temp.relpath("dev_lib.so") f.export_library(path_dso) remote.upload(path_dso) @@ -295,7 +295,7 @@ def check_minrpc(): # export to minrpc temp = utils.tempdir() # system lib prefix will trigger system lib build - f = tvm.build(mod.with_attr("system_lib_prefix", ""), "llvm") + f = tvm.compile(mod.with_attr("system_lib_prefix", ""), "llvm") path_minrpc = temp.relpath("dev_lib.minrpc") f.export_library(path_minrpc, fcompile=rpc.with_minrpc(cc.create_executable)) @@ -338,7 +338,7 @@ def check_remote_link_cl(remote): xo, xi = s.split(x, factors=[None, 32]) s.bind(xo, "blockIdx.x") s.bind(xi, "threadIdx.x") - f = tvm.build(s.mod, "opencl --host=llvm") + f = tvm.compile(s.mod, "opencl --host=llvm") path_tar = temp.relpath("myadd.tar") f.export_library(path_tar) remote.upload(path_tar) @@ -698,7 +698,7 @@ def func_without_arg() -> T.int64: def func_with_arg(unused: T.int64) -> T.int64: return T.int64(42) - built = tvm.build(Module, target="llvm") + built = tvm.compile(Module, target="llvm") server = tvm.rpc.Server(key="x1") client = tvm.rpc.connect("127.0.0.1", server.port, key="x1") diff --git a/tests/python/runtime/test_runtime_trace.py b/tests/python/runtime/test_runtime_trace.py index 58d1a079e46b..5093ce930ec3 100644 --- a/tests/python/runtime/test_runtime_trace.py +++ b/tests/python/runtime/test_runtime_trace.py @@ -23,7 +23,7 @@ def test_trace_default_action(): n = 2 x = te.placeholder((n, n, n), name="X", dtype="float32") y = te.compute(x.shape, lambda i, j, k: tvm.tir.trace([i, j, k, x[i][j][k]])) - f = tvm.build(te.create_prim_func([x, y]), target="llvm") + f = tvm.compile(te.create_prim_func([x, y]), target="llvm") xnd = tvm.nd.array(np.ones((n, n, n), dtype=x.dtype)) ynd = tvm.nd.array(np.zeros((n, n, n), dtype=y.dtype)) f(xnd, ynd) @@ -43,7 +43,7 @@ def check_assign(dtype): z = te.compute( x.shape, lambda i, j, k: tvm.tir.trace([y[i][j][k]], "tvm.tir.trace_callback2") ) - f = tvm.build(te.create_prim_func([x, y, z]), "llvm") + f = tvm.compile(te.create_prim_func([x, y, z]), "llvm") xnd = tvm.nd.array(np.ones((n, n, n), dtype=x.dtype)) ynd = tvm.nd.array(np.zeros((n, n, n), dtype=y.dtype)) @@ -72,7 +72,7 @@ def check_expr_sum(dtype): lambda i, j, k: tvm.tir.trace([a[i][j][k]], "tvm.tir.trace_callback3") + tvm.tir.trace([b[i][j][k]], "tvm.tir.trace_callback3"), ) - f = tvm.build(te.create_prim_func([a, b, c])) + f = tvm.compile(te.create_prim_func([a, b, c])) xnd = tvm.nd.array(np.array(np.ones((n, n, n), dtype=a.dtype))) ynd = tvm.nd.array(np.array(np.ones((n, n, n), dtype=b.dtype))) znd = tvm.nd.array(np.zeros((n, n, n), dtype=c.dtype)) @@ -102,7 +102,7 @@ def check_expr_sum(dtype): + tvm.tir.trace([i, j, k, d[i][j][k]], "tvm.tir.trace_silent") + tvm.tir.trace([i, j, k, e[i][j][k]], "tvm.tir.trace_silent"), ) - f = tvm.build(te.create_prim_func([a, b, d, e, c])) + f = tvm.compile(te.create_prim_func([a, b, d, e, c])) a_nd = tvm.nd.array(np.array(np.ones((n, n, n), dtype=a.dtype))) b_nd = tvm.nd.array(np.array(np.ones((n, n, n), dtype=b.dtype))) d_nd = tvm.nd.array(np.array(np.ones((n, n, n), dtype=d.dtype))) @@ -131,7 +131,7 @@ def check_expr_sum_custom(dtype): lambda i, j: tvm.tir.trace([a[i][j]], "tvm.tir.trace_callback4") + tvm.tir.trace([b[i][j]], "tvm.tir.trace_callback4"), ) - f = tvm.build(te.create_prim_func([a, b, c])) + f = tvm.compile(te.create_prim_func([a, b, c])) npa = np.array([[1, 0, 0, 0], [0, 1, 0, 0], [0, 0, 1, 0], [0, 0, 0, 1]], dtype=a.dtype) npb = np.array([[1, 0, 0, 0], [0, 1, 0, 0], [0, 0, 1, 0], [0, 0, 0, 1]], dtype=a.dtype) xnd = tvm.nd.array(npa) @@ -158,7 +158,7 @@ def check_assign(dtype): x = te.placeholder((n,), name="X", dtype=dtype) y = te.compute(x.shape, lambda i: tvm.tir.trace([x[i]], "tvm.tir.trace_change_int_first")) z = te.compute(x.shape, lambda i: tvm.tir.trace([y[i]], "tvm.tir.trace_change_int_second")) - f = tvm.build(te.create_prim_func([x, y, z])) + f = tvm.compile(te.create_prim_func([x, y, z])) xnd = tvm.nd.array(np.ones((n,), dtype=x.dtype)) ynd = tvm.nd.array(np.zeros((n,), dtype=y.dtype)) @@ -189,7 +189,7 @@ def check_assign(dtype): z = te.compute( x.shape, lambda i: tvm.tir.trace([y[i]], "tvm.tir.trace_change_float_second") ) - f = tvm.build(te.create_prim_func([x, y, z]), target="llvm") + f = tvm.compile(te.create_prim_func([x, y, z]), target="llvm") xnd = tvm.nd.array(np.ones((n,), dtype=x.dtype)) ynd = tvm.nd.array(np.zeros((n,), dtype=y.dtype)) diff --git a/tests/python/target/test_arm_target.py b/tests/python/target/test_arm_target.py index 2fcac847a8e0..686954baade1 100644 --- a/tests/python/target/test_arm_target.py +++ b/tests/python/target/test_arm_target.py @@ -82,7 +82,7 @@ def my_func(a: T.handle): T.func_attr({"global_symbol": "my_module", "tir.noalias": True}) A[0] = T.Div(10000, 4 * T.vscale()) - mod = tvm.build(my_func, target=target) + mod = tvm.compile(my_func, target=target) A_nd = tvm.nd.array(np.empty((1,), dtype="int32"), device=dev) mod(A_nd) @@ -105,7 +105,7 @@ def my_func(a: T.handle, b: T.handle): T.func_attr({"global_symbol": "my_module", "tir.noalias": True}) B[T.ramp(0, 1, 4 * T.vscale())] = A[T.ramp(0, 1, 4 * T.vscale())] - mod = tvm.build(my_func, target=target) + mod = tvm.compile(my_func, target=target) A_np = np.random.uniform(size=(num_elements,)).astype("float32") B_np = np.zeros((num_elements,)).astype("float32") @@ -133,7 +133,7 @@ def my_func(a: T.handle, b: T.handle): for i in T.serial(0, 4 * T.vscale()): B[i] = A[i] - mod = tvm.build(my_func, target=target) + mod = tvm.compile(my_func, target=target) A_np = np.random.uniform(size=(num_elements,)).astype(dtype) B_np = np.zeros((num_elements,)).astype(dtype) @@ -156,7 +156,7 @@ def my_func(a: T.handle): T.func_attr({"global_symbol": "my_module", "tir.noalias": True}) A[T.ramp(0, 1, 4 * T.vscale())] = T.broadcast(1, 4 * T.vscale()) - mod = tvm.build(my_func, target=target) + mod = tvm.compile(my_func, target=target) A_np = np.zeros((num_elements,)).astype("float32") A_nd = tvm.nd.array(A_np, device=dev) diff --git a/tests/python/target/test_target_target.py b/tests/python/target/test_target_target.py index c6908d23f000..eb74629f5a84 100644 --- a/tests/python/target/test_target_target.py +++ b/tests/python/target/test_target_target.py @@ -476,7 +476,7 @@ def func(): func = func.with_attr("Target", target) target2 = tvm.ir.load_json(tvm.ir.save_json(target)) mod = tvm.IRModule({"main": func}) - lib = tvm.build(mod, target=target2) + lib = tvm.compile(mod, target=target2) lib["func"]() diff --git a/tests/python/te/test_te_create_primfunc.py b/tests/python/te/test_te_create_primfunc.py index 486fc0b18c32..eabed1272afc 100644 --- a/tests/python/te/test_te_create_primfunc.py +++ b/tests/python/te/test_te_create_primfunc.py @@ -350,7 +350,7 @@ def test_constant(): ) func = te.create_prim_func([C, A]) - func = tvm.build(func) + func = tvm.compile(func) a_np = np.random.uniform(size=(M,)).astype(A.dtype) c = tvm.nd.array(np.zeros(M, dtype=C.dtype)) x = func(c, tvm.nd.array(a_np)) @@ -363,7 +363,7 @@ def test_data_dependent_access(): C = te.compute((10,), lambda i: A[B[i]]) func = te.create_prim_func([C, A, B]) - func = tvm.build(func) + func = tvm.compile(func) a_np = np.random.uniform(size=(10,)).astype(A.dtype) b_np = np.arange(10, dtype=B.dtype) diff --git a/tests/python/tir-base/test_tir_base.py b/tests/python/tir-base/test_tir_base.py index 3a6750231330..4f6268d6520d 100644 --- a/tests/python/tir-base/test_tir_base.py +++ b/tests/python/tir-base/test_tir_base.py @@ -29,7 +29,7 @@ def build_tir_func(func): if pass_ctx.config.get("tir.noalias", True): func = func.with_attr("tir.noalias", True) mod = tvm.IRModule({"main": func}) - func = tvm.build(mod) + func = tvm.compile(mod) return func diff --git a/tests/python/tir-base/test_tir_imm_values.py b/tests/python/tir-base/test_tir_imm_values.py index 416943c85da6..11213e35364c 100644 --- a/tests/python/tir-base/test_tir_imm_values.py +++ b/tests/python/tir-base/test_tir_imm_values.py @@ -147,7 +147,7 @@ def test_tir_too_large_literal_f64(): def imm_overflow_fp64() -> T.float64: T.evaluate(T.ret(T.float64(1.7976e309), dtype="float64")) - f = tvm.build(imm_overflow_fp64, target="llvm") + f = tvm.compile(imm_overflow_fp64, target="llvm") assert math.isinf(f()) @@ -270,7 +270,7 @@ def float_imm_div(x: T.float32, y: T.float32, z: T.Buffer((), "float32")): z[()] = x / y def __wrap_build(f): - lib = tvm.build(f, target="llvm") + lib = tvm.compile(f, target="llvm") z = tvm.nd.array(np.zeros([]).astype("float32")) def _func(x, y): @@ -331,11 +331,11 @@ def imm_truncdiv(x: T.int8, y: T.int8) -> T.int8: def imm_floordiv(x: T.int8, y: T.int8) -> T.int8: T.evaluate(T.ret(T.floordiv(x, y), dtype="int8")) - fmul = tvm.build(imm_multiply, target="llvm") - fadd = tvm.build(imm_add, target="llvm") - fsub = tvm.build(imm_sub, target="llvm") - ffloordiv = tvm.build(imm_floordiv, target="llvm") - ftruncdiv = tvm.build(imm_truncdiv, target="llvm") + fmul = tvm.compile(imm_multiply, target="llvm") + fadd = tvm.compile(imm_add, target="llvm") + fsub = tvm.compile(imm_sub, target="llvm") + ffloordiv = tvm.compile(imm_floordiv, target="llvm") + ftruncdiv = tvm.compile(imm_truncdiv, target="llvm") # overflow check_tir_const_fold("int8", lambda x, y: x + y, fadd, 127, 1, -128) @@ -387,11 +387,11 @@ def imm_truncdiv(x: T.uint8, y: T.uint8) -> T.uint8: def imm_floordiv(x: T.uint8, y: T.uint8) -> T.uint8: T.evaluate(T.ret(T.floordiv(x, y), dtype="uint8")) - fmul = tvm.build(imm_multiply, target="llvm") - fadd = tvm.build(imm_add, target="llvm") - fsub = tvm.build(imm_sub, target="llvm") - ffloordiv = tvm.build(imm_floordiv, target="llvm") - ftruncdiv = tvm.build(imm_truncdiv, target="llvm") + fmul = tvm.compile(imm_multiply, target="llvm") + fadd = tvm.compile(imm_add, target="llvm") + fsub = tvm.compile(imm_sub, target="llvm") + ffloordiv = tvm.compile(imm_floordiv, target="llvm") + ftruncdiv = tvm.compile(imm_truncdiv, target="llvm") # overflow check_tir_const_fold("uint8", lambda x, y: x + y, fadd, 255, 1, 0) @@ -454,13 +454,13 @@ def imm_floordiv(x: T.int32, y: T.int32) -> T.int32: def imm_floormod(x: T.int32, y: T.int32) -> T.int32: T.evaluate(T.ret(T.floormod(x, y), dtype="int32")) - fmul = tvm.build(imm_multiply, target="llvm") - fadd = tvm.build(imm_add, target="llvm") - fsub = tvm.build(imm_sub, target="llvm") - ffloordiv = tvm.build(imm_floordiv, target="llvm") - ffloormod = tvm.build(imm_floormod, target="llvm") - ftruncdiv = tvm.build(imm_truncdiv, target="llvm") - ftruncmod = tvm.build(imm_truncmod, target="llvm") + fmul = tvm.compile(imm_multiply, target="llvm") + fadd = tvm.compile(imm_add, target="llvm") + fsub = tvm.compile(imm_sub, target="llvm") + ffloordiv = tvm.compile(imm_floordiv, target="llvm") + ffloormod = tvm.compile(imm_floormod, target="llvm") + ftruncdiv = tvm.compile(imm_truncdiv, target="llvm") + ftruncmod = tvm.compile(imm_truncmod, target="llvm") # i32 overflow is not specified, only check for range assert -(2**31) <= int(tir.const(2**31 - 1, "int32") + tir.const(1, "int32")) < 2**31 @@ -534,11 +534,11 @@ def imm_truncdiv(x: T.uint32, y: T.uint32) -> T.uint32: def imm_floordiv(x: T.uint32, y: T.uint32) -> T.uint32: T.evaluate(T.ret(T.floordiv(x, y), dtype="uint32")) - fmul = tvm.build(imm_multiply, target="llvm") - fadd = tvm.build(imm_add, target="llvm") - fsub = tvm.build(imm_sub, target="llvm") - ffloordiv = tvm.build(imm_floordiv, target="llvm") - ftruncdiv = tvm.build(imm_truncdiv, target="llvm") + fmul = tvm.compile(imm_multiply, target="llvm") + fadd = tvm.compile(imm_add, target="llvm") + fsub = tvm.compile(imm_sub, target="llvm") + ffloordiv = tvm.compile(imm_floordiv, target="llvm") + ftruncdiv = tvm.compile(imm_truncdiv, target="llvm") # u32 overflow is not specified, only check for range assert 0 <= int(tir.const(2**32 - 1, "uint32") + tir.const(1, "uint32")) < 2**32 diff --git a/tests/python/tir-base/test_tir_intrin.py b/tests/python/tir-base/test_tir_intrin.py index 8ab18bc84855..d2a73c12e79f 100644 --- a/tests/python/tir-base/test_tir_intrin.py +++ b/tests/python/tir-base/test_tir_intrin.py @@ -37,7 +37,7 @@ def test_nearbyint(): sch = tir.Schedule(mod) # Build from scheduled TIR - func = tvm.build(sch.mod, target="llvm") + func = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) n = 10 @@ -91,7 +91,7 @@ def run_test(tvm_intrin, np_func): sch = tir.Schedule(mod) # Build from scheduled TIR - func = tvm.build(sch.mod, target="llvm") + func = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) n = 10 @@ -125,7 +125,7 @@ def run_test(tvm_intrin, np_func): sch = tir.Schedule(mod) # Build from scheduled TIR - func = tvm.build(sch.mod, target="llvm") + func = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) n = 10 @@ -152,7 +152,7 @@ def test_ldexp(): sch = tir.Schedule(mod) # Build from scheduled TIR - func = tvm.build(sch.mod, target="llvm") + func = tvm.compile(sch.mod, target="llvm") dev = tvm.cpu(0) n = 10 @@ -200,7 +200,7 @@ def clz_np(x, dtype): sch.bind(tx, "threadIdx.x") # Build from scheduled TIR - func = tvm.build(sch.mod, target=target) + func = tvm.compile(sch.mod, target=target) n = 10 highs = [10, 100, 1000, 10000, 100000, 1000000] diff --git a/tests/python/tir-base/test_tir_ptx_cp_async.py b/tests/python/tir-base/test_tir_ptx_cp_async.py index d7600238542d..d5c029c10138 100644 --- a/tests/python/tir-base/test_tir_ptx_cp_async.py +++ b/tests/python/tir-base/test_tir_ptx_cp_async.py @@ -51,7 +51,7 @@ def ptx_cp_async(A: T.Buffer((32, 128), "float16"), B: T.Buffer((32, 128), "floa def test_ptx_cp_async(): f = ptx_cp_async - mod = tvm.build(f, target="cuda") + mod = tvm.compile(f, target="cuda") A_np = np.random.rand(32, 128).astype("float16") B_np = np.zeros((32, 128)).astype("float16") dev = tvm.cuda(0) @@ -98,7 +98,7 @@ def ptx_cp_async_barrier( def test_ptx_cp_async_barrier(): f = ptx_cp_async_barrier - mod = tvm.build(f, target="cuda") + mod = tvm.compile(f, target="cuda") A_np = np.random.rand(32, 128).astype("float16") B_np = np.zeros((32, 128)).astype("float16") dev = tvm.cuda(0) @@ -139,7 +139,7 @@ def ptx_cp_async_bulk(A: T.Buffer((32, 128), "float16"), B: T.Buffer((32, 128), def test_ptx_cp_async_bulk(): f = ptx_cp_async_bulk - mod = tvm.build(f, target="cuda") + mod = tvm.compile(f, target="cuda") A_np = np.random.rand(32, 128).astype("float16") B_np = np.zeros((32, 128)).astype("float16") dev = tvm.cuda(0) diff --git a/tests/python/tir-base/test_tir_ptx_ldmatrix.py b/tests/python/tir-base/test_tir_ptx_ldmatrix.py index 615d33ae004e..346f9c393fcd 100644 --- a/tests/python/tir-base/test_tir_ptx_ldmatrix.py +++ b/tests/python/tir-base/test_tir_ptx_ldmatrix.py @@ -63,7 +63,7 @@ def test_ptx_ldmatrix(): for num in [1, 2, 4]: for trans in [False, True]: - mod = tvm.build(f.specialize({param_num: num, param_trans: trans}), target="cuda") + mod = tvm.compile(f.specialize({param_num: num, param_trans: trans}), target="cuda") A_np = np.random.rand(16, 16).astype("float16") A_mask_np = np.zeros_like(A_np) if num == 1: diff --git a/tests/python/tir-base/test_tir_ptx_mma.py b/tests/python/tir-base/test_tir_ptx_mma.py index cc9eec3a69d7..8f221d95da32 100644 --- a/tests/python/tir-base/test_tir_ptx_mma.py +++ b/tests/python/tir-base/test_tir_ptx_mma.py @@ -15,13 +15,11 @@ # specific language governing permissions and limitations # under the License. -import sys -import pytest +import numpy as np import tvm -from tvm.script import tir as T -import numpy as np import tvm.testing +from tvm.script import tir as T @T.prim_func @@ -69,7 +67,7 @@ def gemm_mma_m8n8k4_row_col_fp64pf64fp64(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m8n8k4_row_col_fp64pf64fp64(): sch = tvm.tir.Schedule(gemm_mma_m8n8k4_row_col_fp64pf64fp64) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [8, 4]).astype("float64") B_np = np.random.uniform(-1, 1, [8, 4]).astype("float64") @@ -145,7 +143,7 @@ def gemm_mma_m8n8k4_row_row_fp16fp16fp16(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7) def test_gemm_mma_m8n8k4_row_row_fp16fp16fp16(): sch = tvm.tir.Schedule(gemm_mma_m8n8k4_row_row_fp16fp16fp16) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 4]).astype("float16") B_np = np.random.uniform(-1, 1, [4, 16]).astype("float16") @@ -228,7 +226,7 @@ def gemm_mma_m8n8k4_row_row_fp16fp16fp32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7) def test_gemm_mma_m8n8k4_row_row_fp16fp16fp32(): sch = tvm.tir.Schedule(gemm_mma_m8n8k4_row_row_fp16fp16fp32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 4]).astype("float16") B_np = np.random.uniform(-1, 1, [4, 16]).astype("float16") @@ -299,7 +297,7 @@ def gemm_mma_m8n8k16_row_col_s8s8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7, 5) def test_gemm_mma_m8n8k16_row_col_s8s8s32(): sch = tvm.tir.Schedule(gemm_mma_m8n8k16_row_col_s8s8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [8, 16]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 16]).astype("int8") @@ -370,7 +368,7 @@ def gemm_mma_m8n8k16_row_col_s8u8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7, 5) def test_gemm_mma_m8n8k16_row_col_s8u8s32(): sch = tvm.tir.Schedule(gemm_mma_m8n8k16_row_col_s8u8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [8, 16]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 16]).astype("uint8") @@ -441,7 +439,7 @@ def gemm_mma_m8n8k32_row_col_s4s4s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7, 5) def test_gemm_mma_m8n8k32_row_col_s4s4s32(): sch = tvm.tir.Schedule(gemm_mma_m8n8k32_row_col_s4s4s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") ctx = tvm.cuda() A_tvm = tvm.nd.empty([8, 32], "int4", ctx) @@ -504,7 +502,7 @@ def gemm_mma_m8n8k32_row_col_s4u4s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(7, 5) def test_gemm_mma_m8n8k32_row_col_s4u4s32(): sch = tvm.tir.Schedule(gemm_mma_m8n8k32_row_col_s4u4s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") ctx = tvm.cuda() A_tvm = tvm.nd.empty([8, 32], "int4", ctx) @@ -569,7 +567,7 @@ def gemm_mma_m16n8k8_row_col_fp16fp16fp32(a: T.handle, b: T.handle, c: T.handle) @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k8_row_col_fp16fp16fp32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k8_row_col_fp16fp16fp32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 8]).astype("float16") B_np = np.random.uniform(-1, 1, [8, 8]).astype("float16") @@ -645,7 +643,7 @@ def gemm_mma_m16n8k16_row_col_fp16fp16fp16(a: T.handle, b: T.handle, c: T.handle @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k16_row_col_fp16fp16fp16(): sch = tvm.tir.Schedule(gemm_mma_m16n8k16_row_col_fp16fp16fp16) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 16]).astype("float16") B_np = np.random.uniform(-1, 1, [8, 16]).astype("float16") @@ -721,7 +719,7 @@ def gemm_mma_m16n8k16_row_col_fp16fp16fp32(a: T.handle, b: T.handle, c: T.handle @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k16_row_col_fp16fp16fp32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k16_row_col_fp16fp16fp32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 16]).astype("float16") B_np = np.random.uniform(-1, 1, [8, 16]).astype("float16") @@ -797,7 +795,7 @@ def gemm_mma_m16n8k16_row_col_s8s8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k16_row_col_s8s8s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k16_row_col_s8s8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [16, 16]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 16]).astype("int8") @@ -873,7 +871,7 @@ def gemm_mma_m16n8k16_row_col_s8u8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k16_row_col_s8u8s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k16_row_col_s8u8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [16, 16]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 16]).astype("uint8") @@ -949,7 +947,7 @@ def gemm_mma_m16n8k32_row_col_s8s8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k32_row_col_s8s8s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k32_row_col_s8s8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [16, 32]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 32]).astype("int8") @@ -1025,7 +1023,7 @@ def gemm_mma_m16n8k32_row_col_s8u8s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k32_row_col_s8u8s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k32_row_col_s8u8s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-10, 10, [16, 32]).astype("int8") B_np = np.random.uniform(-10, 10, [8, 32]).astype("uint8") @@ -1101,7 +1099,7 @@ def gemm_mma_m16n8k64_row_col_s4s4s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k64_row_col_s4s4s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k64_row_col_s4s4s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") ctx = tvm.cuda() A_tvm = tvm.nd.empty([16, 64], "int4", ctx) @@ -1169,7 +1167,7 @@ def gemm_mma_m16n8k64_row_col_s4u4s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k64_row_col_s4u4s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k64_row_col_s4u4s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") ctx = tvm.cuda() A_tvm = tvm.nd.empty([16, 64], "int4", ctx) @@ -1238,7 +1236,7 @@ def gemm_mma_m16n8k256_row_col_b1b1s32(a: T.handle, b: T.handle, c: T.handle): @tvm.testing.requires_cuda_compute_version(8) def test_gemm_mma_m16n8k256_row_col_b1b1s32(): sch = tvm.tir.Schedule(gemm_mma_m16n8k256_row_col_b1b1s32) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") ctx = tvm.cuda() A_tvm = tvm.nd.empty([16, 256], "int1", ctx) diff --git a/tests/python/tir-base/test_tir_ptx_mma_sp.py b/tests/python/tir-base/test_tir_ptx_mma_sp.py index 0b5073864a43..d5c6c9a03b45 100644 --- a/tests/python/tir-base/test_tir_ptx_mma_sp.py +++ b/tests/python/tir-base/test_tir_ptx_mma_sp.py @@ -273,7 +273,7 @@ def get_meta_m16n8k16_half(mask): for out_dtype in ["float16", "float32"]: func = mma_sp_m16n8k16_f16f16f16 if out_dtype == "float16" else mma_sp_m16n8k16_f16f16f32 sch = tvm.tir.Schedule(func) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 8]).astype("float16") B_np = np.random.uniform(-1, 1, [16, 8]).astype("float16") @@ -312,7 +312,7 @@ def get_meta_m16n8k32_half(mask): for out_dtype in ["float16", "float32"]: func = mma_sp_m16n8k32_f16f16f16 if out_dtype == "float16" else mma_sp_m16n8k32_f16f16f32 sch = tvm.tir.Schedule(func) - cuda_mod = tvm.build(sch.mod, target="cuda") + cuda_mod = tvm.compile(sch.mod, target="cuda") A_np = np.random.uniform(-1, 1, [16, 16]).astype("float16") B_np = np.random.uniform(-1, 1, [32, 8]).astype("float16") diff --git a/tests/python/tir-base/test_tir_te_extern_primfunc.py b/tests/python/tir-base/test_tir_te_extern_primfunc.py index 16bc0b0ae2fc..9c375481fe45 100644 --- a/tests/python/tir-base/test_tir_te_extern_primfunc.py +++ b/tests/python/tir-base/test_tir_te_extern_primfunc.py @@ -181,7 +181,7 @@ class TestPrimFuncs: def test_primfunc_call(self, func, verify): target = tvm.target.Target("llvm") - func = tvm.build(func, target=target) + func = tvm.compile(func, target=target) verify(func) def test_te_extern_call(self, func, params, verify): @@ -194,7 +194,7 @@ def test_te_extern_call(self, func, params, verify): rt_prim_func = te.create_prim_func(tensors_from_extern_op(output, prim_func)) target = tvm.target.Target("llvm") - func = tvm.build(rt_prim_func, target=target) + func = tvm.compile(rt_prim_func, target=target) verify(func) diff --git a/tests/python/tir-base/test_tir_texture_scope.py b/tests/python/tir-base/test_tir_texture_scope.py index fb98b6536e66..4b759bb0477d 100644 --- a/tests/python/tir-base/test_tir_texture_scope.py +++ b/tests/python/tir-base/test_tir_texture_scope.py @@ -56,7 +56,7 @@ def schedule_block(block): schedule_block(sch.get_block("C")) target = tvm.target.Target("opencl") - mod = tvm.build(sch.mod["main"], target=target) + mod = tvm.compile(sch.mod["main"], target=target) if __name__ == "__main__": diff --git a/tests/python/tir-schedule/test_tir_schedule_decompose_padding.py b/tests/python/tir-schedule/test_tir_schedule_decompose_padding.py index e8ba0b4e21e0..c8679843dda6 100644 --- a/tests/python/tir-schedule/test_tir_schedule_decompose_padding.py +++ b/tests/python/tir-schedule/test_tir_schedule_decompose_padding.py @@ -35,8 +35,8 @@ def check_decompose_padding(origin, scheduled, expected, check_run=False): x = tvm.nd.array(np.random.uniform(0, 64, in_shape).astype(in_buffer.dtype)) y0 = tvm.nd.array(np.zeros(out_shape).astype(out_buffer.dtype)) y1 = tvm.nd.array(np.zeros(out_shape).astype(out_buffer.dtype)) - f_origin = tvm.build(origin) - f_scheduled = tvm.build(scheduled) + f_origin = tvm.compile(origin) + f_scheduled = tvm.compile(scheduled) f_origin(x, y0) f_scheduled(x, y1) tvm.testing.assert_allclose(y0.numpy(), y1.numpy()) diff --git a/tests/python/tir-schedule/test_tir_schedule_rolling_buffer.py b/tests/python/tir-schedule/test_tir_schedule_rolling_buffer.py index 9d19dd877b75..0ea51aaf83aa 100644 --- a/tests/python/tir-schedule/test_tir_schedule_rolling_buffer.py +++ b/tests/python/tir-schedule/test_tir_schedule_rolling_buffer.py @@ -41,8 +41,8 @@ def check_rolling_buffer( x = tvm.nd.array(np.random.uniform(0, 64, in_shape).astype(in_buffer.dtype)) y0 = tvm.nd.array(np.zeros(out_shape).astype(out_buffer.dtype)) y1 = tvm.nd.array(np.zeros(out_shape).astype(out_buffer.dtype)) - f_origin = tvm.build(origin) - f_scheduled = tvm.build(scheduled) + f_origin = tvm.compile(origin) + f_scheduled = tvm.compile(scheduled) f_origin(x, y0) f_scheduled(x, y1) tvm.testing.assert_allclose(y0.numpy(), y1.numpy()) diff --git a/tests/python/tir-schedule/test_tir_schedule_tensorize_ldmatrix_mma_numeric.py b/tests/python/tir-schedule/test_tir_schedule_tensorize_ldmatrix_mma_numeric.py index 5a80e3e4f6c4..d5646f60fb7a 100644 --- a/tests/python/tir-schedule/test_tir_schedule_tensorize_ldmatrix_mma_numeric.py +++ b/tests/python/tir-schedule/test_tir_schedule_tensorize_ldmatrix_mma_numeric.py @@ -118,7 +118,7 @@ def run_test( mma_store_intrin, ) - f = tvm.build(sch.mod["main"], target="cuda") + f = tvm.compile(sch.mod["main"], target="cuda") dev = tvm.device("cuda", 0) diff --git a/tests/python/tir-schedule/test_tir_schedule_tensorize_mfma_numeric.py b/tests/python/tir-schedule/test_tir_schedule_tensorize_mfma_numeric.py index 2b3e6ce39bfb..c8edaf30fca9 100644 --- a/tests/python/tir-schedule/test_tir_schedule_tensorize_mfma_numeric.py +++ b/tests/python/tir-schedule/test_tir_schedule_tensorize_mfma_numeric.py @@ -109,7 +109,7 @@ def run_test( mma_store_intrin, ) - f = tvm.build(sch.mod["main"], target="rocm") + f = tvm.compile(sch.mod["main"], target="rocm") dev = tvm.device("rocm", 0) if in_dtype == "float32": diff --git a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py index c160e4a31dc3..46487f969b96 100644 --- a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py +++ b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py @@ -142,7 +142,7 @@ def test_inject_async_copy(): continue with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - mod = tvm.build(tvm.IRModule.from_expr(f), target="cuda") + mod = tvm.compile(tvm.IRModule.from_expr(f), target="cuda") A_np = np.random.rand(32, 128).astype(dtype) B_np = np.zeros((32, 128)).astype(dtype) @@ -170,7 +170,7 @@ def test_inject_async_copy_shared_dyn(): return with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - mod = tvm.build(tvm.IRModule.from_expr(f), target="cuda") + mod = tvm.compile(tvm.IRModule.from_expr(f), target="cuda") A_np = np.random.rand(32, 128).astype("float16") B_np = np.random.rand(32, 128).astype("float16") @@ -228,7 +228,7 @@ def test_inject_async_copy_barrier(): if tvm.testing.is_ampere_or_newer(): with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - mod = tvm.build(tvm.IRModule.from_expr(f), target="cuda") + mod = tvm.compile(tvm.IRModule.from_expr(f), target="cuda") A_np = np.random.rand(32, 128).astype(dtype) B_np = np.zeros((32, 128)).astype(dtype) @@ -477,7 +477,7 @@ def simple_compute( mod = tvm.IRModule.from_expr(simple_compute) with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") generated_code = postproc_if_missing_async_support() assert generated_code == expected_cuda_script @@ -938,7 +938,7 @@ def complex_compute( mod = tvm.IRModule.from_expr(complex_compute) with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - tvm.build(mod, target="cuda") + tvm.compile(mod, target="cuda") generated_code = postproc_if_missing_async_support() # generated_code must contain " setp.ne.b32 p, %0, 0;" assert "setp.ne.b32" in generated_code diff --git a/tests/python/tir-transform/test_tir_transform_inject_software_pipeline.py b/tests/python/tir-transform/test_tir_transform_inject_software_pipeline.py index f332eaccc7cf..c4f2756251c5 100644 --- a/tests/python/tir-transform/test_tir_transform_inject_software_pipeline.py +++ b/tests/python/tir-transform/test_tir_transform_inject_software_pipeline.py @@ -1532,7 +1532,7 @@ def index_map(i, j): def build_and_run(sch): if tvm.testing.is_ampere_or_newer(): with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - f = tvm.build(sch.mod["main"], target="cuda") + f = tvm.compile(sch.mod["main"], target="cuda") dev = tvm.device("cuda", 0) a_np = np.random.uniform(size=(N, K)).astype("float16") diff --git a/tests/python/tir-transform/test_tir_transform_lower_intrin.py b/tests/python/tir-transform/test_tir_transform_lower_intrin.py index 3eb642fb51b3..f31cf559764d 100644 --- a/tests/python/tir-transform/test_tir_transform_lower_intrin.py +++ b/tests/python/tir-transform/test_tir_transform_lower_intrin.py @@ -47,7 +47,7 @@ def make_binds(i): return x C = te.compute((n,), make_binds) - f = tvm.build(te.create_prim_func([A, B, C]), "llvm") + f = tvm.compile(te.create_prim_func([A, B, C]), "llvm") a = tvm.nd.array(np.array([x for x, y in data], dtype=expr.dtype)) b = tvm.nd.array(np.array([y for x, y in data], dtype=expr.dtype)) c = tvm.nd.array(np.zeros(len(data), dtype=expr.dtype)) diff --git a/tests/python/tir-transform/test_tir_transform_lower_tvm_builtin.py b/tests/python/tir-transform/test_tir_transform_lower_tvm_builtin.py index 0a040b0eeadb..7d7f610123fe 100644 --- a/tests/python/tir-transform/test_tir_transform_lower_tvm_builtin.py +++ b/tests/python/tir-transform/test_tir_transform_lower_tvm_builtin.py @@ -179,7 +179,7 @@ def build_tir(): ) mod = build_tir() - f = tvm.build(mod, None) + f = tvm.compile(mod, None) a = tvm.nd.array(np.zeros(2, dtype="float32")) f(a) tvm.testing.assert_allclose(a.numpy(), expected_value) @@ -199,7 +199,7 @@ def variance4(rxplaceholder: T.Buffer((T.int64(1), T.int64(32), T.int64(25690112 T_subtract_1[cse_var_1] = rxplaceholder_1[cse_var_1] - rxplaceholder_red_1[ax1] func = variance4 - tvm.build(func, target="llvm") # should not crash + tvm.compile(func, target="llvm") # should not crash class TestLowerDeviceAllocate(tvm.testing.CompareBeforeAfter): diff --git a/tests/python/tir-transform/test_tir_transform_make_packed_api.py b/tests/python/tir-transform/test_tir_transform_make_packed_api.py index 8605d5185d90..c6763300629b 100644 --- a/tests/python/tir-transform/test_tir_transform_make_packed_api.py +++ b/tests/python/tir-transform/test_tir_transform_make_packed_api.py @@ -270,7 +270,7 @@ def func( ): pass - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") with pytest.raises(tvm.TVMError): built() @@ -283,7 +283,7 @@ def test_function_call_with_wrong_type_code(): def func(A: T.Buffer([16, 16], "int32")): pass - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") with pytest.raises(tvm.TVMError): built(0) @@ -297,7 +297,7 @@ def func(A: T.Buffer([16, 16], "int32"), B: T.Buffer([16, 16], "int32")): for i, j in T.grid(16, 16): B[i, j] = A[i, j] - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") A = tvm.nd.empty([16, 16], "int32", tvm.cpu()) B = tvm.nd.empty([16, 16], "int32", tvm.cpu()) @@ -316,7 +316,7 @@ def func(A: T.Buffer([16, 16], "int32"), B: T.Buffer([16, 16], "int32")): for i, j in T.grid(16, 16): B[i, j] = A[i, j] - built = tvm.build(func, target="llvm") + built = tvm.compile(func, target="llvm") A = tvm.nd.empty([16], "int32", tvm.cpu()) B = tvm.nd.empty([16], "int32", tvm.cpu()) diff --git a/tests/python/tir-transform/test_tir_transform_vectorize.py b/tests/python/tir-transform/test_tir_transform_vectorize.py index c5569c829ad5..5c25a1decace 100644 --- a/tests/python/tir-transform/test_tir_transform_vectorize.py +++ b/tests/python/tir-transform/test_tir_transform_vectorize.py @@ -363,7 +363,7 @@ def test_ir(A, B, C): ) try: - tvm.build(te.create_prim_func([A, B, C]), target="llvm") + tvm.compile(te.create_prim_func([A, B, C]), target="llvm") assert False except tvm.error.TVMError as e: error_msg = str(e).split("\n")[-1] @@ -796,7 +796,7 @@ def main(A: T.Buffer((25,), "float32"), B: T.Buffer((25,), "float32")): with tvm.target.Target(target): mod = tvm.tir.transform.VectorizeLoop()(Before) tvm.ir.assert_structural_equal(mod, After) - mod = tvm.build(mod, target=target) + mod = tvm.compile(mod, target=target) @pytest.mark.parametrize( @@ -824,7 +824,7 @@ def main(A: T.Buffer((25,), "int32"), B: T.Buffer((25,), "float32")): with pytest.raises(Exception) as e_info: with tvm.target.Target(target): mod = tvm.tir.transform.VectorizeLoop()(Before) - ex = tvm.build(mod, target=target) + ex = tvm.compile(mod, target=target) tvm.ir.assert_structural_equal(mod, After) assert "Intrinsic does not support vectors" in e_info.value.args[0] diff --git a/tests/python/tvmscript/test_tvmscript_ops.py b/tests/python/tvmscript/test_tvmscript_ops.py index dc539b9327c8..ba56ad20e6fd 100644 --- a/tests/python/tvmscript/test_tvmscript_ops.py +++ b/tests/python/tvmscript/test_tvmscript_ops.py @@ -100,7 +100,7 @@ def test_get_valid_counts_script_func(): mod = tvm.ir.IRModule({"get_valid_counts": get_valid_counts}) print(mod.script()) # check building - f = tvm.build(mod["get_valid_counts"], target=device) + f = tvm.compile(mod["get_valid_counts"], target=device) _check_get_valid_counts_with_numpy(f, (1, 2500, 6), 0.0, 0, 1) @@ -154,8 +154,8 @@ def test_alloc_zero_dim_buffer_round_trip(): func_with_block = alloc_zero_dim_buffer_block rt_func = tvm.script.from_source(func.script()) rt_func_with_block = tvm.script.from_source(func_with_block.script()) - rt_mod = tvm.build(rt_func, "llvm") - rt_mod_with_block = tvm.build(rt_func_with_block, "llvm") + rt_mod = tvm.compile(rt_func, "llvm") + rt_mod_with_block = tvm.compile(rt_func_with_block, "llvm") tvm.ir.assert_structural_equal( func.with_attr("global_symbol", "main"), func_with_block.with_attr("global_symbol", "main") ) @@ -175,7 +175,7 @@ def ceildiv_test(A: T.Buffer(16, "int32")): @tvm.testing.requires_llvm def test_ceildiv(): - f = tvm.build(ceildiv_test, "llvm") + f = tvm.compile(ceildiv_test, "llvm") a = tvm.nd.array(np.arange(16).astype("int32")) f(a) ref = (np.arange(16) + 3) // 4