[sycl-rel] Cherry-pick patches for sycl_ext_intel_device_info (#18886)

Cherry-pick commits that reached the internal branch between intel/llvm
cutoff and release branch pulldown.

Patches included:

---

[SYCL][Doc] Add new device descriptors to sycl_ext_intel_device_info extension (#17386)

Patch-by: Artur Gainullin <[email protected]>

---

[SYCL][CUDA] Add implementation of new device descriptors (#17590)

Tests were not added because there are existing conformance tests which
cover this functionality.

Patch-by: Artur Gainullin <[email protected]>

---

[UR][CUDA] Use different throttle reasons API based on CUDA version (#17719)

Our pre-commit CI uses CUDA 12.6 but nightly uses CUDA 12.1, it turns
out nvml which is part of CUDA 12.6 has
nvmlDeviceGetCurrentClocksEventReasons API, but nvml which is part of
CUDA 12.1 doesn't have it, but only supports older deprecated
nvmlDeviceGetCurrentClocksThrottleReasons.

NVML doesn't provide a version macro to check the support for that API,
so use new API nvmlDeviceGetCurrentClocksEventReasons based on cuda
version.

Patch-by: Artur Gainullin <[email protected]>

Author: KornevNikita

sycl/doc/extensions/supported/sycl_ext_intel_device_info.md

@@ -19,6 +19,7 @@ The Feature Test Macro SYCL\_EXT\_INTEL\_DEVICE\_INFO will be defined as one of
 | 4     | Free device memory query is supported |
 | 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 |
 
 
 
@@ -489,6 +490,144 @@ Then the memory bus width can be obtained using the standard get\_info() interfa
       auto MemoryBusWidth = dev.get_info<ext::intel::info::device::memory_bus_width>();
     }
 
+# Throttle reason #
+
+A new device descriptor is added which provides the current clock throttle reasons.
+A new enum is added with the list of possible throttle reasons.
+
+## Version ##
+
+The extension supports this query in version 7 and later.
+
+## Throttle reasons ##
+
+| Reason             | Description |
+| ------------------ | ----------- |
+| `power_cap` | The clock frequency is throttled due to hitting the power limit. |
+| `current_limit` | The clock frequency is throttled due to hitting the current limit. |
+| `thermal_limit` | The clock frequency is throttled due to hitting the thermal limit. |
+| `psu_alert` | The clock frequency is throttled due to power supply assertion. |
+| `sw_range` | The clock frequency is throttled due to software supplied frequency range. |
+| `hw_range` | The clock frequency is throttled because there is a sub block that has a lower frequency when it receives clocks. |
+| `other` | The clock frequency is throttled due to other reason. |
+
+
+```
+namespace sycl::ext::intel {
+
+  enum class throttle_reason {
+    power_cap,
+    current_limit,
+    thermal_limit,
+    psu_alert,
+    sw_range,
+    hw_range,
+    other
+  }
+
+}
+```
+
+## Device Information Descriptors ##
+
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+| `ext::intel::info::device::current_clock_throttle_reasons` | `std::vector<ext::intel::throttle_reason>` | Returns the set of throttle reasons describing why the frequency is being limited by the hardware. Returns empty set if frequency is not throttled. |
+
+
+## Aspects ##
+
+A new aspect, `ext_intel_current_clock_throttle_reasons`, 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_current_clock_throttle_reasons`.
+
+## Example Usage ##
+
+Then the current clock throttle reasons can be obtained using the standard `get_info()` interface.
+
+```
+if (dev.has(aspect::ext_intel_current_clock_throttle_reasons)) {
+  std::vector<ext::inte::info::throttle_reason> Reasons = dev.get_info<ext::intel::info::device::current_clock_throttle_reasons<>();
+}
+```
+
+
+# Fan speed #
+
+A new device descriptor is added which provides the fan speed for the device.
+
+## Version ##
+
+The extension supports this query in version 7 and later.
+
+## Device Information Descriptors ##
+
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+| `ext::intel::info::device::fan_speed` | `int32_t` | Returns the current speed of device's fan (as a percentage of the maximum speed of the fan). If fan speed can't be measured then returns -1. If there are multiple fans, then returns maximum value. |
+
+
+## Aspects ##
+
+A new aspect, `ext_intel_fan_speed`, 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_fan_speed`.
+
+## Example Usage ##
+
+Then the fan speed can be obtained using the standard `get_info()` interface.
+
+```
+    if (dev.has(aspect::ext_intel_fan_speed)) {
+      auto FanSpeed = dev.get_info<ext::intel::info::device::fan_speed>();
+    }
+```
+
+# Power limits #
+
+New device descriptors are added which provide the maximum and minimum power limits for the device.
+
+## Version ##
+
+The extension supports this query in version 7 and later.
+
+## Device Information Descriptors ##
+
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+|`ext::intel::info::device::min_power_limit` |`int32_t` | Returns the minimum power limit of the device in milliwatts. Returns -1 if the limit is not known. |
+|`ext::intel::info::device::max_power_limit` |`int32_t` | Returns the maximum power limit of the device in milliwatts. Returns -1 if the limit is not known. |
+
+
+## Aspects ##
+
+A new aspect, `ext_intel_power_limits`, 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_power_limits`.
+
+## Example Usage ##
+
+Then the power limits can be obtained using the standard `get_info()` interface.
+
+```
+    if (dev.has(aspect::ext_intel_power_limits)) {
+      auto Min = dev.get_info<ext::intel::info::device::min_power_limit>();
+      auto Max = dev.get_info<ext::intel::info::device::max_power_limit>();
+    }
+```
+
+
+
+
 # Deprecated queries #
 
 The table below lists deprecated, that would soon be removed and their replacements:

unified-runtime/source/adapters/cuda/CMakeLists.txt

@@ -119,6 +119,7 @@ target_link_libraries(${TARGET_NAME} PRIVATE
     ${PROJECT_NAME}::umf
     Threads::Threads
     cudadrv
+    CUDA::nvml
 )
 
 target_include_directories(${TARGET_NAME} PRIVATE

unified-runtime/source/adapters/cuda/common.cpp

@@ -12,6 +12,7 @@
 #include "logger/ur_logger.hpp"
 
 #include <cuda.h>
+#include <nvml.h>
 
 #include <sstream>
 
@@ -36,6 +37,23 @@ ur_result_t mapErrorUR(CUresult Result) {
   }
 }
 
+ur_result_t mapErrorUR(nvmlReturn_t Result) {
+  switch (Result) {
+  case NVML_SUCCESS:
+    return UR_RESULT_SUCCESS;
+  case NVML_ERROR_NOT_SUPPORTED:
+    return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
+  case NVML_ERROR_GPU_IS_LOST:
+    return UR_RESULT_ERROR_DEVICE_LOST;
+  case NVML_ERROR_MEMORY:
+    return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY;
+  case NVML_ERROR_INSUFFICIENT_RESOURCES:
+    return UR_RESULT_ERROR_OUT_OF_RESOURCES;
+  default:
+    return UR_RESULT_ERROR_UNKNOWN;
+  }
+}
+
 void checkErrorUR(CUresult Result, const char *Function, int Line,
                   const char *File) {
   if (Result == CUDA_SUCCESS || Result == CUDA_ERROR_DEINITIALIZED) {
@@ -63,6 +81,30 @@ void checkErrorUR(CUresult Result, const char *Function, int Line,
   throw mapErrorUR(Result);
 }
 
+void checkErrorUR(nvmlReturn_t Result, const char *Function, int Line,
+                  const char *File) {
+  if (Result == NVML_SUCCESS) {
+    return;
+  }
+
+  const char *ErrorString = nullptr;
+  ErrorString = nvmlErrorString(Result);
+  std::stringstream SS;
+  SS << "\nUR NVML ERROR:"
+     << "\n\tValue:           " << Result
+     << "\n\tDescription:     " << ErrorString
+     << "\n\tFunction:        " << Function << "\n\tSource Location: " << File
+     << ":" << Line << "\n";
+  logger::error("{}", SS.str());
+
+  if (std::getenv("PI_CUDA_ABORT") != nullptr ||
+      std::getenv("UR_CUDA_ABORT") != nullptr) {
+    std::abort();
+  }
+
+  throw mapErrorUR(Result);
+}
+
 void checkErrorUR(ur_result_t Result, const char *Function, int Line,
                   const char *File) {
   if (Result == UR_RESULT_SUCCESS) {

unified-runtime/source/adapters/cuda/common.hpp

@@ -10,6 +10,7 @@
 #pragma once
 
 #include <cuda.h>
+#include <nvml.h>
 #include <ur/ur.hpp>
 
 #include <umf/base.h>
@@ -35,6 +36,9 @@ ur_result_t mapErrorUR(CUresult Result);
 void checkErrorUR(CUresult Result, const char *Function, int Line,
                   const char *File);
 
+void checkErrorUR(nvmlReturn_t Result, const char *Function, int Line,
+                  const char *File);
+
 void checkErrorUR(ur_result_t Result, const char *Function, int Line,
                   const char *File);
 

unified-runtime/source/adapters/cuda/device.cpp

@@ -18,6 +18,7 @@
 #include "logger/ur_logger.hpp"
 #include "platform.hpp"
 #include "ur_util.hpp"
+#include <nvml.h>
 
 int getAttribute(ur_device_handle_t device, CUdevice_attribute attribute) {
   int value;
@@ -1085,11 +1086,69 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE:
   case UR_DEVICE_INFO_GPU_HW_THREADS_PER_EU:
   case UR_DEVICE_INFO_IP_VERSION:
-  case UR_DEVICE_INFO_CURRENT_CLOCK_THROTTLE_REASONS:
-  case UR_DEVICE_INFO_FAN_SPEED:
-  case UR_DEVICE_INFO_MIN_POWER_LIMIT:
-  case UR_DEVICE_INFO_MAX_POWER_LIMIT:
     return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
+  case UR_DEVICE_INFO_CURRENT_CLOCK_THROTTLE_REASONS: {
+    unsigned long long ClocksEventReasons;
+#if (CUDA_VERSION >= 12060)
+    UR_CHECK_ERROR(nvmlDeviceGetCurrentClocksEventReasons(hDevice->getNVML(),
+                                                          &ClocksEventReasons));
+#else
+    UR_CHECK_ERROR(nvmlDeviceGetCurrentClocksThrottleReasons(
+        hDevice->getNVML(), &ClocksEventReasons));
+#endif
+    ur_device_throttle_reasons_flags_t ThrottleReasons = 0;
+    constexpr unsigned long long NVMLThrottleFlags[] = {
+        nvmlClocksThrottleReasonSwPowerCap,
+        nvmlClocksThrottleReasonHwThermalSlowdown ||
+            nvmlClocksThrottleReasonSwThermalSlowdown,
+        nvmlClocksThrottleReasonHwPowerBrakeSlowdown,
+        nvmlClocksThrottleReasonApplicationsClocksSetting};
+
+    constexpr ur_device_throttle_reasons_flags_t UrThrottleFlags[] = {
+        UR_DEVICE_THROTTLE_REASONS_FLAG_POWER_CAP,
+        UR_DEVICE_THROTTLE_REASONS_FLAG_THERMAL_LIMIT,
+        UR_DEVICE_THROTTLE_REASONS_FLAG_PSU_ALERT,
+        UR_DEVICE_THROTTLE_REASONS_FLAG_SW_RANGE};
+
+    for (size_t i = 0;
+         i < sizeof(NVMLThrottleFlags) / sizeof(NVMLThrottleFlags[0]); ++i) {
+      if (ClocksEventReasons & NVMLThrottleFlags[i]) {
+        ThrottleReasons |= UrThrottleFlags[i];
+        ClocksEventReasons &= ~NVMLThrottleFlags[i];
+      }
+    }
+    if (ClocksEventReasons) {
+      ThrottleReasons |= UR_DEVICE_THROTTLE_REASONS_FLAG_OTHER;
+    }
+    return ReturnValue(ThrottleReasons);
+  }
+  case UR_DEVICE_INFO_MIN_POWER_LIMIT:
+  case UR_DEVICE_INFO_MAX_POWER_LIMIT: {
+    unsigned int minLimit, maxLimit;
+    auto NVMLHandle = hDevice->getNVML();
+    auto NVMLError = nvmlDeviceGetPowerManagementLimitConstraints(
+        NVMLHandle, &minLimit, &maxLimit);
+    if (NVMLError == NVML_ERROR_NOT_SUPPORTED) {
+      if (propName == UR_DEVICE_INFO_MAX_POWER_LIMIT) {
+        UR_CHECK_ERROR(
+            nvmlDeviceGetPowerManagementLimit(NVMLHandle, &maxLimit));
+        return ReturnValue(static_cast<int32_t>(maxLimit));
+      } else if (propName == UR_DEVICE_INFO_MIN_POWER_LIMIT) {
+        return ReturnValue(static_cast<int32_t>(-1));
+      }
+    }
+    if (propName == UR_DEVICE_INFO_MAX_POWER_LIMIT) {
+      return ReturnValue(static_cast<int32_t>(maxLimit));
+    } else if (propName == UR_DEVICE_INFO_MIN_POWER_LIMIT) {
+      return ReturnValue(static_cast<int32_t>(minLimit));
+    }
+    break;
+  }
+  case UR_DEVICE_INFO_FAN_SPEED: {
+    unsigned int Speed;
+    UR_CHECK_ERROR(nvmlDeviceGetFanSpeed(hDevice->getNVML(), &Speed));
+    return ReturnValue(static_cast<int32_t>(Speed));
+  }
   case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP:
     return ReturnValue(
         static_cast<ur_exp_device_2d_block_array_capability_flags_t>(0));

unified-runtime/source/adapters/cuda/device.hpp

@@ -36,13 +36,14 @@ struct ur_device_handle_t_ {
   int MaxChosenLocalMem{0};
   bool MaxLocalMemSizeChosen{false};
   uint32_t NumComputeUnits{0};
+  std::once_flag NVMLInitFlag;
+  std::optional<nvmlDevice_t> NVMLDevice;
 
 public:
   ur_device_handle_t_(native_type cuDevice, CUcontext cuContext, CUevent evBase,
                       ur_platform_handle_t platform, uint32_t DevIndex)
       : CuDevice(cuDevice), CuContext(cuContext), EvBase(evBase), RefCount{1},
         Platform(platform), DeviceIndex{DevIndex} {
-
     UR_CHECK_ERROR(cuDeviceGetAttribute(
         &MaxRegsPerBlock, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK,
         cuDevice));
@@ -102,11 +103,28 @@ struct ur_device_handle_t_ {
     if (MemoryProviderShared) {
       umfMemoryProviderDestroy(MemoryProviderShared);
     }
+    if (NVMLDevice.has_value()) {
+      UR_CHECK_ERROR(nvmlShutdown());
+    }
     cuDevicePrimaryCtxRelease(CuDevice);
   }
 
   native_type get() const noexcept { return CuDevice; };
 
+  nvmlDevice_t getNVML() {
+    // Initialization happens lazily once per device object. Call to nvmlInit by
+    // different objects will just increase the reference count. Each object's
+    // destructor calls shutdown method, so once there will be no NVML users
+    // left, resources will be released.
+    std::call_once(NVMLInitFlag, [this]() {
+      UR_CHECK_ERROR(nvmlInit());
+      nvmlDevice_t Handle;
+      UR_CHECK_ERROR(nvmlDeviceGetHandleByIndex(DeviceIndex, &Handle));
+      NVMLDevice = Handle;
+    });
+    return NVMLDevice.value();
+  };
+
   CUcontext getNativeContext() const noexcept { return CuContext; };
 
   uint32_t getReferenceCount() const noexcept { return RefCount; }

unified-runtime/test/adapters/cuda/CMakeLists.txt

@@ -29,4 +29,6 @@ target_include_directories(test-adapter-cuda PRIVATE
     ${PROJECT_SOURCE_DIR}/source/adapters/cuda
 )
 
-target_link_libraries(test-adapter-cuda PRIVATE cudadrv ${PROJECT_NAME}::umf)
+find_package(CUDAToolkit 10.1 REQUIRED)
+
+target_link_libraries(test-adapter-cuda PRIVATE cudadrv CUDA::nvml ${PROJECT_NAME}::umf)