[SYCL][ABI-Break] Add SYCL 2020 kernel_device_specific::max_sub_group_size (#6782)

SYCL 2020 promotes the info::kernel_device_specific::max_sub_group_size
query on kernels, but removes the additional argument. This commit adds
an overload with no additional argument and deprecates the old variant.

Author: steffenlarsen

sycl/include/sycl/detail/info_desc_helpers.hpp

@@ -81,20 +81,7 @@ struct IsSubGroupInfo<info::kernel_device_specific::compile_sub_group_size>
   struct is_##DescType##_info_desc<info::DescType::Desc> : std::true_type {    \
     using return_type = info::DescType::Desc::return_type;                     \
   };
-#define __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT(DescType, Desc, ReturnT, InputT,   \
-                                            PiCode)                            \
-  template <> struct PiInfoCode<info::DescType::Desc> {                        \
-    static constexpr                                                           \
-        typename std::conditional<IsSubGroupInfo<info::DescType::Desc>::value, \
-                                  pi_kernel_sub_group_info,                    \
-                                  pi_kernel_group_info>::type value = PiCode;  \
-  };                                                                           \
-  template <>                                                                  \
-  struct is_##DescType##_info_desc<info::DescType::Desc> : std::true_type {    \
-    using with_input_return_type = info::DescType::Desc::return_type;          \
-  };
 #include <sycl/info/kernel_device_specific_traits.def>
-#undef __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT
 #undef __SYCL_PARAM_TRAITS_SPEC
 // Need a static_cast here since piDeviceGetInfo can also accept
 // pi_usm_capability_query values.

sycl/include/sycl/info/info_desc.hpp

@@ -118,12 +118,8 @@ namespace kernel {
 } // namespace kernel
 
 namespace kernel_device_specific {
-#define __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT(DescType, Desc, ReturnT, InputT,   \
-                                            PiCode)                            \
-  __SYCL_PARAM_TRAITS_SPEC(DescType, Desc, ReturnT, PiCode)
 #include <sycl/info/kernel_device_specific_traits.def>
 } // namespace kernel_device_specific
-#undef __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT
 
 // A.6 Event information desctiptors
 enum class event_command_status : pi_int32 {

sycl/include/sycl/info/kernel_device_specific_traits.def

@@ -7,7 +7,6 @@ __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific,
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, private_mem_size, size_t, PI_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, max_num_sub_groups, uint32_t, PI_KERNEL_MAX_NUM_SUB_GROUPS)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, compile_num_sub_groups, uint32_t, PI_KERNEL_COMPILE_NUM_SUB_GROUPS)
-__SYCL_PARAM_TRAITS_SPEC_WITH_INPUT(kernel_device_specific, max_sub_group_size,
-                             uint32_t, sycl::range<3>, PI_KERNEL_MAX_SUB_GROUP_SIZE)
+__SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, max_sub_group_size, uint32_t, PI_KERNEL_MAX_SUB_GROUP_SIZE)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, compile_sub_group_size, uint32_t, PI_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, ext_codeplay_num_regs, uint32_t, PI_KERNEL_GROUP_INFO_NUM_REGS)

sycl/include/sycl/kernel.hpp

