From 24ed565802f527224882c7bec04508abb525ce2d Mon Sep 17 00:00:00 2001
From: Aaron Greig <aaron.greig@codeplay.com>
Date: Tue, 1 Oct 2024 15:49:55 +0100
Subject: [PATCH] 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.
---
 source/adapters/opencl/enqueue.cpp            | 25 +++++++++++++++++--
 .../enqueue/enqueue_adapter_opencl.match      |  1 -
 2 files changed, 23 insertions(+), 3 deletions(-)

diff --git a/source/adapters/opencl/enqueue.cpp b/source/adapters/opencl/enqueue.cpp
index 9fb4239c38..6596a01317 100644
--- a/source/adapters/opencl/enqueue.cpp
+++ b/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;
diff --git a/test/conformance/enqueue/enqueue_adapter_opencl.match b/test/conformance/enqueue/enqueue_adapter_opencl.match
index d67d35d9b0..c9c582b41c 100644
--- a/test/conformance/enqueue/enqueue_adapter_opencl.match
+++ b/test/conformance/enqueue/enqueue_adapter_opencl.match
@@ -1,5 +1,4 @@
 {{NONDETERMINISTIC}}
 {{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/Intel_R__OpenCL___{{.*}}_
-urEnqueueKernelLaunchKernelWgSizeTest.Success/Intel_R__OpenCL___{{.*}}_
 urEnqueueKernelLaunchKernelSubGroupTest.Success/Intel_R__OpenCL___{{.*}}_
 {{OPT}}urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__OpenCL___{{.*}}_UsePoolEnabled