[SYCL] Add support for getting device LUID on windows

DBDuncan · DBDuncan · commit b49a6c6c313d · 2025-07-08T16:54:55.000+01:00
Adds a new aspect to get device LUID. This feature is only available on Windows and allows for device matching when performing SYCL/DirectX interop
diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
@@ -92,6 +92,7 @@ def AspectExt_intel_current_clock_throttle_reasons : Aspect<"ext_intel_current_c
 def AspectExt_intel_fan_speed : Aspect<"ext_intel_fan_speed">;
 def AspectExt_intel_power_limits : Aspect<"ext_intel_power_limits">;
 def AspectExt_oneapi_async_memory_alloc : Aspect<"ext_oneapi_async_memory_alloc">;
+def AspectExt_intel_device_info_luid : Aspect<"ext_intel_device_info_luid">;
 
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
@@ -163,7 +164,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_intel_current_clock_throttle_reasons,
     AspectExt_intel_fan_speed,
     AspectExt_intel_power_limits,
-    AspectExt_oneapi_async_memory_alloc],
+    AspectExt_oneapi_async_memory_alloc,
+    AspectExt_intel_device_info_luid],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.
diff --git a/sycl/doc/extensions/supported/sycl_ext_intel_device_info.md b/sycl/doc/extensions/supported/sycl_ext_intel_device_info.md
@@ -20,6 +20,7 @@ The Feature Test Macro SYCL\_EXT\_INTEL\_DEVICE\_INFO will be defined as one of
 | 5     | Device ID is supported |
 | 6     | Memory clock rate and bus width queries are supported |
 | 7     | Throttle reasons, fan speed and power limits queries are supported |
+| 8     | Device LUID is supported |
 
 
 
