Query out and use local size set in program IL in CL adapter.

The CL spec wording on this is kind of fuzzy but every CL driver I
tested (across intel, amd, nvidia cpu + gpu) returns an error when you
have a local size set in the program source/IL and you don't specify any
local size in your clEnqueueNDRangeKernel call (i.e. you leave it as
NULL).

Our spec does allow you to leave local size as null if you have a size
specified in your program, so this change adds some logic to query out
the size set in the program and passes it to the enqueue call.

Initially I was concerned this might impact performance of current users
but it looks like SYCL always passes a local size when calling
urEnqueueKernelLaunch so it won't hit the path with the extra query.

Author: Aaron Greig

source/adapters/opencl/enqueue.cpp

@@ -30,12 +30,33 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     const size_t *pGlobalWorkOffset, const size_t *pGlobalWorkSize,
     const size_t *pLocalWorkSize, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
+  std::vector<size_t> compiledLocalWorksize;
+  if (!pLocalWorkSize) {
+    cl_device_id device = nullptr;
+    CL_RETURN_ON_FAILURE(clGetCommandQueueInfo(
+        cl_adapter::cast<cl_command_queue>(hQueue), CL_QUEUE_DEVICE,
+        sizeof(device), &device, nullptr));
+    // This query always returns size_t[3], if nothing was specified it returns
+    // all zeroes.
+    size_t queriedLocalWorkSize[3] = {0, 0, 0};
+    CL_RETURN_ON_FAILURE(clGetKernelWorkGroupInfo(
+        cl_adapter::cast<cl_kernel>(hKernel), device,
+        CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof(size_t[3]),
+        queriedLocalWorkSize, nullptr));
+    if (queriedLocalWorkSize[0] != 0) {
+      for (uint32_t i = 0; i < workDim; i++) {
+        compiledLocalWorksize.push_back(queriedLocalWorkSize[i]);
+      }
+    }
+  }
 
   CL_RETURN_ON_FAILURE(clEnqueueNDRangeKernel(
       cl_adapter::cast<cl_command_queue>(hQueue),
       cl_adapter::cast<cl_kernel>(hKernel), workDim, pGlobalWorkOffset,
-      pGlobalWorkSize, pLocalWorkSize, numEventsInWaitList,
-      cl_adapter::cast<const cl_event *>(phEventWaitList),
+      pGlobalWorkSize,
+      compiledLocalWorksize.empty() ? pLocalWorkSize
+                                    : compiledLocalWorksize.data(),
+      numEventsInWaitList, cl_adapter::cast<const cl_event *>(phEventWaitList),
       cl_adapter::cast<cl_event *>(phEvent)));
 
   return UR_RESULT_SUCCESS;

test/conformance/enqueue/enqueue_adapter_opencl.match

@@ -1,3 +1,2 @@
 {{NONDETERMINISTIC}}
-urEnqueueKernelLaunchKernelWgSizeTest.Success/Intel_R__OpenCL___{{.*}}_
 {{OPT}}urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__OpenCL___{{.*}}_UsePoolEnabled