[SYCL] Add support for getting device LUID on windows (#19349)

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

Author: DBDuncan

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -92,6 +92,8 @@ 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">;
+def AspectExt_intel_device_info_node_mask : Aspect<"ext_intel_device_info_node_mask">;
 
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
@@ -163,7 +165,9 @@ 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,
+    AspectExt_intel_device_info_node_mask],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

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 and device node mask is supported |
 
 
 
@@ -626,7 +627,66 @@ 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>();
+    }
+```
+
+
+# Device Node Mask #
+
+A new device descriptor is added which will provide the device node mask.
+
+## Version ##
+
+The extension supports this query in version 8 and later.
+
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+|`ext::intel::info::device::node_mask` |`unsigned int` | Returns the device node mask. |
+
+## Aspects ##
+
+A new aspect, `ext_intel_device_info_node_mask`, 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_node_mask`.
+
+## Example Usage ##
+
+The device node mask can be obtained using the standard `get_info()` interface.
+
+```
+    if (dev.has(aspect::ext_intel_device_info_node_mask)) {
+      auto node_mask = dev.get_info<ext::intel::info::device::node_mask>();
+    }
+```
 
 # Deprecated queries #
 

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;
 

sycl/include/sycl/info/aspects.def

@@ -78,4 +78,6 @@ __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)
+__SYCL_ASPECT(ext_intel_device_info_node_mask, 89)
 

sycl/include/sycl/info/ext_intel_device_traits.def

@@ -21,6 +21,8 @@ __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, __SYCL_TRAIT_HANDLED_IN_RT)
+__SYCL_PARAM_TRAITS_SPEC(ext::intel, device, node_mask, uint32_t, __SYCL_TRAIT_HANDLED_IN_RT)
 #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF

sycl/source/detail/device_impl.hpp

@@ -1194,6 +1194,20 @@ 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");
+      return get_info_impl<UR_DEVICE_INFO_LUID>();
+    }
+    CASE(ext::intel::info::device::node_mask) {
+      if (!has(aspect::ext_intel_device_info_node_mask))
+        throw exception(make_error_code(errc::feature_not_supported),
+                        "The device does not have the "
+                        "ext_intel_device_info_node_mask aspect");
+      return get_info_impl<UR_DEVICE_INFO_NODE_MASK>();
+    }
     else {
       constexpr auto Desc = UrInfoCode<Param>::value;
       return static_cast<typename Param::return_type>(get_info_impl<Desc>());
@@ -1304,6 +1318,12 @@ 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_device_info_node_mask) {
+      return has_info_desc(UR_DEVICE_INFO_NODE_MASK);
+    }
     CASE(ext_intel_max_mem_bandwidth) {
       // currently not supported
       return false;

sycl/source/detail/ur_device_info_ret_types.inc

@@ -154,6 +154,10 @@ MAP(UR_DEVICE_INFO_FAN_SPEED, int32_t)
 MAP(UR_DEVICE_INFO_MIN_POWER_LIMIT, int32_t)
 MAP(UR_DEVICE_INFO_MAX_POWER_LIMIT, int32_t)
 MAP(UR_DEVICE_INFO_BFLOAT16_CONVERSIONS_NATIVE, ur_bool_t)
+MAP(UR_DEVICE_INFO_KERNEL_LAUNCH_CAPABILITIES, ur_kernel_launch_properties_flags_t)
+// Manually changed std::vector<uint8_t> -> std::array<uint8_t, 8>
+MAP(UR_DEVICE_INFO_LUID, std::array<uint8_t, 8>)
+MAP(UR_DEVICE_INFO_NODE_MASK, uint32_t)
 
 // These aren't present in the specification, extracted from ur_api.h
 // instead.
@@ -187,5 +191,4 @@ MAP(UR_DEVICE_INFO_MIPMAP_LEVEL_REFERENCE_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY_EXP, uint32_t)
 MAP(UR_DEVICE_INFO_MIPMAP_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_TIMESTAMP_RECORDING_SUPPORT_EXP, ur_bool_t)
-MAP(UR_DEVICE_INFO_KERNEL_LAUNCH_CAPABILITIES, ur_kernel_launch_properties_flags_t)
 // clang-format on

sycl/test-e2e/Adapters/level_zero/luid.cpp

@@ -0,0 +1,53 @@
+// 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
+
+// Test that the LUID is read correctly from Level Zero.
+
+#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;
+}

sycl/test-e2e/Adapters/level_zero/node_mask.cpp

@@ -0,0 +1,47 @@
+// REQUIRES: aspect-ext_intel_device_info_node_mask
+// REQUIRES: gpu, level_zero, level_zero_dev_kit, windows
+
+// RUN: %{build} %level_zero_options -o %t.out
+// RUN: %{run} %t.out
+
+// Test that the node mask is read correctly from Level Zero.
+
+#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 nodeMaskSYCL = dev.get_info<sycl::ext::intel::info::device::node_mask>();
+
+  std::cout << "SYCL: " << nodeMaskSYCL << 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);
+
+  uint32_t nodeMaskL0 = luid_dev_prop->nodeMask;
+
+  std::cout << "L0  : " << nodeMaskL0 << std::endl;
+
+  if (nodeMaskSYCL != nodeMaskL0) {
+    std::cout << "FAILED" << std::endl;
+    return -1;
+  }
+
+  std::cout << "PASSED" << std::endl;
+  return 0;
+}

sycl/test-e2e/Adapters/luid-cuda.cpp

@@ -0,0 +1,51 @@
+// REQUIRES: aspect-ext_intel_device_info_luid
+// REQUIRES: gpu, target-nvidia, cuda_dev_kit, windows
+
+// RUN: %{build} %cuda_options -o %t.out
+// RUN: %{run} %t.out
+
+// Test that the LUID is read correctly from CUDA.
+
+#include <iomanip>
+#include <iostream>
+#include <sstream>
+#define SYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL 1
+#include <cuda.h>
+#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 luidSYCLHex;
+  for (int i = 0; i < luid.size(); ++i) {
+    luidSYCLHex << std::hex << std::setw(2) << std::setfill('0')
+                << int(luid[i]);
+  }
+  std::cout << "SYCL: " << luidSYCLHex.str() << std::endl;
+
+  CUdevice cudaDevice = sycl::get_native<sycl::backend::ext_oneapi_cuda>(dev);
+
+  std::array<char, 8> luidCuda{};
+
+  cuDeviceGetLuid(luidCuda.data(), nullptr, cudaDevice);
+
+  // Cuda returns luid as char, not unsigned char so convert that here.
+  std::array<unsigned char, 8> luidCudaConverted{};
+  std::copy(luidCuda.begin(), luidCuda.end(), luidCudaConverted.begin());
+
+  std::stringstream luidCudaHex;
+  for (int i = 0; i < 8; ++i)
+    luidCudaHex << std::hex << std::setw(2) << std::setfill('0')
+                << int(luidCudaConverted[i]);
+  std::cout << "CUDA  : " << luidCudaHex.str() << std::endl;
+
+  if (luidSYCLHex.str() != luidCudaHex.str()) {
+    std::cout << "FAILED" << std::endl;
+    return -1;
+  }
+
+  std::cout << "PASSED" << std::endl;
+  return 0;
+}