[SYCL][PI][CUDA] Kernel group private memory size

Signed-off-by: Steffen Larsen <[email protected]>

Author: Steffen Larsen

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

@@ -246,7 +246,8 @@ typedef enum {
     CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
   PI_KERNEL_LOCAL_MEM_SIZE                     = CL_KERNEL_LOCAL_MEM_SIZE,
   PI_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE =
-      CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
+      CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
+  PI_KERNEL_PRIVATE_MEM_SIZE                   = CL_KERNEL_PRIVATE_MEM_SIZE
 } _pi_kernel_group_info;
 
 typedef enum {

sycl/plugins/cuda/pi_cuda.cpp

@@ -2074,6 +2074,15 @@ pi_result cuda_piKernelGetGroupInfo(pi_kernel kernel, pi_device device,
       return getInfo(param_value_size, param_value, param_value_size_ret,
                      static_cast<size_t>(warpSize));
     }
+    case PI_KERNEL_PRIVATE_MEM_SIZE: {
+      // OpenCL PRIVATE == CUDA LOCAL
+      int bytes = 0;
+      cl::sycl::detail::pi::assertion(cuFuncGetAttribute(&bytes,
+                                       CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
+                                       kernel->get()) == CUDA_SUCCESS);
+      return getInfo(param_value_size, param_value, param_value_size_ret,
+                     pi_uint64(bytes));
+    }
     default:
       PI_HANDLE_UNKNOWN_PARAM_NAME(param_name);
     }