Add 'number of registers' to kernel_device_specific info

Author: joeatodd

sycl/include/CL/sycl/detail/cl.h

@@ -16,3 +16,5 @@
 
 #include <CL/cl.h>
 #include <CL/cl_ext.h>
+
+#define CL_KERNEL_NUM_REGS                  0x11BB

sycl/include/CL/sycl/detail/pi.h

@@ -346,7 +346,8 @@ typedef enum {
   PI_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE = CL_KERNEL_LOCAL_MEM_SIZE,
   PI_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE =
       CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
-  PI_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE = CL_KERNEL_PRIVATE_MEM_SIZE
+  PI_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE = CL_KERNEL_PRIVATE_MEM_SIZE,
+  PI_KERNEL_GROUP_INFO_NUM_REGS = CL_KERNEL_NUM_REGS
 } _pi_kernel_group_info;
 
 typedef enum {

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

@@ -250,6 +250,7 @@ enum class kernel_device_specific : cl_kernel_work_group_info {
   preferred_work_group_size_multiple =
       CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
   private_mem_size = CL_KERNEL_PRIVATE_MEM_SIZE,
+  num_regs = CL_KERNEL_NUM_REGS,
   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,

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

@@ -4,6 +4,7 @@ __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, global_work_size, cl::sycl::ran
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific,
                   preferred_work_group_size_multiple, size_t)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, private_mem_size, cl_ulong)
+__SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, num_regs, uint32_t)
 __SYCL_PARAM_TRAITS_SPEC(kernel_device_specific, work_group_size, size_t)
 __SYCL_PARAM_TRAITS_SPEC_WITH_INPUT(kernel_device_specific, max_sub_group_size,
                              uint32_t, cl::sycl::range<3>)

sycl/plugins/cuda/pi_cuda.cpp

@@ -2548,6 +2548,15 @@ pi_result cuda_piKernelGetGroupInfo(pi_kernel kernel, pi_device device,
       return getInfo(param_value_size, param_value, param_value_size_ret,
                      pi_uint64(bytes));
     }
+    case PI_KERNEL_GROUP_INFO_NUM_REGS: {
+      // OpenCL PRIVATE == CUDA LOCAL
+      int bytes = 0;
+      cl::sycl::detail::pi::assertion(
+          cuFuncGetAttribute(&bytes, CU_FUNC_ATTRIBUTE_NUM_REGS,
+                             kernel->get()) == CUDA_SUCCESS);
+      return getInfo(param_value_size, param_value, param_value_size_ret,
+                     pi_uint32(bytes));
+    }
     default:
       __SYCL_PI_HANDLE_UNKNOWN_PARAM_NAME(param_name);
     }

sycl/source/detail/kernel_info.hpp

@@ -71,6 +71,9 @@ struct IsWorkGroupInfo<
 template <>
 struct IsWorkGroupInfo<info::kernel_device_specific::private_mem_size>
     : std::true_type {};
+template <>
+struct IsWorkGroupInfo<info::kernel_device_specific::num_regs>
+    : std::true_type {};
 
 template <typename T, info::kernel_device_specific Param>
 struct get_kernel_device_specific_info {
@@ -143,6 +146,12 @@ inline cl_ulong get_kernel_device_specific_info_host<
   return 0;
 }
 
+template <>
+inline uint32_t get_kernel_device_specific_info_host<
+    info::kernel_device_specific::num_regs>(const cl::sycl::device &) {
+  return 0;
+}
+
 template <>
 inline uint32_t get_kernel_device_specific_info_host<
     info::kernel_device_specific::max_num_sub_groups>(