diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 72841724fa01d..5f8caf7f43c6d 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/0x12CC/unified-runtime.git") include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index e640fe9bdba6e..61ed630738768 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -4,4 +4,4 @@ # Date: Tue Nov 26 13:32:11 2024 +0000 # Merge pull request #2351 from nrspruit/mcl_1_1 # [L0] Add support for the MCL 1.1 apis thru the spec extensions -set(UNIFIED_RUNTIME_TAG e0f22b54080411537e8a12fe42f06dbac5e7a2d2) +set(UNIFIED_RUNTIME_TAG ca0a0e3e3b3ad933605a35850ebc11bbf5d5f367) diff --git a/sycl/include/sycl/info/ext_oneapi_kernel_queue_specific_traits.def b/sycl/include/sycl/info/ext_oneapi_kernel_queue_specific_traits.def index 0ec11af0bb6b1..f1c92ccd1bea1 100644 --- a/sycl/include/sycl/info/ext_oneapi_kernel_queue_specific_traits.def +++ b/sycl/include/sycl/info/ext_oneapi_kernel_queue_specific_traits.def @@ -1,4 +1 @@ -// TODO: Revisit 'max_num_work_group_sync' and align it with the -// 'sycl_ext_oneapi_forward_progress' extension once #7598 is merged. -__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, kernel_queue_specific, max_num_work_group_sync, size_t,) __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, kernel_queue_specific, max_num_work_groups, size_t,) diff --git a/sycl/include/sycl/kernel.hpp b/sycl/include/sycl/kernel.hpp index 654373c104c85..c0346ceee5610 100644 --- a/sycl/include/sycl/kernel.hpp +++ b/sycl/include/sycl/kernel.hpp @@ -168,6 +168,32 @@ class __SYCL_EXPORT kernel : public detail::OwnerLessBase { typename detail::is_kernel_queue_specific_info_desc::return_type ext_oneapi_get_info(queue Queue) const; + /// Query queue/launch-specific information from a kernel using the + /// info::kernel_queue_specific descriptor for a specific Queue and values. + /// max_num_work_groups is the only valid descriptor for this function. + /// + /// \param Queue is a valid SYCL queue. + /// \param WorkGroupSize is the work-group size the number of work-groups is + /// requested for. + /// \return depends on information being queried. + template + typename detail::is_kernel_queue_specific_info_desc::return_type + ext_oneapi_get_info(queue Queue, const range<1> &WorkGroupSize, + size_t DynamicLocalMemorySize) const; + + /// Query queue/launch-specific information from a kernel using the + /// info::kernel_queue_specific descriptor for a specific Queue and values. + /// max_num_work_groups is the only valid descriptor for this function. + /// + /// \param Queue is a valid SYCL queue. + /// \param WorkGroupSize is the work-group size the number of work-groups is + /// requested for. + /// \return depends on information being queried. + template + typename detail::is_kernel_queue_specific_info_desc::return_type + ext_oneapi_get_info(queue Queue, const range<2> &WorkGroupSize, + size_t DynamicLocalMemorySize) const; + /// Query queue/launch-specific information from a kernel using the /// info::kernel_queue_specific descriptor for a specific Queue and values. /// max_num_work_groups is the only valid descriptor for this function. diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp index 986e78aa21530..f89ef979f7c9e 100644 --- a/sycl/source/detail/kernel_impl.cpp +++ b/sycl/source/detail/kernel_impl.cpp @@ -106,38 +106,6 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const { "interoperability function or to query a device built-in kernel"); } -bool kernel_impl::exceedsOccupancyResourceLimits( - const device &Device, const range<3> &WorkGroupSize, - size_t DynamicLocalMemorySize) const { - // Respect occupancy limits for WorkGroupSize and DynamicLocalMemorySize. - // Generally, exceeding hardware resource limits will yield in an error when - // the kernel is launched. - const size_t MaxWorkGroupSize = - get_info(Device); - const size_t MaxLocalMemorySizeInBytes = - Device.get_info(); - - if (WorkGroupSize.size() > MaxWorkGroupSize) - return true; - - if (DynamicLocalMemorySize > MaxLocalMemorySizeInBytes) - return true; - - // It will be impossible to launch a kernel for Cuda when the hardware limit - // for the 32-bit registers page file size is exceeded. - if (Device.get_backend() == backend::ext_oneapi_cuda) { - const uint32_t RegsPerWorkItem = - get_info(Device); - const uint32_t MaxRegsPerWorkGroup = - Device.get_info(); - if ((MaxWorkGroupSize * RegsPerWorkItem) > MaxRegsPerWorkGroup) - return true; - } - - return false; -} - template <> typename info::platform::version::return_type kernel_impl::get_backend_info() const { diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index 67fe2b155ea1e..0ba62bd642efb 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -123,6 +123,32 @@ class kernel_impl { template typename Param::return_type ext_oneapi_get_info(queue Queue) const; + /// Query queue/launch-specific information from a kernel using the + /// info::kernel_queue_specific descriptor for a specific Queue and values. + /// max_num_work_groups is the only valid descriptor for this function. + /// + /// \param Queue is a valid SYCL queue. + /// \param WorkGroupSize is the work-group size the number of work-groups is + /// requested for. + /// \return depends on information being queried. + template + typename Param::return_type + ext_oneapi_get_info(queue Queue, const range<1> &MaxWorkGroupSize, + size_t DynamicLocalMemorySize) const; + + /// Query queue/launch-specific information from a kernel using the + /// info::kernel_queue_specific descriptor for a specific Queue and values. + /// max_num_work_groups is the only valid descriptor for this function. + /// + /// \param Queue is a valid SYCL queue. + /// \param WorkGroupSize is the work-group size the number of work-groups is + /// requested for. + /// \return depends on information being queried. + template + typename Param::return_type + ext_oneapi_get_info(queue Queue, const range<2> &MaxWorkGroupSize, + size_t DynamicLocalMemorySize) const; + /// Query queue/launch-specific information from a kernel using the /// info::kernel_queue_specific descriptor for a specific Queue and values. /// max_num_work_groups is the only valid descriptor for this function. @@ -193,11 +219,49 @@ class kernel_impl { /// Check if the occupancy limits are exceeded for the given kernel launch /// configuration. + template bool exceedsOccupancyResourceLimits(const device &Device, - const range<3> &WorkGroupSize, + const range &WorkGroupSize, size_t DynamicLocalMemorySize) const; + template + size_t queryMaxNumWorkGroups(queue Queue, + const range &WorkGroupSize, + size_t DynamicLocalMemorySize) const; }; +template +bool kernel_impl::exceedsOccupancyResourceLimits( + const device &Device, const range &WorkGroupSize, + size_t DynamicLocalMemorySize) const { + // Respect occupancy limits for WorkGroupSize and DynamicLocalMemorySize. + // Generally, exceeding hardware resource limits will yield in an error when + // the kernel is launched. + const size_t MaxWorkGroupSize = + get_info(Device); + const size_t MaxLocalMemorySizeInBytes = + Device.get_info(); + + if (WorkGroupSize.size() > MaxWorkGroupSize) + return true; + + if (DynamicLocalMemorySize > MaxLocalMemorySizeInBytes) + return true; + + // It will be impossible to launch a kernel for Cuda when the hardware limit + // for the 32-bit registers page file size is exceeded. + if (Device.get_backend() == backend::ext_oneapi_cuda) { + const uint32_t RegsPerWorkItem = + get_info(Device); + const uint32_t MaxRegsPerWorkGroup = + Device.get_info(); + if ((MaxWorkGroupSize * RegsPerWorkItem) > MaxRegsPerWorkGroup) + return true; + } + + return false; +} + template inline typename Param::return_type kernel_impl::get_info() const { static_assert(is_kernel_info_desc::value, @@ -244,13 +308,11 @@ kernel_impl::get_info(const device &Device, namespace syclex = ext::oneapi::experimental; -template <> -inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: - return_type - kernel_impl::ext_oneapi_get_info< - syclex::info::kernel_queue_specific::max_num_work_groups>( - queue Queue, const range<3> &WorkGroupSize, - size_t DynamicLocalMemorySize) const { +template +size_t +kernel_impl::queryMaxNumWorkGroups(queue Queue, + const range &WorkGroupSize, + size_t DynamicLocalMemorySize) const { if (WorkGroupSize.size() == 0) throw exception(sycl::make_error_code(errc::invalid), "The launch work-group size cannot be zero."); @@ -259,12 +321,21 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: const auto &Handle = getHandleRef(); auto Device = Queue.get_device(); + size_t WG[Dimensions]; + WG[0] = WorkGroupSize[0]; + if constexpr (Dimensions >= 2) + WG[1] = WorkGroupSize[1]; + if constexpr (Dimensions == 3) + WG[2] = WorkGroupSize[2]; + uint32_t GroupCount{0}; if (auto Result = Adapter->call_nocheck< UrApiKind::urKernelSuggestMaxCooperativeGroupCountExp>( - Handle, WorkGroupSize.size(), DynamicLocalMemorySize, &GroupCount); - Result != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - // The feature is supported. Check for other errors and throw if any. + Handle, Dimensions, WG, DynamicLocalMemorySize, &GroupCount); + Result != UR_RESULT_ERROR_UNSUPPORTED_FEATURE && + Result != UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE) { + // The feature is supported and the group size is valid. Check for other + // errors and throw if any. Adapter->checkUrResult(Result); return GroupCount; } @@ -278,30 +349,33 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: } template <> -inline typename syclex::info::kernel_queue_specific::max_num_work_group_sync:: +inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: return_type kernel_impl::ext_oneapi_get_info< - syclex::info::kernel_queue_specific::max_num_work_group_sync>( - queue Queue, const range<3> &WorkGroupSize, + syclex::info::kernel_queue_specific::max_num_work_groups>( + queue Queue, const range<1> &WorkGroupSize, size_t DynamicLocalMemorySize) const { - return ext_oneapi_get_info< - syclex::info::kernel_queue_specific::max_num_work_groups>( - Queue, WorkGroupSize, DynamicLocalMemorySize); + return queryMaxNumWorkGroups(Queue, WorkGroupSize, DynamicLocalMemorySize); } template <> -inline typename syclex::info::kernel_queue_specific::max_num_work_group_sync:: +inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: return_type kernel_impl::ext_oneapi_get_info< - syclex::info::kernel_queue_specific::max_num_work_group_sync>( - queue Queue) const { - auto Device = Queue.get_device(); - const auto MaxWorkGroupSize = - get_info(Device); - const sycl::range<3> WorkGroupSize{MaxWorkGroupSize, 1, 1}; - return ext_oneapi_get_info< - syclex::info::kernel_queue_specific::max_num_work_group_sync>( - Queue, WorkGroupSize, /* DynamicLocalMemorySize */ 0); + syclex::info::kernel_queue_specific::max_num_work_groups>( + queue Queue, const range<2> &WorkGroupSize, + size_t DynamicLocalMemorySize) const { + return queryMaxNumWorkGroups(Queue, WorkGroupSize, DynamicLocalMemorySize); +} + +template <> +inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: + return_type + kernel_impl::ext_oneapi_get_info< + syclex::info::kernel_queue_specific::max_num_work_groups>( + queue Queue, const range<3> &WorkGroupSize, + size_t DynamicLocalMemorySize) const { + return queryMaxNumWorkGroups(Queue, WorkGroupSize, DynamicLocalMemorySize); } } // namespace detail diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index 634f22b09bafb..2b61c9dd22d05 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -113,26 +113,41 @@ kernel::ext_oneapi_get_info(queue Queue) const { template typename detail::is_kernel_queue_specific_info_desc::return_type -kernel::ext_oneapi_get_info(queue Queue, const range<3> &WorkGroupSize, +kernel::ext_oneapi_get_info(queue Queue, const range<1> &WorkGroupSize, size_t DynamicLocalMemorySize) const { return impl->ext_oneapi_get_info(Queue, WorkGroupSize, DynamicLocalMemorySize); } -template __SYCL_EXPORT typename ext::oneapi::experimental::info:: - kernel_queue_specific::max_num_work_group_sync::return_type - kernel::ext_oneapi_get_info< - ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(queue Queue) const; +template +typename detail::is_kernel_queue_specific_info_desc::return_type +kernel::ext_oneapi_get_info(queue Queue, const range<2> &WorkGroupSize, + size_t DynamicLocalMemorySize) const { + return impl->ext_oneapi_get_info(Queue, WorkGroupSize, + DynamicLocalMemorySize); +} + +template +typename detail::is_kernel_queue_specific_info_desc::return_type +kernel::ext_oneapi_get_info(queue Queue, const range<3> &WorkGroupSize, + size_t DynamicLocalMemorySize) const { + return impl->ext_oneapi_get_info(Queue, WorkGroupSize, + DynamicLocalMemorySize); +} #define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT) \ template __SYCL_EXPORT ReturnT \ + kernel::ext_oneapi_get_info( \ + queue, const range<1> &, size_t) const; \ + template __SYCL_EXPORT ReturnT \ + kernel::ext_oneapi_get_info( \ + queue, const range<2> &, size_t) const; \ + template __SYCL_EXPORT ReturnT \ kernel::ext_oneapi_get_info( \ queue, const range<3> &, size_t) const; // Not including "ext_oneapi_kernel_queue_specific_traits.def" because not all // kernel_queue_specific queries require the above-defined get_info interface. // clang-format off -__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, kernel_queue_specific, max_num_work_group_sync, size_t) __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, kernel_queue_specific, max_num_work_groups, size_t) // clang-format on #undef __SYCL_PARAM_TRAITS_SPEC diff --git a/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp b/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp index fe8246bb2f56d..06a8a52b56cf0 100644 --- a/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp +++ b/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp @@ -161,14 +161,10 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) { // It cannot be possible to launch a kernel successfully with a configuration // that exceeds the available resources as in the above defined workGroupSize. // workGroupSize is larger than maxWorkGroupSize, hence maxWGs must equal 0. - // Note: Level-Zero currently always returns a non-zero value. - // TODO: Remove the backend condition once the Level-Zero API issue is fixed. - if (dev.get_backend() != sycl::backend::ext_oneapi_level_zero) { - assert(maxWGs == 0 && - "max_num_work_groups query failed.\n" - "It should return 0 possible groups when the requested resources " - "by the lanuch config exceed those available in the hardware."); - } + assert(maxWGs == 0 && + "max_num_work_groups query failed.\n" + "It should return 0 possible groups when the requested resources " + "by the lanuch config exceed those available in the hardware."); // As we ensured that the 'max_num_work_groups' query correctly // returns 0 possible work-groups, test that the kernel launch will fail. diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 92e5d69ffcab4..369057d63adac 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -27,13 +27,18 @@ void testQueriesAndProperties() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); - const auto wgRange = sycl::range{WorkGroupSize, 1, 1}; - const auto maxWGsWithLimits = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q, wgRange, wgRange.size() * sizeof(int)); + const auto local_range = sycl::range<1>(1); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, local_range, 0); + const auto wgRange = sycl::range<3>{WorkGroupSize, 1, 1}; + const auto maxWGsWithLimits = + kernel + .ext_oneapi_get_info( + q, wgRange, wgRange.size() * sizeof(int)); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; q.single_task(props, []() {}); @@ -52,9 +57,11 @@ void testRootGroup() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, WorkGroupSize, 0); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; sycl::buffer dataBuf{sycl::range{maxWGs * WorkGroupSize}}; @@ -94,9 +101,11 @@ void testRootGroupFunctions() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, WorkGroupSize, 0); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 853ac28bad1d5..d24e795ee20ec 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3941,9 +3941,9 @@ _ZNK4sycl3_V16kernel16get_backend_infoINS0_4info6device15backend_versionEEENS0_6 _ZNK4sycl3_V16kernel16get_backend_infoINS0_4info6device7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv _ZNK4sycl3_V16kernel16get_backend_infoINS0_4info8platform7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv _ZNK4sycl3_V16kernel17get_kernel_bundleEv +_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific19max_num_work_groupsEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi1EEEm +_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific19max_num_work_groupsEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi2EEEm _ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific19max_num_work_groupsEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi3EEEm -_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific23max_num_work_group_syncEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueE -_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific23max_num_work_group_syncEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi3EEEm _ZNK4sycl3_V16kernel3getEv _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific15work_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific16global_work_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 8e29aba2726c9..282c5d1b47d16 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -13,8 +13,8 @@ ??$create_sub_devices@$0BAIH@@device@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@_KV?$allocator@_K@std@@@4@@Z ??$create_sub_devices@$0BAII@@device@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@W4partition_affinity_domain@info@12@@Z ??$create_sub_devices@$0BAIJ@@device@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ -??$ext_oneapi_get_info@Umax_num_work_group_sync@kernel_queue_specific@info@experimental@oneapi@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KVqueue@12@@Z -??$ext_oneapi_get_info@Umax_num_work_group_sync@kernel_queue_specific@info@experimental@oneapi@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KVqueue@12@AEBV?$range@$02@12@_K@Z +??$ext_oneapi_get_info@Umax_num_work_groups@kernel_queue_specific@info@experimental@oneapi@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KVqueue@12@AEBV?$range@$00@12@_K@Z +??$ext_oneapi_get_info@Umax_num_work_groups@kernel_queue_specific@info@experimental@oneapi@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KVqueue@12@AEBV?$range@$01@12@_K@Z ??$ext_oneapi_get_info@Umax_num_work_groups@kernel_queue_specific@info@experimental@oneapi@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KVqueue@12@AEBV?$range@$02@12@_K@Z ??$get_backend_info@Ubackend_version@device@info@_V1@sycl@@@context@_V1@sycl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ ??$get_backend_info@Ubackend_version@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