Merge branch 'main' into except-san-layer

Author: hdelan

.github/codeql/codeql-config.yml

@@ -27,7 +27,6 @@ jobs:
       uses: github/codeql-action/init@f079b8493333aace61c81488f8bd40919487bd9f # v3.25.7
       with:
         languages: cpp, python
-        config-file: ./.github/codeql/codeql-config.yml
 
     - name: Install pip packages
       run: pip install -r third_party/requirements.txt

.github/workflows/codeql.yml

@@ -35,7 +35,6 @@ jobs:
           format: 'sarif'
           output: 'trivy-results.sarif'
           exit-code: 1  # Fail if issue found
-          skip-dirs: '**/_deps/**'
           # file with suppressions: .trivyignore (in root dir)
 
       - name: Print report and trivyignore file

.github/workflows/trivy.yml

@@ -40,7 +40,7 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR)
         set(UR_LEVEL_ZERO_LOADER_REPO "https://github.com/oneapi-src/level-zero.git")
     endif()
     if (UR_LEVEL_ZERO_LOADER_TAG STREQUAL "")
-        set(UR_LEVEL_ZERO_LOADER_TAG v1.18.3)
+        set(UR_LEVEL_ZERO_LOADER_TAG v1.17.39)
     endif()
 
     # Disable due to a bug https://github.com/oneapi-src/level-zero/issues/104

cmake/FetchLevelZero.cmake

@@ -627,32 +627,6 @@ urCommandBufferFinalizeExp(ur_exp_command_buffer_handle_t CommandBuffer) {
   return UR_RESULT_SUCCESS;
 }
 
