Add target device arch string in device properties

Signed-off-by: Ettore Tiotto <[email protected]>

Author: etiotto

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()

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

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):