[SYCL] Implement SYCL_INTEL_device_specific_kernel_queries

sycl/include/CL/sycl/info/info_desc.hpp

@@ -214,6 +214,19 @@ enum class kernel_sub_group : cl_kernel_sub_group_info {
   compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
 };
 
+enum class kernel_device_specific : cl_kernel_work_group_info {
+  global_work_size = CL_KERNEL_GLOBAL_WORK_SIZE,
+  work_group_size = CL_KERNEL_WORK_GROUP_SIZE,
+  compile_work_group_size = CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
+  preferred_work_group_size_multiple =
+      CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
+  private_mem_size = CL_KERNEL_PRIVATE_MEM_SIZE,
+  max_sub_group_size = CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
+  max_num_sub_groups = CL_KERNEL_MAX_NUM_SUB_GROUPS,
+  compile_num_sub_groups = CL_KERNEL_COMPILE_NUM_SUB_GROUPS,
+  compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
+};
+
 // A.6 Program information desctiptors
 enum class program : cl_program_info {
   context = CL_PROGRAM_CONTEXT,
@@ -242,6 +255,8 @@ enum class event_profiling : cl_profiling_info {
 // Provide an alias to the return type for each of the info parameters
 template <typename T, T param> class param_traits {};
 
+template <typename T, T param> struct compatibility_param_traits {};
+
 #define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
   template <> class param_traits<param_type, param_type::param> {              \
   public:                                                                      \
@@ -263,6 +278,7 @@ template <typename T, T param> class param_traits {};
 
 #include <CL/sycl/info/event_profiling_traits.def>
 
+#include <CL/sycl/info/kernel_device_specific_traits.def>
 #include <CL/sycl/info/kernel_sub_group_traits.def>
 #include <CL/sycl/info/kernel_traits.def>
 #include <CL/sycl/info/kernel_work_group_traits.def>
@@ -276,6 +292,24 @@ template <typename T, T param> class param_traits {};
 #undef PARAM_TRAITS_SPEC
 #undef PARAM_TRAITS_SPEC_WITH_INPUT
 
+#define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
+  template <>                                                                  \
+  struct compatibility_param_traits<param_type, param_type::param> {           \
+    static constexpr auto value = kernel_device_specific::param;               \
+  };
+
+#define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type)     \
+  template <>                                                                  \
+  struct compatibility_param_traits<param_type, param_type::param> {           \
+    static constexpr auto value = kernel_device_specific::param;               \
+  };
+
+#include <CL/sycl/info/kernel_sub_group_traits.def>
+#include <CL/sycl/info/kernel_work_group_traits.def>
+
+#undef PARAM_TRAITS_SPEC
+#undef PARAM_TRAITS_SPEC_WITH_INPUT
+
 } // namespace info
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/info/kernel_device_specific_traits.def

@@ -0,0 +1,12 @@
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_work_group_size,
+                  cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, global_work_size, cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, 
+                  preferred_work_group_size_multiple, size_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, private_mem_size, cl_ulong)
+PARAM_TRAITS_SPEC(kernel_device_specific, work_group_size, size_t)
+PARAM_TRAITS_SPEC_WITH_INPUT(kernel_device_specific, max_sub_group_size,
+                             uint32_t, cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, max_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_sub_group_size, uint32_t)

sycl/include/CL/sycl/kernel.hpp

@@ -92,6 +92,27 @@ class __SYCL_EXPORT kernel {
   typename info::param_traits<info::kernel, param>::return_type
   get_info() const;
 
+  /// Query device-specific information from the kernel object using the
+  /// info::kernel_device_specific descriptor.
+  ///
+  /// \param Device is a valid SYCL device to query info for.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device) const;
+
+  /// Query device-specific information from a kernel using the
+  /// info::kernel_device_specific descriptor for a specific device and value.
+  ///
+  /// \param Device is a valid SYCL device.
+  /// \param Value depends on information being queried.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device,
+           typename info::param_traits<info::kernel_device_specific,
+                                       param>::input_type Value) const;
+
   /// Query work-group information from a kernel using the
   /// info::kernel_work_group descriptor for a specific device.
   ///
@@ -107,8 +128,11 @@ class __SYCL_EXPORT kernel {
   /// \param Device is a valid SYCL device.
   /// \return depends on information being queried.
   template <info::kernel_sub_group param>
+  // clang-format off
   typename info::param_traits<info::kernel_sub_group, param>::return_type
+  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
   get_sub_group_info(const device &Device) const;
+  // clang-format on
 
   /// Query sub-group information from a kernel using the
   /// info::kernel_sub_group descriptor for a specific device and value.
@@ -117,11 +141,13 @@ class __SYCL_EXPORT kernel {
   /// \param Value depends on information being queried.
   /// \return depends on information being queried.
   template <info::kernel_sub_group param>
+  // clang-format off
   typename info::param_traits<info::kernel_sub_group, param>::return_type
-  get_sub_group_info(
-      const device &Device,
-      typename info::param_traits<info::kernel_sub_group, param>::input_type
-          Value) const;
+  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
+  get_sub_group_info(const device &Device,
+                     typename info::param_traits<info::kernel_sub_group,
+                     param>::input_type Value) const;
+  // clang-format on
 
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.

sycl/source/detail/kernel_impl.cpp

@@ -82,43 +82,59 @@ template <> program kernel_impl::get_info<info::kernel::program>() const {
   return createSyclObjFromImpl<program>(MProgramImpl);
 }
 
-template <info::kernel_work_group param>
-typename info::param_traits<info::kernel_work_group, param>::return_type
-kernel_impl::get_work_group_info(const device &Device) const {
+template <info::kernel_device_specific param>
+typename info::param_traits<info::kernel_device_specific, param>::return_type
+kernel_impl::get_info(const device &Device) const {
   if (is_host()) {
-    return get_kernel_work_group_info_host<param>(Device);
+    return get_kernel_device_specific_info_host<param>(Device);
   }
-  return get_kernel_work_group_info<
-      typename info::param_traits<info::kernel_work_group, param>::return_type,
+  return get_kernel_device_specific_info<
+      typename info::param_traits<info::kernel_device_specific,
+                                  param>::return_type,
       param>::get(this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
                   getPlugin());
 }
 
-template <info::kernel_sub_group param>
-typename info::param_traits<info::kernel_sub_group, param>::return_type
-kernel_impl::get_sub_group_info(const device &Device) const {
+template <info::kernel_device_specific param>
+typename info::param_traits<info::kernel_device_specific, param>::return_type
+kernel_impl::get_info(
+    const device &Device,
+    typename info::param_traits<info::kernel_device_specific, param>::input_type
+        Value) const {
   if (is_host()) {
     throw runtime_error("Sub-group feature is not supported on HOST device.",
                         PI_INVALID_DEVICE);
   }
-  return get_kernel_sub_group_info<param>::get(
-      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
+  return get_kernel_device_specific_info_with_input<param>::get(
+      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), Value,
       getPlugin());
 }
 
+template <info::kernel_work_group param>
+typename info::param_traits<info::kernel_work_group, param>::return_type
+kernel_impl::get_work_group_info(const device &Device) const {
+  return get_info<
+      info::compatibility_param_traits<info::kernel_work_group, param>::value>(
+      Device);
+}
+
+template <info::kernel_sub_group param>
+typename info::param_traits<info::kernel_sub_group, param>::return_type
+kernel_impl::get_sub_group_info(const device &Device) const {
+  return get_info<
+      info::compatibility_param_traits<info::kernel_sub_group, param>::value>(
+      Device);
+}
+
 template <info::kernel_sub_group param>
 typename info::param_traits<info::kernel_sub_group, param>::return_type
 kernel_impl::get_sub_group_info(
     const device &Device,
     typename info::param_traits<info::kernel_sub_group, param>::input_type
         Value) const {
-  if (is_host()) {
-    throw runtime_error("Sub-group feature is not supported on HOST device.",
-                        PI_INVALID_DEVICE);
-  }
-  return get_kernel_sub_group_info_with_input<param>::get(
-      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), Value,
-      getPlugin());
+  return get_info<
+      info::compatibility_param_traits<info::kernel_sub_group, param>::value>(
+      Device, Value);
 }
 
 #define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
@@ -128,6 +144,18 @@ kernel_impl::get_sub_group_info(
 
 #undef PARAM_TRAITS_SPEC
 
+#define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
+  template ret_type kernel_impl::get_info<info::param_type::param>(            \
+      const device &) const;
+#define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type)     \
+  template ret_type kernel_impl::get_info<info::param_type::param>(            \
+      const device &, in_type) const;
+
+#include <CL/sycl/info/kernel_device_specific_traits.def>
+
+#undef PARAM_TRAITS_SPEC
+#undef PARAM_TRAITS_SPEC_WITH_INPUT
+
 #define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
   template ret_type kernel_impl::get_work_group_info<info::param_type::param>( \
       const device &) const;

sycl/source/detail/kernel_impl.hpp

@@ -103,6 +103,27 @@ class kernel_impl {
   typename info::param_traits<info::kernel, param>::return_type
   get_info() const;
 
+  /// Query device-specific information from a kernel object using the
+  /// info::kernel_device_specific descriptor.
+  ///
+  /// \param Device is a valid SYCL device to query info for.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device) const;
+
+  /// Query device-specific information from a kernel using the
+  /// info::kernel_device_specific descriptor for a specific device and value.
+  ///
+  /// \param Device is a valid SYCL device.
+  /// \param Value depends on information being queried.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device,
+           typename info::param_traits<info::kernel_device_specific,
+                                       param>::input_type Value) const;
+
   /// Query work-group information from a kernel using the
   /// info::kernel_work_group descriptor for a specific device.
   ///