-/**
- * Sets the global offset for a kernel command that will be appended to the
- * command buffer.
- * @param[in] CommandBuffer The CommandBuffer where the command will be
- * appended.
- * @param[in] Kernel The handle to the kernel that will be appended.
- * @param[in] GlobalWorkOffset The global offset value.
- * @return UR_RESULT_SUCCESS or an error code on failure
- */
-ur_result_t setKernelGlobalOffset(ur_exp_command_buffer_handle_t CommandBuffer,
-                                  ur_kernel_handle_t Kernel,
-                                  const size_t *GlobalWorkOffset) {
-
-  if (!CommandBuffer->Context->getPlatform()
-           ->ZeDriverGlobalOffsetExtensionFound) {
-    logger::debug("No global offset extension found on this driver");
-    return UR_RESULT_ERROR_INVALID_VALUE;
-  }
-
-  ZE2UR_CALL(zeKernelSetGlobalOffsetExp,
-             (Kernel->ZeKernel, GlobalWorkOffset[0], GlobalWorkOffset[1],
-              GlobalWorkOffset[2]));
-
-  return UR_RESULT_SUCCESS;
-}
-
 /**
  * Sets the kernel arguments for a kernel command that will be appended to the
  * command buffer.
@@ -754,7 +728,8 @@ ur_result_t urCommandBufferAppendKernelLaunchExp(
       Kernel->Mutex, Kernel->Program->Mutex, CommandBuffer->Mutex);
 
   if (GlobalWorkOffset != NULL) {
-    UR_CALL(setKernelGlobalOffset(CommandBuffer, Kernel, GlobalWorkOffset));
+    UR_CALL(setKernelGlobalOffset(CommandBuffer->Context, Kernel->ZeKernel,
+                                  WorkDim, GlobalWorkOffset));
   }
 
   // If there are any pending arguments set them now.

source/adapters/level_zero/command_buffer.cpp

@@ -67,16 +67,17 @@ ur_result_t getSuggestedLocalWorkSize(ur_device_handle_t hDevice,
 }
 
 ur_result_t setKernelGlobalOffset(ur_context_handle_t Context,
-                                  ze_kernel_handle_t Kernel,
+                                  ze_kernel_handle_t Kernel, uint32_t WorkDim,
                                   const size_t *GlobalWorkOffset) {
   if (!Context->getPlatform()->ZeDriverGlobalOffsetExtensionFound) {
     logger::debug("No global offset extension found on this driver");
     return UR_RESULT_ERROR_INVALID_VALUE;
   }
 
-  ZE2UR_CALL(
-      zeKernelSetGlobalOffsetExp,
-      (Kernel, GlobalWorkOffset[0], GlobalWorkOffset[1], GlobalWorkOffset[2]));
+  auto OffsetX = GlobalWorkOffset[0];
+  auto OffsetY = WorkDim > 1 ? GlobalWorkOffset[1] : 0;
+  auto OffsetZ = WorkDim > 2 ? GlobalWorkOffset[2] : 0;
+  ZE2UR_CALL(zeKernelSetGlobalOffsetExp, (Kernel, OffsetX, OffsetY, OffsetZ));
 
   return UR_RESULT_SUCCESS;
 }

source/adapters/level_zero/helpers/kernel_helpers.cpp

@@ -36,11 +36,12 @@ ur_result_t calculateKernelWorkDimensions(
  * command buffer.
  * @param[in] Context Context associated with the queue.
  * @param[in] Kernel The handle to the kernel that will be appended.
- * @param[in] GlobalWorkOffset The global offset value.
+ * @param[in] WorkDim The number of work dimensions.
+ * @param[in] GlobalWorkOffset Array of size WorkDim.
  * @return UR_RESULT_SUCCESS or an error code on failure
  */
 ur_result_t setKernelGlobalOffset(ur_context_handle_t Context,
-                                  ze_kernel_handle_t Kernel,
+                                  ze_kernel_handle_t Kernel, uint32_t WorkDim,
                                   const size_t *GlobalWorkOffset);
 
 /**

source/adapters/level_zero/helpers/kernel_helpers.hpp

@@ -96,14 +96,8 @@ ur_result_t urEnqueueKernelLaunch(
   std::scoped_lock<ur_shared_mutex, ur_shared_mutex, ur_shared_mutex> Lock(
       Queue->Mutex, Kernel->Mutex, Kernel->Program->Mutex);
   if (GlobalWorkOffset != NULL) {
-    if (!Queue->Device->Platform->ZeDriverGlobalOffsetExtensionFound) {
-      logger::error("No global offset extension found on this driver");
-      return UR_RESULT_ERROR_INVALID_VALUE;
-    }
-
-    ZE2UR_CALL(zeKernelSetGlobalOffsetExp,
-               (ZeKernel, GlobalWorkOffset[0], GlobalWorkOffset[1],
-                GlobalWorkOffset[2]));
+    UR_CALL(setKernelGlobalOffset(Queue->Context, ZeKernel, WorkDim,
+                                  GlobalWorkOffset));
   }
 
   // If there are any pending arguments set them now.
@@ -257,14 +251,8 @@ ur_result_t urEnqueueCooperativeKernelLaunchExp(
   std::scoped_lock<ur_shared_mutex, ur_shared_mutex, ur_shared_mutex> Lock(
       Queue->Mutex, Kernel->Mutex, Kernel->Program->Mutex);
   if (GlobalWorkOffset != NULL) {
-    if (!Queue->Device->Platform->ZeDriverGlobalOffsetExtensionFound) {
-      logger::error("No global offset extension found on this driver");
-      return UR_RESULT_ERROR_INVALID_VALUE;
-    }
-
-    ZE2UR_CALL(zeKernelSetGlobalOffsetExp,
-               (ZeKernel, GlobalWorkOffset[0], GlobalWorkOffset[1],
-                GlobalWorkOffset[2]));
+    UR_CALL(setKernelGlobalOffset(Queue->Context, ZeKernel, WorkDim,
+                                  GlobalWorkOffset));
   }
 
   // If there are any pending arguments set them now.

source/adapters/level_zero/kernel.cpp

@@ -252,7 +252,8 @@ ur_result_t ur_queue_immediate_in_order_t::enqueueKernelLaunch(
       hKernel->Mutex, hKernel->getProgramHandle()->Mutex, this->Mutex);
 
   if (pGlobalWorkOffset != NULL) {
-    UR_CALL(setKernelGlobalOffset(hContext, hZeKernel, pGlobalWorkOffset));
+    UR_CALL(
+        setKernelGlobalOffset(hContext, hZeKernel, workDim, pGlobalWorkOffset));
   }
 
   ze_group_count_t zeThreadGroupDimensions{1, 1, 1};

source/adapters/level_zero/v2/queue_immediate_in_order.cpp

@@ -37,6 +37,7 @@ add_ur_adapter(${TARGET_NAME} SHARED
     ${CMAKE_CURRENT_SOURCE_DIR}/program.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/queue.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/sampler.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/usm.hpp
     ${CMAKE_CURRENT_SOURCE_DIR}/usm.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/usm_p2p.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/virtual_mem.cpp

source/adapters/opencl/CMakeLists.txt

@@ -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;

source/adapters/opencl/enqueue.cpp

@@ -81,7 +81,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithIL(
 
     *phProgram = cl_adapter::cast<ur_program_handle_t>(clCreateProgramWithIL(
         cl_adapter::cast<cl_context>(hContext), pIL, length, &Err));
-    CL_RETURN_ON_FAILURE(Err);
   } else {
 
     /* If none of the devices conform with CL 2.1 or newer make sure they all
@@ -109,6 +108,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithIL(
 
     *phProgram = cl_adapter::cast<ur_program_handle_t>(
         FuncPtr(cl_adapter::cast<cl_context>(hContext), pIL, length, &Err));
+  }
+
+  // INVALID_VALUE is only returned in three circumstances according to the cl
+  // spec:
+  // * pIL == NULL
+  // * length == 0
+  // * pIL is not a well-formed binary
+  // UR has a unique error code for each of these, so here we figure out which
+  // to return
+  if (Err == CL_INVALID_VALUE) {
+    if (pIL == nullptr) {
+      return UR_RESULT_ERROR_INVALID_NULL_POINTER;
+    }
+    if (length == 0) {
+      return UR_RESULT_ERROR_INVALID_SIZE;
+    }
+    return UR_RESULT_ERROR_INVALID_BINARY;
+  } else {
     CL_RETURN_ON_FAILURE(Err);
   }