[SYCL][PI][CUDA] Kernel group preferred size multiple

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

Author: Steffen Larsen

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

@@ -241,9 +241,12 @@ typedef enum {
 } _pi_kernel_info;
 
 typedef enum {
-  PI_KERNEL_GROUP_INFO_SIZE         = CL_KERNEL_WORK_GROUP_SIZE,
-  PI_KERNEL_COMPILE_GROUP_INFO_SIZE = CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
-  PI_KERNEL_LOCAL_MEM_SIZE          = CL_KERNEL_LOCAL_MEM_SIZE
+  PI_KERNEL_GROUP_INFO_SIZE                    = CL_KERNEL_WORK_GROUP_SIZE,
+  PI_KERNEL_COMPILE_GROUP_INFO_SIZE            =
+    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
 } _pi_kernel_group_info;
 
 typedef enum {

sycl/plugins/cuda/pi_cuda.cpp

@@ -2065,6 +2065,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_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: {
+      // Work groups should be multiples of the warp size
+      int warpSize = 0;
+      cl::sycl::detail::pi::assertion(cuDeviceGetAttribute(&warpSize,
+                                       CU_DEVICE_ATTRIBUTE_WARP_SIZE,
+                                       device->get()) == CUDA_SUCCESS);
+      return getInfo(param_value_size, param_value, param_value_size_ret,
+                     static_cast<size_t>(warpSize));
+    }
     default:
       PI_HANDLE_UNKNOWN_PARAM_NAME(param_name);
     }