From 8e57c5e7bbbb06120f086c98cde823909500753d Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 11 Aug 2026 07:16:26 -0700 Subject: [PATCH 1/4] [SYCL] Implement sycl_khr_max_work_group_queries Signed-off-by: Hu, Peisen --- .../sycl/khr/max_work_group_queries.hpp | 51 +++++++++++++++++++ sycl/include/sycl/sycl.hpp | 1 + sycl/source/detail/device_impl.hpp | 24 +++++++++ sycl/source/feature_test.hpp.in | 1 + .../test-e2e/Basic/max_work_group_queries.cpp | 50 ++++++++++++++++++ sycl/test/basic_tests/Inputs/khr_all.hpp | 1 + 6 files changed, 128 insertions(+) create mode 100644 sycl/include/sycl/khr/max_work_group_queries.hpp create mode 100644 sycl/test-e2e/Basic/max_work_group_queries.cpp diff --git a/sycl/include/sycl/khr/max_work_group_queries.hpp b/sycl/include/sycl/khr/max_work_group_queries.hpp new file mode 100644 index 0000000000000..ae724e2e3e2ff --- /dev/null +++ b/sycl/include/sycl/khr/max_work_group_queries.hpp @@ -0,0 +1,51 @@ +//==-- max_work_group_queries.hpp - oneapi max work groups info traits ------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace khr::info::device { + +template struct max_work_group_range; + +// max_work_group_range<1> and <2> are RT-only; only <3> dispatches via UR. +template <> +struct max_work_group_range<1> + : sycl::detail::rt_traits_base { + using return_type = sycl::range<1>; +}; + +template <> +struct max_work_group_range<2> + : sycl::detail::rt_traits_base { + using return_type = sycl::range<2>; +}; + +template <> +struct max_work_group_range<3> + : sycl::detail::ur_traits_base { + using return_type = sycl::range<3>; +}; + +struct max_work_group_range_size + : sycl::detail::ur_traits_base { + using return_type = size_t; +}; + +} // namespace khr::info::device +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index d133824184470..ba320874c0e94 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -165,5 +165,6 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include +#include #include #include diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index ea45d51058753..02f1b1f723f24 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -26,6 +26,7 @@ #include #include #include +#include #include #include @@ -1219,6 +1220,29 @@ class device_impl { "ext_intel_max_lanes_per_hw_thread aspect"); return get_info_impl(); } + + // khr device traits (defined under sycl/ext/oneapi/...). + + CASE(khr::info::device::max_work_group_range_size) { + return get_info_impl(); + } + CASE(khr::info::device::max_work_group_range<3>) { + size_t result[3] = {}; + getAdapter().call( + getHandleRef(), UR_DEVICE_INFO_MAX_WORK_GROUPS_3D, sizeof(result), + &result, nullptr); + return range<3>(result[2], result[1], result[0]); + } + CASE(khr::info::device::max_work_group_range<2>) { + range<3> max_3d = + get_info, DependentFalse>(); + return range<2>{max_3d[1], max_3d[2]}; + } + CASE(khr::info::device::max_work_group_range<1>) { + range<3> max_3d = + get_info, DependentFalse>(); + return range<1>{max_3d[2]}; + } else { constexpr auto Desc = UrInfoCode::value; return static_cast(get_info_impl()); diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 6aebc773b514c..398db031b60b0 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -107,6 +107,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_CURRENT_DEVICE 1 #define SYCL_KHR_QUEUE_EMPTY_QUERY 1 #define SYCL_KHR_QUEUE_FLUSH 1 +#define SYCL_KHR_MAX_WORK_GROUP_QUERIES 1 #define SYCL_EXT_ONEAPI_MEMORY_EXPORT 1 #define SYCL_EXT_ONEAPI_CLOCK 1 #define SYCL_EXT_ONEAPI_DEVICE_IS_INTEGRATED_GPU 1 diff --git a/sycl/test-e2e/Basic/max_work_group_queries.cpp b/sycl/test-e2e/Basic/max_work_group_queries.cpp new file mode 100644 index 0000000000000..ae078ba03ddb6 --- /dev/null +++ b/sycl/test-e2e/Basic/max_work_group_queries.cpp @@ -0,0 +1,50 @@ +// RUN: %{build} -o %t.out +// REQUIRES: cuda || hip || level_zero +// RUN: %{run} %t.out + +#include +#include + +#include +#include + +using namespace sycl; + +int main() { + queue q; + device dev = q.get_device(); + +#if !defined(SYCL_KHR_MAX_WORK_GROUP_QUERIES) +#error SYCL_KHR_MAX_WORK_GROUP_QUERIES is not defined! +#endif + + sycl::id<1> groupD = dev.get_info< + sycl::khr::info::device::max_work_group_range<1>>(); + std::cout << "Max work group size in 1D \n"; + std::cout << "Dimension 1:" << groupD[0] << std::endl; + + sycl::id<2> group2D = dev.get_info< + sycl::khr::info::device::max_work_group_range<2>>(); + std::cout << "Max work group size in 2D \n"; + std::cout << "Dimension 1:" << group2D[0] << "\n" + << "Dimension 2:" << group2D[1] << std::endl; + + sycl::id<3> group3D = dev.get_info< + sycl::khr::info::device::max_work_group_range<3>>(); + std::cout << "Max work group size in 3D \n"; + std::cout << "Dimension 1:" << group3D[0] << "\n" + << "Dimension 2:" << group3D[1] << "\n" + << "Dimension 3:" << group3D[2] << std::endl; + + size_t group_max = dev.get_info< + sycl::khr::info::device::max_work_group_range_size>(); + std::cout << "Max global work group size:" << group_max << "\n"; + + assert((group3D[0] <= group_max && group3D[1] <= group_max && + group3D[2] <= group_max) && + "Max work-group size of each dimension must be smaller than " + "global work-group size"); + + std::cout << "Passed!" << std::endl; + return 0; +} diff --git a/sycl/test/basic_tests/Inputs/khr_all.hpp b/sycl/test/basic_tests/Inputs/khr_all.hpp index d94a19f8692ca..8abc6a239e880 100644 --- a/sycl/test/basic_tests/Inputs/khr_all.hpp +++ b/sycl/test/basic_tests/Inputs/khr_all.hpp @@ -21,6 +21,7 @@ #include #include #include +#include #include #include #include From 7f171efa603fe75f8b491f2a28ecf44f986fe708 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 11 Aug 2026 07:57:26 -0700 Subject: [PATCH 2/4] [SYCL] Fix formatting issue Signed-off-by: Hu, Peisen --- sycl/include/sycl/khr/max_work_group_queries.hpp | 5 +++-- sycl/source/detail/device_impl.hpp | 8 ++++---- sycl/test-e2e/Basic/max_work_group_queries.cpp | 16 ++++++++-------- 3 files changed, 15 insertions(+), 14 deletions(-) diff --git a/sycl/include/sycl/khr/max_work_group_queries.hpp b/sycl/include/sycl/khr/max_work_group_queries.hpp index ae724e2e3e2ff..b0488ea192fc2 100644 --- a/sycl/include/sycl/khr/max_work_group_queries.hpp +++ b/sycl/include/sycl/khr/max_work_group_queries.hpp @@ -1,4 +1,5 @@ -//==-- max_work_group_queries.hpp - oneapi max work groups info traits ------------==// +//==-- max_work_group_queries.hpp - oneapi max work groups info traits +//------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -18,7 +19,7 @@ namespace sycl { inline namespace _V1 { namespace khr::info::device { -template struct max_work_group_range; +template struct max_work_group_range; // max_work_group_range<1> and <2> are RT-only; only <3> dispatches via UR. template <> diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 02f1b1f723f24..2f967ac5f2ce7 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1234,13 +1234,13 @@ class device_impl { return range<3>(result[2], result[1], result[0]); } CASE(khr::info::device::max_work_group_range<2>) { - range<3> max_3d = - get_info, DependentFalse>(); + range<3> max_3d = get_info, + DependentFalse>(); return range<2>{max_3d[1], max_3d[2]}; } CASE(khr::info::device::max_work_group_range<1>) { - range<3> max_3d = - get_info, DependentFalse>(); + range<3> max_3d = get_info, + DependentFalse>(); return range<1>{max_3d[2]}; } else { diff --git a/sycl/test-e2e/Basic/max_work_group_queries.cpp b/sycl/test-e2e/Basic/max_work_group_queries.cpp index ae078ba03ddb6..0918990d915b1 100644 --- a/sycl/test-e2e/Basic/max_work_group_queries.cpp +++ b/sycl/test-e2e/Basic/max_work_group_queries.cpp @@ -18,26 +18,26 @@ int main() { #error SYCL_KHR_MAX_WORK_GROUP_QUERIES is not defined! #endif - sycl::id<1> groupD = dev.get_info< - sycl::khr::info::device::max_work_group_range<1>>(); + sycl::id<1> groupD = + dev.get_info>(); std::cout << "Max work group size in 1D \n"; std::cout << "Dimension 1:" << groupD[0] << std::endl; - sycl::id<2> group2D = dev.get_info< - sycl::khr::info::device::max_work_group_range<2>>(); + sycl::id<2> group2D = + dev.get_info>(); std::cout << "Max work group size in 2D \n"; std::cout << "Dimension 1:" << group2D[0] << "\n" << "Dimension 2:" << group2D[1] << std::endl; - sycl::id<3> group3D = dev.get_info< - sycl::khr::info::device::max_work_group_range<3>>(); + sycl::id<3> group3D = + dev.get_info>(); std::cout << "Max work group size in 3D \n"; std::cout << "Dimension 1:" << group3D[0] << "\n" << "Dimension 2:" << group3D[1] << "\n" << "Dimension 3:" << group3D[2] << std::endl; - size_t group_max = dev.get_info< - sycl::khr::info::device::max_work_group_range_size>(); + size_t group_max = + dev.get_info(); std::cout << "Max global work group size:" << group_max << "\n"; assert((group3D[0] <= group_max && group3D[1] <= group_max && From 2966287ba349a091afab8ad6e65bb5a20bb69753 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 13 Aug 2026 07:18:23 -0700 Subject: [PATCH 3/4] [SYCL] Fix test issue and improve OpenCL Implementation Signed-off-by: Hu, Peisen --- sycl/source/device.cpp | 9 +++++++++ unified-runtime/source/adapters/opencl/device.cpp | 14 ++++++++++++-- 2 files changed, 21 insertions(+), 2 deletions(-) diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 5d399abea7237..4ae0c689cc454 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -394,6 +394,15 @@ __SYCL_ONEAPI_DEVICE_INST(ext::oneapi::experimental::info::device, __SYCL_ONEAPI_DEVICE_INST(ext::oneapi::info::device, num_compute_units, size_t) #undef __SYCL_ONEAPI_DEVICE_INST +#define __SYCL_KHR_DEVICE_INST(NS, NAME, RETURN_T) \ + template __SYCL_EXPORT detail::ABINeutralT_t \ + device::get_info_impl() const; +__SYCL_KHR_DEVICE_INST(khr::info::device, max_work_group_range_size, size_t) +__SYCL_KHR_DEVICE_INST(khr::info::device, max_work_group_range<1>, range<1>) +__SYCL_KHR_DEVICE_INST(khr::info::device, max_work_group_range<2>, range<2>) +__SYCL_KHR_DEVICE_INST(khr::info::device, max_work_group_range<3>, range<3>) +#undef __SYCL_KHR_DEVICE_INST + #define __SYCL_ONEAPI_PROGRESS_INST(NAME, SCOPE) \ template __SYCL_EXPORT detail::ABINeutralT_t< \ std::vector> \ diff --git a/unified-runtime/source/adapters/opencl/device.cpp b/unified-runtime/source/adapters/opencl/device.cpp index acccf0fee01b0..4dce963b25e91 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -258,7 +258,12 @@ ur_result_t urDeviceGetInfo(ur_device_handle_t hDevice, * use to submit a kernel. There is no such query defined in OpenCL. So * we'll return the maximum value. */ static constexpr uint32_t MaxWorkItemDimensions = 3u; - static constexpr size_t Max = (std::numeric_limits::max)(); + + size_t Max = 0; + CL_RETURN_ON_FAILURE(clGetDeviceInfo(Device->CLDevice, + CL_DEVICE_MAX_WORK_GROUP_SIZE, + sizeof(Max), &Max, nullptr)); + assert(Max >= 0); struct { size_t sizes[MaxWorkItemDimensions]; @@ -270,7 +275,12 @@ ur_result_t urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(ReturnSizes); } case UR_DEVICE_INFO_MAX_WORK_GROUPS: { - return ReturnValue(std::numeric_limits::max()); + size_t Max = 0; + CL_RETURN_ON_FAILURE(clGetDeviceInfo(Device->CLDevice, + CL_DEVICE_MAX_WORK_GROUP_SIZE, + sizeof(Max), &Max, nullptr)); + assert(Max >= 0); + return ReturnValue(Max); } case UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES: { return ReturnValue(static_cast(1u)); From 2e4afaf5eaecd79734389df04aa0ebfff4630cb2 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 13 Aug 2026 09:19:04 -0700 Subject: [PATCH 4/4] [SYCL] Fix OpenCL UR Call Signed-off-by: Hu, Peisen --- .../source/adapters/opencl/device.cpp | 36 ++++++++++--------- 1 file changed, 20 insertions(+), 16 deletions(-) diff --git a/unified-runtime/source/adapters/opencl/device.cpp b/unified-runtime/source/adapters/opencl/device.cpp index 4dce963b25e91..73c8abfa37c7f 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -254,25 +254,29 @@ ur_result_t urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(URValue.data(), URValue.size()); } case UR_DEVICE_INFO_MAX_WORK_GROUPS_3D: { - /* Returns the maximum sizes of a work group for each dimension one could - * use to submit a kernel. There is no such query defined in OpenCL. So - * we'll return the maximum value. */ - static constexpr uint32_t MaxWorkItemDimensions = 3u; + constexpr size_t ReturnBufferSize = 3; + cl_uint MaxWorkItemDimensions = 0; - size_t Max = 0; - CL_RETURN_ON_FAILURE(clGetDeviceInfo(Device->CLDevice, - CL_DEVICE_MAX_WORK_GROUP_SIZE, - sizeof(Max), &Max, nullptr)); - assert(Max >= 0); + CL_RETURN_ON_FAILURE(clGetDeviceInfo( + Device->CLDevice, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, + sizeof(MaxWorkItemDimensions), &MaxWorkItemDimensions, nullptr)); + assert(MaxWorkItemDimensions >= 3); - struct { - size_t sizes[MaxWorkItemDimensions]; - } ReturnSizes; + size_t ReturnBuffer[ReturnBufferSize]; + size_t *ClCallBuffer = new size_t[MaxWorkItemDimensions]; - ReturnSizes.sizes[0] = Max; - ReturnSizes.sizes[1] = Max; - ReturnSizes.sizes[2] = Max; - return ReturnValue(ReturnSizes); + oclv::OpenCLVersion DevVer; + UR_RETURN_ON_FAILURE(Device->getDeviceVersion(DevVer)); + + CL_RETURN_ON_FAILURE( + clGetDeviceInfo(Device->CLDevice, CL_DEVICE_MAX_WORK_ITEM_SIZES, + sizeof(ClCallBuffer), &ClCallBuffer, nullptr)); + + ReturnBuffer[0] = ClCallBuffer[0]; + ReturnBuffer[1] = ClCallBuffer[1]; + ReturnBuffer[2] = ClCallBuffer[2]; + delete[] ClCallBuffer; + return ReturnValue(ReturnBuffer); } case UR_DEVICE_INFO_MAX_WORK_GROUPS: { size_t Max = 0;