[SYCL] Added support for kernel descriptor compile_sub_group_size

Signed-off-by: Tian, Shilei <[email protected]>

Author: shileiti

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

@@ -203,7 +203,8 @@ enum class kernel_sub_group : cl_kernel_sub_group_info {
   sub_group_count_for_ndrange = CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE,
   local_size_for_sub_group_count = CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
   max_num_sub_groups = CL_KERNEL_MAX_NUM_SUB_GROUPS,
-  compile_num_sub_groups = CL_KERNEL_COMPILE_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
@@ -364,6 +365,7 @@ PARAM_TRAITS_SPEC_WITH_INPUT(kernel_sub_group, local_size_for_sub_group_count,
                              cl::sycl::range<3>, size_t)
 PARAM_TRAITS_SPEC(kernel_sub_group, max_num_sub_groups, size_t)
 PARAM_TRAITS_SPEC(kernel_sub_group, compile_num_sub_groups, size_t)
+PARAM_TRAITS_SPEC(kernel_sub_group, compile_sub_group_size, size_t)
 
 PARAM_TRAITS_SPEC(platform, profile, string_class)
 PARAM_TRAITS_SPEC(platform, version, string_class)

sycl/test/sub_group/info.cpp

@@ -72,6 +72,18 @@ int main() {
     /* Sub-group size is not specified in kernel or IL*/
     exit_if_not_equal<size_t>(Res, 0, "compile_num_sub_groups");
 
+    // According to specification, this kernel query requires `cl_khr_subgroups`
+    // or `cl_intel_subgroups`
+    if ((Device.has_extension("cl_khr_subgroups") ||
+         Device.has_extension("cl_intel_subgroups")) &&
+        Device.has_extension("cl_intel_required_subgroup_size")) {
+      Res = Kernel.get_sub_group_info<
+          info::kernel_sub_group::compile_sub_group_size>(Device);
+
+      /* Required sub-group size is not specified in kernel or IL*/
+      exit_if_not_equal<size_t>(Res, 0, "compile_sub_group_size");
+    }
+
     /* Check work-group sizea which can accommodate the requested number of
      * sub-groups*/
     for (auto s : {(size_t)200, (size_t)1, (size_t)3, (size_t)5, (size_t)7,
@@ -103,4 +115,3 @@ int main() {
   std::cout << "Test passed.\n";
   return 0;
 }
-