@@ -626,7 +627,35 @@ Then the power limits can be obtained using the standard `get_info()` interface.
 ```
 
 
+# Device LUID #
 
+A new device descriptor is added which will provide the device Locally Unique ID (LUID).
+
+## Version ##
+
+The extension supports this query in version 8 and later.
+
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+|`ext::intel::info::device::luid` |`std::array<unsigned char, 8>` | Returns the device LUID. |
+
+## Aspects ##
+
+A new aspect, `ext_intel_device_info_luid`, is added.
+
+## Error Condition ##
+
+Throws a synchronous `exception` with the `errc::feature_not_supported` error code if the device does not have `aspect::ext_intel_device_info_luid`.
+
+## Example Usage ##
+
+The LUID can be obtained using the standard `get_info()` interface.
+
+```
+    if (dev.has(aspect::ext_intel_device_info_luid)) {
+      auto LUID = dev.get_info<ext::intel::info::device::luid>();
+    }
+```
 
 # Deprecated queries #
 
diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp
@@ -110,10 +110,11 @@ inline constexpr bool is_group_helper_v =
 } // namespace ext::oneapi::experimental
 
 namespace detail {
-// Type for Intel device UUID extension.
+// Types for Intel's device UUID and device LUID extension.
 // For details about this extension, see
 // sycl/doc/extensions/supported/sycl_ext_intel_device_info.md
 using uuid_type = std::array<unsigned char, 16>;
+using luid_type = std::array<unsigned char, 8>;
 
 template <typename T, typename R> struct copy_cv_qualifiers;
 
diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def
@@ -78,4 +78,5 @@ __SYCL_ASPECT(ext_intel_current_clock_throttle_reasons, 84)
 __SYCL_ASPECT(ext_intel_fan_speed, 85)
 __SYCL_ASPECT(ext_intel_power_limits, 86)
 __SYCL_ASPECT(ext_oneapi_async_memory_alloc, 87)
+__SYCL_ASPECT(ext_intel_device_info_luid, 88)
 
diff --git a/sycl/include/sycl/info/device_traits.def b/sycl/include/sycl/info/device_traits.def
@@ -237,6 +237,9 @@ __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_max_work_groups_3d, id<3>,
 __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_max_global_work_groups, size_t, __SYCL_TRAIT_HANDLED_IN_RT)
 __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_cuda_cluster_group, bool, __SYCL_TRAIT_HANDLED_IN_RT)
 
+__SYCL_PARAM_TRAITS_SPEC(device, ext_intel_device_info_luid, detail::luid_type,
+                         UR_DEVICE_INFO_LUID)
+
 #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
diff --git a/sycl/include/sycl/info/ext_intel_device_traits.def b/sycl/include/sycl/info/ext_intel_device_traits.def
@@ -21,6 +21,7 @@ __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, current_clock_throttle_reasons, std
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, fan_speed, int32_t, UR_DEVICE_INFO_FAN_SPEED)
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, min_power_limit, int32_t, UR_DEVICE_INFO_MIN_POWER_LIMIT)
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, max_power_limit, int32_t, UR_DEVICE_INFO_MAX_POWER_LIMIT)
+__SYCL_PARAM_TRAITS_SPEC(ext::intel, device, luid, detail::luid_type, UR_DEVICE_INFO_LUID)
 #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
@@ -510,6 +510,7 @@ EXPORT_GET_INFO(ext::intel::info::device::current_clock_throttle_reasons)
 EXPORT_GET_INFO(ext::intel::info::device::fan_speed)
 EXPORT_GET_INFO(ext::intel::info::device::min_power_limit)
 EXPORT_GET_INFO(ext::intel::info::device::max_power_limit)
+EXPORT_GET_INFO(ext::intel::info::device::luid)
 
 EXPORT_GET_INFO(ext::codeplay::experimental::info::device::supports_fusion)
 EXPORT_GET_INFO(ext::codeplay::experimental::info::device::max_registers_per_work_group)
diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp
@@ -1195,6 +1195,15 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
             "The device does not have the ext_intel_power_limits aspect");
       return get_info_impl<UR_DEVICE_INFO_MIN_POWER_LIMIT>();
     }
+    CASE(ext::intel::info::device::luid) {
+      if (!has(aspect::ext_intel_device_info_luid))
+        throw exception(
+            make_error_code(errc::feature_not_supported),
+            "The device does not have the ext_intel_device_info_luid aspect");
+      // TODO: we're essentially memcpy'ing here...
+      static_assert(std::is_same_v<luid_type, std::array<unsigned char, 8>>);
+      return get_info_impl<UR_DEVICE_INFO_LUID>();
+    }
     else {
       constexpr auto Desc = UrInfoCode<Param>::value;
       return static_cast<typename Param::return_type>(get_info_impl<Desc>());
@@ -1305,6 +1314,9 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     CASE(ext_intel_device_info_uuid) {
       return has_info_desc(UR_DEVICE_INFO_UUID);
     }
+    CASE(ext_intel_device_info_luid) {
+      return has_info_desc(UR_DEVICE_INFO_LUID);
+    }
     CASE(ext_intel_max_mem_bandwidth) {
       // currently not supported
       return false;
@@ -2313,6 +2325,7 @@ EXPORT_GET_INFO(ext::intel::info::device::current_clock_throttle_reasons)
 EXPORT_GET_INFO(ext::intel::info::device::fan_speed)
 EXPORT_GET_INFO(ext::intel::info::device::min_power_limit)
 EXPORT_GET_INFO(ext::intel::info::device::max_power_limit)
+EXPORT_GET_INFO(ext::intel::info::device::luid)
 
 EXPORT_GET_INFO(ext::codeplay::experimental::info::device::supports_fusion)
 EXPORT_GET_INFO(ext::codeplay::experimental::info::device::max_registers_per_work_group)
diff --git a/sycl/source/detail/ur_device_info_ret_types.inc b/sycl/source/detail/ur_device_info_ret_types.inc
@@ -129,6 +129,7 @@ MAP(UR_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES, ur_memory_order_capability_
 MAP(UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES, ur_memory_scope_capability_flags_t)
 MAP(UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES, ur_memory_order_capability_flags_t)
 MAP(UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES, ur_memory_scope_capability_flags_t)
+MAP(UR_DEVICE_INFO_LUID, std::array<uint8_t, 8>)
 // Deprecated, we're not using it, so comment out to avoid warnings:
 // MAP(UR_DEVICE_INFO_BFLOAT16, ur_bool_t)
 MAP(UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES, uint32_t)
diff --git a/sycl/test-e2e/Adapters/level_zero/luid.cpp b/sycl/test-e2e/Adapters/level_zero/luid.cpp
@@ -0,0 +1,54 @@
+// REQUIRES: aspect-ext_intel_device_info_luid
+// REQUIRES: gpu, level_zero, level_zero_dev_kit, windows
+
+// RUN: %{build} %level_zero_options -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %s
+
+// Test that the LUID is read correctly from Level Zero.
+
+// CHECK: PASSED
+#include <iomanip>
+#include <iostream>
+#include <level_zero/ze_api.h>
+#include <sstream>
+#include <sycl/backend.hpp>
+#include <sycl/detail/core.hpp>
+
+int main() {
+  sycl::device dev;
+  auto luid = dev.get_info<sycl::ext::intel::info::device::luid>();
+
+  std::stringstream luid_sycl;
+  for (int i = 0; i < luid.size(); ++i) {
+    luid_sycl << std::hex << std::setw(2) << std::setfill('0') << int(luid[i]);
+  }
+  std::cout << "SYCL: " << luid_sycl.str() << std::endl;
+
+  auto zedev = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(dev);
+  ze_device_properties_t device_properties{};
+  device_properties.stype = ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES;
+
+  ze_device_luid_ext_properties_t luid_device_properties{};
+  luid_device_properties.stype = ZE_STRUCTURE_TYPE_DEVICE_LUID_EXT_PROPERTIES;
+
+  device_properties.pNext = &luid_device_properties;
+
+  zeDeviceGetProperties(zedev, &device_properties);
+
+  ze_device_luid_ext_properties_t *luid_dev_prop =
+      static_cast<ze_device_luid_ext_properties_t *>(device_properties.pNext);
+
+  std::stringstream luid_l0;
+  for (int i = 0; i < ZE_MAX_DEVICE_LUID_SIZE_EXT; ++i)
+    luid_l0 << std::hex << std::setw(2) << std::setfill('0')
+            << int(luid_dev_prop->luid.id[i]);
+  std::cout << "L0  : " << luid_l0.str() << std::endl;
+
+  if (luid_sycl.str() != luid_l0.str()) {
+    std::cout << "FAILED" << std::endl;
+    return -1;
+  }
+
+  std::cout << "PASSED" << std::endl;
+  return 0;
+}
diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h
diff --git a/unified-runtime/include/ur_print.hpp b/unified-runtime/include/ur_print.hpp
diff --git a/unified-runtime/scripts/core/device.yml b/unified-runtime/scripts/core/device.yml
@@ -464,6 +464,8 @@ etors:
       desc: "[$x_bool_t] support for native bfloat16 conversions"
     - name: KERNEL_LAUNCH_CAPABILITIES
       desc: "[$x_kernel_launch_properties_flags_t] Bitfield of supported kernel launch properties."
+    - name: LUID
+      desc: "[uint8_t[]][optional-query] return device LUID"
 --- #--------------------------------------------------------------------------
 type: function
 desc: "Retrieves various information about device"
diff --git a/unified-runtime/source/adapters/cuda/device.cpp b/unified-runtime/source/adapters/cuda/device.cpp
@@ -1186,6 +1186,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
 
     return ReturnValue(0);
   }
+  case UR_DEVICE_INFO_LUID: {
+    // LUID is only available on Windows.
+    // Intel extension for device LUID. This returns the LUID as
+    // std::array<std::byte, 8>. For details about this extension,
+    // see sycl/doc/extensions/supported/sycl_ext_intel_device_info.md.
+    char *LUID = nullptr;
+    cuDeviceGetLuid(LUID, nullptr, hDevice->get());
+
+    std::array<unsigned char, 8> Name = {};
+
+    if (LUID == nullptr) {
+      return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    }
+
+    std::copy(LUID, LUID + 8, Name.begin());
+    return ReturnValue(Name.data(), 8);
+  }
   default:
     break;
   }
diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp
@@ -1346,6 +1346,26 @@ ur_result_t urDeviceGetInfo(
   }
   case UR_DEVICE_INFO_KERNEL_LAUNCH_CAPABILITIES:
     return ReturnValue(UR_KERNEL_LAUNCH_PROPERTIES_FLAG_COOPERATIVE);
+  case UR_DEVICE_INFO_LUID: {
+    // LUID is only available on Windows.
+    // Intel extension for device LUID. This returns the LUID as
+    // std::array<std::byte, 8>. For details about this extension,
+    // see sycl/doc/extensions/supported/sycl_ext_intel_device_info.md.
+    if (Device->Platform->ZeLUIDSupported) {
+      ze_device_properties_t DeviceProp = {};
+      DeviceProp.stype = ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES;
+      ze_device_luid_ext_properties_t LuidDesc = {};
+      LuidDesc.stype = ZE_STRUCTURE_TYPE_DEVICE_LUID_EXT_PROPERTIES;
+      DeviceProp.pNext = (void *)&LuidDesc;
+
+      ZE2UR_CALL(zeDeviceGetProperties, (ZeDevice, &DeviceProp));
+
+      const auto &LUID = LuidDesc.luid.id;
+      return ReturnValue(LUID, sizeof(LUID));
+    } else {
+      return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    }
+  }
   default:
     UR_LOG(ERR, "Unsupported ParamName in urGetDeviceInfo");
     UR_LOG(ERR, "ParamNameParamName={}(0x{})", ParamName,
diff --git a/unified-runtime/source/adapters/level_zero/platform.cpp b/unified-runtime/source/adapters/level_zero/platform.cpp
@@ -300,6 +300,12 @@ ur_result_t ur_platform_handle_t_::initialize() {
         ZeBindlessImagesExtensionSupported = true;
       }
     }
+    if (strncmp(extension.name, ZE_DEVICE_LUID_EXT_NAME,
+                strlen(ZE_DEVICE_LUID_EXT_NAME) + 1) == 0) {
+      if (extension.version == ZE_DEVICE_LUID_EXT_VERSION_1_0) {
+        ZeLUIDSupported = true;
+      }
+    }
     zeDriverExtensionMap[extension.name] = extension.version;
   }
 
diff --git a/unified-runtime/source/adapters/level_zero/platform.hpp b/unified-runtime/source/adapters/level_zero/platform.hpp
@@ -69,6 +69,7 @@ struct ur_platform_handle_t_ : ur::handle_base<ur::level_zero::ddi_getter>,
   bool ZeDriverEuCountExtensionFound{false};
   bool ZeCopyOffloadExtensionSupported{false};
   bool ZeBindlessImagesExtensionSupported{false};
+  bool ZeLUIDSupported{false};
 
   // Cache UR devices for reuse
   std::vector<std::unique_ptr<ur_device_handle_t_>> URDevicesCache;
diff --git a/unified-runtime/test/conformance/testing/include/uur/optional_queries.h b/unified-runtime/test/conformance/testing/include/uur/optional_queries.h
diff --git a/unified-runtime/tools/urinfo/urinfo.hpp b/unified-runtime/tools/urinfo/urinfo.hpp
