From ae056a1eafa7af2fc3cf09e55ba14efac793acb8 Mon Sep 17 00:00:00 2001 From: Ettore Tiotto Date: Fri, 5 Dec 2025 22:21:16 +0000 Subject: [PATCH 1/3] Add target device arch string in device properties Signed-off-by: Ettore Tiotto --- python/triton/_internal_testing.py | 77 +++++++++++++++++++++++++ third_party/intel/backend/arch_parser.c | 2 +- third_party/intel/backend/driver.py | 10 ++++ 3 files changed, 88 insertions(+), 1 deletion(-) diff --git a/python/triton/_internal_testing.py b/python/triton/_internal_testing.py index 562e8df721..6ed2505f8b 100644 --- a/python/triton/_internal_testing.py +++ b/python/triton/_internal_testing.py @@ -105,6 +105,83 @@ def is_xpu(): target = get_current_target() return False if target is None else target.backend == "xpu" +def is_xpu_arl_h(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'arl_h' in target.arch + + +def is_xpu_arl_s(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'arl_s' in target.arch + + +def is_xpu_bmg(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'bmg' in target.arch + + +def is_xpu_dg2(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'dg2' in target.arch + + +def is_xpu_lnl(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'lnl' in target.arch + + +def is_xpu_mtl(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'mtl' in target.arch + + +def is_xpu_pvc(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'pvc' in target.arch + + +def is_xpu_ptl_h(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'ptl_h' in target.arch + + +def is_xpu_ptl_u(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'ptl_u' in target.arch + + +def is_xpu_cri(): + target = get_current_target() + return target is not None and target.backend == 'xpu' and 'cri' in target.arch + + +def is_xpu_lpg(): + return is_xpu_arl_s() + + +def is_xpu_lpgP(): + return is_xpu_arl_h() + + +def is_xpu_hpg(): + return is_xpu_dg2() + + +def is_xpu_hpc(): + return is_xpu_pvc() + + +def is_xpu_xe2(): + return is_xpu_bmg() + + +def is_xpu_xe3(): + return is_xpu_ptl_h() or is_xpu_ptl_u() + + +def is_xpu_xe3P(): + return is_xpu_cri() + def get_arch(): target = get_current_target() diff --git a/third_party/intel/backend/arch_parser.c b/third_party/intel/backend/arch_parser.c index eaa3846cb5..c9f00671bf 100644 --- a/third_party/intel/backend/arch_parser.c +++ b/third_party/intel/backend/arch_parser.c @@ -52,7 +52,7 @@ extern "C" EXPORT_FUNC const char *parse_device_arch(uint64_t dev_arch) { break; #endif case sycl::ext::oneapi::experimental::architecture::unknown: - std::cerr << "unknown sycl_arch" << std::endl; + std::cerr << "unknown" << std::endl; break; default: std::cerr << "sycl_arch not recognized: " << (uint64_t)sycl_arch diff --git a/third_party/intel/backend/driver.py b/third_party/intel/backend/driver.py index be59664675..bd259bac56 100644 --- a/third_party/intel/backend/driver.py +++ b/third_party/intel/backend/driver.py @@ -942,7 +942,17 @@ def update_advanced_features(device, dev_property): dev_property["has_subgroup_2d_block_io"] = check(device, b"cl_intel_subgroup_2d_block_io") dev_property["has_bfloat16_conversions"] = check(device, b"cl_intel_bfloat16_conversions") + def update_device_arch(dev_property): + if not (arch := knobs.intel.device_arch): + dirname = os.path.dirname(os.path.realpath(__file__)) + parser = compile_module_from_src(src=Path(os.path.join(dirname, "arch_parser.c")).read_text(), + name="arch_utils") + arch_name = parser.parse_device_arch(dev_property["architecture"]) + dev_property["arch"] = arch + update_advanced_features(device, dev_property) + update_device_arch(dev_property) + return GPUTarget("xpu", dev_property, warp_size=32) def build_proton_help_lib(self): From d0fa7dad800ae82be7801ee06cf9e74ed7b296e1 Mon Sep 17 00:00:00 2001 From: Ettore Tiotto Date: Mon, 8 Dec 2025 18:45:00 +0000 Subject: [PATCH 2/3] Address review comments Signed-off-by: Ettore Tiotto --- third_party/intel/backend/arch_parser.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/third_party/intel/backend/arch_parser.c b/third_party/intel/backend/arch_parser.c index c9f00671bf..814f1cab81 100644 --- a/third_party/intel/backend/arch_parser.c +++ b/third_party/intel/backend/arch_parser.c @@ -52,10 +52,10 @@ extern "C" EXPORT_FUNC const char *parse_device_arch(uint64_t dev_arch) { break; #endif case sycl::ext::oneapi::experimental::architecture::unknown: - std::cerr << "unknown" << std::endl; + arch = "unknown"; break; default: - std::cerr << "sycl_arch not recognized: " << (uint64_t)sycl_arch + std::cerr << "device architecture not recognized: " << (uint64_t)sycl_arch << std::endl; } From 774810a5edfb43918f7d932473e2dbdf0801bef9 Mon Sep 17 00:00:00 2001 From: Ettore Tiotto Date: Mon, 8 Dec 2025 19:12:15 +0000 Subject: [PATCH 3/3] Fix pre-commit Signed-off-by: Ettore Tiotto --- python/triton/_internal_testing.py | 1 + third_party/intel/backend/driver.py | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/python/triton/_internal_testing.py b/python/triton/_internal_testing.py index 6ed2505f8b..2e3d84ee54 100644 --- a/python/triton/_internal_testing.py +++ b/python/triton/_internal_testing.py @@ -105,6 +105,7 @@ def is_xpu(): target = get_current_target() return False if target is None else target.backend == "xpu" + def is_xpu_arl_h(): target = get_current_target() return target is not None and target.backend == 'xpu' and 'arl_h' in target.arch diff --git a/third_party/intel/backend/driver.py b/third_party/intel/backend/driver.py index bd259bac56..3ccde3116c 100644 --- a/third_party/intel/backend/driver.py +++ b/third_party/intel/backend/driver.py @@ -947,7 +947,7 @@ def update_device_arch(dev_property): dirname = os.path.dirname(os.path.realpath(__file__)) parser = compile_module_from_src(src=Path(os.path.join(dirname, "arch_parser.c")).read_text(), name="arch_utils") - arch_name = parser.parse_device_arch(dev_property["architecture"]) + arch = parser.parse_device_arch(dev_property["architecture"]) dev_property["arch"] = arch update_advanced_features(device, dev_property)