[SYCL][UR] Add max_threads_per_compute_unit device query

sycl/include/sycl/ext/oneapi/info/device.hpp

@@ -23,6 +23,13 @@ struct num_compute_units
   using return_type = size_t;
 };
 
+struct max_threads_per_compute_unit
+    : sycl::detail::ur_traits_base<
+          sycl::detail::info_class::device,
+          UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT> {
+  using return_type = size_t;
+};
+
 } // namespace ext::oneapi::info::device
 } // namespace _V1
 } // namespace sycl

sycl/source/detail/device_impl.hpp

@@ -1006,6 +1006,10 @@ class device_impl {
       return static_cast<size_t>(
           get_info_impl<UR_DEVICE_INFO_NUM_COMPUTE_UNITS>());
     }
+    CASE(ext::oneapi::info::device::max_threads_per_compute_unit) {
+      return static_cast<size_t>(
+          get_info_impl<UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT>());
+    }
 
     // ext::intel device traits (defined under sycl/ext/intel/info/device.hpp).
 

sycl/source/detail/ur_device_info_ret_types.inc

@@ -167,6 +167,7 @@ MAP(UR_DEVICE_INFO_XE_CLUSTERS_PER_REGION, uint32_t)
 MAP(UR_DEVICE_INFO_XE_CORES_PER_CLUSTER, uint32_t)
 MAP(UR_DEVICE_INFO_EUS_PER_XE_CORE, uint32_t)
 MAP(UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD, uint32_t)
+MAP(UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT, uint32_t)
 
 // These aren't present in the specification, extracted from ur_api.h
 // instead.

sycl/source/device.cpp

@@ -391,6 +391,8 @@ __SYCL_ONEAPI_DEVICE_INST(ext::oneapi::experimental::info::device,
 __SYCL_ONEAPI_DEVICE_INST(ext::oneapi::experimental::info::device,
                           composite_device, sycl::device)
 __SYCL_ONEAPI_DEVICE_INST(ext::oneapi::info::device, num_compute_units, size_t)
+__SYCL_ONEAPI_DEVICE_INST(ext::oneapi::info::device,
+                          max_threads_per_compute_unit, size_t)
 #undef __SYCL_ONEAPI_DEVICE_INST
 
 #define __SYCL_ONEAPI_PROGRESS_INST(NAME, SCOPE)                               \

sycl/test/abi/sycl_symbols_linux.dump

@@ -3848,6 +3848,7 @@ _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31w
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device32work_group_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi4info6device17num_compute_unitsEEENS0_6detail11ABINeutralTINS8_19is_device_info_descIT_E11return_typeEE4typeEv
+_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi4info6device28max_threads_per_compute_unitEEENS0_6detail11ABINeutralTINS8_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext8codeplay12experimental4info6device28max_registers_per_work_groupEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device10extensionsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device11device_typeEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv

sycl/test/abi/sycl_symbols_windows.dump

@@ -151,6 +151,7 @@
 ??$get_info_impl@Umax_read_image_args@device@info@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ
 ??$get_info_impl@Umax_registers_per_work_group@device@info@experimental@codeplay@ext@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ
 ??$get_info_impl@Umax_samplers@device@info@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ
+??$get_info_impl@Umax_threads_per_compute_unit@device@info@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ
 ??$get_info_impl@Umax_work_group_size@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ
 ??$get_info_impl@Umax_work_item_dimensions@device@info@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ
 ??$get_info_impl@Umax_write_image_args@device@info@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ

unified-runtime/include/unified-runtime/ur_api.h

@@ -2427,6 +2427,12 @@ typedef enum ur_device_info_t {
   /// [uint32_t][optional-query] return Intel GPU maximal number of lanes
   /// (virtual SIMD size) per hardware thread
   UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD = 139,
+  /// [uint32_t][optional-query] the maximum number of work-items (threads)
+  /// that can be resident on a single compute unit (CUDA/HIP streaming
+  /// multiprocessor). Corresponds to
+  /// CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR /
+  /// hipDeviceAttributeMaxThreadsPerMultiProcessor.
+  UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT = 140,
   /// [::ur_bool_t] Returns true if the device supports the use of
   /// command-buffers.
   UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP = 0x1000,

unified-runtime/include/unified-runtime/ur_print.hpp

@@ -3265,6 +3265,9 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_device_info_t value) {
   case UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD:
     os << "UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD";
     break;
+  case UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT:
+    os << "UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT";
+    break;
   case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP:
     os << "UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP";
     break;
@@ -5208,6 +5211,19 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr,
 
     os << ")";
   } break;
+  case UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT: {
+    const uint32_t *tptr = (const uint32_t *)ptr;
+    if (sizeof(uint32_t) > size) {
+      os << "invalid size (is: " << size << ", expected: >=" << sizeof(uint32_t)
+         << ")";
+      return UR_RESULT_ERROR_INVALID_SIZE;
+    }
+    os << (const void *)(tptr) << " (";
+
+    os << *tptr;
+
+    os << ")";
+  } break;
   case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: {
     const ur_bool_t *tptr = (const ur_bool_t *)ptr;
     if (sizeof(ur_bool_t) > size) {

unified-runtime/scripts/core/device.yml

@@ -487,6 +487,8 @@ etors:
       desc: "[uint32_t][optional-query] return Intel GPU number of execution engines (EUs) per XE Core"
     - name: MAX_LANES_PER_HW_THREAD
       desc: "[uint32_t][optional-query] return Intel GPU maximal number of lanes (virtual SIMD size) per hardware thread"
+    - name: MAX_THREADS_PER_COMPUTE_UNIT
+      desc: "[uint32_t][optional-query] the maximum number of work-items (threads) that can be resident on a single compute unit (CUDA/HIP streaming multiprocessor). Corresponds to CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR / hipDeviceAttributeMaxThreadsPerMultiProcessor."
 --- #--------------------------------------------------------------------------
 type: function
 desc: "Retrieves various information about device"

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

@@ -60,6 +60,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_MAX_COMPUTE_UNITS: {
     return ReturnValue(hDevice->getNumComputeUnits());
   }
+  case UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT: {
+    int MaxThreads = 0;
+    UR_CHECK_ERROR(cuDeviceGetAttribute(
+        &MaxThreads, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR,
+        hDevice->get()));
+    assert(MaxThreads >= 0);
+    return ReturnValue(static_cast<uint32_t>(MaxThreads));
+  }
   case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: {
     return ReturnValue(MaxWorkItemDimensions);
   }

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

@@ -66,6 +66,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
     assert(ComputeUnits >= 0);
     return ReturnValue(static_cast<uint32_t>(ComputeUnits));
   }
+  case UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT: {
+    int MaxThreads = 0;
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &MaxThreads, hipDeviceAttributeMaxThreadsPerMultiProcessor,
+        hDevice->get()));
+    assert(MaxThreads >= 0);
+    return ReturnValue(static_cast<uint32_t>(MaxThreads));
+  }
   case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: {
     return ReturnValue(MaxWorkItemDimensions);
   }