@@ -162,6 +162,39 @@ class kernel_impl {
162162 ext_oneapi_get_info (queue Queue, const range<3 > &MaxWorkGroupSize,
163163 size_t DynamicLocalMemorySize) const ;
164164
165+ // / Query queue/launch-specific information from a kernel using the
166+ // / info::kernel_queue_specific descriptor for a specific Queue and values.
167+ // / max_num_work_groups is the only valid descriptor for this function.
168+ // /
169+ // / \param Queue is a valid SYCL queue.
170+ // / \param WG is a work group size
171+ // / \return depends on information being queried.
172+ template <typename Param>
173+ typename Param::return_type ext_oneapi_get_info (queue Queue,
174+ const range<3 > &WG) const ;
175+
176+ // / Query queue/launch-specific information from a kernel using the
177+ // / info::kernel_queue_specific descriptor for a specific Queue and values.
178+ // / max_num_work_groups is the only valid descriptor for this function.
179+ // /
180+ // / \param Queue is a valid SYCL queue.
181+ // / \param WG is a work group size
182+ // / \return depends on information being queried.
183+ template <typename Param>
184+ typename Param::return_type ext_oneapi_get_info (queue Queue,
185+ const range<2 > &WG) const ;
186+
187+ // / Query queue/launch-specific information from a kernel using the
188+ // / info::kernel_queue_specific descriptor for a specific Queue and values.
189+ // / max_num_work_groups is the only valid descriptor for this function.
190+ // /
191+ // / \param Queue is a valid SYCL queue.
192+ // / \param WG is a work group size
193+ // / \return depends on information being queried.
194+ template <typename Param>
195+ typename Param::return_type ext_oneapi_get_info (queue Queue,
196+ const range<1 > &WG) const ;
197+
165198 // / Get a constant reference to a raw kernel object.
166199 // /
167200 // / \return a constant reference to a valid UrKernel instance with raw
@@ -383,6 +416,91 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups::
383416 DynamicLocalMemorySize);
384417}
385418
419+ template <>
420+ inline typename syclex::info::kernel_queue_specific::max_work_group_size::
421+ return_type
422+ kernel_impl::ext_oneapi_get_info<
423+ syclex::info::kernel_queue_specific::max_work_group_size>(
424+ queue Queue) const {
425+ const auto &Adapter = getAdapter ();
426+ const auto DeviceNativeHandle =
427+ getSyclObjImpl (Queue.get_device ())->getHandleRef ();
428+
429+ size_t KernelWGSize = 0 ;
430+ Adapter->call <UrApiKind::urKernelGetGroupInfo>(
431+ MKernel, DeviceNativeHandle, UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE,
432+ sizeof (size_t ), &KernelWGSize, nullptr );
433+ return KernelWGSize;
434+ }
435+
436+ template <int Dimensions>
437+ inline sycl::id<Dimensions>
438+ generate_id (const sycl::range<Dimensions> &DevMaxWorkItemSizes,
439+ const size_t DevWgSize) {
440+ sycl::id<Dimensions> Ret;
441+ for (int i = 0 ; i < Dimensions; i++) {
442+ // DevMaxWorkItemSizes values are inverted, see
443+ // sycl/source/detail/device_info.hpp:582
444+ Ret[i] = std::min (DevMaxWorkItemSizes[i], DevWgSize);
445+ }
446+ return Ret;
447+ }
448+
449+ #define ADD_TEMPLATE_METHOD_SPEC (Num ) \
450+ template <> \
451+ inline typename syclex::info::kernel_queue_specific::max_work_item_sizes< \
452+ Num>::return_type \
453+ kernel_impl::ext_oneapi_get_info< \
454+ syclex::info::kernel_queue_specific::max_work_item_sizes<Num>>( \
455+ queue Queue) const { \
456+ const auto Dev = Queue.get_device (); \
457+ const auto DeviceWgSize = \
458+ get_info<info::kernel_device_specific::work_group_size>(Dev); \
459+ const auto DeviceMaxWorkItemSizes = \
460+ Dev.get_info <info::device::max_work_item_sizes<Num>>(); \
461+ return generate_id<Num>(DeviceMaxWorkItemSizes, DeviceWgSize); \
462+ } // namespace detail
463+
464+ ADD_TEMPLATE_METHOD_SPEC (1 )
465+ ADD_TEMPLATE_METHOD_SPEC (2 )
466+ ADD_TEMPLATE_METHOD_SPEC (3 )
467+
468+ #undef ADD_TEMPLATE_METHOD_SPEC
469+
470+ #define ADD_TEMPLATE_METHOD_SPEC (QueueSpec, Num, Kind, Reg ) \
471+ template <> \
472+ inline typename syclex::info::kernel_queue_specific::QueueSpec::return_type \
473+ kernel_impl::ext_oneapi_get_info< \
474+ syclex::info::kernel_queue_specific::QueueSpec>( \
475+ queue Queue, const range<Num> &WG) const { \
476+ if (WG.size () == 0 ) \
477+ throw exception (sycl::make_error_code (errc::invalid), \
478+ " The work-group size cannot be zero."
479+ const auto &Adapter = getAdapter (); \
480+ const auto DeviceNativeHandle = \
481+ getSyclObjImpl (Queue.get_device ())->getHandleRef (); \
482+ uint32_t KernelSubWGSize = 0 ; \
483+ Adapter->call <UrApiKind::Kind>(MKernel, DeviceNativeHandle, Reg, \
484+ sizeof (uint32_t ), &KernelSubWGSize, \
485+ nullptr ); \
486+ return KernelSubWGSize; \
487+ }
488+
489+ ADD_TEMPLATE_METHOD_SPEC (max_sub_group_size, 3 , urKernelGetSubGroupInfo,
490+ UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE)
491+ ADD_TEMPLATE_METHOD_SPEC (max_sub_group_size, 2 , urKernelGetSubGroupInfo,
492+ UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE)
493+ ADD_TEMPLATE_METHOD_SPEC (max_sub_group_size, 1 , urKernelGetSubGroupInfo,
494+ UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE)
495+
496+ ADD_TEMPLATE_METHOD_SPEC (num_sub_groups, 3 , urKernelGetSubGroupInfo,
497+ UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS)
498+ ADD_TEMPLATE_METHOD_SPEC (num_sub_groups, 2 , urKernelGetSubGroupInfo,
499+ UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS)
500+ ADD_TEMPLATE_METHOD_SPEC (num_sub_groups, 1 , urKernelGetSubGroupInfo,
501+ UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS)
502+
503+ #undef ADD_TEMPLATE_METHOD_SPEC
386504} // namespace detail
387505} // namespace _V1
388506} // namespace sycl
0 commit comments