@@ -149,9 +149,9 @@ class __SYCL_EXPORT kernel {
   /// \param WGSize is the work-group size the sub-group size is requested for.
   /// \return depends on information being queried.
   template <typename Param>
-  typename detail::is_kernel_device_specific_info_desc<
-      Param>::with_input_return_type
-  get_info(const device &Device, const range<3> &WGSize) const;
+  __SYCL2020_DEPRECATED("Use the overload without the second parameter")
+  typename detail::is_kernel_device_specific_info_desc<Param>::return_type
+      get_info(const device &Device, const range<3> &WGSize) const;
 
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.

sycl/plugins/opencl/pi_opencl.cpp

@@ -26,6 +26,7 @@
 #include <cstring>
 #include <limits>
 #include <map>
+#include <memory>
 #include <sstream>
 #include <string>
 #include <vector>
@@ -860,6 +861,21 @@ pi_result piKernelGetSubGroupInfo(pi_kernel kernel, pi_device device,
   (void)param_value_size;
   size_t ret_val;
   cl_int ret_err;
+
+  std::shared_ptr<void> implicit_input_value;
+  if (param_name == PI_KERNEL_MAX_SUB_GROUP_SIZE && !input_value) {
+    // OpenCL needs an input value for PI_KERNEL_MAX_SUB_GROUP_SIZE so if no
+    // value is given we use the max work item sizes of the device to avoid
+    // truncation of max sub-group size.
+    implicit_input_value = std::shared_ptr<size_t[]>(new size_t[3]);
+    pi_result pi_ret_err = piDeviceGetInfo(
+        device, PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES, 3 * sizeof(size_t),
+        implicit_input_value.get(), nullptr);
+    if (pi_ret_err != PI_SUCCESS)
+      return pi_ret_err;
+    input_value = implicit_input_value.get();
+  }
+
   ret_err = cast<pi_result>(clGetKernelSubGroupInfo(
       cast<cl_kernel>(kernel), cast<cl_device_id>(device),
       cast<cl_kernel_sub_group_info>(param_name), input_value_size, input_value,

sycl/source/detail/kernel_info.hpp

@@ -105,6 +105,9 @@ get_kernel_device_specific_info(RT::PiKernel Kernel, RT::PiDevice Device,
   return sycl::range<3>(Result[0], Result[1], Result[2]);
 }
 
+// TODO: This is used by a deprecated version of
+// info::kernel_device_specific::max_sub_group_size taking an input paramter.
+// This should be removed when the deprecated info query is removed.
 template <typename Param>
 uint32_t get_kernel_device_specific_info_with_input(RT::PiKernel Kernel,
                                                     RT::PiDevice Device,
@@ -178,6 +181,13 @@ inline uint32_t get_kernel_device_specific_info_host<
                              PI_ERROR_INVALID_KERNEL);
 }
 
+template <>
+inline uint32_t get_kernel_device_specific_info_host<
+    info::kernel_device_specific::max_sub_group_size>(const sycl::device &) {
+  throw invalid_object_error("This instance of kernel is a host instance",
+                             PI_ERROR_INVALID_KERNEL);
+}
+
 template <>
 inline uint32_t get_kernel_device_specific_info_host<
     info::kernel_device_specific::compile_num_sub_groups>(

sycl/source/kernel.cpp

@@ -56,25 +56,28 @@ kernel::get_info(const device &Dev) const {
   return impl->get_info<Param>(Dev);
 }
 
+// Deprecated overload for kernel_device_specific::max_sub_group_size taking
+// an extra argument.
 template <typename Param>
-typename detail::is_kernel_device_specific_info_desc<
-    Param>::with_input_return_type
+typename detail::is_kernel_device_specific_info_desc<Param>::return_type
 kernel::get_info(const device &Device, const range<3> &WGSize) const {
+  static_assert(
+      std::is_same_v<Param, info::kernel_device_specific::max_sub_group_size>,
+      "Unexpected param for kernel::get_info with range argument.");
   return impl->get_info<Param>(Device, WGSize);
 }
 
 #define __SYCL_PARAM_TRAITS_SPEC(DescType, Desc, ReturnT, PiCode)              \
   template __SYCL_EXPORT ReturnT kernel::get_info<info::DescType::Desc>(       \
       const device &) const;
-#define __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT(DescType, Desc, ReturnT, InputT,   \
-                                            PiCode)                            \
-  template __SYCL_EXPORT ReturnT kernel::get_info<info::DescType::Desc>(       \
-      const device &, const InputT &) const;
 
 #include <sycl/info/kernel_device_specific_traits.def>
 
 #undef __SYCL_PARAM_TRAITS_SPEC
-#undef __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT
+
+template __SYCL_EXPORT uint32_t
+kernel::get_info<info::kernel_device_specific::max_sub_group_size>(
+    const device &, const sycl::range<3> &) const;
 
 kernel::kernel(std::shared_ptr<detail::kernel_impl> Impl) : impl(Impl) {}
 

sycl/test/abi/sycl_symbols_linux.dump

@@ -4282,7 +4282,8 @@ _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific15work_group_size
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific16global_work_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific16private_mem_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific18max_num_sub_groupsEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
-_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific18max_sub_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E22with_input_return_typeERKNS0_6deviceERKNS0_5rangeILi3EEE
+_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific18max_sub_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
+_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific18max_sub_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceERKNS0_5rangeILi3EEE
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific21ext_codeplay_num_regsEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific22compile_num_sub_groupsEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
 _ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific22compile_sub_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE