From 7abe2d635134ba4aeb5f456f09e2aa8229745fc2 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 05:56:09 -0700 Subject: [PATCH] [SYCL][UR] Add max_threads_per_compute_unit device query Expose the maximum number of work-items (threads) that can be resident on a single compute unit, which CUDA and HIP report as CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR and hipDeviceAttributeMaxThreadsPerMultiProcessor respectively. This complements the existing ext::oneapi::info::device::num_compute_units query and lets applications reason about per-SM/CU occupancy in a backend-portable way. Plumbing (mirrors num_compute_units): - UR: new optional-query UR_DEVICE_INFO_MAX_THREADS_PER_COMPUTE_UNIT (device.yml + generated ur_api.h/ur_print.hpp). - Adapters: implemented in the CUDA and HIP adapters; other backends fall through to the existing UNSUPPORTED_ENUMERATION default, as befits an optional query. - SYCL: new ext::oneapi::info::device::max_threads_per_compute_unit descriptor (return_type size_t), runtime dispatch in device_impl.hpp, return-type map entry, and explicit template instantiation in device.cpp. - Updated the Linux/Windows ABI symbol dumps for the new instantiation. Co-authored-by: Cursor --- sycl/include/sycl/ext/oneapi/info/device.hpp | 7 +++++++ sycl/source/detail/device_impl.hpp | 4 ++++ sycl/source/detail/ur_device_info_ret_types.inc | 1 + sycl/source/device.cpp | 2 ++ sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/abi/sycl_symbols_windows.dump | 1 + unified-runtime/include/unified-runtime/ur_api.h | 6 ++++++ .../include/unified-runtime/ur_print.hpp | 16 ++++++++++++++++ unified-runtime/scripts/core/device.yml | 2 ++ unified-runtime/source/adapters/cuda/device.cpp | 8 ++++++++ unified-runtime/source/adapters/hip/device.cpp | 8 ++++++++ 11 files changed, 56 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/info/device.hpp b/sycl/include/sycl/ext/oneapi/info/device.hpp index d4c6660482510..d8b32e7a0ee82 100644 --- a/sycl/include/sycl/ext/oneapi/info/device.hpp +++ b/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 diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 3715e0c0ae1ce..ed660685a7059 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1006,6 +1006,10 @@ class device_impl { return static_cast( get_info_impl()); } + CASE(ext::oneapi::info::device::max_threads_per_compute_unit) { + return static_cast( + get_info_impl()); + } // ext::intel device traits (defined under sycl/ext/intel/info/device.hpp). diff --git a/sycl/source/detail/ur_device_info_ret_types.inc b/sycl/source/detail/ur_device_info_ret_types.inc index 54e729dd8ff65..e7b88ddd0bcf6 100644 --- a/sycl/source/detail/ur_device_info_ret_types.inc +++ b/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. diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 1a8e6e9f199f4..9c73669c4c4ba 100644 --- a/sycl/source/device.cpp +++ b/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) \ diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index c4987b6733b7b..21700f886a8ff 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/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 diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 29870ad4f5e92..2de5f43cbae95 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/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 diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index fac4fcbc64e75..77f8f26de849d 100644 --- a/unified-runtime/include/unified-runtime/ur_api.h +++ b/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, diff --git a/unified-runtime/include/unified-runtime/ur_print.hpp b/unified-runtime/include/unified-runtime/ur_print.hpp index 10c01a2952e13..7408a731025a5 100644 --- a/unified-runtime/include/unified-runtime/ur_print.hpp +++ b/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) { diff --git a/unified-runtime/scripts/core/device.yml b/unified-runtime/scripts/core/device.yml index 0f183a6576b58..532997b876493 100644 --- a/unified-runtime/scripts/core/device.yml +++ b/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" diff --git a/unified-runtime/source/adapters/cuda/device.cpp b/unified-runtime/source/adapters/cuda/device.cpp index 452e55984339e..27e772103ce99 100644 --- a/unified-runtime/source/adapters/cuda/device.cpp +++ b/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(MaxThreads)); + } case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: { return ReturnValue(MaxWorkItemDimensions); } diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index 42ed5cbeaae2c..f3b0cbe5f2880 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/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(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(MaxThreads)); + } case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: { return ReturnValue(MaxWorkItemDimensions); }