This is an automated email from the ASF dual-hosted git repository.
mshr-h pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 15b607d6bf [CUDA] Add NVIDIA Jetson AGX Thor target tag (#20259)
15b607d6bf is described below
commit 15b607d6bf1d3a64d45f7227a3b0ec55afd5c015
Author: Shushi Hong <[email protected]>
AuthorDate: Wed Sep 2 08:49:47 2026 -0400
[CUDA] Add NVIDIA Jetson AGX Thor target tag (#20259)
This PR adds a CUDA target tag for NVIDIA Jetson AGX Thor:
```python
tvm.target.Target("nvidia/jetson-agx-thor")
```
The target configuration includes:
- CUDA architecture: `sm_110a`
- Maximum shared memory per block: `232448` bytes
- Registers per block: `65536`
- L2 cache: `32 MiB`
- Host CPU: Arm Neoverse V3AE with 14 cores
The change is limited to the new target tag and its configuration test.
It does not change generic CUDA target detection, CUDA runtime behavior,
or existing targets.
Validation performed on NVIDIA Thor:
- CUDA kernel compilation and execution
- NVCC and NVRTC compilation
- TMA, mbarrier, TMEM, and tcgen05 instructions
- BF16 tcgen05 GEMM correctness
- Dynamic shared-memory launches above 48 KiB
---
python/tvm/backend/cuda/target_tags.py | 18 ++++++++++++++++++
tests/python/target/test_target_target.py | 15 +++++++++++++++
2 files changed, 33 insertions(+)
diff --git a/python/tvm/backend/cuda/target_tags.py
b/python/tvm/backend/cuda/target_tags.py
index 87c83b378f..7d9033a836 100644
--- a/python/tvm/backend/cuda/target_tags.py
+++ b/python/tvm/backend/cuda/target_tags.py
@@ -372,3 +372,21 @@ _register_jetson_tag("nvidia/jetson-agx-xavier", "sm_72",
"carmel", 8)
_register_jetson_tag("nvidia/jetson-orin-nano", "sm_87", "carmel", 6)
_register_jetson_tag("nvidia/jetson-agx-orin-32gb", "sm_87", "cortex-a78", 8)
_register_jetson_tag("nvidia/jetson-agx-orin-64gb", "sm_87", "cortex-a78", 12)
+register_tag(
+ "nvidia/jetson-agx-thor",
+ {
+ "kind": "cuda",
+ "arch": "sm_110a",
+ "max_shared_memory_per_block": 232448,
+ "max_threads_per_block": 1024,
+ "thread_warp_size": 32,
+ "registers_per_block": 65536,
+ "l2_cache_size_bytes": 33554432,
+ "host": {
+ "kind": "llvm",
+ "mtriple": "aarch64-linux-gnu",
+ "mcpu": "neoverse-v3ae",
+ "num-cores": 14,
+ },
+ },
+)
diff --git a/tests/python/target/test_target_target.py
b/tests/python/target/test_target_target.py
index 5f264e23b9..ed0b4d0142 100644
--- a/tests/python/target/test_target_target.py
+++ b/tests/python/target/test_target_target.py
@@ -178,6 +178,21 @@ def test_target_tag_1():
assert tgt.attrs["registers_per_block"] == 32768
+def test_target_tag_jetson_agx_thor():
+ tgt = tvm.target.Target("nvidia/jetson-agx-thor")
+ assert tgt.kind.name == "cuda"
+ assert tgt.attrs["arch"] == "sm_110a"
+ assert tgt.attrs["max_shared_memory_per_block"] == 232448
+ assert tgt.attrs["max_threads_per_block"] == 1024
+ assert tgt.attrs["thread_warp_size"] == 32
+ assert tgt.attrs["registers_per_block"] == 65536
+ assert tgt.attrs["l2_cache_size_bytes"] == 33554432
+ assert tgt.host.kind.name == "llvm"
+ assert tgt.host.attrs["mtriple"] == "aarch64-linux-gnu"
+ assert tgt.host.attrs["mcpu"] == "neoverse-v3ae"
+ assert tgt.host.attrs["num-cores"] == 14
+
+
def test_target_tag_override():
"""Test creating a target from a tag with attribute overrides."""
tgt = tvm.target.Target({"tag": "nvidia/nvidia-a100",
"l2_cache_size_bytes": 12345})
