You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by tq...@apache.org on 2026/02/16 00:57:38 UTC
(tvm) branch target created (now 3c471c53eb)
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a change to branch target
in repository https://gitbox.apache.org/repos/asf/tvm.git
at 3c471c53eb [REFACTOR][TARGET] Phase out legacy target string format
This branch includes the following new commits:
new 3c471c53eb [REFACTOR][TARGET] Phase out legacy target string format
The 1 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
(tvm) 01/01: [REFACTOR][TARGET] Phase out legacy target string format
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch target
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 3c471c53eb52d40525ad781fedf209eacc928236
Author: tqchen <ti...@gmail.com>
AuthorDate: Sun Feb 15 19:55:42 2026 -0500
[REFACTOR][TARGET] Phase out legacy target string format
This PR phases out legacy target string format in favor of
the json style format that is more well formed. It also simplfies
our overall code in handling multiple formats.
---
apps/android_rpc/tests/android_rpc_test.py | 2 +-
apps/ios_rpc/tests/ios_rpc_test.py | 2 +-
docs/how_to/tutorials/cross_compilation_and_rpc.py | 32 +-
.../how_to/tutorials/export_and_load_executable.py | 2 +-
include/tvm/target/target.h | 5 +-
python/tvm/contrib/hexagon/pytest_plugin.py | 16 +-
python/tvm/exec/gpu_memory_bandwidth.py | 2 +-
python/tvm/relax/frontend/torch/dynamo.py | 4 +-
python/tvm/relax/pipeline.py | 2 +-
python/tvm/runtime/module.py | 13 +-
python/tvm/s_tir/dlight/benchmark/bench.py | 2 +-
.../meta_schedule/testing/validate_database.py | 2 +-
python/tvm/target/target.py | 503 +++++++++++----------
python/tvm/testing/plugin.py | 18 +-
python/tvm/testing/utils.py | 21 +-
src/target/llvm/llvm_instance.cc | 136 ++++--
src/target/target.cc | 456 ++++---------------
tests/cpp/target/virtual_device_test.cc | 8 +-
tests/python/arith/test_arith_rewrite_simplify.py | 2 +-
tests/python/arith/test_arith_simplify.py | 4 +-
tests/python/codegen/test_target_codegen.py | 8 +-
.../python/codegen/test_target_codegen_aarch64.py | 30 +-
tests/python/codegen/test_target_codegen_arm.py | 14 +-
.../codegen/test_target_codegen_cross_llvm.py | 2 +-
tests/python/codegen/test_target_codegen_device.py | 2 +-
.../codegen/test_target_codegen_gpu_common.py | 4 +-
tests/python/codegen/test_target_codegen_llvm.py | 11 +-
.../python/codegen/test_target_codegen_llvm_vla.py | 50 +-
tests/python/codegen/test_target_codegen_opencl.py | 8 +-
tests/python/codegen/test_target_codegen_riscv.py | 48 +-
tests/python/codegen/test_target_codegen_vulkan.py | 8 +-
tests/python/codegen/test_target_codegen_x86.py | 22 +-
.../python/contrib/test_android/infrastructure.py | 2 +-
tests/python/nightly/test_nnapi/infrastructure.py | 2 +-
.../adreno/test_transform_annotate_custom_scope.py | 2 +-
tests/python/relax/backend/clml/conftest.py | 2 +-
tests/python/relax/backend/clml/utils.py | 4 +-
.../relax/test_backend_dispatch_sort_scan.py | 8 +-
tests/python/relax/test_frontend_nn_op.py | 4 +-
.../relax/test_meta_schedule_relax_integration.py | 2 +-
.../python/relax/test_transform_few_shot_tuning.py | 2 +-
.../test_transform_meta_schedule_apply_database.py | 2 +-
.../relax/test_transform_meta_schedule_tuning.py | 2 +-
.../python/relax/test_transform_realize_vdevice.py | 14 +-
.../python/relax/test_transform_update_vdevice.py | 4 +-
tests/python/relax/test_tvmscript_parser.py | 8 +-
tests/python/relax/texture/adreno_utils.py | 2 +-
tests/python/relax/texture/test_texture_nd.py | 7 +-
tests/python/runtime/test_runtime_module_load.py | 2 +-
tests/python/runtime/test_runtime_rpc.py | 4 +-
tests/python/s_tir/dlight/test_benchmark.py | 8 +-
tests/python/s_tir/dlight/test_gpu_gemv.py | 4 +-
tests/python/s_tir/dlight/test_gpu_matmul.py | 4 +-
.../meta_schedule/test_meta_schedule_cost_model.py | 2 +-
.../test_meta_schedule_measure_callback.py | 4 +-
.../test_meta_schedule_mutator_mutate_parallel.py | 2 +-
.../test_meta_schedule_mutator_mutate_tile_size.py | 4 +-
.../test_meta_schedule_mutator_mutate_unroll.py | 2 +-
.../test_meta_schedule_post_order_apply.py | 34 +-
...est_meta_schedule_postproc_rewrite_tensorize.py | 2 +-
...meta_schedule_postproc_rewrite_unbound_block.py | 2 +-
...test_meta_schedule_schedule_rule_add_rfactor.py | 4 +-
...eta_schedule_schedule_rule_apply_custom_rule.py | 2 +-
.../test_meta_schedule_schedule_rule_mlt_intrin.py | 8 +-
.../test_meta_schedule_schedule_rule_mlt_tc.py | 22 +-
...dule_schedule_rule_parallel_vectorize_unroll.py | 4 +-
.../test_meta_schedule_space_post_opt.py | 2 +-
.../test_meta_schedule_trace_apply.py | 2 +-
.../meta_schedule/test_meta_schedule_tune_tir.py | 4 +-
.../s_tir/schedule/test_tir_schedule_split_fuse.py | 6 +-
tests/python/target/test_arm_target.py | 76 +++-
tests/python/target/test_llvm_features_info.py | 4 +-
tests/python/target/test_riscv_features.py | 20 +-
tests/python/target/test_target_parser_mprofile.py | 8 +-
tests/python/target/test_target_target.py | 76 +++-
tests/python/target/test_x86_features.py | 161 ++++---
tests/python/tir-base/test_tir_intrin.py | 2 +-
.../tir-transform/test_tir_transform_helpers.py | 21 +-
.../test_tir_transform_split_host_device.py | 12 +-
.../tir-transform/test_tir_transform_vectorize.py | 15 +-
web/tests/python/prepare_test_libs.py | 4 +-
web/tests/python/relax_rpc_test.py | 4 +-
web/tests/python/webgpu_rpc_test.py | 4 +-
83 files changed, 1035 insertions(+), 1006 deletions(-)
diff --git a/apps/android_rpc/tests/android_rpc_test.py b/apps/android_rpc/tests/android_rpc_test.py
index 77e7cb565c..48c42a6e65 100644
--- a/apps/android_rpc/tests/android_rpc_test.py
+++ b/apps/android_rpc/tests/android_rpc_test.py
@@ -36,7 +36,7 @@ key = "android"
# Change target configuration.
# Run `adb shell cat /proc/cpuinfo` to find the arch.
arch = "arm64"
-target = "llvm -mtriple=%s-linux-android" % arch
+target = {"kind": "llvm", "mtriple": f"{arch}-linux-android"}
# whether enable to execute test on OpenCL target
test_opencl = False
diff --git a/apps/ios_rpc/tests/ios_rpc_test.py b/apps/ios_rpc/tests/ios_rpc_test.py
index 6757dfb47b..73d53c99e4 100644
--- a/apps/ios_rpc/tests/ios_rpc_test.py
+++ b/apps/ios_rpc/tests/ios_rpc_test.py
@@ -33,7 +33,7 @@ from tvm.contrib import utils, xcode
# Change target configuration, this is setting for iphone6s
arch = "arm64"
sdk = "iphoneos"
-target = "llvm -mtriple=%s-apple-darwin" % arch
+target = {"kind": "llvm", "mtriple": f"{arch}-apple-darwin"}
MODES = {"proxy": rpc.connect, "tracker": rpc.connect_tracker, "standalone": rpc.connect}
diff --git a/docs/how_to/tutorials/cross_compilation_and_rpc.py b/docs/how_to/tutorials/cross_compilation_and_rpc.py
index 229f40f938..b3af2dfd8d 100644
--- a/docs/how_to/tutorials/cross_compilation_and_rpc.py
+++ b/docs/how_to/tutorials/cross_compilation_and_rpc.py
@@ -108,7 +108,7 @@ mod = tvm.IRModule.from_expr(te.create_prim_func([A, B]).with_attr("global_symbo
######################################################################
# Then we cross compile the kernel.
-# The target should be 'llvm -mtriple=armv7l-linux-gnueabihf' for
+# The target should be {"kind": "llvm", "mtriple": "armv7l-linux-gnueabihf"} for
# Raspberry Pi 3B, but we use 'llvm' here to make this tutorial runnable
# on our webpage building server. See the detailed note in the following block.
@@ -117,7 +117,7 @@ local_demo = True
if local_demo:
target = "llvm"
else:
- target = "llvm -mtriple=armv7l-linux-gnueabihf"
+ target = {"kind": "llvm", "mtriple": "armv7l-linux-gnueabihf"}
func = tvm.compile(mod, target=target)
# save the lib at a local temp folder
@@ -132,8 +132,8 @@ func.export_library(path)
# to False and replace :code:`target` in :code:`build` with the appropriate
# target triple for your device. The target triple which might be
# different for different devices. For example, it is
-# :code:`'llvm -mtriple=armv7l-linux-gnueabihf'` for Raspberry Pi 3B and
-# :code:`'llvm -mtriple=aarch64-linux-gnu'` for RK3399.
+# :code:`{"kind": "llvm", "mtriple": "armv7l-linux-gnueabihf"}` for Raspberry Pi 3B and
+# :code:`{"kind": "llvm", "mtriple": "aarch64-linux-gnu"}` for RK3399.
#
# Usually, you can query the target by running :code:`gcc -v` on your
# device, and looking for the line starting with :code:`Target:`
@@ -228,7 +228,7 @@ def run_opencl():
# them according to your environment.
opencl_device_host = "10.77.1.145"
opencl_device_port = 9090
- target = tvm.target.Target("opencl", host="llvm -mtriple=aarch64-linux-gnu")
+ target = tvm.target.Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
# create schedule for the above "add one" compute declaration
mod = tvm.IRModule.from_expr(te.create_prim_func([A, B]))
@@ -356,23 +356,23 @@ def run_pytorch_model_via_rpc():
# Choose the appropriate target for your device:
#
# ARM devices:
- # - Raspberry Pi 3/4 (32-bit): "llvm -mtriple=armv7l-linux-gnueabihf"
- # - Raspberry Pi 4 (64-bit) / Jetson: "llvm -mtriple=aarch64-linux-gnu"
- # - Android: "llvm -mtriple=aarch64-linux-android"
+ # - Raspberry Pi 3/4 (32-bit): {"kind": "llvm", "mtriple": "armv7l-linux-gnueabihf"}
+ # - Raspberry Pi 4 (64-bit) / Jetson: {"kind": "llvm", "mtriple": "aarch64-linux-gnu"}
+ # - Android: {"kind": "llvm", "mtriple": "aarch64-linux-android"}
#
# x86 servers:
- # - Linux x86_64: "llvm -mtriple=x86_64-linux-gnu"
- # - With AVX-512: "llvm -mtriple=x86_64-linux-gnu -mcpu=skylake-avx512"
+ # - Linux x86_64: {"kind": "llvm", "mtriple": "x86_64-linux-gnu"}
+ # - With AVX-512: {"kind": "llvm", "mtriple": "x86_64-linux-gnu", "mcpu": "skylake-avx512"}
#
# RISC-V:
- # - RV64: "llvm -mtriple=riscv64-unknown-linux-gnu"
+ # - RV64: {"kind": "llvm", "mtriple": "riscv64-unknown-linux-gnu"}
#
# GPU targets:
- # - CUDA: tvm.target.Target("cuda", host="llvm -mtriple=x86_64-linux-gnu")
- # - OpenCL: tvm.target.Target("opencl", host="llvm -mtriple=aarch64-linux-gnu")
+ # - CUDA: tvm.target.Target("cuda", host={"kind": "llvm", "mtriple": "x86_64-linux-gnu"})
+ # - OpenCL: tvm.target.Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
#
# For this example, we use ARM 64-bit
- target = tvm.target.Target("llvm -mtriple=aarch64-linux-gnu")
+ target = tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
print(f"Cross-compiling for target: {target}")
# Apply optimization pipeline
@@ -527,12 +527,12 @@ def run_pytorch_model_via_rpc():
#
# # Example: ARM with NEON
# target = tvm.target.Target(
- # "llvm -mtriple=aarch64-linux-gnu -mattr=+neon"
+ # {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": "+neon"}
# )
#
# # Example: x86 with AVX-512
# target = tvm.target.Target(
- # "llvm -mtriple=x86_64-linux-gnu -mcpu=skylake-avx512"
+ # {"kind": "llvm", "mtriple": "x86_64-linux-gnu", "mcpu": "skylake-avx512"}
# )
#
# See :doc:`e2e_opt_model </how_to/tutorials/e2e_opt_model>` for detailed
diff --git a/docs/how_to/tutorials/export_and_load_executable.py b/docs/how_to/tutorials/export_and_load_executable.py
index 9f1225646c..4074dac4a6 100644
--- a/docs/how_to/tutorials/export_and_load_executable.py
+++ b/docs/how_to/tutorials/export_and_load_executable.py
@@ -336,7 +336,7 @@ if RUN_EXAMPLE:
# from tvm import relax
#
# # Step 1: Cross-compile for ARM target (on local machine)
-# TARGET = tvm.target.Target("llvm -mtriple=aarch64-linux-gnu")
+# TARGET = tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
# executable = tvm.compile(built_mod, target=TARGET)
# executable.export_library("mlp_arm.so")
#
diff --git a/include/tvm/target/target.h b/include/tvm/target/target.h
index 78d4d102f4..50839dc1ab 100644
--- a/include/tvm/target/target.h
+++ b/include/tvm/target/target.h
@@ -62,9 +62,8 @@ class TargetNode : public Object {
ffi::Map<ffi::String, Any> features;
/*!
- * \brief The raw string representation of the target
- * \return the full device string to pass to codegen::Build
- * \note It will be deprecated after the Target RFC is fully landed.
+ * \brief The JSON string representation of the target
+ * \return JSON string of the target configuration (e.g. {"kind": "llvm", "mcpu": "cortex-a53"})
*/
TVM_DLL const std::string& str() const;
/*! \return Export target to JSON-like configuration */
diff --git a/python/tvm/contrib/hexagon/pytest_plugin.py b/python/tvm/contrib/hexagon/pytest_plugin.py
index 303414e391..c289e2cfe6 100644
--- a/python/tvm/contrib/hexagon/pytest_plugin.py
+++ b/python/tvm/contrib/hexagon/pytest_plugin.py
@@ -39,11 +39,13 @@ ANDROID_SERIAL_NUMBER = "ANDROID_SERIAL_NUMBER"
ADB_SERVER_SOCKET = "ADB_SERVER_SOCKET"
RNG_SEEDED = False
-HEXAGON_AOT_LLVM_TARGET = (
- "llvm -keys=hexagon,cpu "
- "-mattr=+hvxv68,+hvx-length128b,+hvx-qfloat,-hvx-ieee-fp "
- "-mcpu=hexagonv68 -mtriple=hexagon"
-)
+HEXAGON_AOT_LLVM_TARGET = {
+ "kind": "llvm",
+ "keys": ["hexagon", "cpu"],
+ "mattr": ["+hvxv68", "+hvx-length128b", "+hvx-qfloat", "-hvx-ieee-fp"],
+ "mcpu": "hexagonv68",
+ "mtriple": "hexagon",
+}
@tvm.testing.fixture
@@ -316,7 +318,9 @@ aot_host_target = tvm.testing.parameter(HEXAGON_AOT_LLVM_TARGET)
def aot_target(aot_host_target):
if aot_host_target == "c":
yield tvm.target.hexagon("v68")
- elif aot_host_target.startswith("llvm"):
+ elif isinstance(aot_host_target, dict) and aot_host_target.get("kind") == "llvm":
+ yield aot_host_target
+ elif isinstance(aot_host_target, str) and aot_host_target.startswith("llvm"):
yield aot_host_target
else:
assert False, "Incorrect AoT host target: {aot_host_target}. Options are [c, llvm]."
diff --git a/python/tvm/exec/gpu_memory_bandwidth.py b/python/tvm/exec/gpu_memory_bandwidth.py
index ac462f85e3..0720d1d9aa 100644
--- a/python/tvm/exec/gpu_memory_bandwidth.py
+++ b/python/tvm/exec/gpu_memory_bandwidth.py
@@ -40,7 +40,7 @@ def _parse_args() -> argparse.Namespace:
--vec "1,2,4" \
Example for Android GPU: \
- python -m tvm.exec.gpu_memory_bandwidth "opencl" --target_host "llvm -mtriple=arm64-linux-android" \
+ python -m tvm.exec.gpu_memory_bandwidth "opencl" --target_host '{"kind": "llvm", "mtriple": "arm64-linux-android"}' \
--rpc_host "127.0.0.1" \
--rpc_port 9190 \
--rpc_key "android" \
diff --git a/python/tvm/relax/frontend/torch/dynamo.py b/python/tvm/relax/frontend/torch/dynamo.py
index 21388dbef7..668543fc8a 100644
--- a/python/tvm/relax/frontend/torch/dynamo.py
+++ b/python/tvm/relax/frontend/torch/dynamo.py
@@ -200,8 +200,8 @@ def llvm_target():
import platform
import subprocess
- AVX512_TARGET = "llvm -mcpu=skylake-avx512"
- AVX2_TARGET = "llvm -mcpu=core-avx2"
+ AVX512_TARGET = {"kind": "llvm", "mcpu": "skylake-avx512"}
+ AVX2_TARGET = {"kind": "llvm", "mcpu": "core-avx2"}
DEFAULT_TARGET = "llvm"
system = platform.system()
diff --git a/python/tvm/relax/pipeline.py b/python/tvm/relax/pipeline.py
index 3b613db03c..67b025d00d 100644
--- a/python/tvm/relax/pipeline.py
+++ b/python/tvm/relax/pipeline.py
@@ -151,7 +151,7 @@ def static_shape_tuning_pipeline(
mod = relax.pipeline.static_shape_tuning_pipeline(
total_trials=1000,
- target="llvm -num-cores 16",
+ target={"kind": "llvm", "num-cores": 16},
work_dir="tuning_logs",
cpu_weight_prepack=True,
max_trials_per_task=64,
diff --git a/python/tvm/runtime/module.py b/python/tvm/runtime/module.py
index 71b3bdd94b..d21e650eb1 100644
--- a/python/tvm/runtime/module.py
+++ b/python/tvm/runtime/module.py
@@ -17,6 +17,7 @@
# pylint: disable=invalid-name, unused-import, import-outside-toplevel, inconsistent-return-statements
"""Runtime Module namespace."""
+import json
import os
import struct
from typing import Sequence
@@ -215,7 +216,7 @@ class Module(_Module):
is_system_lib = False
has_c_module = False
system_lib_prefix = None
- llvm_target_string = None
+ llvm_target = None
global_object_format = "o"
def get_source_format_from_module(module):
@@ -250,7 +251,7 @@ class Module(_Module):
files.append(path_obj)
if module.kind == "llvm":
is_system_lib = module.get_function("__tvm_is_system_module")()
- llvm_target_string = module.get_function("_get_target_string")()
+ llvm_target = module.get_function("_get_target_string")()
system_lib_prefix = module.get_function("__tvm_get_system_lib_prefix")()
if not fcompile:
@@ -261,10 +262,10 @@ class Module(_Module):
else:
fcompile = _cc.create_shared
- if llvm_target_string is None and hasattr(fcompile, "get_target_triple"):
+ if llvm_target is None and hasattr(fcompile, "get_target_triple"):
triple = fcompile.get_target_triple()
assert triple, "Target triple should not be empty"
- llvm_target_string = "llvm -mtriple " + triple
+ llvm_target = json.dumps({"kind": "llvm", "mtriple": triple.strip()})
if getattr(fcompile, "need_system_lib", False) and not is_system_lib:
raise ValueError(f"{str(fcompile)} need --system-lib option")
@@ -275,12 +276,12 @@ class Module(_Module):
if fpack_imports is not None:
path_out = fpack_imports(self, is_system_lib, pack_lib_prefix, workspace_dir)
files.append(path_out)
- elif _ffi_api.RuntimeEnabled("llvm") and llvm_target_string:
+ elif _ffi_api.RuntimeEnabled("llvm") and llvm_target:
path_obj = os.path.join(
workspace_dir, f"{pack_lib_prefix}devc.{global_object_format}"
)
m = _ffi_api.ModulePackImportsToLLVM(
- self, is_system_lib, llvm_target_string, pack_lib_prefix
+ self, is_system_lib, llvm_target, pack_lib_prefix
)
m.write_to_file(path_obj)
files.append(path_obj)
diff --git a/python/tvm/s_tir/dlight/benchmark/bench.py b/python/tvm/s_tir/dlight/benchmark/bench.py
index 0968c3ad2b..da35541118 100644
--- a/python/tvm/s_tir/dlight/benchmark/bench.py
+++ b/python/tvm/s_tir/dlight/benchmark/bench.py
@@ -242,7 +242,7 @@ def benchmark_relax_func(
[Dict[str, str]],
Dict[str, int],
] = default_dym_var_sample_func,
- target: Union[str, tvm.target.Target] = "llvm -num-cores=4",
+ target: Union[str, dict, tvm.target.Target] = {"kind": "llvm", "num-cores": 4},
evaluator_config: Optional["EvaluatorConfig"] = None,
rpc_config: Optional["RPCConfig"] = None,
) -> None:
diff --git a/python/tvm/s_tir/meta_schedule/testing/validate_database.py b/python/tvm/s_tir/meta_schedule/testing/validate_database.py
index 3ed1887272..c315203fef 100644
--- a/python/tvm/s_tir/meta_schedule/testing/validate_database.py
+++ b/python/tvm/s_tir/meta_schedule/testing/validate_database.py
@@ -56,7 +56,7 @@ def _parse_args():
args.add_argument(
"--baseline-target",
type=Target,
- default="llvm -num-cores=1",
+ default='{"kind": "llvm", "num-cores": 1}',
required=False,
help="The baseline target to compile the original module.",
)
diff --git a/python/tvm/target/target.py b/python/tvm/target/target.py
index 31788d5f06..8671331dec 100644
--- a/python/tvm/target/target.py
+++ b/python/tvm/target/target.py
@@ -118,12 +118,6 @@ class Target(Object):
When using a dictionary or json string to configure target, the possible values are
same as target.
"""
- if isinstance(target, str) and "-libs=mkldnn" in target:
- target = target.replace("mkldnn", "dnnl")
- warnings.warn(
- "Legacy support of mkldnn is going to be deprecated. "
- "Please use -libs=dnnl instead.",
- )
if isinstance(target, (dict, str)):
target = convert(target)
if isinstance(host, (dict, str)):
@@ -302,7 +296,7 @@ class Target(Object):
Can convert from:
- None (to None).
- An existing TVM Target object.
- - A string, eg "cuda" or "cuda -arch=sm_80"
+ - A string, eg "cuda"
- A Python dictionary, eg {"kind": "cuda", "arch": "sm_80" }
"""
if target is None:
@@ -415,19 +409,7 @@ class Target(Object):
return target
-# TODO(@tvm-team): Deprecate the helper functions below. Encourage the usage of config dict instead.
-def _merge_opts(opts, new_opts):
- """Helper function to merge options"""
- if isinstance(new_opts, str):
- new_opts = new_opts.split()
- if new_opts:
- opt_set = set(opts)
- new_opts = [opt for opt in new_opts if opt not in opt_set]
- return opts + new_opts
- return opts
-
-
-def cuda(model="unknown", arch=None, options=None):
+def cuda(model="unknown", arch=None, **kwargs):
"""Returns a cuda target.
Parameters
@@ -436,81 +418,104 @@ def cuda(model="unknown", arch=None, options=None):
The model of cuda device (e.g. 1080ti)
arch: str
The cuda architecture (e.g. sm_61)
- options : str or list of str
- Additional options
+ kwargs: dict
+ Additional target attributes
"""
- opts = _merge_opts(["-model=%s" % model], options)
+ config = {"kind": "cuda", "model": model}
if arch:
- opts = _merge_opts(["-arch=%s" % arch], opts)
- if not any(["-arch" in opt for opt in opts]):
- warnings.warn("Try specifying cuda arch by adding 'arch=sm_xx' to your target.")
- return Target(" ".join(["cuda"] + opts))
+ config["arch"] = arch
+ else:
+ warnings.warn("Try specifying cuda arch by adding arch='sm_xx' to your target.")
+ config.update(kwargs)
+ return Target(config)
-def rocm(model="unknown", options=None):
+def rocm(model="unknown", **kwargs):
"""Returns a ROCM target.
Parameters
----------
model: str
The model of this device
- options : str or list of str
- Additional options
+ kwargs: dict
+ Additional target attributes
"""
- opts = _merge_opts(["-model=%s" % model], options)
- return Target(" ".join(["rocm"] + opts))
+ config = {"kind": "rocm", "model": model}
+ config.update(kwargs)
+ return Target(config)
-def mali(model="unknown", options=None):
+def mali(model="unknown", **kwargs):
"""Returns a ARM Mali GPU target.
Parameters
----------
model: str
The model of this device
- options : str or list of str
- Additional options
+ kwargs: dict
+ Additional target attributes
"""
- opts = ["-device=mali", "-model=%s" % model]
- opts = _merge_opts(opts, options)
- return Target(" ".join(["opencl"] + opts))
+ config = {"kind": "opencl", "device": "mali", "model": model}
+ config.update(kwargs)
+ return Target(config)
-def intel_graphics(model="unknown", options=None):
+def intel_graphics(model="unknown", **kwargs):
"""Returns an Intel Graphics target.
Parameters
----------
model: str
The model of this device
- options : str or list of str
- Additional options
+ kwargs: dict
+ Additional target attributes
"""
- opts = ["-device=intel_graphics", "-model=%s" % model, "-thread_warp_size=16"]
- opts = _merge_opts(opts, options)
- return Target(" ".join(["opencl"] + opts))
+ config = {
+ "kind": "opencl",
+ "device": "intel_graphics",
+ "model": model,
+ "thread_warp_size": 16,
+ }
+ config.update(kwargs)
+ return Target(config)
MICRO_SUPPORTED_MODELS = {
- "host": [],
- "atsamd51": ["-mcpu=cortex-m4"],
- "cxd5602gg": ["-mcpu=cortex-m4"],
- "esp32": [],
- "imxrt10xx": ["-mcpu=cortex-m7"],
- "mps2_an521": ["-mcpu=cortex-m33"],
- "mps3_an547": ["-mcpu=cortex-m55"],
- "nrf52840": ["-mcpu=cortex-m4+nodsp"],
- "nrf5340dk": ["-mcpu=cortex-m33"],
- "rp2040": ["-mcpu=cortex-m0"],
- "sam3x8e": ["-mcpu=cortex-m3"],
- "stm32f746xx": ["-mcpu=cortex-m7", "-march=armv7e-m"],
- "stm32h7xx": ["-mcpu=cortex-m7"],
- "stm32l4r5zi": ["-mcpu=cortex-m4"],
- "stm32u5xx": ["-mcpu=cortex-m33"],
- "zynq_mp_r5": ["-mcpu=cortex-r5"],
+ "host": {},
+ "atsamd51": {"mcpu": "cortex-m4"},
+ "cxd5602gg": {"mcpu": "cortex-m4"},
+ "esp32": {},
+ "imxrt10xx": {"mcpu": "cortex-m7"},
+ "mps2_an521": {"mcpu": "cortex-m33"},
+ "mps3_an547": {"mcpu": "cortex-m55"},
+ "nrf52840": {"mcpu": "cortex-m4+nodsp"},
+ "nrf5340dk": {"mcpu": "cortex-m33"},
+ "rp2040": {"mcpu": "cortex-m0"},
+ "sam3x8e": {"mcpu": "cortex-m3"},
+ "stm32f746xx": {"mcpu": "cortex-m7", "march": "armv7e-m"},
+ "stm32h7xx": {"mcpu": "cortex-m7"},
+ "stm32l4r5zi": {"mcpu": "cortex-m4"},
+ "stm32u5xx": {"mcpu": "cortex-m33"},
+ "zynq_mp_r5": {"mcpu": "cortex-r5"},
}
+def _parse_cli_opts_to_dict(opts):
+ """Convert a list of CLI-style options (e.g. ['-mcpu=cortex-a72']) to a dict."""
+ result = {}
+ for opt in opts:
+ opt = opt.lstrip("-")
+ if "=" in opt:
+ key, val = opt.split("=", 1)
+ # Handle comma-separated values as lists
+ if "," in val:
+ val = val.split(",")
+ result[key] = val
+ else:
+ result[opt] = True
+ return result
+
+
def arm_cpu(model="unknown", options=None):
"""Returns a ARM CPU target.
This function will also download pre-tuned op parameters when there is none.
@@ -523,51 +528,63 @@ def arm_cpu(model="unknown", options=None):
Additional options
"""
trans_table = {
- "pixel2": ["-model=snapdragon835", "-mtriple=arm64-linux-android", "-mattr=+neon"],
- "mate10": ["-model=kirin970", "-mtriple=arm64-linux-android", "-mattr=+neon"],
- "mate10pro": ["-model=kirin970", "-mtriple=arm64-linux-android", "-mattr=+neon"],
- "p20": ["-model=kirin970", "-mtriple=arm64-linux-android", "-mattr=+neon"],
- "p20pro": ["-model=kirin970", "-mtriple=arm64-linux-android", "-mattr=+neon"],
- "rasp3b": ["-model=bcm2837", "-mtriple=armv7l-linux-gnueabihf", "-mattr=+neon"],
- "rasp4b": [
- "-model=bcm2711",
- "-mtriple=armv8l-linux-gnueabihf",
- "-mattr=+neon",
- "-mcpu=cortex-a72",
- ],
- "rasp4b64": [
- "-model=bcm2711",
- "-mtriple=aarch64-linux-gnu",
- "-mattr=+neon",
- "-mcpu=cortex-a72",
- ],
- "rk3399": ["-model=rk3399", "-mtriple=aarch64-linux-gnu", "-mattr=+neon"],
- "pynq": ["-model=pynq", "-mtriple=armv7a-linux-eabi", "-mattr=+neon"],
- "ultra96": ["-model=ultra96", "-mtriple=aarch64-linux-gnu", "-mattr=+neon"],
- "beagleai": [
- "-model=beagleai",
- "-mtriple=armv7a-linux-gnueabihf",
- "-mattr=+neon,+vfp4,+thumb2",
- "-mcpu=cortex-a15",
- ],
- "stm32mp1": [
- "-model=stm32mp1",
- "-mtriple=armv7a-linux-gnueabihf",
- "-mattr=+neon,+vfp4,+thumb2",
- "-mcpu=cortex-a7",
- ],
- "thunderx": [
- "-model=thunderx",
- "-mtriple=aarch64-linux-gnu",
- "-mattr=+neon,+crc,+lse",
- "-mcpu=thunderxt88",
- ],
+ "pixel2": {"model": "snapdragon835", "mtriple": "arm64-linux-android", "mattr": ["+neon"]},
+ "mate10": {"model": "kirin970", "mtriple": "arm64-linux-android", "mattr": ["+neon"]},
+ "mate10pro": {"model": "kirin970", "mtriple": "arm64-linux-android", "mattr": ["+neon"]},
+ "p20": {"model": "kirin970", "mtriple": "arm64-linux-android", "mattr": ["+neon"]},
+ "p20pro": {"model": "kirin970", "mtriple": "arm64-linux-android", "mattr": ["+neon"]},
+ "rasp3b": {
+ "model": "bcm2837",
+ "mtriple": "armv7l-linux-gnueabihf",
+ "mattr": ["+neon"],
+ },
+ "rasp4b": {
+ "model": "bcm2711",
+ "mtriple": "armv8l-linux-gnueabihf",
+ "mattr": ["+neon"],
+ "mcpu": "cortex-a72",
+ },
+ "rasp4b64": {
+ "model": "bcm2711",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon"],
+ "mcpu": "cortex-a72",
+ },
+ "rk3399": {"model": "rk3399", "mtriple": "aarch64-linux-gnu", "mattr": ["+neon"]},
+ "pynq": {"model": "pynq", "mtriple": "armv7a-linux-eabi", "mattr": ["+neon"]},
+ "ultra96": {"model": "ultra96", "mtriple": "aarch64-linux-gnu", "mattr": ["+neon"]},
+ "beagleai": {
+ "model": "beagleai",
+ "mtriple": "armv7a-linux-gnueabihf",
+ "mattr": ["+neon", "+vfp4", "+thumb2"],
+ "mcpu": "cortex-a15",
+ },
+ "stm32mp1": {
+ "model": "stm32mp1",
+ "mtriple": "armv7a-linux-gnueabihf",
+ "mattr": ["+neon", "+vfp4", "+thumb2"],
+ "mcpu": "cortex-a7",
+ },
+ "thunderx": {
+ "model": "thunderx",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon", "+crc", "+lse"],
+ "mcpu": "thunderxt88",
+ },
}
- pre_defined_opt = trans_table.get(model, ["-model=%s" % model])
+ pre_defined = trans_table.get(model, {"model": model})
- opts = ["-keys=arm_cpu,cpu", "-device=arm_cpu"] + pre_defined_opt
- opts = _merge_opts(opts, options)
- return Target(" ".join(["llvm"] + opts))
+ config = {
+ "kind": "llvm",
+ "keys": ["arm_cpu", "cpu"],
+ "device": "arm_cpu",
+ }
+ config.update(pre_defined)
+ if options:
+ if isinstance(options, str):
+ options = options.split()
+ config.update(_parse_cli_opts_to_dict(options))
+ return Target(config)
def rasp(options=None):
@@ -584,17 +601,19 @@ def rasp(options=None):
return arm_cpu("rasp3b", options)
-def bifrost(model="unknown", options=None):
+def bifrost(model="unknown", **kwargs):
"""Return an ARM Mali GPU target (Bifrost architecture).
Parameters
----------
- options : str or list of str
- Additional options
+ model: str
+ The model of this device
+ kwargs: dict
+ Additional target attributes
"""
- opts = ["-device=bifrost", "-model=%s" % model]
- opts = _merge_opts(opts, options)
- return Target(" ".join(["opencl"] + opts))
+ config = {"kind": "opencl", "device": "bifrost", "model": model}
+ config.update(kwargs)
+ return Target(config)
def riscv_cpu(model="sifive-u54", options=None):
@@ -609,48 +628,51 @@ def riscv_cpu(model="sifive-u54", options=None):
Additional options
"""
trans_table = {
- "sifive-e31": [
- "-model=sifive-e31",
- "-mtriple=riscv32-unknown-linux-gnu",
- "-mcpu=sifive-e31",
- "-mabi=ilp32",
- # cc: riscv64-unknown-linux-gnu-g++ -march=rv32imac -mabi=ilp32 -mcpu=sifive-e31
- ],
- "sifive-e76": [
- "-model=sifive-e76",
- "-mtriple=riscv32-unknown-linux-gnu",
- "-mcpu=sifive-e76",
- "-mabi=ilp32",
- # cc: riscv64-unknown-linux-gnu-g++ -march=rv32imafc -mabi=ilp32 -mcpu=sifive-e76
- ],
- "sifive-u54": [
- "-model=sifive-u54",
- "-mtriple=riscv64-unknown-linux-gnu",
- "-mcpu=sifive-u54",
- "-mabi=lp64d",
- # cc: riscv64-unknown-linux-gnu-g++ -march=rv64gc -mabi=lp64d -mcpu=sifive-u54
- ],
- "sifive-u74": [
- "-model=sifive-u74",
- "-mtriple=riscv64-unknown-linux-gnu",
- "-mcpu=sifive-u74",
- "-mabi=lp64d",
- # cc: riscv64-unknown-linux-gnu-g++ -march=rv64gc -mabi=lp64d -mcpu=sifive-u74
- ],
- "licheepi3a": [
- "-num-cores=8",
- "-mtriple=riscv64-unknown-linux-gnu",
- "-mcpu=spacemit-x60",
- "-mfloat-abi=hard",
- "-mabi=lp64d",
- # cc: riscv64-unknown-linux-gnu-g++ -march=rv64gcv -mabi=lp64d -mcpu=spacemit-x60
- ],
+ "sifive-e31": {
+ "model": "sifive-e31",
+ "mtriple": "riscv32-unknown-linux-gnu",
+ "mcpu": "sifive-e31",
+ "mabi": "ilp32",
+ },
+ "sifive-e76": {
+ "model": "sifive-e76",
+ "mtriple": "riscv32-unknown-linux-gnu",
+ "mcpu": "sifive-e76",
+ "mabi": "ilp32",
+ },
+ "sifive-u54": {
+ "model": "sifive-u54",
+ "mtriple": "riscv64-unknown-linux-gnu",
+ "mcpu": "sifive-u54",
+ "mabi": "lp64d",
+ },
+ "sifive-u74": {
+ "model": "sifive-u74",
+ "mtriple": "riscv64-unknown-linux-gnu",
+ "mcpu": "sifive-u74",
+ "mabi": "lp64d",
+ },
+ "licheepi3a": {
+ "num-cores": 8,
+ "mtriple": "riscv64-unknown-linux-gnu",
+ "mcpu": "spacemit-x60",
+ "mfloat-abi": "hard",
+ "mabi": "lp64d",
+ },
}
- pre_defined_opt = trans_table.get(model, ["-model=%s" % model])
+ pre_defined = trans_table.get(model, {"model": model})
- opts = ["-keys=arm_cpu,cpu", "-device=arm_cpu"] + pre_defined_opt
- opts = _merge_opts(opts, options)
- return Target(" ".join(["llvm"] + opts))
+ config = {
+ "kind": "llvm",
+ "keys": ["arm_cpu", "cpu"],
+ "device": "arm_cpu",
+ }
+ config.update(pre_defined)
+ if options:
+ if isinstance(options, str):
+ options = options.split()
+ config.update(_parse_cli_opts_to_dict(options))
+ return Target(config)
def hexagon(cpu_ver="v68", **kwargs):
@@ -680,14 +702,6 @@ def hexagon(cpu_ver="v68", **kwargs):
Note: Floating point support in HVX requires LLVM 14+.
"""
- # Some of the target parameters correspond to target kind attributes
- # listed in src/target/target_kind.cc. For those parameters, their
- # names follow the attribute names with the exception of '_' being used
- # in place of '-'.
-
- # Example compiler arguments
- # llvm -mtriple=hexagon -mcpu=hexagonv68 -mattr=+hvxv68,+hvx-length128b
-
def get_arch_version(cpu_ver):
m = re.match(r"v([0-9]+).*", cpu_ver)
assert m
@@ -714,110 +728,89 @@ def hexagon(cpu_ver="v68", **kwargs):
}
return default_vtcm_sizes.get(cpu_ver, 0)
- # Target configuration:
arch_version = get_arch_version(cpu_ver)
- config = {
+ local_config = {
"hvx": 128,
"llvm_options": None,
"use_qfloat": arch_version >= 68,
"use_ieee_fp": False,
"vtcm_capacity": get_vtcm_capacity(cpu_ver),
}
- config.update(kwargs)
+ local_config.update(kwargs)
# Warn about obsolete parameter names.
- if config.get("sim_args") or config.get("sim_options"):
+ if local_config.get("sim_args") or local_config.get("sim_options"):
msg = (
"Setting simulator options in target is deprecated, set environment variable "
"HEXAGON_SIM_ARGS instead"
)
warnings.warn(msg, stacklevel=2)
- if config.get("llvm_args"):
+ if local_config.get("llvm_args"):
msg = "The keyword parameter 'llvm_args' is deprecated, use 'llvm_options' instead"
warnings.warn(msg, stacklevel=2)
- config.update({"llvm_options": config["llvm_args"]})
-
- # LLVM target string
- def create_llvm_target(cpu_ver, config):
- """Create LLVM target string."""
-
- target = " -mtriple=hexagon"
- mcpu = " -mcpu=hexagon" + cpu_ver
-
- # Process the options that affect target features and return the
- # target feature string.
- def create_target_features(config):
- features = {
- "use_qfloat": "hvx-qfloat",
- "use_ieee_fp": "hvx-ieee-fp",
- }
- tfs = []
- if config["hvx"] > 0:
- valid_hvx = [0, 64, 128]
- if not config["hvx"] in valid_hvx:
- raise ValueError("Invalid hvx value, should be one of " + str(valid_hvx))
- tfs += ["+hvx" + cpu_ver, "+hvx-length" + str(config["hvx"]) + "b"]
- else:
- tfs += ["-hvx"]
- # All the additional features happen to only apply to v68+.
- # Don't bother applying them (even with '-') to lower versions.
- if arch_version >= 68:
- tfs += ["-+"[config[f]] + features[f] for f in features]
-
- return "-mattr=" + ",".join(tfs) if tfs else ""
-
- return target + mcpu + " " + create_target_features(config)
-
- # LLVM options string
- def create_llvm_options(cpu_ver, config): # pylint: disable=unused-argument
- """Create LLVM options string."""
-
- llvm_options = config["llvm_options"]
-
- # To enable auto-vectorization for v68 target added the below llvm-option by default
- if arch_version == 68:
- if not llvm_options:
- llvm_options = ""
- llvm_options += " -force-hvx-float"
-
- # TVM's option parser doesn't allow '=' in values, but '=' can
- # appear in LLVM flags. Replace it with '@', since it's unlikely
- # that '@' will be used in another context.
- if llvm_options is None or len(llvm_options.strip()) == 0:
- return ""
- args = [s.replace("=", "@") for s in llvm_options.split()]
- return "--llvm-options=" + ",".join(args)
-
- target_str = create_llvm_target(cpu_ver, config)
- llvm_str = create_llvm_options(cpu_ver, config)
-
- args_list = target_str.split() + llvm_str.split()
-
- num_cores = config["num_cores"] if "num_cores" in kwargs else 4
- args_list.append("--num-cores=%d" % num_cores)
- args_list.append("--vtcm-capacity=%d" % config["vtcm_capacity"])
-
- return Target(" ".join(["hexagon"] + args_list))
+ local_config.update({"llvm_options": local_config["llvm_args"]})
+
+ # Build mattr list from config
+ features_map = {
+ "use_qfloat": "hvx-qfloat",
+ "use_ieee_fp": "hvx-ieee-fp",
+ }
+ mattr = []
+ if local_config["hvx"] > 0:
+ valid_hvx = [0, 64, 128]
+ if local_config["hvx"] not in valid_hvx:
+ raise ValueError("Invalid hvx value, should be one of " + str(valid_hvx))
+ mattr += ["+hvx" + cpu_ver, "+hvx-length" + str(local_config["hvx"]) + "b"]
+ else:
+ mattr += ["-hvx"]
+ if arch_version >= 68:
+ for f, feat_name in features_map.items():
+ mattr.append(("-", "+")[local_config[f]] + feat_name)
+
+ # Build llvm-options list
+ llvm_options_list = []
+ llvm_options = local_config["llvm_options"]
+ if arch_version == 68:
+ if not llvm_options:
+ llvm_options = ""
+ llvm_options += " -force-hvx-float"
+ if llvm_options and len(llvm_options.strip()) > 0:
+ llvm_options_list = [s.replace("=", "@") for s in llvm_options.split()]
+
+ num_cores = local_config["num_cores"] if "num_cores" in kwargs else 4
+
+ target_config = {
+ "kind": "hexagon",
+ "mtriple": "hexagon",
+ "mcpu": "hexagon" + cpu_ver,
+ "mattr": mattr,
+ "num-cores": num_cores,
+ "vtcm-capacity": local_config["vtcm_capacity"],
+ }
+ if llvm_options_list:
+ target_config["llvm-options"] = llvm_options_list
+
+ return Target(target_config)
STM32_SUPPORTED_SERIES = {
# High-Performance
- "stm32H7xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m7", "-march=armv7e-m"],
- "stm32F7xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m7"],
- "stm32F4xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m4"],
- "stm32F2xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m3"],
+ "stm32H7xx": {"mcpu": "cortex-m7", "march": "armv7e-m"},
+ "stm32F7xx": {"mcpu": "cortex-m7"},
+ "stm32F4xx": {"mcpu": "cortex-m4"},
+ "stm32F2xx": {"mcpu": "cortex-m3"},
# Mainstream
- "stm32G0xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m0+"],
- "stm32F0xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m0"],
- "stm32F1xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m3"],
- "stm32G4xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m4"],
- "stm32F3xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m4"],
+ "stm32G0xx": {"mcpu": "cortex-m0+"},
+ "stm32F0xx": {"mcpu": "cortex-m0"},
+ "stm32F1xx": {"mcpu": "cortex-m3"},
+ "stm32G4xx": {"mcpu": "cortex-m4"},
+ "stm32F3xx": {"mcpu": "cortex-m4"},
# Low-power
- "stm32U5xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m33"],
- "stm32L5xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m33"],
- "stm32L4xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m4"],
- "stm32L1xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m3"],
- "stm32L0xx": ["-keys=arm_cpu,cpu", "-device=arm_cpu", "-mcpu=cortex-m0+"],
+ "stm32U5xx": {"mcpu": "cortex-m33"},
+ "stm32L5xx": {"mcpu": "cortex-m33"},
+ "stm32L4xx": {"mcpu": "cortex-m4"},
+ "stm32L1xx": {"mcpu": "cortex-m3"},
+ "stm32L0xx": {"mcpu": "cortex-m0+"},
}
@@ -831,11 +824,19 @@ def stm32(series="unknown", options=None):
options : str or list of str
Additional options
"""
-
if series not in STM32_SUPPORTED_SERIES:
raise ValueError(f"Series {series} is not supported by tvm.target.stm32.")
- opts = _merge_opts(STM32_SUPPORTED_SERIES[series], options)
- return Target(" ".join(["c"] + opts))
+ config = {
+ "kind": "c",
+ "keys": ["arm_cpu", "cpu"],
+ "device": "arm_cpu",
+ }
+ config.update(STM32_SUPPORTED_SERIES[series])
+ if options:
+ if isinstance(options, str):
+ options = options.split()
+ config.update(_parse_cli_opts_to_dict(options))
+ return Target(config)
def adreno(model="unknown", options=None, cfg=None, backend="opencl"):
@@ -855,13 +856,21 @@ def adreno(model="unknown", options=None, cfg=None, backend="opencl"):
if backend not in ["opencl", "vulkan"]:
raise ValueError(f"Unsupported API: {backend}. Must be 'opencl' or 'vulkan'.")
- keys = f"adreno,{backend},gpu"
+ keys = ["adreno", backend, "gpu"]
if cfg:
- keys += f",{cfg}"
+ keys.append(cfg)
- opts = ["-device=adreno", f"--keys={keys}", f"-model={model}"]
- opts = _merge_opts(opts, options)
- return Target(" ".join([backend] + opts))
+ config = {
+ "kind": backend,
+ "device": "adreno",
+ "keys": keys,
+ "model": model,
+ }
+ if options:
+ if isinstance(options, str):
+ options = options.split()
+ config.update(_parse_cli_opts_to_dict(options))
+ return Target(config)
def create(target):
diff --git a/python/tvm/testing/plugin.py b/python/tvm/testing/plugin.py
index 6963eab292..a2995b9489 100644
--- a/python/tvm/testing/plugin.py
+++ b/python/tvm/testing/plugin.py
@@ -65,7 +65,12 @@ def pytest_configure(config):
for feature in utils.Feature._all_features.values():
feature._register_marker(config)
- print("enabled targets:", "; ".join(map(lambda x: x[0], utils.enabled_targets())))
+ print(
+ "enabled targets:",
+ "; ".join(
+ map(lambda x: str(x[0]) if isinstance(x[0], dict) else x[0], utils.enabled_targets())
+ ),
+ )
print("pytest marker:", config.option.markexpr)
@@ -97,6 +102,8 @@ def pytest_collection_modifyitems(config, items):
@pytest.fixture
def dev(target):
"""Give access to the device to tests that need it."""
+ if isinstance(target, dict):
+ return tvm.device(target["kind"])
return tvm.device(target)
@@ -182,7 +189,12 @@ def _add_target_specific_marks(metafunc):
if mark in metafunc.definition.own_markers:
xfail_targets = getattr(metafunc.function, "tvm_known_failing_targets", [])
- target_kind = target.split()[0] if isinstance(target, str) else target.kind.name
+ if isinstance(target, str):
+ target_kind = target.split()[0]
+ elif isinstance(target, dict):
+ target_kind = target["kind"]
+ else:
+ target_kind = target.kind.name
if target_kind in xfail_targets:
additional_marks.append(
pytest.mark.xfail(
@@ -279,7 +291,7 @@ def _sort_tests(items):
def _target_to_requirement(target):
- if isinstance(target, str):
+ if isinstance(target, (str, dict)):
target = tvm.target.Target(target)
# mapping from target to decorator
diff --git a/python/tvm/testing/utils.py b/python/tvm/testing/utils.py
index 9b511fd849..e30865c62e 100644
--- a/python/tvm/testing/utils.py
+++ b/python/tvm/testing/utils.py
@@ -407,7 +407,10 @@ def _get_targets(target_names=None):
targets = []
for target in target_names:
- target_kind = target.split()[0]
+ if isinstance(target, dict):
+ target_kind = target["kind"]
+ else:
+ target_kind = target.split()[0]
if target_kind == "cuda" and "cudnn" in tvm.target.Target(target).attrs.get("libs", []):
is_enabled = tvm.support.libinfo()["USE_CUDNN"].lower() in ["on", "true", "1"]
@@ -451,10 +454,10 @@ DEFAULT_TEST_TARGETS = [
"llvm",
"cuda",
"nvptx",
- "vulkan -from_device=0",
+ {"kind": "vulkan", "from_device": 0},
"opencl",
- "opencl -device=mali",
- "opencl -device=intel_graphics",
+ {"kind": "opencl", "device": "mali"},
+ {"kind": "opencl", "device": "intel_graphics"},
"metal",
"rocm",
"hexagon",
@@ -519,7 +522,7 @@ def enabled_targets():
A list of pairs of all enabled devices and the associated context
"""
- return [(t["target"], tvm.device(t["target"])) for t in _get_targets() if t["is_runnable"]]
+ return [(t["target"], tvm.device(t["target_kind"])) for t in _get_targets() if t["is_runnable"]]
class Feature:
@@ -658,8 +661,12 @@ class Feature:
if self.target_kind_enabled is not None:
target_kind = self.target_kind_enabled.split()[0]
+
+ def _get_target_kind(t):
+ return t["kind"] if isinstance(t, dict) else t.split()[0]
+
yield pytest.mark.skipif(
- all(enabled.split()[0] != target_kind for enabled in _tvm_test_targets()),
+ all(_get_target_kind(enabled) != target_kind for enabled in _tvm_test_targets()),
reason=(
f"{self.target_kind_enabled} tests disabled "
f"by TVM_TEST_TARGETS environment variable"
@@ -1032,7 +1039,7 @@ requires_aprofile_aem_fvp = Feature(
def _has_cpu_feat(features):
cpu = codegen.llvm_get_system_cpu()
triple = codegen.llvm_get_system_triple()
- target = "llvm -mtriple=%s -mcpu=%s" % (triple, cpu)
+ target = {"kind": "llvm", "mtriple": triple, "mcpu": cpu}
has_feat = codegen.target_has_features(features, tvm.target.Target(target))
return has_feat
diff --git a/src/target/llvm/llvm_instance.cc b/src/target/llvm/llvm_instance.cc
index ecb887243d..b634140af8 100644
--- a/src/target/llvm/llvm_instance.cc
+++ b/src/target/llvm/llvm_instance.cc
@@ -455,113 +455,170 @@ std::string LLVMTargetInfo::GetTargetFeatureString() const { //
return Join(",", attrs_);
}
+static void JSONStr(std::ostringstream& os, const std::string& s) {
+ os << '"';
+ for (char c : s) {
+ if (c == '"')
+ os << "\\\"";
+ else if (c == '\\')
+ os << "\\\\";
+ else
+ os << c;
+ }
+ os << '"';
+}
+
std::string LLVMTargetInfo::str() const {
std::ostringstream os;
- os << "llvm";
+ // Emit a sorted JSON dict for LLVM target representation.
+ // We build key-value pairs, sort by key, then emit.
+ std::vector<std::pair<std::string, std::string>> kvs;
+
+ // kind is always "llvm"
+ {
+ std::ostringstream v;
+ v << "\"llvm\"";
+ kvs.push_back({"kind", v.str()});
+ }
+
if (!triple_.empty()) {
- os << " -mtriple=" << triple_;
+ std::ostringstream v;
+ JSONStr(v, triple_);
+ kvs.push_back({"mtriple", v.str()});
}
if (!cpu_.empty() && cpu_ != defaults::cpu) {
- os << " -mcpu=" << cpu_;
+ std::ostringstream v;
+ JSONStr(v, cpu_);
+ kvs.push_back({"mcpu", v.str()});
}
if (!attrs_.empty()) {
- os << " -mattr=" << GetTargetFeatureString();
+ std::ostringstream v;
+ v << '[';
+ bool first = true;
+ for (const auto& attr : attrs_) {
+ if (!first) v << ", ";
+ first = false;
+ JSONStr(v, attr);
+ }
+ v << ']';
+ kvs.push_back({"mattr", v.str()});
}
switch (target_options_.FloatABIType) {
case llvm::FloatABI::Soft:
- os << " -mfloat-abi=soft";
+ kvs.push_back({"mfloat-abi", "\"soft\""});
break;
case llvm::FloatABI::Hard:
- os << " -mfloat-abi=hard";
+ kvs.push_back({"mfloat-abi", "\"hard\""});
break;
case llvm::FloatABI::Default:
break;
}
if (!target_options_.MCOptions.ABIName.empty()) {
- os << " -mabi=" << target_options_.MCOptions.ABIName;
+ std::ostringstream v;
+ JSONStr(v, target_options_.MCOptions.ABIName);
+ kvs.push_back({"mabi", v.str()});
}
bool do_individual = true;
#if TVM_LLVM_VERSION >= 60
if (fast_math_flags_.isFast()) {
- os << " -fast-math";
+ kvs.push_back({"fast-math", "true"});
do_individual = false;
}
#else
if (fast_math_flags_.unsafeAlgebra()) {
- os << " -fast-math";
+ kvs.push_back({"fast-math", "true"});
do_individual = false;
}
#endif
if (do_individual) {
- if (fast_math_flags_.noNaNs()) os << " -fast-math-nnan";
- if (fast_math_flags_.noInfs()) os << " -fast-math-ninf";
- if (fast_math_flags_.noSignedZeros()) os << " -fast-math-nsz";
- if (fast_math_flags_.allowReciprocal()) os << " -fast-math-arcp";
+ if (fast_math_flags_.noNaNs()) kvs.push_back({"fast-math-nnan", "true"});
+ if (fast_math_flags_.noInfs()) kvs.push_back({"fast-math-ninf", "true"});
+ if (fast_math_flags_.noSignedZeros()) kvs.push_back({"fast-math-nsz", "true"});
+ if (fast_math_flags_.allowReciprocal()) kvs.push_back({"fast-math-arcp", "true"});
#if TVM_LLVM_VERSION >= 50
- if (fast_math_flags_.allowContract()) os << " -fast-math-contract";
+ if (fast_math_flags_.allowContract()) kvs.push_back({"fast-math-contract", "true"});
#endif
#if TVM_LLVM_VERSION >= 60
- if (fast_math_flags_.allowReassoc()) os << " -fast-math-reassoc";
- if (fast_math_flags_.approxFunc()) os << " -fast-math-afn";
+ if (fast_math_flags_.allowReassoc()) kvs.push_back({"fast-math-reassoc", "true"});
+ if (fast_math_flags_.approxFunc()) kvs.push_back({"fast-math-afn", "true"});
#endif
}
#if TVM_LLVM_VERSION <= 170
if (opt_level_ != defaults::opt_level) {
- os << " -opt-level=";
+ std::string level;
switch (opt_level_) {
case llvm::CodeGenOpt::None:
- os << "0";
+ level = "0";
break;
case llvm::CodeGenOpt::Less:
- os << "1";
+ level = "1";
break;
case llvm::CodeGenOpt::Default:
- os << "2";
+ level = "2";
break;
case llvm::CodeGenOpt::Aggressive:
- os << "3";
+ level = "3";
break;
}
+ kvs.push_back({"opt-level", level});
}
#else
if (opt_level_ != defaults::opt_level) {
- os << " -opt-level=";
+ std::string level;
switch (opt_level_) {
case llvm::CodeGenOptLevel::None:
- os << "0";
+ level = "0";
break;
case llvm::CodeGenOptLevel::Less:
- os << "1";
+ level = "1";
break;
case llvm::CodeGenOptLevel::Default:
- os << "2";
+ level = "2";
break;
case llvm::CodeGenOptLevel::Aggressive:
- os << "3";
+ level = "3";
break;
}
+ kvs.push_back({"opt-level", level});
}
#endif
- if (size_t num = llvm_options_.size(); num > 0) {
- os << " -cl-opt=";
- std::vector<std::string> opts;
+ if (!llvm_options_.empty()) {
+ std::ostringstream v;
+ v << '[';
+ bool first = true;
for (const Option& opt : llvm_options_) {
- std::stringstream os;
- os << opt;
- opts.emplace_back(os.str());
+ if (!first) v << ", ";
+ first = false;
+ std::stringstream opt_s;
+ opt_s << opt;
+ JSONStr(v, opt_s.str());
}
- auto* quote = num > 1 ? "'" : "";
- os << quote << Join(",", opts) << quote;
+ v << ']';
+ kvs.push_back({"cl-opt", v.str()});
}
if (jit_engine_ != "orcjit") {
- os << " -jit=" << jit_engine_;
+ std::ostringstream v;
+ JSONStr(v, jit_engine_);
+ kvs.push_back({"jit", v.str()});
+ }
+
+ // Sort by key for deterministic output
+ std::sort(kvs.begin(), kvs.end());
+
+ os << '{';
+ bool first = true;
+ for (const auto& kv : kvs) {
+ if (!first) os << ", ";
+ first = false;
+ os << '"' << kv.first << "\": " << kv.second;
}
+ os << '}';
return os.str();
}
@@ -973,18 +1030,19 @@ std::string LLVMTarget::GetTargetMetadata(const llvm::Module& module) {
if (llvm::Metadata* tvm_target = module.getModuleFlag("tvm_target")) {
auto* mdstr = llvm::cast<llvm::MDString>(tvm_target);
llvm::StringRef meta = mdstr->getString();
+ // Accept both JSON form (starts with '{') and legacy CLI form (starts with 'llvm')
#if TVM_LLVM_VERSION >= 180
- if (meta.starts_with("llvm")) {
+ if (meta.starts_with("{") || meta.starts_with("llvm")) {
#else
- if (meta.startswith("llvm")) {
+ if (meta.startswith("{") || meta.startswith("llvm")) {
#endif
return meta.str();
}
}
#if TVM_LLVM_VERSION >= 210
- return "llvm -mtriple " + module.getTargetTriple().str();
+ return "{\"kind\": \"llvm\", \"mtriple\": \"" + module.getTargetTriple().str() + "\"}";
#else
- return "llvm -mtriple " + module.getTargetTriple();
+ return "{\"kind\": \"llvm\", \"mtriple\": \"" + module.getTargetTriple() + "\"}";
#endif
}
diff --git a/src/target/target.cc b/src/target/target.cc
index 1f7e0a2868..113388623d 100644
--- a/src/target/target.cc
+++ b/src/target/target.cc
@@ -31,13 +31,10 @@
#include <tvm/tir/expr.h>
#include <algorithm>
-#include <cctype>
-#include <ios>
#include <sstream>
#include <stack>
#include <string>
#include <unordered_map>
-#include <utility>
#include <vector>
namespace tvm {
@@ -51,15 +48,12 @@ class TargetInternal {
static ffi::Map<ffi::String, ffi::Any> Export(Target target) { return target->Export(); }
static const TargetKindNode::ValueTypeInfo& FindTypeInfo(const TargetKind& kind,
const std::string& key);
- static ffi::Optional<ffi::String> StringifyAttrsToRaw(
- const ffi::Map<ffi::String, ffi::Any>& attrs);
- static Any ParseType(const std::string& str, const TargetKindNode::ValueTypeInfo& info);
static Any ParseType(const Any& obj, const TargetKindNode::ValueTypeInfo& info);
static ObjectPtr<TargetNode> FromString(const ffi::String& tag_or_config_or_target_str);
static ObjectPtr<TargetNode> FromConfigString(const ffi::String& config_str);
- static ObjectPtr<TargetNode> FromRawString(const ffi::String& target_str);
static ObjectPtr<TargetNode> FromConfig(ffi::Map<ffi::String, ffi::Any> config);
static void ConstructorDispatcher(ffi::PackedArgs args, ffi::Any* rv);
+ static std::string ExportToJSON(const ffi::Map<ffi::String, ffi::Any>& data);
static Target WithHost(const Target& target, const Target& target_host) {
ObjectPtr<TargetNode> n = ffi::make_object<TargetNode>(*target.get());
n->host = target_host;
@@ -69,17 +63,8 @@ class TargetInternal {
private:
static std::unordered_map<ffi::String, ffi::Any> QueryDevice(int device_id,
const TargetNode* target);
- static bool IsQuoted(const std::string& str);
- static std::string Quote(const std::string& str);
- static std::string JoinString(const std::vector<std::string>& array, char separator);
- static std::vector<std::string> SplitString(const std::string& str, char separator);
- static std::string Interpret(const std::string& str);
- static std::string Uninterpret(const std::string& str);
- static std::string StringifyAtomicType(const Any& obj);
- static std::string StringifyArray(const ffi::ArrayObj& array);
-
- static constexpr char quote = '\'';
- static constexpr char escape = '\\';
+ static void JSONSerializeValue(std::ostringstream& os, const Any& obj);
+ static void JSONSerializeString(std::ostringstream& os, const std::string& str);
};
/********** Helper functions **********/
@@ -127,211 +112,6 @@ static TargetKind GetTargetKind(const ffi::String& name) {
return kind.value();
}
-static std::string RemovePrefixDashes(const std::string& s) {
- int n_dashes = 0;
- int len = s.length();
- for (; n_dashes < len && s[n_dashes] == '-'; ++n_dashes) {
- }
- if (n_dashes == 0) {
- TVM_FFI_THROW(ValueError) << "Attribute keys should start with '-', not an attribute key: " + s;
- }
- if (n_dashes >= len) {
- TVM_FFI_THROW(ValueError) << "Not an attribute key: " + s;
- }
- return s.substr(n_dashes);
-}
-
-bool TargetInternal::IsQuoted(const std::string& str) {
- std::string::size_type start = 0, end = str.size();
- if (end < 2 || str[start] != quote || str[end - 1] != quote) {
- return false;
- }
- bool escaping = false;
- for (auto i = start + 1, e = end - 1; i < e; ++i) {
- if (escaping) {
- escaping = false;
- } else if (str[i] == escape) {
- escaping = true;
- } else if (str[i] == quote) {
- return false;
- }
- }
- // If the reduced string ends with \, then the terminating quote is escaped.
- return !escaping;
-}
-
-std::string TargetInternal::Quote(const std::string& str) {
- std::string result(1, quote);
- result.append(str);
- result.push_back(quote);
- return result;
-}
-
-std::string TargetInternal::JoinString(const std::vector<std::string>& array, char separator) {
- std::string result;
- ICHECK(separator != quote && separator != escape)
- << "string join separator cannot be " << quote << " or " << escape;
-
- bool is_first = true;
- for (const auto& s : array) {
- if (!is_first) {
- result.push_back(separator);
- }
- result.append(s);
- is_first = false;
- }
-
- return result;
-}
-
-std::vector<std::string> TargetInternal::SplitString(const std::string& str, char separator) {
- std::vector<std::string> output;
-
- const char* start = str.data();
- const char* end = start + str.size();
- const char* pos = start;
-
- std::stringstream current_word;
-
- auto finish_word = [&]() {
- std::string word = current_word.str();
- if (word.size()) {
- output.push_back(word);
- current_word.str("");
- }
- };
-
- bool pos_quoted = false;
-
- while (pos < end) {
- if ((*pos == separator) && !pos_quoted) {
- finish_word();
- pos++;
- } else if (*pos == escape && pos + 1 < end) {
- current_word << escape;
- current_word << pos[1];
- pos += 2;
- } else if (*pos == quote) {
- current_word << quote;
- pos_quoted = !pos_quoted;
- pos++;
- } else {
- current_word << *pos;
- pos++;
- }
- }
-
- ICHECK(!pos_quoted) << "Mismatched quotes '' in string";
-
- finish_word();
-
- return output;
-}
-
-std::string TargetInternal::Interpret(const std::string& str) {
- // ffi::String interpretation deals with quotes (') and escapes(\).
- // - An escape character must be followed by another character forming an
- // "escape sequence". (Trailing escape is not allowed.) An escape prevents
- // interpretation of the character that follows. This happens regardless of
- // whether the escape sequence appears within quoted substring or not.
- // - A quote character, when interpreted, marks the beginning or the end of a
- // quoted substring. (A quoted substring cannot contain unescaped quotes.)
- // - Any other character, when interpreted, represents itself.
- //
- // Interpretation happens in two steps:
- // 1. If the entire string is quoted, the quotes are removed first, and the
- // resulting string is treated as unquoted.
- // 2. Each character or escape sequence is interpreted, and the result is copied
- // to the result. When not inside a quoted substring, the interpretation of an
- // escape sequence is the escaped character, otherwise it is the entire escape
- // sequence.
- //
- // Examples:
- // blah -> blah Nothing happened
- // 'blah' -> blah Enclosing quotes removed
- // 'bl'ah -> 'bl'ah Non-enclosing quotes remain
- // '\'blah\'' -> 'blah' Enclosing quotes removed, escaped quotes
- // interpreted.
- // '\'\\\'blah\\\'\'' -> '\'blah\'' Same as above.
- //
- // Note that
- // '\'\\\'blah\\\'\'' -> '\'blah\'' -> 'blah'
-
- std::string result;
- if (str.empty()) {
- return result;
- }
-
- // Check if the entire string is enclosed in quotes ''. If so, strip the quotes
- // and treat the string as unquoted (so that escapes are interpreted). Doing that
- // will allow '\'foo\'' to become 'foo', instead of \'foo\'.
- std::string::size_type start = 0, end = str.size();
- if (IsQuoted(str)) {
- start++;
- end--;
- }
-
- bool inside_quote = false;
- bool escaping = false;
-
- for (auto i = start, e = end; i < e; ++i) {
- std::string::value_type c = str[i];
- if (escaping) {
- escaping = false;
- } else if (c == escape) {
- escaping = true;
- if (!inside_quote) {
- continue;
- }
- } else if (c == quote) {
- inside_quote = !inside_quote;
- }
- result.push_back(c);
- }
-
- return result;
-}
-
-std::string TargetInternal::Uninterpret(const std::string& str) {
- // Do the opposite to `Interpret`, so that Interpret(Uninterpret(str)) == str.
- std::string result;
-
- for (std::string::size_type i = 0, e = str.size(); i < e; ++i) {
- std::string::value_type c = str[i];
- if (c == escape || c == quote) {
- result.push_back(escape);
- }
- result.push_back(c);
- }
-
- return result;
-}
-
-static int ParseKVPair(const std::string& s, const std::string& s_next, std::string* key,
- std::string* value) {
- std::string::size_type pos;
- std::string& result_k = *key;
- std::string& result_v = *value;
- if ((pos = s.find_first_of('=')) != std::string::npos) {
- // case 1. --key=value
- result_k = s.substr(0, pos);
- result_v = s.substr(pos + 1);
- if (result_k.empty() || result_v.empty()) {
- TVM_FFI_THROW(ValueError) << "Empty attribute key or value in \"" + s + "\"";
- }
- return 1;
- } else if (!s_next.empty() && s_next[0] != '-') {
- // case 2. --key value
- result_k = s;
- result_v = s_next;
- return 2;
- }
- // case 3. --boolean-key
- result_k = s;
- result_v = "1";
- return 1;
-}
-
const TargetKindNode::ValueTypeInfo& TargetInternal::FindTypeInfo(const TargetKind& kind,
const std::string& key) {
auto it = kind->key2vtype_.find(key);
@@ -354,64 +134,6 @@ const TargetKindNode::ValueTypeInfo& TargetInternal::FindTypeInfo(const TargetKi
/********** Parsing **********/
-Any TargetInternal::ParseType(const std::string& str, const TargetKindNode::ValueTypeInfo& info) {
- std::string interp_str = Interpret(str);
- if (info.type_index == ffi::TypeIndex::kTVMFFIInt ||
- info.type_index == ffi::TypeIndex::kTVMFFIBool) {
- // Parsing integer or boolean
- std::istringstream is(interp_str);
- int64_t v;
- if (!(is >> v)) {
- std::string lower(interp_str.size(), '\x0');
- std::transform(interp_str.begin(), interp_str.end(), lower.begin(),
- [](unsigned char c) { return std::tolower(c); });
- // Mimic C++ automatic conversions, allowing bool to be used for
- // integer parameters.
- if (lower == "true") {
- v = 1;
- } else if (lower == "false") {
- v = 0;
- } else {
- TVM_FFI_THROW(ValueError) << "Cannot parse integer from string: " + interp_str;
- }
- }
-
- if (info.type_index == ffi::TypeIndex::kTVMFFIInt) {
- return static_cast<int64_t>(v);
- } else {
- return static_cast<bool>(v);
- }
- } else if (info.type_index == ffi::TypeIndex::kTVMFFIStr) {
- // Parsing string, strip leading/trailing spaces, and enclosing quotes if any
- auto start = interp_str.find_first_not_of(' ');
- auto end = interp_str.find_last_not_of(' ');
- if (start == std::string::npos || end == std::string::npos) {
- // The whole string is made of spaces.
- return ffi::String();
- }
- return ffi::String(interp_str.substr(start, (end - start + 1)));
-
- } else if (info.type_index == Target::ContainerType::RuntimeTypeIndex()) {
- // Parsing target
- return Target(TargetInternal::FromString(interp_str));
- } else if (info.type_index == ffi::ArrayObj::RuntimeTypeIndex()) {
- // Parsing array
- std::vector<ObjectRef> result;
- for (const std::string& substr : SplitString(interp_str, ',')) {
- try {
- ObjectRef parsed = TargetInternal::ParseType(substr, *info.key).cast<ObjectRef>();
- result.push_back(parsed);
- } catch (const Error& e) {
- std::string index = "[" + std::to_string(result.size()) + "]";
- throw Error(e.kind(), e.message() + index, e.backtrace());
- }
- }
- return ffi::Array<ObjectRef>(result);
- }
- TVM_FFI_THROW(TypeError) << "Unsupported type \"" + info.type_key
- << "\" for parsing from string: " + interp_str;
-}
-
Any TargetInternal::ParseType(const Any& obj, const TargetKindNode::ValueTypeInfo& info) {
if (info.type_index == ffi::TypeIndex::kTVMFFIInt) {
// Parsing integer
@@ -483,88 +205,88 @@ Any TargetInternal::ParseType(const Any& obj, const TargetKindNode::ValueTypeInf
return obj;
}
-/********** Stringifying **********/
+/********** JSON Serialization **********/
-std::string TargetInternal::StringifyAtomicType(const Any& obj) {
- if (obj.type_index() == ffi::TypeIndex::kTVMFFIBool) {
- return std::to_string(obj.cast<bool>());
- } else if (obj.type_index() == ffi::TypeIndex::kTVMFFIInt) {
- return std::to_string(obj.cast<int64_t>());
- } else if (auto opt_str = obj.as<ffi::String>()) {
- std::string s = opt_str.value();
- auto u = Uninterpret(s);
- if (u.find_first_of(' ') != std::string::npos && !IsQuoted(u)) {
- u = Quote(u);
+void TargetInternal::JSONSerializeString(std::ostringstream& os, const std::string& str) {
+ os << '"';
+ for (char c : str) {
+ switch (c) {
+ case '"':
+ os << "\\\"";
+ break;
+ case '\\':
+ os << "\\\\";
+ break;
+ case '\n':
+ os << "\\n";
+ break;
+ case '\r':
+ os << "\\r";
+ break;
+ case '\t':
+ os << "\\t";
+ break;
+ default:
+ os << c;
+ break;
}
- return u;
}
- LOG(FATAL) << "Cannot stringify object of type " << obj.GetTypeKey();
- TVM_FFI_UNREACHABLE();
+ os << '"';
}
-std::string TargetInternal::StringifyArray(const ffi::ArrayObj& array) {
- std::vector<std::string> elements;
-
- for (const Any& item : array) {
- std::string s = StringifyAtomicType(item);
- std::string u = Uninterpret(s);
- if (u.find_first_of(',') != std::string::npos && !IsQuoted(u)) {
- u = Quote(u);
+void TargetInternal::JSONSerializeValue(std::ostringstream& os, const Any& obj) {
+ if (obj.type_index() == ffi::TypeIndex::kTVMFFIBool) {
+ os << (obj.cast<bool>() ? "true" : "false");
+ } else if (obj.type_index() == ffi::TypeIndex::kTVMFFIInt) {
+ os << obj.cast<int64_t>();
+ } else if (auto opt_str = obj.as<ffi::String>()) {
+ JSONSerializeString(os, std::string(opt_str.value()));
+ } else if (const auto* array = obj.as<ffi::ArrayObj>()) {
+ os << '[';
+ bool first = true;
+ for (const Any& item : *array) {
+ if (!first) os << ", ";
+ first = false;
+ JSONSerializeValue(os, item);
}
- elements.push_back(u);
+ os << ']';
+ } else if (const auto* map = obj.as<ffi::MapObj>()) {
+ // Nested map (e.g. host target's Export)
+ ffi::Map<ffi::String, ffi::Any> typed_map = ffi::GetRef<ffi::Map<ffi::String, ffi::Any>>(map);
+ os << ExportToJSON(typed_map);
+ } else {
+ // Fallback: use the object's string representation
+ os << obj;
}
-
- return JoinString(elements, ',');
}
-ffi::Optional<ffi::String> TargetInternal::StringifyAttrsToRaw(
- const ffi::Map<ffi::String, ffi::Any>& attrs) {
- std::ostringstream os;
- std::vector<ffi::String> keys;
- for (const auto& kv : attrs) {
- keys.push_back(kv.first);
+std::string TargetInternal::ExportToJSON(const ffi::Map<ffi::String, ffi::Any>& data) {
+ // Collect and sort keys for deterministic output
+ std::vector<std::string> keys;
+ for (const auto& kv : data) {
+ keys.push_back(std::string(kv.first));
}
std::sort(keys.begin(), keys.end());
- std::vector<std::string> result;
+ std::ostringstream os;
+ os << '{';
+ bool first = true;
for (const auto& key : keys) {
- const Any& obj = attrs[key];
- std::string value;
- // skip undefined attrs
- if (obj == nullptr) continue;
- if (const auto* array = obj.as<ffi::ArrayObj>()) {
- value = ffi::String(StringifyArray(*array));
- } else {
- value = StringifyAtomicType(obj);
- }
- if (!value.empty()) {
- result.push_back("-" + key + "=" + value);
- }
- }
- return ffi::String(JoinString(result, ' '));
+ const Any& val = data[ffi::String(key)];
+ if (val == nullptr) continue;
+ if (!first) os << ", ";
+ first = false;
+ JSONSerializeString(os, key);
+ os << ": ";
+ JSONSerializeValue(os, val);
+ }
+ os << '}';
+ return os.str();
}
const std::string& TargetNode::str() const {
if (str_repr_.empty()) {
- std::ostringstream os;
- os << kind->name;
- if (!this->keys.empty()) {
- os << " -keys=";
- bool is_first = true;
- for (const ffi::String& s : keys) {
- if (is_first) {
- is_first = false;
- } else {
- os << ',';
- }
- os << s;
- }
- }
- if (ffi::Optional<ffi::String> attrs_str = TargetInternal::StringifyAttrsToRaw(attrs)) {
- os << ' ' << attrs_str.value();
- }
-
- str_repr_ = os.str();
+ str_repr_ = TargetInternal::ExportToJSON(Export());
}
return str_repr_;
}
@@ -781,7 +503,15 @@ ObjectPtr<TargetNode> TargetInternal::FromString(const ffi::String& tag_or_confi
if (!tag_or_config_or_target_str.empty() && tag_or_config_or_target_str.data()[0] == '{') {
return TargetInternal::FromConfigString(tag_or_config_or_target_str);
}
- return TargetInternal::FromRawString(tag_or_config_or_target_str);
+ // Treat as bare kind name (e.g. "llvm", "cuda"). Reject strings with spaces.
+ std::string s(tag_or_config_or_target_str);
+ if (s.find(' ') != std::string::npos) {
+ TVM_FFI_THROW(ValueError)
+ << "Cannot parse target string \"" << s
+ << "\". CLI target string form (e.g. \"llvm -mcpu=xxx\") is no longer supported. "
+ << "Please use JSON dict form (e.g. {\"kind\": \"llvm\", \"mcpu\": \"xxx\"}) instead.";
+ }
+ return TargetInternal::FromConfig({{"kind", ffi::String(s)}});
}
ObjectPtr<TargetNode> TargetInternal::FromConfigString(const ffi::String& config_str) {
@@ -796,38 +526,6 @@ ObjectPtr<TargetNode> TargetInternal::FromConfigString(const ffi::String& config
return TargetInternal::FromConfig({config.value().begin(), config.value().end()});
}
-ObjectPtr<TargetNode> TargetInternal::FromRawString(const ffi::String& target_str) {
- ICHECK_GT(target_str.length(), 0) << "Cannot parse empty target string";
- // Split the string by empty spaces
- std::vector<std::string> options = SplitString(std::string(target_str), ' ');
- std::string name = options[0];
- // Create the target config
- std::unordered_map<ffi::String, ffi::Any> config = {{"kind", ffi::String(name)}};
- TargetKind kind = GetTargetKind(name);
- for (size_t iter = 1, end = options.size(); iter < end;) {
- std::string key, value;
- try {
- // Parse key-value pair
- std::string s_next = (iter + 1 < options.size()) ? options[iter + 1] : "";
- iter += ParseKVPair(RemovePrefixDashes(options[iter]), s_next, &key, &value);
- } catch (const Error& e) {
- throw Error(e.kind(), e.message() + ", during parsing target `" + target_str + "`",
- e.backtrace());
- }
- try {
- // check if `key` has been used
- if (config.count(key)) {
- TVM_FFI_THROW(ValueError) << "The key \"" + key + "\" appears more than once";
- }
- config[key] = TargetInternal::ParseType(value, TargetInternal::FindTypeInfo(kind, key));
- } catch (const Error& e) {
- throw Error(e.kind(), std::string(e.message()) + ", during parsing target[\"" + key + "\"]",
- e.backtrace());
- }
- }
- return TargetInternal::FromConfig(config);
-}
-
ObjectPtr<TargetNode> TargetInternal::FromConfig(ffi::Map<ffi::String, ffi::Any> config) {
const ffi::String kKind = "kind";
const ffi::String kTag = "tag";
diff --git a/tests/cpp/target/virtual_device_test.cc b/tests/cpp/target/virtual_device_test.cc
index 4f4b945cae..60b643396c 100644
--- a/tests/cpp/target/virtual_device_test.cc
+++ b/tests/cpp/target/virtual_device_test.cc
@@ -117,8 +117,12 @@ TEST(VirtualDeviceCache, Memoized) {
EXPECT_NE(cache.Make(kDLCPU, 3, target_b, "local"), virtual_device_a);
EXPECT_NE(cache.Make(kDLCUDA, 3, target_a, "global"), virtual_device_a);
EXPECT_EQ(cache.Make(kDLCUDA, 3, Target("cuda"), "local"), virtual_device_a);
- EXPECT_NE(cache.Make(kDLCUDA, 3, Target("cuda -max_threads_per_block=4096"), "local"),
- virtual_device_a);
+ EXPECT_NE(
+ cache.Make(kDLCUDA, 3,
+ Target(ffi::Map<ffi::String, ffi::Any>{{"kind", ffi::String("cuda")},
+ {"max_threads_per_block", int64_t(4096)}}),
+ "local"),
+ virtual_device_a);
}
} // namespace
diff --git a/tests/python/arith/test_arith_rewrite_simplify.py b/tests/python/arith/test_arith_rewrite_simplify.py
index 5eaaac68f0..b85a6b758c 100644
--- a/tests/python/arith/test_arith_rewrite_simplify.py
+++ b/tests/python/arith/test_arith_rewrite_simplify.py
@@ -889,7 +889,7 @@ class TestScalableIndex(BaseCompare):
)
def test_simplify(self, test_case):
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu -mattr=+sve"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}):
super().test_simplify(test_case)
diff --git a/tests/python/arith/test_arith_simplify.py b/tests/python/arith/test_arith_simplify.py
index 161548a7a1..11da0a3f0e 100644
--- a/tests/python/arith/test_arith_simplify.py
+++ b/tests/python/arith/test_arith_simplify.py
@@ -99,7 +99,7 @@ def test_simplify_symbolic_comparison():
def test_simplify_vscale_comparison_with_sve_target(expression):
ana = tvm.arith.Analyzer()
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu -mattr=+sve"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}):
assert ana.can_prove(expression)
@@ -108,7 +108,7 @@ def test_simplify_vscale_comparison_without_sve_target(capfd):
vs = tvm.tir.vscale()
with pytest.raises(AssertionError):
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu"}):
assert ana.can_prove(vs * 32 < vs * 64)
warning_msg = (
diff --git a/tests/python/codegen/test_target_codegen.py b/tests/python/codegen/test_target_codegen.py
index 108e212ed3..af6322c1ae 100644
--- a/tests/python/codegen/test_target_codegen.py
+++ b/tests/python/codegen/test_target_codegen.py
@@ -34,7 +34,9 @@ def test_buffer_store_predicate_not_supported(target):
tvm.compile(func)
-@tvm.testing.parametrize_targets("cuda", "opencl", "metal", "rocm", "vulkan -from_device=0")
+@tvm.testing.parametrize_targets(
+ "cuda", "opencl", "metal", "rocm", {"kind": "vulkan", "from_device": 0}
+)
def test_buffer_store_predicate_not_supported_gpu(target):
@T.prim_func
def func(a: T.handle, b: T.handle):
@@ -70,7 +72,9 @@ def test_buffer_load_predicate_not_supported(target):
tvm.compile(func)
-@tvm.testing.parametrize_targets("cuda", "opencl", "metal", "rocm", "vulkan -from_device=0")
+@tvm.testing.parametrize_targets(
+ "cuda", "opencl", "metal", "rocm", {"kind": "vulkan", "from_device": 0}
+)
def test_buffer_load_predicate_not_supported_gpu(target):
@T.prim_func
def func(a: T.handle, b: T.handle):
diff --git a/tests/python/codegen/test_target_codegen_aarch64.py b/tests/python/codegen/test_target_codegen_aarch64.py
index 5517d988a3..f5c25ecded 100644
--- a/tests/python/codegen/test_target_codegen_aarch64.py
+++ b/tests/python/codegen/test_target_codegen_aarch64.py
@@ -35,7 +35,7 @@ from tvm.target.codegen import llvm_version_major
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_mul(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -74,7 +74,7 @@ def test_mul(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_add(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -113,7 +113,7 @@ def test_add(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_sub(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -152,7 +152,7 @@ def test_sub(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_muladd(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -192,7 +192,7 @@ def test_muladd(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_max(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -235,7 +235,7 @@ def test_max(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_min(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -278,7 +278,7 @@ def test_min(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_div(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -316,7 +316,7 @@ def test_div(dtype):
"dtype", ["uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"]
)
def test_mod(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -355,7 +355,7 @@ def test_mod(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_eq(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -394,7 +394,7 @@ def test_eq(dtype):
["float", "float16", "uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"],
)
def test_neq(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -432,7 +432,7 @@ def test_neq(dtype):
"dtype", ["uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"]
)
def test_or(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -470,7 +470,7 @@ def test_or(dtype):
"dtype", ["uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"]
)
def test_and(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -508,7 +508,7 @@ def test_and(dtype):
"dtype", ["uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"]
)
def test_not(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -549,7 +549,7 @@ def test_not(dtype):
"dtype", ["uint8", "uint16", "uint32", "uint64", "int8", "int16", "int32", "int64"]
)
def test_memcpy(dtype):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
@I.ir_module
class Module:
@@ -590,7 +590,7 @@ def test_memcpy(dtype):
],
)
def test_vscale_range_function_attribute(mattr, expect_attr):
- target = f"llvm -mtriple=aarch64-linux-gnu -mattr={mattr}"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": [mattr]}
@I.ir_module
class Module:
diff --git a/tests/python/codegen/test_target_codegen_arm.py b/tests/python/codegen/test_target_codegen_arm.py
index d6791c7859..51d10b3fb4 100644
--- a/tests/python/codegen/test_target_codegen_arm.py
+++ b/tests/python/codegen/test_target_codegen_arm.py
@@ -20,7 +20,12 @@ from tvm.script import tir as T, ir as I
def test_popcount():
- target = "llvm -mtriple=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon"
+ target = {
+ "kind": "llvm",
+ "mtriple": "armv7l-none-linux-gnueabihf",
+ "mcpu": "cortex-a53",
+ "mattr": ["+neon"],
+ }
def check_correct_assembly(type, elements, counts):
@I.ir_module
@@ -51,7 +56,12 @@ def test_popcount():
def test_vmlal_s16():
- target = "llvm -mtriple=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon"
+ target = {
+ "kind": "llvm",
+ "mtriple": "armv7l-none-linux-gnueabihf",
+ "mcpu": "cortex-a53",
+ "mattr": ["+neon"],
+ }
def check_correct_assembly(N):
@I.ir_module
diff --git a/tests/python/codegen/test_target_codegen_cross_llvm.py b/tests/python/codegen/test_target_codegen_cross_llvm.py
index 224924ff72..a0d928f21b 100644
--- a/tests/python/codegen/test_target_codegen_cross_llvm.py
+++ b/tests/python/codegen/test_target_codegen_cross_llvm.py
@@ -56,7 +56,7 @@ def test_llvm_add_pipeline():
assert struct.unpack(endian + "h", arr[0x12:0x14])[0] == e_machine
def build_arm():
- target = "llvm -mtriple=armv7-none-linux-gnueabihf"
+ target = {"kind": "llvm", "mtriple": "armv7-none-linux-gnueabihf"}
if not tvm.runtime.enabled(target):
print("Skip because %s is not enabled.." % target)
return
diff --git a/tests/python/codegen/test_target_codegen_device.py b/tests/python/codegen/test_target_codegen_device.py
index 04e1857c9e..cbdd3f7fde 100644
--- a/tests/python/codegen/test_target_codegen_device.py
+++ b/tests/python/codegen/test_target_codegen_device.py
@@ -49,7 +49,7 @@ def test_large_uint_imm():
assert a.numpy()[0] == value + 3
check_target("cuda")
- check_target("vulkan -from_device=0")
+ check_target({"kind": "vulkan", "from_device": 0})
@tvm.testing.requires_gpu
diff --git a/tests/python/codegen/test_target_codegen_gpu_common.py b/tests/python/codegen/test_target_codegen_gpu_common.py
index bf2fc12083..84fcb6d427 100644
--- a/tests/python/codegen/test_target_codegen_gpu_common.py
+++ b/tests/python/codegen/test_target_codegen_gpu_common.py
@@ -25,7 +25,9 @@ from tvm.script import tir as T, ir as I
@tvm.testing.requires_gpu
-@tvm.testing.parametrize_targets("cuda", "metal", "vulkan -supports_int64=1", "opencl")
+@tvm.testing.parametrize_targets(
+ "cuda", "metal", {"kind": "vulkan", "supports_int64": True}, "opencl"
+)
@pytest.mark.parametrize("dtype", ["int32", "uint32", "int64", "uint64"])
def test_int_intrin(target, dev, dtype):
test_funcs = [
diff --git a/tests/python/codegen/test_target_codegen_llvm.py b/tests/python/codegen/test_target_codegen_llvm.py
index 78f2abf523..ebbf4aeb49 100644
--- a/tests/python/codegen/test_target_codegen_llvm.py
+++ b/tests/python/codegen/test_target_codegen_llvm.py
@@ -781,7 +781,7 @@ def test_dwarf_debug_information():
"fadd2": Module["main"].with_attr("global_symbol", "fadd2"),
}
)
- m = tvm.tir.build(mod, target="llvm -mtriple=aarch64-linux-gnu")
+ m = tvm.tir.build(mod, target={"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
ll = m.inspect_source("ll")
# On non-Darwin OS, don't explicitly specify DWARF version.
@@ -791,7 +791,7 @@ def test_dwarf_debug_information():
assert re.search(r"""llvm.dbg.value""", ll)
# Try Darwin, require DWARF-2
- m = tvm.tir.build(mod, target="llvm -mtriple=x86_64-apple-darwin-macho")
+ m = tvm.tir.build(mod, target={"kind": "llvm", "mtriple": "x86_64-apple-darwin-macho"})
ll = m.inspect_source("ll")
assert re.search(r"""i32 4, !"Dwarf Version", i32 2""", ll)
assert re.search(r"""llvm.dbg.value""", ll)
@@ -994,7 +994,12 @@ def test_llvm_target_attributes():
T.writes(C[v_i])
C[v_i] = B[v_i] + T.float32(1.0)
- target_llvm = "llvm -mtriple=x86_64-linux-gnu -mcpu=skylake -mattr=+avx512f"
+ target_llvm = {
+ "kind": "llvm",
+ "mtriple": "x86_64-linux-gnu",
+ "mcpu": "skylake",
+ "mattr": ["+avx512f"],
+ }
target = tvm.target.Target(target_llvm, host=target_llvm)
module = tvm.tir.build(Module, target=target)
diff --git a/tests/python/codegen/test_target_codegen_llvm_vla.py b/tests/python/codegen/test_target_codegen_llvm_vla.py
index 8930159481..4e03596b1f 100644
--- a/tests/python/codegen/test_target_codegen_llvm_vla.py
+++ b/tests/python/codegen/test_target_codegen_llvm_vla.py
@@ -32,8 +32,14 @@ from tvm.target.codegen import llvm_version_major
llvm_version_major() < 11, reason="Vscale is not supported in earlier versions of LLVM"
)
@tvm.testing.parametrize_targets(
- "llvm -mtriple=aarch64-linux-gnu -mattr=+sve",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]},
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_codegen_vscale(target):
vscale = tvm.tir.vscale()
@@ -54,8 +60,14 @@ def test_codegen_vscale(target):
llvm_version_major() < 11, reason="Vscale is not supported in earlier versions of LLVM"
)
@tvm.testing.parametrize_targets(
- "llvm -mtriple=aarch64-linux-gnu -mattr=+sve",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]},
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_scalable_buffer_load_store(target):
@T.prim_func
@@ -77,8 +89,14 @@ def test_scalable_buffer_load_store(target):
llvm_version_major() < 11, reason="Vscale is not supported in earlier versions of LLVM"
)
@tvm.testing.parametrize_targets(
- "llvm -mtriple=aarch64-linux-gnu -mattr=+sve",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]},
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_scalable_broadcast(target):
@T.prim_func
@@ -101,8 +119,14 @@ def test_scalable_broadcast(target):
reason="Vscale and get.active.lane.mask are not supported in earlier versions of LLVM",
)
@tvm.testing.parametrize_targets(
- "llvm -mtriple=aarch64-linux-gnu -mattr=+sve",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]},
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_get_active_lane_mask(target):
@T.prim_func
@@ -122,8 +146,14 @@ def test_get_active_lane_mask(target):
reason="Vscale and get.active.lane.mask are not supported in earlier versions of LLVM",
)
@tvm.testing.parametrize_targets(
- "llvm -mtriple=aarch64-linux-gnu -mattr=+sve",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]},
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_predicated_scalable_buffer(target):
@T.prim_func
diff --git a/tests/python/codegen/test_target_codegen_opencl.py b/tests/python/codegen/test_target_codegen_opencl.py
index 6493dbc27f..2e0fd0fd31 100644
--- a/tests/python/codegen/test_target_codegen_opencl.py
+++ b/tests/python/codegen/test_target_codegen_opencl.py
@@ -218,10 +218,12 @@ def test_opencl_type_casting():
@tvm.testing.requires_gpu
@tvm.testing.requires_opencl
-@tvm.testing.parametrize_targets("opencl", "opencl -device=adreno")
+@tvm.testing.parametrize_targets("opencl", {"kind": "opencl", "device": "adreno"})
def test_opencl_ceil_log2(target):
def _check(target, n, dtype):
- inter_dtype = "float32" if "adreno" in target else "float64"
+ target_obj = tvm.target.Target(target)
+ is_adreno = "adreno" in target_obj.attrs.get("device", "")
+ inter_dtype = "float32" if is_adreno else "float64"
@I.ir_module
class Module:
@@ -237,7 +239,7 @@ def test_opencl_ceil_log2(target):
fun = tvm.tir.build(Module, target=target)
assembly = fun.imports[0].inspect_source()
- if "adreno" in target:
+ if is_adreno:
pattern = "convert_float"
else:
pattern = "convert_double"
diff --git a/tests/python/codegen/test_target_codegen_riscv.py b/tests/python/codegen/test_target_codegen_riscv.py
index 008c7c2108..e19174db26 100644
--- a/tests/python/codegen/test_target_codegen_riscv.py
+++ b/tests/python/codegen/test_target_codegen_riscv.py
@@ -23,10 +23,34 @@ from tvm.target.codegen import target_has_features
@tvm.testing.requires_llvm_minimum_version(14)
@tvm.testing.parametrize_targets(
- "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+i,+m",
- "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+i,+m,+v",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv32-linux-gnu",
+ "mcpu": "generic-rv32",
+ "mattr": ["+i", "+m"],
+ },
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv32-linux-gnu",
+ "mcpu": "generic-rv32",
+ "mattr": ["+i", "+m", "+v"],
+ },
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m"],
+ },
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_rvv(target):
def check_rvv_presence(N, extent):
@@ -49,8 +73,20 @@ def test_rvv(target):
@tvm.testing.requires_llvm_minimum_version(14)
@tvm.testing.parametrize_targets(
- "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+i,+m,+v",
- "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v",
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv32-linux-gnu",
+ "mcpu": "generic-rv32",
+ "mattr": ["+i", "+m", "+v"],
+ },
+ {
+ "kind": "llvm",
+ "device": "riscv_cpu",
+ "mtriple": "riscv64-linux-gnu",
+ "mcpu": "generic-rv64",
+ "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"],
+ },
)
def test_rvv_vscale_llvm_dbginfo(target):
# fmt: off
diff --git a/tests/python/codegen/test_target_codegen_vulkan.py b/tests/python/codegen/test_target_codegen_vulkan.py
index d010da46da..3ba261608b 100644
--- a/tests/python/codegen/test_target_codegen_vulkan.py
+++ b/tests/python/codegen/test_target_codegen_vulkan.py
@@ -150,7 +150,7 @@ vulkan_parameter_dtype = tvm.testing.parameter("int32", "float32", "int64")
# Only run on vulkan because extremely large numbers of input
# parameters can crash cuda/llvm compiler.
-@tvm.testing.parametrize_targets("vulkan -from_device=0")
+@tvm.testing.parametrize_targets({"kind": "vulkan", "from_device": 0})
def test_vulkan_constant_passing(target, dev, vulkan_parameter_impl, vulkan_parameter_dtype):
target = tvm.target.Target(target)
dtype = vulkan_parameter_dtype
@@ -285,7 +285,7 @@ def test_vulkan_local_threadidx(target, dev):
tvm.testing.assert_allclose(b.numpy(), a_np)
-@tvm.testing.parametrize_targets("vulkan -from_device=0")
+@tvm.testing.parametrize_targets({"kind": "vulkan", "from_device": 0})
def test_vectorized_index_ramp(target, dev):
"""Test vectorized copy with ramp indices (load N values, write to N locations)"""
n = 4
@@ -315,7 +315,7 @@ def test_vectorized_index_ramp(target, dev):
tvm.testing.assert_allclose(b.numpy(), a_np)
-@tvm.testing.parametrize_targets("vulkan -from_device=0")
+@tvm.testing.parametrize_targets({"kind": "vulkan", "from_device": 0})
def test_vectorized_index_broadcast(target, dev):
"""Test broadcast index (load 1 value, write to N locations)"""
n = 4
@@ -478,7 +478,7 @@ def test_cooperative_matrix(out_dtype):
T.tvm_store_matrix_sync(A.data, 16, 16, 16, A.elem_offset // A.strides[0] // 16 * (A.strides[0] // 16) + A.elem_offset % A.strides[0] // 16, T.tvm_access_ptr(T.type_annotation(out_dtype), C.data, C.elem_offset, C.strides[0] * 16, 2), C.strides[0], "row_major")
# fmt: on
- target = "vulkan -from_device=0"
+ target = {"kind": "vulkan", "from_device": 0}
tgt_attrs = tvm.target.Target(target).attrs
if tgt_attrs.get("supports_cooperative_matrix"):
diff --git a/tests/python/codegen/test_target_codegen_x86.py b/tests/python/codegen/test_target_codegen_x86.py
index 3d68155c9f..545411b13d 100644
--- a/tests/python/codegen/test_target_codegen_x86.py
+++ b/tests/python/codegen/test_target_codegen_x86.py
@@ -61,14 +61,20 @@ def test_fp16_to_fp32():
not_matches = [l for l in assembly if re.search(not_match, l)]
assert not not_matches
- fp16_to_fp32("llvm -mcpu=skylake-avx512", 15, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=skylake-avx512", 16, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=skylake-avx512", 17, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=skylake-avx512", 49, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=skylake-avx512 -mattr=-avx512f", 49, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=skylake-avx512 -mattr=-f16c,-avx512f", 49, not_match="vcvtph2ps")
- fp16_to_fp32("llvm -mcpu=core-avx2", 8, match="vcvtph2ps.*mm")
- fp16_to_fp32("llvm -mcpu=core-avx2", 9, match="vcvtph2ps.*mm")
+ fp16_to_fp32({"kind": "llvm", "mcpu": "skylake-avx512"}, 15, match="vcvtph2ps.*mm")
+ fp16_to_fp32({"kind": "llvm", "mcpu": "skylake-avx512"}, 16, match="vcvtph2ps.*mm")
+ fp16_to_fp32({"kind": "llvm", "mcpu": "skylake-avx512"}, 17, match="vcvtph2ps.*mm")
+ fp16_to_fp32({"kind": "llvm", "mcpu": "skylake-avx512"}, 49, match="vcvtph2ps.*mm")
+ fp16_to_fp32(
+ {"kind": "llvm", "mcpu": "skylake-avx512", "mattr": ["-avx512f"]}, 49, match="vcvtph2ps.*mm"
+ )
+ fp16_to_fp32(
+ {"kind": "llvm", "mcpu": "skylake-avx512", "mattr": ["-f16c", "-avx512f"]},
+ 49,
+ not_match="vcvtph2ps",
+ )
+ fp16_to_fp32({"kind": "llvm", "mcpu": "core-avx2"}, 8, match="vcvtph2ps.*mm")
+ fp16_to_fp32({"kind": "llvm", "mcpu": "core-avx2"}, 9, match="vcvtph2ps.*mm")
fp16_to_fp32("llvm", 9, not_match="vcvtph2ps")
diff --git a/tests/python/contrib/test_android/infrastructure.py b/tests/python/contrib/test_android/infrastructure.py
index fca7989cb1..9baf220ba4 100644
--- a/tests/python/contrib/test_android/infrastructure.py
+++ b/tests/python/contrib/test_android/infrastructure.py
@@ -53,5 +53,5 @@ def get_rpc_runner() -> tvm.s_tir.meta_schedule.runner.RPCRunner:
def get_android_gpu_target() -> tvm.target.Target:
"""Creates a Android GPU target"""
target_c = "opencl"
- target_h = "llvm -mtriple=arm64-linux-android"
+ target_h = {"kind": "llvm", "mtriple": "arm64-linux-android"}
return tvm.target.Target(target_c, host=target_h)
diff --git a/tests/python/nightly/test_nnapi/infrastructure.py b/tests/python/nightly/test_nnapi/infrastructure.py
index 07aa0bd752..2f4089fea9 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.compile(mod, target="llvm -mtriple=aarch64-linux-android")
+ ex = tvm.compile(mod, target={"kind": "llvm", "mtriple": "aarch64-linux-android"})
return ex
diff --git a/tests/python/relax/adreno/test_transform_annotate_custom_scope.py b/tests/python/relax/adreno/test_transform_annotate_custom_scope.py
index 6c6dbf0231..bfcead3011 100644
--- a/tests/python/relax/adreno/test_transform_annotate_custom_scope.py
+++ b/tests/python/relax/adreno/test_transform_annotate_custom_scope.py
@@ -72,7 +72,7 @@ class ValidateScope(PyExprVisitor): # pylint: disable=abstract-method
def verify(mod, expected):
- tgt = tvm.target.Target("opencl --device=adreno", host="llvm")
+ tgt = tvm.target.Target({"kind": "opencl", "device": "adreno"}, host="llvm")
skip_ops = [
"relax.nn.conv2d",
"relax.nn.max_pool2d",
diff --git a/tests/python/relax/backend/clml/conftest.py b/tests/python/relax/backend/clml/conftest.py
index 00bad5da21..dd9e444457 100644
--- a/tests/python/relax/backend/clml/conftest.py
+++ b/tests/python/relax/backend/clml/conftest.py
@@ -30,7 +30,7 @@ def rpc():
host = os.getenv("TVM_TRACKER_HOST", "localhost")
port = int(os.getenv("TVM_TRACKER_PORT", 9090))
target = "opencl"
- target_host = "llvm -mtriple=aarch64-linux-gnu"
+ target_host = {"kind": "llvm", "mtriple": "aarch64-linux-gnu"}
device_key = os.getenv("RPC_DEVICE_KEY", "android")
cross_compile = os.getenv("TVM_NDK_CC", "aarch64-linux-android-g++")
tracker = _rpc.connect_tracker(host, port)
diff --git a/tests/python/relax/backend/clml/utils.py b/tests/python/relax/backend/clml/utils.py
index a5b7ebe9cf..2c167a2473 100644
--- a/tests/python/relax/backend/clml/utils.py
+++ b/tests/python/relax/backend/clml/utils.py
@@ -39,7 +39,7 @@ def get_rpc():
host = os.getenv("TVM_TRACKER_HOST", "localhost")
port = int(os.getenv("TVM_TRACKER_PORT", 9090))
target = "opencl"
- target_host = "llvm -mtriple=aarch64-linux-gnu"
+ target_host = {"kind": "llvm", "mtriple": "aarch64-linux-gnu"}
device_key = os.getenv("RPC_DEVICE_KEY", "android")
cross_compile = os.getenv("TVM_NDK_CC", "aarch64-linux-android-g++")
tracker = _rpc.connect_tracker(host, port)
@@ -55,7 +55,7 @@ def build_and_run(
rpc=None,
load_path="vm_library.so",
):
- tgt = tvm.target.Target(target, host="llvm -mtriple=aarch64-linux-gnu")
+ tgt = tvm.target.Target(target, host={"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
relax_pipeline = relax.pipeline.get_default_pipeline(tgt)
tir_pipeline = tvm.tir.get_default_tir_pipeline(tgt)
diff --git a/tests/python/relax/test_backend_dispatch_sort_scan.py b/tests/python/relax/test_backend_dispatch_sort_scan.py
index 564ea05274..0bf1eacb3d 100644
--- a/tests/python/relax/test_backend_dispatch_sort_scan.py
+++ b/tests/python/relax/test_backend_dispatch_sort_scan.py
@@ -169,7 +169,7 @@ def test_dispatch_sort_cuda():
R.output(gv)
return gv
- target = tvm.target.Target("cuda -libs=thrust", host="llvm")
+ target = tvm.target.Target({"kind": "cuda", "libs": ["thrust"]}, host="llvm")
vdevices = [I.vdevice("cuda", 0)]
x = relax.Var("x", R.Tensor((2, 3), "float32", vdevices[0]))
@@ -265,7 +265,7 @@ def test_dispatch_argsort_cuda():
R.output(gv)
return gv
- target = tvm.target.Target("cuda -libs=thrust", host="llvm")
+ target = tvm.target.Target({"kind": "cuda", "libs": ["thrust"]}, host="llvm")
vdevices = [I.vdevice("cuda", 0)]
x = relax.Var("x", R.Tensor((2, 3), "float32", vdevices[0]))
@@ -350,7 +350,7 @@ def test_dispatch_topk_cuda():
R.output(gv)
return gv
- target = tvm.target.Target("cuda -libs=thrust", host="llvm")
+ target = tvm.target.Target({"kind": "cuda", "libs": ["thrust"]}, host="llvm")
vdevices = [I.vdevice("cuda", 0)]
x = relax.Var("x", R.Tensor((2, 3), "float32", vdevices[0]))
@@ -409,7 +409,7 @@ def test_dispatch_topk_gpu():
assert_structural_equal(mod, expected_mod)
-@tvm.testing.parametrize_targets("cuda", "vulkan -supports_int64=1")
+@tvm.testing.parametrize_targets("cuda", {"kind": "vulkan", "supports_int64": True})
def test_dispatch_cumsum_gpu(target, dev):
"""Test cumsum kernel dispatch and numerical correctness"""
diff --git a/tests/python/relax/test_frontend_nn_op.py b/tests/python/relax/test_frontend_nn_op.py
index 958192a479..5f561e06ef 100644
--- a/tests/python/relax/test_frontend_nn_op.py
+++ b/tests/python/relax/test_frontend_nn_op.py
@@ -1107,7 +1107,7 @@ def test_sample_top_p_top_k_from_sorted_prob():
tvm.ir.assert_structural_equal(mod, Expected)
- target = tvm.target.Target("cuda -libs=thrust", host="llvm")
+ target = tvm.target.Target({"kind": "cuda", "libs": ["thrust"]}, host="llvm")
with target:
mod = s_tir.transform.DefaultGPUSchedule()(mod)
@@ -1224,7 +1224,7 @@ def test_renormalize_top_p_top_k_prob():
tvm.ir.assert_structural_equal(mod, Expected)
- target = tvm.target.Target("cuda -libs=thrust", host="llvm")
+ target = tvm.target.Target({"kind": "cuda", "libs": ["thrust"]}, host="llvm")
with target:
mod = relax.transform.LegalizeOps()(mod)
mod = s_tir.transform.DefaultGPUSchedule()(mod)
diff --git a/tests/python/relax/test_meta_schedule_relax_integration.py b/tests/python/relax/test_meta_schedule_relax_integration.py
index 3d5a29dce1..76d4130dab 100644
--- a/tests/python/relax/test_meta_schedule_relax_integration.py
+++ b/tests/python/relax/test_meta_schedule_relax_integration.py
@@ -48,7 +48,7 @@ class Module0:
def test_extracting_tasks():
- target = "llvm -mcpu=core-avx2 -num-cores=1"
+ target = {"kind": "llvm", "mcpu": "core-avx2", "num-cores": 1}
relax_mod = Module0
relax_mod = relax.transform.LegalizeOps()(relax_mod)
diff --git a/tests/python/relax/test_transform_few_shot_tuning.py b/tests/python/relax/test_transform_few_shot_tuning.py
index 61d4a2ab23..0d15ea2b63 100644
--- a/tests/python/relax/test_transform_few_shot_tuning.py
+++ b/tests/python/relax/test_transform_few_shot_tuning.py
@@ -313,7 +313,7 @@ class Module:
def _target() -> tvm.target.Target:
- return tvm.target.Target("llvm -num-cores=4")
+ return tvm.target.Target({"kind": "llvm", "num-cores": 4})
# for local testing only
# return tvm.target.Target("nvidia/geforce-rtx-3070")
diff --git a/tests/python/relax/test_transform_meta_schedule_apply_database.py b/tests/python/relax/test_transform_meta_schedule_apply_database.py
index ab481e75a5..a7b30d05ee 100644
--- a/tests/python/relax/test_transform_meta_schedule_apply_database.py
+++ b/tests/python/relax/test_transform_meta_schedule_apply_database.py
@@ -22,7 +22,7 @@ from tvm.s_tir import meta_schedule as ms
from tvm import relax
from tvm.script import ir as I, tir as T
-target = tvm.target.Target("llvm --num-cores=16")
+target = tvm.target.Target({"kind": "llvm", "num-cores": 16})
def test_apply_to_func_with_different_block_name():
diff --git a/tests/python/relax/test_transform_meta_schedule_tuning.py b/tests/python/relax/test_transform_meta_schedule_tuning.py
index 4a0d2995a1..27595bd333 100644
--- a/tests/python/relax/test_transform_meta_schedule_tuning.py
+++ b/tests/python/relax/test_transform_meta_schedule_tuning.py
@@ -28,7 +28,7 @@ from tvm.ir.transform import PassContext
from tvm.script import relax as R
from tvm.script import tir as T
-target = tvm.target.Target("llvm --num-cores=16")
+target = tvm.target.Target({"kind": "llvm", "num-cores": 16})
@tvm.script.ir_module
diff --git a/tests/python/relax/test_transform_realize_vdevice.py b/tests/python/relax/test_transform_realize_vdevice.py
index fa64282184..8c6476ba7f 100644
--- a/tests/python/relax/test_transform_realize_vdevice.py
+++ b/tests/python/relax/test_transform_realize_vdevice.py
@@ -31,7 +31,7 @@ vdevices = [
VDevice("llvm"),
VDevice("cuda", 0),
VDevice("metal", 0, "global"),
- VDevice("cuda -arch=sm_80", 0),
+ VDevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
@@ -45,7 +45,7 @@ def test_dataflow_binding():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
@@ -76,7 +76,7 @@ def test_dataflow_binding():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
@@ -267,7 +267,7 @@ def test_multi_device():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
@@ -296,7 +296,7 @@ def test_multi_device():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
@@ -329,7 +329,7 @@ def test_insert_to_vdevice():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
@@ -359,7 +359,7 @@ def test_insert_to_vdevice():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
diff --git a/tests/python/relax/test_transform_update_vdevice.py b/tests/python/relax/test_transform_update_vdevice.py
index 76caccb536..ee9acd7f1b 100644
--- a/tests/python/relax/test_transform_update_vdevice.py
+++ b/tests/python/relax/test_transform_update_vdevice.py
@@ -31,7 +31,7 @@ def test_update():
VDevice("llvm"),
VDevice("cuda", 0),
VDevice("metal", 0, "global"),
- VDevice("cuda -arch=sm_80", 0),
+ VDevice({"kind": "cuda", "arch": "sm_80"}, 0),
VDevice("metal", 1, "global"),
VDevice("llvm", 1),
]
@@ -45,7 +45,7 @@ def test_update():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
diff --git a/tests/python/relax/test_tvmscript_parser.py b/tests/python/relax/test_tvmscript_parser.py
index e01f214d8b..e6bcd88bf4 100644
--- a/tests/python/relax/test_tvmscript_parser.py
+++ b/tests/python/relax/test_tvmscript_parser.py
@@ -314,7 +314,7 @@ def test_global_info_vdevice():
vdevices = [
VDevice("llvm"),
VDevice("cuda", 0),
- VDevice("cuda -arch=sm_80", 0),
+ VDevice({"kind": "cuda", "arch": "sm_80"}, 0),
VDevice("metal", 0, "global"),
]
@@ -326,7 +326,7 @@ def test_global_info_vdevice():
"vdevice": [
I.vdevice("llvm"),
I.vdevice("cuda", 0),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
I.vdevice("metal", 0, "global"),
]
}
@@ -774,7 +774,7 @@ def test_tensor_with_vdevice():
VDevice("llvm"),
VDevice("cuda", 0),
VDevice("metal", 0, "global"),
- VDevice("cuda -arch=sm_80", 0),
+ VDevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
@I.ir_module
@@ -786,7 +786,7 @@ def test_tensor_with_vdevice():
I.vdevice("llvm"),
I.vdevice("cuda", 0),
I.vdevice("metal", 0, "global"),
- I.vdevice("cuda -arch=sm_80", 0),
+ I.vdevice({"kind": "cuda", "arch": "sm_80"}, 0),
]
}
)
diff --git a/tests/python/relax/texture/adreno_utils.py b/tests/python/relax/texture/adreno_utils.py
index b6ffd83714..61f8c41ff1 100644
--- a/tests/python/relax/texture/adreno_utils.py
+++ b/tests/python/relax/texture/adreno_utils.py
@@ -114,7 +114,7 @@ def build_run(mod, inputs, backend, is_adreno=False):
if remote is None:
tgt = tvm.target.Target(target, host="llvm")
else:
- tgt = tvm.target.Target(target, host="llvm -mtriple=aarch64-linux-gnu")
+ tgt = tvm.target.Target(target, host={"kind": "llvm", "mtriple": "aarch64-linux-gnu"})
relax_pipeline = relax.pipeline.get_default_pipeline(tgt)
tir_pipeline = tvm.tir.get_default_tir_pipeline(tgt)
mod = relax_pipeline(mod)
diff --git a/tests/python/relax/texture/test_texture_nd.py b/tests/python/relax/texture/test_texture_nd.py
index 058885e1d6..fb7a6a4ec5 100644
--- a/tests/python/relax/texture/test_texture_nd.py
+++ b/tests/python/relax/texture/test_texture_nd.py
@@ -146,9 +146,12 @@ def test_texture_copy(backend, dtype, channel_size, read_width):
mod = TextureCopy
if remote is None:
- target = Target(backend + " -device=adreno")
+ target = Target({"kind": backend, "device": "adreno"})
else:
- target = Target(backend + " -device=adreno", "llvm -mtriple=aarch64-linux-android")
+ target = Target(
+ {"kind": backend, "device": "adreno"},
+ {"kind": "llvm", "mtriple": "aarch64-linux-android"},
+ )
with target:
mod = preprocess_pipeline(mod)
diff --git a/tests/python/runtime/test_runtime_module_load.py b/tests/python/runtime/test_runtime_module_load.py
index fcf4180fb5..62de1239f0 100644
--- a/tests/python/runtime/test_runtime_module_load.py
+++ b/tests/python/runtime/test_runtime_module_load.py
@@ -42,7 +42,7 @@ print("Finish runtime checking...")
@tvm.testing.requires_llvm
-@pytest.mark.parametrize("target", ["llvm", "llvm -jit=mcjit"])
+@pytest.mark.parametrize("target", ["llvm", {"kind": "llvm", "jit": "mcjit"}])
def test_dso_module_load(target):
dtype = "int64"
temp = utils.tempdir()
diff --git a/tests/python/runtime/test_runtime_rpc.py b/tests/python/runtime/test_runtime_rpc.py
index 0d14fa14b6..94d4f0c616 100644
--- a/tests/python/runtime/test_runtime_rpc.py
+++ b/tests/python/runtime/test_runtime_rpc.py
@@ -88,7 +88,7 @@ def test_bigendian_rpc():
print("Test RPC connection to PowerPC...")
remote = rpc.connect(host, port)
- target = "llvm -mtriple=powerpc-linux-gnu"
+ target = {"kind": "llvm", "mtriple": "powerpc-linux-gnu"}
for dtype in ["float32", "float64", "int32", "int8"]:
verify_rpc(remote, target, (10,), dtype)
@@ -320,7 +320,7 @@ def test_rpc_remote_module():
xo, xi = s.split(x, factors=[None, 32])
s.bind(xo, "blockIdx.x")
s.bind(xi, "threadIdx.x")
- f = tvm.compile(s.mod, "opencl --host=llvm")
+ f = tvm.compile(s.mod, tvm.target.Target("opencl", host="llvm"))
path_tar = temp.relpath("myadd.tar")
f.export_library(path_tar)
remote.upload(path_tar)
diff --git a/tests/python/s_tir/dlight/test_benchmark.py b/tests/python/s_tir/dlight/test_benchmark.py
index 410d06ed75..c02ec30cf8 100644
--- a/tests/python/s_tir/dlight/test_benchmark.py
+++ b/tests/python/s_tir/dlight/test_benchmark.py
@@ -251,7 +251,7 @@ def test_benchmark_prim_func_full_rpc():
def test_benchmark_relax_func():
- with tvm.target.Target("llvm -num-cores=4"):
+ with tvm.target.Target({"kind": "llvm", "num-cores": 4}):
benchmark_relax_func(Module, "test")
@@ -266,7 +266,7 @@ def test_extract_prim_func_full1():
dym_var_dict={"n": "int32"},
weight=2,
sample_number=10,
- target="llvm -num-cores=4",
+ target={"kind": "llvm", "num-cores": 4},
)
)
@@ -280,13 +280,13 @@ def test_extract_prim_func_matmul1():
func=Module["matmul1"], # type: ignore
weight=2,
sample_number=10,
- target="llvm -num-cores=4",
+ target={"kind": "llvm", "num-cores": 4},
)
)
def test_extract_from_relax():
- with tvm.target.Target("llvm -num-cores=4"):
+ with tvm.target.Target({"kind": "llvm", "num-cores": 4}):
with tempfile.TemporaryDirectory() as filepath:
extract_from_relax(
Module,
diff --git a/tests/python/s_tir/dlight/test_gpu_gemv.py b/tests/python/s_tir/dlight/test_gpu_gemv.py
index a23521c7e6..a8fe8fbd10 100644
--- a/tests/python/s_tir/dlight/test_gpu_gemv.py
+++ b/tests/python/s_tir/dlight/test_gpu_gemv.py
@@ -771,7 +771,7 @@ def test_outer_reduction_adreno():
p_output0_intermediate[0, 0, v0] = lv570[0, 0, v0] + var_matmul_intermediate_local[0, 0, v0]
# fmt: on
mod = tvm.IRModule({"main": before})
- with Target("opencl", host="llvm -mtriple=aarch64-linux-android"):
+ with Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-android"}):
mod = dl.ApplyDefaultSchedule(dl.gpu.GEMV())(mod)
tvm.ir.assert_structural_equal(mod["main"], expected)
@@ -913,7 +913,7 @@ def test_outer_reduction_adreno_dynamic():
# fmt: on
mod = tvm.IRModule({"main": before})
- with Target("opencl", host="llvm -mtriple=aarch64-linux-android"):
+ with Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-android"}):
mod = dl.ApplyDefaultSchedule(dl.gpu.GEMV())(mod)
tvm.ir.assert_structural_equal(mod["main"], expected)
diff --git a/tests/python/s_tir/dlight/test_gpu_matmul.py b/tests/python/s_tir/dlight/test_gpu_matmul.py
index 30f999ac72..7cb01b44e4 100644
--- a/tests/python/s_tir/dlight/test_gpu_matmul.py
+++ b/tests/python/s_tir/dlight/test_gpu_matmul.py
@@ -702,7 +702,7 @@ def test_matmul_android():
matmul[v0, v1, v2] = matmul_pad_local[v0, v1, v2]
mod = tvm.IRModule({"main": before})
- with Target("opencl", host="llvm -mtriple=aarch64-linux-android"):
+ with Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-android"}):
mod = dl.ApplyDefaultSchedule(dl.gpu.Matmul())(mod)
tvm.ir.assert_structural_equal(mod["main"], expected)
@@ -848,7 +848,7 @@ def test_fused_dequant_matmul_android():
# fmt: on
mod = tvm.IRModule({"main": before})
- with Target("opencl", host="llvm -mtriple=aarch64-linux-android"):
+ with Target("opencl", host={"kind": "llvm", "mtriple": "aarch64-linux-android"}):
mod = dl.ApplyDefaultSchedule(dl.gpu.Matmul())(mod)
tvm.ir.assert_structural_equal(mod["main"], expected)
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_cost_model.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_cost_model.py
index 2030cd119a..2355e0362b 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_cost_model.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_cost_model.py
@@ -182,7 +182,7 @@ def test_meta_schedule_xgb_model_no_feature():
model = XGBModel(num_warmup_samples=0)
tune_ctx = TuneContext(
FullModule,
- target="llvm --num-cores 16",
+ target={"kind": "llvm", "num-cores": 16},
space_generator="post-order-apply",
search_strategy="evolutionary",
)
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_measure_callback.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_measure_callback.py
index e80ad5c6d8..e96fc85580 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_measure_callback.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_measure_callback.py
@@ -140,7 +140,7 @@ def test_meta_schedule_measure_callback_update_cost_model_with_zero():
with tempfile.TemporaryDirectory() as work_dir:
ms.tune_tir(
mod=Matmul,
- target="llvm -num-cores=1",
+ target={"kind": "llvm", "num-cores": 1},
work_dir=work_dir,
max_trials_global=10,
runner=AllZeroRunner(),
@@ -165,7 +165,7 @@ def test_meta_schedule_measure_callback_update_cost_model_with_runtime_error():
with tempfile.TemporaryDirectory() as work_dir:
ms.tune_tir(
mod=Matmul,
- target="llvm -num-cores=1",
+ target={"kind": "llvm", "num-cores": 1},
work_dir=work_dir,
max_trials_global=10,
runner=EmptyRunner(),
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_parallel.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_parallel.py
index 2583f67dbc..d692ac8104 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_parallel.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_parallel.py
@@ -93,7 +93,7 @@ def _make_mutator(target: Target, max_jobs_per_core: int) -> ms.Mutator:
def test_mutate_parallel_matmul():
mutator = _make_mutator(
- target=Target("llvm --num-cores=16"),
+ target=Target({"kind": "llvm", "num-cores": 16}),
max_jobs_per_core=256,
)
sch = _sch(
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_tile_size.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_tile_size.py
index dfe6e858a1..c6af099ff3 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_tile_size.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_tile_size.py
@@ -81,7 +81,7 @@ def _make_mutator(target: Target) -> ms.Mutator:
def test_mutate_tile_size_matmul():
mutator = _make_mutator(
- target=Target("llvm --num-cores=16"),
+ target=Target({"kind": "llvm", "num-cores": 16}),
)
results = {}
sch = _sch(decisions=[[4, 32, 4, 1]])
@@ -97,7 +97,7 @@ def test_mutate_tile_size_matmul():
def test_mutate_sample_categorical_single_candidate():
mutator = _make_mutator(
- target=Target("llvm --num-cores=16"),
+ target=Target({"kind": "llvm", "num-cores": 16}),
)
sch = Schedule(matmul, debug_mask="all")
sch.sample_categorical(candidates=[1], probs=[1.0], decision=0)
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_unroll.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_unroll.py
index 8241a30fec..3cc89aa38f 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_unroll.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_mutator_mutate_unroll.py
@@ -97,7 +97,7 @@ def _make_mutator(target: Target) -> ms.Mutator:
def test_mutate_unroll_matmul():
- mutator = _make_mutator(target=Target("llvm --num-cores=16"))
+ mutator = _make_mutator(target=Target({"kind": "llvm", "num-cores": 16}))
sch = _sch(
decisions=[
[4, 32, 4, 1],
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_post_order_apply.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_post_order_apply.py
index 1c9fcf8734..84e5915585 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_post_order_apply.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_post_order_apply.py
@@ -432,27 +432,53 @@ def test_target_sblocks_search_space():
"target,mod,expected_intr",
[
(
- Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+neon -num-cores 2"),
+ Target(
+ {
+ "kind": "llvm",
+ "device": "arm_cpu",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon"],
+ "num-cores": 2,
+ }
+ ),
IRModule({"main": get_matmul_packed(128, 128, 128, "int8", "int8", "int32")}),
"dot_4x4_i8i8s32_neon",
),
(
Target(
- "llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+neon,+v8.2a,+dotprod -num-cores 2"
+ {
+ "kind": "llvm",
+ "device": "arm_cpu",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon", "+v8.2a", "+dotprod"],
+ "num-cores": 2,
+ }
),
IRModule({"main": get_matmul_packed(128, 128, 128, "int8", "int8", "int32")}),
"dot_4x4_i8i8s32_sdot",
),
(
Target(
- "llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+neon,+v8.2a,+dotprod -num-cores 2"
+ {
+ "kind": "llvm",
+ "device": "arm_cpu",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon", "+v8.2a", "+dotprod"],
+ "num-cores": 2,
+ }
),
IRModule({"main": get_matmul_packed(128, 128, 128, "uint8", "uint8", "uint32")}),
"dot_4x4_u8u8u32_udot",
),
(
Target(
- "llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+neon,+v8.2a,+dotprod -num-cores 2"
+ {
+ "kind": "llvm",
+ "device": "arm_cpu",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+neon", "+v8.2a", "+dotprod"],
+ "num-cores": 2,
+ }
),
IRModule({"main": get_matmul_packed(128, 128, 128, "uint8", "uint8", "int32")}),
"dot_4x4_u8u8i32_hdot",
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_tensorize.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_tensorize.py
index 2d3bbc429c..bedb22a4f9 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_tensorize.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_tensorize.py
@@ -468,7 +468,7 @@ def _create_context(mod, target, postprocs) -> ms.TuneContext:
def test_rewrite_tensorize_conv2d_nchwc_vnni():
mod = Conv2dNCHWcVNNIModuleTiled
- target = tvm.target.Target("llvm -mcpu=cascadelake -num-cores 4")
+ target = tvm.target.Target({"kind": "llvm", "mcpu": "cascadelake", "num-cores": 4})
ctx = _create_context(
mod,
target,
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_unbound_block.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_unbound_block.py
index 5dd8b0ad36..afbbd95a7d 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_unbound_block.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_postproc_rewrite_unbound_block.py
@@ -24,7 +24,7 @@ from tvm.target import Target
def _target() -> Target:
- return Target("cuda --max_threads_per_block=1024", host="llvm")
+ return Target({"kind": "cuda", "max_threads_per_block": 1024}, host="llvm")
def _create_context(mod, target) -> ms.TuneContext:
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_add_rfactor.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_add_rfactor.py
index d9f41bd601..3320236100 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_add_rfactor.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_add_rfactor.py
@@ -110,7 +110,7 @@ def test_cpu_matmul():
actual = generate_design_space(
kind="llvm",
mod=mod,
- target=Target("llvm --num-cores=32"),
+ target=Target({"kind": "llvm", "num-cores": 32}),
types=ms.schedule_rule.AddRFactor,
)
check_sketches(
@@ -274,7 +274,7 @@ def test_cpu_argmax():
actual = generate_design_space(
kind="llvm",
mod=mod,
- target=Target("llvm --num-cores=32"),
+ target=Target({"kind": "llvm", "num-cores": 32}),
types=ms.schedule_rule.AddRFactor,
)
check_sketches(
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_apply_custom_rule.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_apply_custom_rule.py
index abd6df1623..33bd9c4e74 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_apply_custom_rule.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_apply_custom_rule.py
@@ -54,7 +54,7 @@ def test_custom_rule():
space_gen = ms.space_generator.PostOrderApply(sch_rules=sch_rules)
ms.tune_tir(
mod=Matmul,
- target="llvm -num-cores=1",
+ target={"kind": "llvm", "num-cores": 1},
work_dir=tmpdir,
max_trials_global=10,
space=space_gen,
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_intrin.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_intrin.py
index 221ada7c0e..d9507500dc 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_intrin.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_intrin.py
@@ -30,7 +30,9 @@ from tvm.s_tir.tensor_intrin.x86 import AVX512_DOT_16x4_INTRIN as AVX512_INTRIN
from tvm.s_tir.tensor_intrin.x86 import VNNI_DOT_16x4_INTRIN as VNNI_INTRIN
-def test_x86_conv2d_nchwc(intrin=VNNI_INTRIN, target="llvm -mcpu=cascadelake -num-cores=4"):
+def test_x86_conv2d_nchwc(
+ intrin=VNNI_INTRIN, target={"kind": "llvm", "mcpu": "cascadelake", "num-cores": 4}
+):
@T.prim_func
def conv2d_nchwc(
placeholder: T.Buffer((1, 4, 56, 56, 16), "uint8"),
@@ -277,7 +279,7 @@ def _check_dp4a_dense(m, n, k, in_dtype, out_dtype, expected_mods, expected_deci
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=Target("cuda --arch=sm_70"),
+ target=Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
ms.schedule_rule.MultiLevelTilingWithIntrin(
@@ -418,7 +420,7 @@ def test_dp4a_dense_no_tensorize_2():
if __name__ == "__main__":
test_x86_conv2d_nchwc()
- test_x86_conv2d_nchwc(AVX512_INTRIN, "llvm -mcpu=skylake-avx512 -num-cores=4")
+ test_x86_conv2d_nchwc(AVX512_INTRIN, {"kind": "llvm", "mcpu": "skylake-avx512", "num-cores": 4})
test_dp4a_dense()
test_dp4a_dense_no_tensorize_1()
test_dp4a_dense_no_tensorize_2()
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_tc.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_tc.py
index 5c069909bb..32197e7910 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_tc.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_mlt_tc.py
@@ -213,7 +213,7 @@ def test_matmul_relu(shared_scope):
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
multi_level_tiling_tensor_core(
@@ -364,7 +364,7 @@ def test_matmul_relu_with_fallback():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
multi_level_tiling_tensor_core(),
@@ -527,7 +527,7 @@ def test_conv2d(shared_scope):
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
multi_level_tiling_tensor_core(
@@ -547,7 +547,7 @@ def test_conv2d(shared_scope):
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
multi_level_tiling_tensor_core(
@@ -711,7 +711,7 @@ def test_matmul_relu_pipeline(shared_scope):
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[
multi_level_tiling_tensor_core(
@@ -741,7 +741,7 @@ def test_matmul_relu_non_tensorizable():
(sch,) = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core(write_reuse_scope="shared")]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
@@ -884,7 +884,7 @@ def test_padded_matmul_relu():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core(write_reuse_scope="shared")]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
@@ -1042,7 +1042,7 @@ def test_conv_1x1():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core(write_reuse_scope="shared")]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
@@ -1194,7 +1194,7 @@ def test_padded_conv():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core(write_reuse_scope="shared")]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
@@ -1342,7 +1342,7 @@ def test_padded_matmul_single_padded_input():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core()]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
@@ -1489,7 +1489,7 @@ def test_padded_matmul_no_padded_output():
actual = generate_design_space(
kind="cuda",
mod=mod,
- target=tvm.target.Target("cuda --arch=sm_70"),
+ target=tvm.target.Target({"kind": "cuda", "arch": "sm_70"}),
types=None,
sch_rules=[multi_level_tiling_tensor_core()]
+ get_rules("cuda", ms.schedule_rule.AutoInline),
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_parallel_vectorize_unroll.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_parallel_vectorize_unroll.py
index 2cbba875f5..20c57ba66d 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_parallel_vectorize_unroll.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_schedule_rule_parallel_vectorize_unroll.py
@@ -258,7 +258,7 @@ def test_parallel_vectorize_unroll():
actual = generate_design_space(
kind="llvm",
mod=mod,
- target=Target("llvm --num-cores=32"),
+ target=Target({"kind": "llvm", "num-cores": 32}),
types=None,
sch_rules=[
ms.schedule_rule.ParallelizeVectorizeUnroll(
@@ -282,7 +282,7 @@ def test_parallel_vectorize_unroll_spatial():
actual = generate_design_space(
kind="llvm",
mod=mod,
- target=Target("llvm --num-cores=32"),
+ target=Target({"kind": "llvm", "num-cores": 32}),
types=None,
sch_rules=[
ms.schedule_rule.ParallelizeVectorizeUnroll(
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_space_post_opt.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_space_post_opt.py
index 28b08d19e7..c76eea7254 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_space_post_opt.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_space_post_opt.py
@@ -49,7 +49,7 @@ def matmul(a: T.handle, b: T.handle, c: T.handle) -> None:
@tvm.testing.requires_llvm
def test_tune_matmul_cpu():
with tempfile.TemporaryDirectory() as work_dir:
- target = Target("llvm --num-cores=16")
+ target = Target({"kind": "llvm", "num-cores": 16})
database = ms.tir_integration.tune_tir(
mod=matmul,
target=target,
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_trace_apply.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_trace_apply.py
index 869f052186..4e30969943 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_trace_apply.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_trace_apply.py
@@ -2720,7 +2720,7 @@ def test_conv2d_int8_vnni():
Conv2dInt8_NCHWc,
apply_trace,
Conv2dInt8_NCHWc_target,
- "llvm -mcpu=cascadelake",
+ {"kind": "llvm", "mcpu": "cascadelake"},
get_conv2d_vnni_mod(vnni_id),
)
diff --git a/tests/python/s_tir/meta_schedule/test_meta_schedule_tune_tir.py b/tests/python/s_tir/meta_schedule/test_meta_schedule_tune_tir.py
index 17bc1f90e9..51e93d58e6 100644
--- a/tests/python/s_tir/meta_schedule/test_meta_schedule_tune_tir.py
+++ b/tests/python/s_tir/meta_schedule/test_meta_schedule_tune_tir.py
@@ -66,7 +66,7 @@ def two_step(a: T.handle, c: T.handle) -> None:
@tvm.testing.requires_llvm
def test_tune_matmul_cpu():
with tempfile.TemporaryDirectory() as work_dir:
- target = Target("llvm --num-cores=16")
+ target = Target({"kind": "llvm", "num-cores": 16})
database = ms.tir_integration.tune_tir(
mod=matmul,
target=target,
@@ -163,7 +163,7 @@ def test_tune_block_cpu():
return RemoveBlock()
with tempfile.TemporaryDirectory() as work_dir:
- target = Target("llvm --num-cores=16")
+ target = Target({"kind": "llvm", "num-cores": 16})
database = ms.tir_integration.tune_tir(
mod=two_step,
target=target,
diff --git a/tests/python/s_tir/schedule/test_tir_schedule_split_fuse.py b/tests/python/s_tir/schedule/test_tir_schedule_split_fuse.py
index 9a610c9e58..0810f259cf 100644
--- a/tests/python/s_tir/schedule/test_tir_schedule_split_fuse.py
+++ b/tests/python/s_tir/schedule/test_tir_schedule_split_fuse.py
@@ -697,7 +697,7 @@ def test_sve_scalable_split_predicated(num_elements):
compile-time, we don't know if vscale is a multiple of the extent of the
loop to be split.
"""
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu -mattr=+sve"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}):
outer_extent = tvm.arith.Analyzer().simplify(T.ceildiv(num_elements, 4 * T.vscale()))
@T.prim_func
@@ -733,7 +733,7 @@ def test_sve_scalable_split_assume_exact_multiple():
a predicate is not created. This can be used to ensure predication is not
inserted.
"""
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu -mattr=+sve"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}):
outer_extent = tvm.arith.Analyzer().simplify(T.ceildiv(128, 4 * T.vscale()))
@T.prim_func
@@ -785,7 +785,7 @@ def test_sve_split_over_scalable_loop():
T.where(i_0 * (T.vscale() * 2) + i_1 < T.vscale() * 4)
A[v_i] = 1.0
- with tvm.target.Target("llvm -mtriple=aarch64-linux-gnu -mattr=+sve"):
+ with tvm.target.Target({"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}):
sch = tvm.s_tir.Schedule(before)
(a,) = sch.get_loops("A")
sch.split(
diff --git a/tests/python/target/test_arm_target.py b/tests/python/target/test_arm_target.py
index d656031ad9..2b60827094 100644
--- a/tests/python/target/test_arm_target.py
+++ b/tests/python/target/test_arm_target.py
@@ -28,23 +28,59 @@ from tvm.target import codegen
llvm_version, arm_target, input_dtype, kernel_dtype, is_supported = tvm.testing.parameters(
# Testing mcpu type
- (8, "c -mcpu=cortex-m4", "int8", "int8", False),
- (8, "c -mcpu=cortex-m7", "int8", "int8", False),
- (8, "c -mcpu=cortex-m33", "int8", "int8", False),
- (8, "c -mcpu=cortex-m55", "int8", "int8", False),
- (8, "c -mcpu=cortex-m3", "int8", "int8", False),
- (7, "llvm -mtriple=arm-linux-gnueabi -mattr=+neon", "int8", "int8", False),
- (8, "llvm -mtriple=arm-linux-gnueabi -mattr=+neon", "int8", "int8", True),
- (9, "llvm -mtriple=arm-linux-gnueabi -mattr=+neon", "int8", "int8", True),
- (8, "llvm -mtriple=arm-linux-gnueabi", "int8", "int8", False),
- (7, "llvm -mtriple=aarch64-linux-gnu -mattr=+v8.4a,+dotprod", "int8", "int8", False),
- (8, "llvm -mtriple=aarch64-linux-gnu -mattr=+v8.4a,+dotprod", "int8", "int8", True),
- (9, "llvm -mtriple=arm-linux-gnueabi -mattr=+neon", "int8", "int8", True),
- (8, "llvm -mtriple=aarch64-linux-gnu", "int8", "int8", True),
+ (8, {"kind": "c", "mcpu": "cortex-m4"}, "int8", "int8", False),
+ (8, {"kind": "c", "mcpu": "cortex-m7"}, "int8", "int8", False),
+ (8, {"kind": "c", "mcpu": "cortex-m33"}, "int8", "int8", False),
+ (8, {"kind": "c", "mcpu": "cortex-m55"}, "int8", "int8", False),
+ (8, {"kind": "c", "mcpu": "cortex-m3"}, "int8", "int8", False),
+ (
+ 7,
+ {"kind": "llvm", "mtriple": "arm-linux-gnueabi", "mattr": ["+neon"]},
+ "int8",
+ "int8",
+ False,
+ ),
+ (8, {"kind": "llvm", "mtriple": "arm-linux-gnueabi", "mattr": ["+neon"]}, "int8", "int8", True),
+ (9, {"kind": "llvm", "mtriple": "arm-linux-gnueabi", "mattr": ["+neon"]}, "int8", "int8", True),
+ (8, {"kind": "llvm", "mtriple": "arm-linux-gnueabi"}, "int8", "int8", False),
+ (
+ 7,
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+v8.4a", "+dotprod"]},
+ "int8",
+ "int8",
+ False,
+ ),
+ (
+ 8,
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+v8.4a", "+dotprod"]},
+ "int8",
+ "int8",
+ True,
+ ),
+ (9, {"kind": "llvm", "mtriple": "arm-linux-gnueabi", "mattr": ["+neon"]}, "int8", "int8", True),
+ (8, {"kind": "llvm", "mtriple": "aarch64-linux-gnu"}, "int8", "int8", True),
# Testing dtype
- (8, "llvm -mtriple=aarch64-linux-gnu -mattr=+neon", "int16", "int8", False),
- (8, "llvm -mtriple=aarch64-linux-gnu -mattr=+neon", "int8", "int16", False),
- (8, "llvm -mtriple=aarch64-linux-gnu -mattr=+neon", "int16", "int16", False),
+ (
+ 8,
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+neon"]},
+ "int16",
+ "int8",
+ False,
+ ),
+ (
+ 8,
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+neon"]},
+ "int8",
+ "int16",
+ False,
+ ),
+ (
+ 8,
+ {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+neon"]},
+ "int16",
+ "int16",
+ False,
+ ),
)
@@ -73,7 +109,7 @@ def sve_device_vector_length():
@tvm.testing.requires_aarch64_sve
def test_scalable_div(sve_device_vector_length):
np.random.seed(0)
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
dev = tvm.cpu(0)
@T.prim_func
@@ -94,7 +130,7 @@ def test_scalable_div(sve_device_vector_length):
@tvm.testing.requires_aarch64_sve
def test_scalable_buffer_load_store(sve_device_vector_length):
np.random.seed(0)
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
num_elements = sve_device_vector_length // 32
dev = tvm.cpu(0)
@@ -122,7 +158,7 @@ def test_scalable_loop_bound(sve_device_vector_length):
dtype = "float32"
num_elements = sve_device_vector_length // 32
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
dev = tvm.cpu(0)
@T.prim_func
@@ -146,7 +182,7 @@ def test_scalable_loop_bound(sve_device_vector_length):
@tvm.testing.requires_aarch64_sve
def test_scalable_broadcast(sve_device_vector_length):
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+sve"
+ target = {"kind": "llvm", "mtriple": "aarch64-linux-gnu", "mattr": ["+sve"]}
num_elements = sve_device_vector_length // 32
dev = tvm.cpu(0)
diff --git a/tests/python/target/test_llvm_features_info.py b/tests/python/target/test_llvm_features_info.py
index e627c89aa8..7677cf27fc 100644
--- a/tests/python/target/test_llvm_features_info.py
+++ b/tests/python/target/test_llvm_features_info.py
@@ -39,7 +39,7 @@ def test_llvm_targets(capfd):
assert str(codegen.llvm_get_targets()) == str(_ffi_api.llvm_get_targets())
tvm.target.codegen.llvm_get_cpu_features(
- tvm.target.Target("llvm -mtriple=x86_64-linux-gnu -mcpu=dummy")
+ tvm.target.Target({"kind": "llvm", "mtriple": "x86_64-linux-gnu", "mcpu": "dummy"})
)
expected_str = (
" with `-mcpu=dummy` is not valid in "
@@ -68,7 +68,7 @@ min_llvm_version, llvm_target, cpu_arch, cpu_features, is_supported = tvm.testin
def test_target_features(min_llvm_version, llvm_target, cpu_arch, cpu_features, is_supported):
- target = Target("llvm -mtriple=%s-- -mcpu=%s" % (llvm_target, cpu_arch))
+ target = Target({"kind": "llvm", "mtriple": f"{llvm_target}--", "mcpu": cpu_arch})
##
## legalize llvm_target
diff --git a/tests/python/target/test_riscv_features.py b/tests/python/target/test_riscv_features.py
index 17452a86db..cefdc0903f 100644
--- a/tests/python/target/test_riscv_features.py
+++ b/tests/python/target/test_riscv_features.py
@@ -25,19 +25,19 @@ LLVM_VERSION = codegen.llvm_version_major()
# fmt: off
min_llvm_version, tvm_target, vec_width = tvm.testing.parameters(
# generic, no vector -> (default 128)
- (-1, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+i,+m", 128),
- (-1, "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+64bit,+a,+c,+d,+f,+m", 128),
+ (-1, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "generic-rv64", "mattr": ["+i", "+m"]}, 128),
+ (-1, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv32-linux-gnu", "mcpu": "generic-rv32", "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m"]}, 128),
# generic, with vector -> (default zvl128b)
- (-1, "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+i,+m,+v", 128),
- (-1, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v", 128),
+ (-1, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv32-linux-gnu", "mcpu": "generic-rv32", "mattr": ["+i", "+m", "+v"]}, 128),
+ (-1, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "generic-rv64", "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v"]}, 128),
# explicit +zvlXXXb
- (14, "llvm -device=riscv_cpu -mtriple=riscv32-linux-gnu -mcpu=generic-rv32 -mattr=+i,+m,+v,+zvl64b", 128),
- (14, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v,+zvl256b", 256),
- (14, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=generic-rv64 -mattr=+64bit,+a,+c,+d,+f,+m,+v,+zvl512b", 512),
+ (14, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv32-linux-gnu", "mcpu": "generic-rv32", "mattr": ["+i", "+m", "+v", "+zvl64b"]}, 128),
+ (14, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "generic-rv64", "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v", "+zvl256b"]}, 256),
+ (14, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "generic-rv64", "mattr": ["+64bit", "+a", "+c", "+d", "+f", "+m", "+v", "+zvl512b"]}, 512),
# vendor CPU
- (17, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=sifive-x280", 512),
- (18, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=sifive-p670", 128),
- (19, "llvm -device=riscv_cpu -mtriple=riscv64-linux-gnu -mcpu=spacemit-x60", 256),
+ (17, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "sifive-x280"}, 512),
+ (18, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "sifive-p670"}, 128),
+ (19, {"kind": "llvm", "device": "riscv_cpu", "mtriple": "riscv64-linux-gnu", "mcpu": "spacemit-x60"}, 256),
) # fmt: on
diff --git a/tests/python/target/test_target_parser_mprofile.py b/tests/python/target/test_target_parser_mprofile.py
index 75a0a93566..a5ddcb9ff3 100644
--- a/tests/python/target/test_target_parser_mprofile.py
+++ b/tests/python/target/test_target_parser_mprofile.py
@@ -25,7 +25,7 @@ from tvm.target import Target
@pytest.mark.parametrize(["cpu_target"], [["c"], ["llvm"]])
def test_target_parser_mprofile(cpu_target):
- parsed_target = Target(f"{cpu_target} -mcpu=cortex-m55")
+ parsed_target = Target({"kind": cpu_target, "mcpu": "cortex-m55"})
assert len(parsed_target.keys) == 2
assert parsed_target.keys[0] == "arm_cpu"
assert parsed_target.keys[1] == "cpu"
@@ -36,7 +36,7 @@ def test_target_parser_mprofile(cpu_target):
@pytest.mark.parametrize(["cpu_target"], [["c"], ["llvm"]])
def test_target_parser_mprofile_no_mve(cpu_target):
- parsed_target = Target(f"{cpu_target} -mcpu=cortex-m7")
+ parsed_target = Target({"kind": cpu_target, "mcpu": "cortex-m7"})
assert len(parsed_target.keys) == 2
assert parsed_target.keys[0] == "arm_cpu"
assert parsed_target.keys[1] == "cpu"
@@ -47,7 +47,7 @@ def test_target_parser_mprofile_no_mve(cpu_target):
@pytest.mark.parametrize(["cpu_target"], [["c"], ["llvm"]])
def test_target_parser_mprofile_no_dsp(cpu_target):
- parsed_target = Target(f"{cpu_target} -mcpu=cortex-m3")
+ parsed_target = Target({"kind": cpu_target, "mcpu": "cortex-m3"})
assert len(parsed_target.keys) == 2
assert parsed_target.keys[0] == "arm_cpu"
assert parsed_target.keys[1] == "cpu"
@@ -58,7 +58,7 @@ def test_target_parser_mprofile_no_dsp(cpu_target):
@pytest.mark.parametrize(["cpu_target"], [["llvm"]])
def test_target_parser_mprofile_mattr(cpu_target):
- parsed_target = Target(f"{cpu_target} -mcpu=cortex-m55 -mattr=+nomve,+woof")
+ parsed_target = Target({"kind": cpu_target, "mcpu": "cortex-m55", "mattr": ["+nomve", "+woof"]})
assert len(parsed_target.keys) == 2
assert parsed_target.keys[0] == "arm_cpu"
assert parsed_target.keys[1] == "cpu"
diff --git a/tests/python/target/test_target_target.py b/tests/python/target/test_target_target.py
index 8aa314bd62..d9c31c6ae8 100644
--- a/tests/python/target/test_target_target.py
+++ b/tests/python/target/test_target_target.py
@@ -40,13 +40,13 @@ def test_all_targets_device_type_verify():
def test_target_string_parse():
- target = tvm.target.Target("cuda -model=unknown -libs=cublas,cudnn")
+ target = tvm.target.Target({"kind": "cuda", "model": "unknown", "libs": ["cublas", "cudnn"]})
assert target.kind.name == "cuda"
assert target.model == "unknown"
assert set(target.keys) == set(["cuda", "gpu"])
assert set(target.libs) == set(["cublas", "cudnn"])
- assert str(target) == str(tvm.target.cuda(options="-libs=cublas,cudnn"))
+ assert str(target) == str(tvm.target.cuda(libs=["cublas", "cudnn"]))
assert tvm.target.intel_graphics().device_name == "intel_graphics"
assert tvm.target.mali().device_name == "mali"
@@ -55,7 +55,7 @@ def test_target_string_parse():
def test_target_string_with_spaces():
target = tvm.target.Target(
- "vulkan -device_name='Name of GPU with spaces' -device_type=discrete"
+ {"kind": "vulkan", "device_name": "Name of GPU with spaces", "device_type": "discrete"}
)
assert target.attrs["device_name"] == "Name of GPU with spaces"
assert target.attrs["device_type"] == "discrete"
@@ -67,23 +67,25 @@ def test_target_string_with_spaces():
def test_target_llvm_options():
- target = tvm.target.Target("llvm -cl-opt='-unroll-threshold:uint=100,-unroll-count:uint=3'")
+ target = tvm.target.Target(
+ {"kind": "llvm", "cl-opt": ["-unroll-threshold:uint=100", "-unroll-count:uint=3"]}
+ )
assert sorted(target.attrs["cl-opt"]) == sorted(
["-unroll-threshold:uint=100", "-unroll-count:uint=3"]
)
def test_target_llvm_jit_options():
- target = tvm.target.Target("llvm -jit=mcjit")
+ target = tvm.target.Target({"kind": "llvm", "jit": "mcjit"})
assert target.attrs["jit"] == "mcjit"
- target = tvm.target.Target("llvm -jit=orcjit")
+ target = tvm.target.Target({"kind": "llvm", "jit": "orcjit"})
assert target.attrs["jit"] == "orcjit"
def test_target_llvm_vector_width():
- target = tvm.target.Target("llvm -vector-width=256")
+ target = tvm.target.Target({"kind": "llvm", "vector-width": 256})
assert target.attrs["vector-width"] == 256
- target = tvm.target.Target("llvm -vector-width=1024")
+ target = tvm.target.Target({"kind": "llvm", "vector-width": 1024})
assert target.attrs["vector-width"] == 1024
@@ -129,7 +131,9 @@ def test_config_map():
def test_composite_target():
- tgt = tvm.target.Target("composite --host=llvm --devices=cuda,opencl")
+ tgt = tvm.target.Target(
+ {"kind": "composite", "host": {"kind": "llvm"}, "devices": ["cuda", "opencl"]}
+ )
assert tgt.kind.name == "composite"
assert tgt.host.kind.name == "llvm"
assert len(tgt.attrs["devices"]) == 2
@@ -204,13 +208,13 @@ def test_target_host_single_dict():
def test_target_host_single_string():
- tgt = tvm.target.Target("cuda --host llvm")
+ tgt = tvm.target.Target({"kind": "cuda", "host": {"kind": "llvm"}})
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "llvm"
def test_target_host_single_string_with_tag():
- tgt = tvm.target.Target("cuda --host nvidia/jetson-nano")
+ tgt = tvm.target.Target({"kind": "cuda", "host": "nvidia/jetson-nano"})
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "cuda"
assert tgt.host.attrs["arch"] == "sm_53"
@@ -221,7 +225,7 @@ def test_target_host_single_string_with_tag():
def test_target_host_merge_0():
- tgt = tvm.target.Target(tvm.target.Target("cuda --host nvidia/jetson-nano"), None)
+ tgt = tvm.target.Target(tvm.target.Target({"kind": "cuda", "host": "nvidia/jetson-nano"}), None)
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "cuda"
assert tgt.host.attrs["arch"] == "sm_53"
@@ -232,7 +236,7 @@ def test_target_host_merge_0():
def test_target_host_merge_1():
- tgt = tvm.target.Target("cuda --host llvm")
+ tgt = tvm.target.Target({"kind": "cuda", "host": {"kind": "llvm"}})
tgt = tvm.target.Target(tgt, tgt.host)
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "llvm"
@@ -240,7 +244,10 @@ def test_target_host_merge_1():
def test_target_host_merge_2():
"""Test picking the same host is ok."""
- tgt = tvm.target.Target(tvm.target.Target("cuda --host llvm"), tvm.target.Target("llvm"))
+ tgt = tvm.target.Target(
+ tvm.target.Target({"kind": "cuda", "host": {"kind": "llvm"}}),
+ tvm.target.Target("llvm"),
+ )
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "llvm"
@@ -248,7 +255,7 @@ def test_target_host_merge_2():
def test_target_tvm_object():
"""Test creating Target by using TVM Objects"""
String = tvm.runtime.container.String
- tgt = tvm.target.Target(target=String("cuda --host llvm"))
+ tgt = tvm.target.Target(target={"kind": "cuda", "host": {"kind": "llvm"}})
assert tgt.kind.name == "cuda"
assert tgt.host.kind.name == "llvm"
tgt = tvm.target.Target(target=String("cuda"), host=String("llvm"))
@@ -259,7 +266,7 @@ def test_target_tvm_object():
@pytest.mark.skip(reason="Causing infinite loop because of pytest and handle issue")
def test_target_host_merge_3():
with pytest.raises(ValueError, match=r"target host has to be a string or dictionary."):
- tvm.target.Target(tvm.target.Target("cuda --host llvm"), 12.34)
+ tvm.target.Target(tvm.target.Target({"kind": "cuda", "host": {"kind": "llvm"}}), 12.34)
def test_target_with_host():
@@ -384,13 +391,13 @@ def test_canon_target_map_and_host():
def test_target_attr_bool_value():
- target0 = Target("vulkan --supports_float16=True")
+ target0 = Target({"kind": "vulkan", "supports_float16": True})
assert target0.attrs["supports_float16"] == 1
- target1 = Target("vulkan --supports_float16=true")
+ target1 = Target({"kind": "vulkan", "supports_float16": True})
assert target1.attrs["supports_float16"] == 1
- target2 = Target("vulkan --supports_float16=False")
+ target2 = Target({"kind": "vulkan", "supports_float16": False})
assert target2.attrs["supports_float16"] == 0
- target3 = Target("vulkan --supports_float16=false")
+ target3 = Target({"kind": "vulkan", "supports_float16": False})
assert target3.attrs["supports_float16"] == 0
@@ -484,5 +491,34 @@ def test_module_dict_from_deserialized_targets():
lib["func"]()
+def test_json_roundtrip():
+ """Test that Target(str(target)) roundtrips correctly."""
+ target = Target({"kind": "llvm", "mcpu": "cortex-a53"})
+ target2 = Target(str(target))
+ assert target2.kind.name == "llvm"
+ assert target2.attrs["mcpu"] == "cortex-a53"
+
+ # Test with more complex target
+ target = Target({"kind": "cuda", "arch": "sm_80", "max_threads_per_block": 1024})
+ target2 = Target(str(target))
+ assert target2.kind.name == "cuda"
+ assert target2.attrs["arch"] == "sm_80"
+
+
+def test_str_is_json():
+ """Test that str() output is valid JSON."""
+ target = Target({"kind": "llvm", "mcpu": "cortex-a53"})
+ s = str(target)
+ parsed = json.loads(s)
+ assert parsed["kind"] == "llvm"
+ assert parsed["mcpu"] == "cortex-a53"
+
+
+def test_cli_string_rejected():
+ """Test that CLI string form is rejected."""
+ with pytest.raises(ValueError):
+ Target("llvm -mcpu=cortex-a53")
+
+
if __name__ == "__main__":
tvm.testing.main()
diff --git a/tests/python/target/test_x86_features.py b/tests/python/target/test_x86_features.py
index 41ac844eea..cb810f7cab 100644
--- a/tests/python/target/test_x86_features.py
+++ b/tests/python/target/test_x86_features.py
@@ -24,83 +24,108 @@ LLVM_VERSION = codegen.llvm_version_major()
min_llvm_version, tvm_target, x86_feature, is_supported = tvm.testing.parameters(
# sse4.1
- (-1, "llvm -mtriple=x86_64-- -mcpu=btver2", "sse4a", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=penryn", "sse4.1", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=silvermont", "sse4.2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=slm", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=goldmont", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=goldmont-plus", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=tremont", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=nehalem", "sse4.2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=corei7", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=westmere", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=bdver1", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=bdver2", "sse4.2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=bdver3", "sse4.2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=x86-64-v2", "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "btver2"}, "sse4a", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "penryn"}, "sse4.1", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "silvermont"}, "sse4.2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "slm"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "goldmont"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "goldmont-plus"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "tremont"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "nehalem"}, "sse4.2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "corei7"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "westmere"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "bdver1"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "bdver2"}, "sse4.2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "bdver3"}, "sse4.2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64-v2"}, "sse4.2", True),
# avx
- (-1, "llvm -mtriple=x86_64-- -mcpu=sandybridge", "avx", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=corei7-avx", "avx", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=ivybridge", "avx", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=core-avx-i", "avx", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "sandybridge"}, "avx", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "corei7-avx"}, "avx", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "ivybridge"}, "avx", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "core-avx-i"}, "avx", True),
# avx2
- (-1, "llvm -mtriple=x86_64-- -mcpu=haswell", "avx2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=core-avx2", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=broadwell", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=skylake", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=bdver4", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=znver1", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=znver2", "avx2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=znver3", "avx2", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=x86-64-v3", "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "haswell"}, "avx2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "core-avx2"}, "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "broadwell"}, "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "skylake"}, "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "bdver4"}, "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "znver1"}, "avx2", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "znver2"}, "avx2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "znver3"}, "avx2", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64-v3"}, "avx2", True),
# avx512bw
- (-1, "llvm -mtriple=x86_64-- -mcpu=skylake-avx512", "avx512bw", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=skx", "avx512bw", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=knl", "avx512bw", False),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knl", "avx512f", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=knl", ["avx512bw", "avx512f"], False),
- (11, "llvm -mtriple=x86_64-- -mcpu=knl", ("avx512bw", "avx512f"), False),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knl", "avx512cd", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=knl", ["avx512cd", "avx512f"], True),
- (11, "llvm -mtriple=x86_64-- -mcpu=knl", ("avx512cd", "avx512f"), True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knl", "avx512er", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knl", "avx512pf", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=knm", "avx512bw", False),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knm", "avx512f", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knm", "avx512cd", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knm", "avx512er", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=knm", "avx512pf", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=x86-64-v4", "avx512bw", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=cannonlake", "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "skylake-avx512"}, "avx512bw", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "skx"}, "avx512bw", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, "avx512bw", False),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, "avx512f", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, ["avx512bw", "avx512f"], False),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, ("avx512bw", "avx512f"), False),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, "avx512cd", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, ["avx512cd", "avx512f"], True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, ("avx512cd", "avx512f"), True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, "avx512er", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knl"}, "avx512pf", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knm"}, "avx512bw", False),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knm"}, "avx512f", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knm"}, "avx512cd", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knm"}, "avx512er", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "knm"}, "avx512pf", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64-v4"}, "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cannonlake"}, "avx512bw", True),
# explicit enumeration of VNNI capable due to collision with alderlake
- (11, "llvm -mtriple=x86_64-- -mcpu=alderlake", "avx512bw", False),
- (-1, "llvm -mtriple=x86_64-- -mcpu=cascadelake", "avx512bw", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=icelake-client", "avx512bw", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=icelake-server", "avx512bw", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=rocketlake", "avx512bw", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=tigerlake", "avx512bw", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=cooperlake", "avx512bw", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=sapphirerapids", "avx512bw", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "alderlake"}, "avx512bw", False),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cascadelake"}, "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "icelake-client"}, "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "icelake-server"}, "avx512bw", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "rocketlake"}, "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "tigerlake"}, "avx512bw", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cooperlake"}, "avx512bw", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "sapphirerapids"}, "avx512bw", True),
# avx512vnni
- (11, "llvm -mtriple=x86_64-- -mcpu=alderlake", "avx512vnni", False),
- (11, "llvm -mtriple=x86_64-- -mcpu=alderlake", "avxvnni", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=cascadelake", "avx512vnni", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=icelake-client", "avx512vnni", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=icelake-server", "avx512vnni", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=rocketlake", "avx512vnni", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=tigerlake", "avx512vnni", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=cooperlake", "avx512vnni", True),
- (11, "llvm -mtriple=x86_64-- -mcpu=sapphirerapids", "avx512vnni", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "alderlake"}, "avx512vnni", False),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "alderlake"}, "avxvnni", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cascadelake"}, "avx512vnni", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "icelake-client"}, "avx512vnni", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "icelake-server"}, "avx512vnni", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "rocketlake"}, "avx512vnni", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "tigerlake"}, "avx512vnni", True),
+ (-1, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cooperlake"}, "avx512vnni", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "sapphirerapids"}, "avx512vnni", True),
# amx-int8
- (11, "llvm -mtriple=x86_64-- -mcpu=sapphirerapids", "amx-int8", True),
+ (11, {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "sapphirerapids"}, "amx-int8", True),
# generic CPU (no features) but with extra -mattr
- (-1, "llvm -mtriple=x86_64-- -mcpu=x86-64 -mattr=+sse4.1,+avx2", "avx2", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=x86-64 -mattr=+sse4.1,+avx2", "sse4.1", True),
+ (
+ -1,
+ {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64", "mattr": ["+sse4.1", "+avx2"]},
+ "avx2",
+ True,
+ ),
+ (
+ -1,
+ {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64", "mattr": ["+sse4.1", "+avx2"]},
+ "sse4.1",
+ True,
+ ),
# enabling +sse4.1 implies ssse3 presence in LLVM
- (-1, "llvm -mtriple=x86_64-- -mcpu=x86-64 -mattr=+sse4.1,+avx2", "ssse3", True),
- (-1, "llvm -mtriple=x86_64-- -mcpu=ivybridge -mattr=-ssse3", "ssse3", False),
+ (
+ -1,
+ {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "x86-64", "mattr": ["+sse4.1", "+avx2"]},
+ "ssse3",
+ True,
+ ),
+ (
+ -1,
+ {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "ivybridge", "mattr": ["-ssse3"]},
+ "ssse3",
+ False,
+ ),
# disabling avx512f (foundation) also disables avx512bw
- (-1, "llvm -mtriple=x86_64-- -mcpu=cascadelake -mattr=-avx512f", "avx512bw", False),
+ (
+ -1,
+ {"kind": "llvm", "mtriple": "x86_64--", "mcpu": "cascadelake", "mattr": ["-avx512f"]},
+ "avx512bw",
+ False,
+ ),
)
diff --git a/tests/python/tir-base/test_tir_intrin.py b/tests/python/tir-base/test_tir_intrin.py
index c17e872f27..9f199a2b63 100644
--- a/tests/python/tir-base/test_tir_intrin.py
+++ b/tests/python/tir-base/test_tir_intrin.py
@@ -249,7 +249,7 @@ def test_ldexp():
dtype = tvm.testing.parameter("int32", "int64")
-@tvm.testing.parametrize_targets("llvm", "vulkan -from_device=0")
+@tvm.testing.parametrize_targets("llvm", {"kind": "vulkan", "from_device": 0})
def test_clz(target, dev, dtype):
target = tvm.target.Target(target)
if (
diff --git a/tests/python/tir-transform/test_tir_transform_helpers.py b/tests/python/tir-transform/test_tir_transform_helpers.py
index b8bb2ff822..dedc60af87 100644
--- a/tests/python/tir-transform/test_tir_transform_helpers.py
+++ b/tests/python/tir-transform/test_tir_transform_helpers.py
@@ -186,14 +186,14 @@ def test_bind_target_updates_host():
T.func_attr(
{
"global_symbol": "func",
- "target": T.target("nvptx", host="llvm -opt-level=0"),
+ "target": T.target("nvptx", host={"kind": "llvm", "opt-level": 0}),
}
)
T.evaluate(0)
- After = tvm.tir.transform.BindTarget(tvm.target.Target("cuda", host="llvm -opt-level=0"))(
- Before
- )
+ After = tvm.tir.transform.BindTarget(
+ tvm.target.Target("cuda", host={"kind": "llvm", "opt-level": 0})
+ )(Before)
tvm.ir.assert_structural_equal(After, Expected)
@@ -256,7 +256,7 @@ def test_bind_target_with_device_host_call_same_func():
@T.prim_func(private=True)
def add_host(a: T.int32, b: T.int32) -> T.int32:
- T.func_attr({"target": T.target("llvm -opt-level=0")})
+ T.func_attr({"target": T.target({"kind": "llvm", "opt-level": 0})})
return a + b
@T.prim_func
@@ -266,16 +266,19 @@ def test_bind_target_with_device_host_call_same_func():
C: T.Buffer((128, 128), "int32"),
):
T.func_attr(
- {"global_symbol": "main", "target": T.target("cuda", host="llvm -opt-level=0")}
+ {
+ "global_symbol": "main",
+ "target": T.target("cuda", host={"kind": "llvm", "opt-level": 0}),
+ }
)
length: T.int32 = Expected.add_host(64, 64) # Call from host
for bx in T.thread_binding(length, "blockIdx.x"):
for tx in T.thread_binding(length, "threadIdx.x"):
C[bx, tx] = Expected.add(A[bx, tx], B[bx, tx]) # Call from device
- After = tvm.tir.transform.BindTarget(tvm.target.Target("cuda", host="llvm -opt-level=0"))(
- Before
- )
+ After = tvm.tir.transform.BindTarget(
+ tvm.target.Target("cuda", host={"kind": "llvm", "opt-level": 0})
+ )(Before)
tvm.ir.assert_structural_equal(After, Expected)
diff --git a/tests/python/tir-transform/test_tir_transform_split_host_device.py b/tests/python/tir-transform/test_tir_transform_split_host_device.py
index 5fb59326a7..e5bc6c31b9 100644
--- a/tests/python/tir-transform/test_tir_transform_split_host_device.py
+++ b/tests/python/tir-transform/test_tir_transform_split_host_device.py
@@ -58,7 +58,7 @@ def test_split_host_device():
class Before:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
T.attr(T.target("cuda"), "target", 0)
T.evaluate(n)
@@ -66,7 +66,7 @@ def test_split_host_device():
class Expected:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
Expected.main_kernel(n)
@T.prim_func(private=True)
@@ -91,7 +91,7 @@ def test_split_host_device_on_cpu():
class Before:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
T.attr(T.target("llvm"), "target", 0)
T.evaluate(n)
@@ -99,7 +99,7 @@ def test_split_host_device_on_cpu():
class Expected:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
err = Expected.main_kernel(n)
assert err == 0, "Error executing compute kernel"
@@ -187,7 +187,7 @@ def test_split_host_device_name_collision():
class Before:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
T.attr(T.target("cuda"), "target", 0)
T.evaluate(n)
@@ -200,7 +200,7 @@ def test_split_host_device_name_collision():
class Expected:
@T.prim_func
def main(n: T.int32):
- T.func_attr({"target": T.target("cuda", host="llvm -opt-level=0")})
+ T.func_attr({"target": T.target("cuda", host={"kind": "llvm", "opt-level": 0})})
Expected.main_kernel_1(n)
@T.prim_func(private=True)
diff --git a/tests/python/tir-transform/test_tir_transform_vectorize.py b/tests/python/tir-transform/test_tir_transform_vectorize.py
index bba0a500bc..898b9428e5 100644
--- a/tests/python/tir-transform/test_tir_transform_vectorize.py
+++ b/tests/python/tir-transform/test_tir_transform_vectorize.py
@@ -21,8 +21,15 @@ import tvm.testing
from tvm.script import ir as I
from tvm.script import tir as T
-simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
-sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
+simple_target = tvm.target.Target({"kind": "llvm", "mtriple": "x86_64-linux-gnu"})
+sve_target = tvm.target.Target(
+ {
+ "kind": "llvm",
+ "device": "arm_cpu",
+ "mtriple": "aarch64-linux-gnu",
+ "mattr": ["+v8.2a", "+sve"],
+ }
+)
@pytest.mark.parametrize("extent, target", [(4, simple_target), (T.vscale() * 4, sve_target)])
@@ -505,8 +512,8 @@ def test_illegal_vscale_in_non_sve_compilation():
A[j] = 13
msg = (
- f"Failed to vectorize loop with extent T.vscale\\(\\) \\* 4 for target "
- f"llvm -keys=cpu -mtriple=x86_64-linux-gnu"
+ "Failed to vectorize loop with extent T.vscale\\(\\) \\* 4 for target "
+ '\\{"keys": \\["cpu"\\], "kind": "llvm", "mtriple": "x86_64-linux-gnu"'
)
with tvm.target.Target(simple_target):
with pytest.raises(tvm.error.InternalError, match=msg):
diff --git a/web/tests/python/prepare_test_libs.py b/web/tests/python/prepare_test_libs.py
index 7198c16341..79546c47aa 100644
--- a/web/tests/python/prepare_test_libs.py
+++ b/web/tests/python/prepare_test_libs.py
@@ -34,7 +34,7 @@ def prepare_relax_lib(base_path):
lv0 = R.add(x, y)
return lv0
- target = tvm.target.Target("llvm -mtriple=wasm32-unknown-unknown-wasm")
+ target = tvm.target.Target({"kind": "llvm", "mtriple": "wasm32-unknown-unknown-wasm"})
mod = pipeline(Mod)
ex = relax.build(mod, target)
@@ -43,7 +43,7 @@ def prepare_relax_lib(base_path):
def prepare_tir_lib(base_path):
- target = "llvm -mtriple=wasm32-unknown-unknown-wasm"
+ target = {"kind": "llvm", "mtriple": "wasm32-unknown-unknown-wasm"}
if not tvm.runtime.enabled(target):
raise RuntimeError("Target %s is not enbaled" % target)
n = te.var("n")
diff --git a/web/tests/python/relax_rpc_test.py b/web/tests/python/relax_rpc_test.py
index e407d6ba63..b50c570eb7 100644
--- a/web/tests/python/relax_rpc_test.py
+++ b/web/tests/python/relax_rpc_test.py
@@ -54,7 +54,9 @@ def test_rpc():
dtype = "float32"
temp = utils.tempdir()
wasm_path = temp.relpath("relax.wasm")
- target = tvm.target.Target("webgpu", host="llvm -mtriple=wasm32-unknown-unknown-wasm")
+ target = tvm.target.Target(
+ "webgpu", host={"kind": "llvm", "mtriple": "wasm32-unknown-unknown-wasm"}
+ )
mod = get_model()
ex = relax.build(mod, target)
diff --git a/web/tests/python/webgpu_rpc_test.py b/web/tests/python/webgpu_rpc_test.py
index bd84d153b8..ff392a4128 100644
--- a/web/tests/python/webgpu_rpc_test.py
+++ b/web/tests/python/webgpu_rpc_test.py
@@ -34,7 +34,9 @@ def test_rpc():
if not tvm.runtime.enabled("rpc"):
return
# generate the wasm library
- target = tvm.target.Target("webgpu", host="llvm -mtriple=wasm32-unknown-unknown-wasm")
+ target = tvm.target.Target(
+ "webgpu", host={"kind": "llvm", "mtriple": "wasm32-unknown-unknown-wasm"}
+ )
n = te.var("n")
A = te.placeholder((n,), name="A")