Merge branch 'main' into tensormap-exp-api

Author: Hugh Delaney

.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

@@ -38,6 +38,7 @@ add_ur_adapter(${TARGET_NAME} SHARED
     ${CMAKE_CURRENT_SOURCE_DIR}/queue.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/sampler.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/tensor_map.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);
   }
 

source/adapters/opencl/program.cpp

@@ -11,6 +11,14 @@
 #include <ur/ur.hpp>
 
 #include "common.hpp"
+#include "usm.hpp"
+
+template <class T>
+void AllocDeleterCallback(cl_event event, cl_int, void *pUserData) {
+  clReleaseEvent(event);
+  auto Info = static_cast<T *>(pUserData);
+  delete Info;
+}
 
 namespace umf {
 ur_result_t getProviderNativeError(const char *, int32_t) {
@@ -312,32 +320,19 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill(
       numEventsInWaitList, cl_adapter::cast<const cl_event *>(phEventWaitList),
       &CopyEvent));
 
-  struct DeleteCallbackInfo {
-    DeleteCallbackInfo(clMemBlockingFreeINTEL_fn USMFree, cl_context CLContext,
-                       void *HostBuffer)
-        : USMFree(USMFree), CLContext(CLContext), HostBuffer(HostBuffer) {
-      clRetainContext(CLContext);
-    }
-    ~DeleteCallbackInfo() {
-      USMFree(CLContext, HostBuffer);
-      clReleaseContext(CLContext);
-    }
-    DeleteCallbackInfo(const DeleteCallbackInfo &) = delete;
-    DeleteCallbackInfo &operator=(const DeleteCallbackInfo &) = delete;
-
-    clMemBlockingFreeINTEL_fn USMFree;
-    cl_context CLContext;
-    void *HostBuffer;
-  };
-
-  auto Info = new DeleteCallbackInfo(USMFree, CLContext, HostBuffer);
+  if (phEvent) {
+    // Since we're releasing this in the callback above we need to retain it
+    // here to keep the user copy alive.
+    CL_RETURN_ON_FAILURE(clRetainEvent(CopyEvent));
+    *phEvent = cl_adapter::cast<ur_event_handle_t>(CopyEvent);
+  }
 
-  auto DeleteCallback = [](cl_event, cl_int, void *pUserData) {
-    auto Info = static_cast<DeleteCallbackInfo *>(pUserData);
-    delete Info;
-  };
+  // This self destructs taking the event and allocation with it.
+  auto Info = new AllocDeleterCallbackInfo(USMFree, CLContext, HostBuffer);
 
-  ClErr = clSetEventCallback(CopyEvent, CL_COMPLETE, DeleteCallback, Info);
+  ClErr =
+      clSetEventCallback(CopyEvent, CL_COMPLETE,
+                         AllocDeleterCallback<AllocDeleterCallbackInfo>, Info);
   if (ClErr != CL_SUCCESS) {
     // We can attempt to recover gracefully by attempting to wait for the copy
     // to finish and deleting the info struct here.
@@ -346,11 +341,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill(
     clReleaseEvent(CopyEvent);
     CL_RETURN_ON_FAILURE(ClErr);
   }
-  if (phEvent) {
-    *phEvent = cl_adapter::cast<ur_event_handle_t>(CopyEvent);
-  } else {
-    CL_RETURN_ON_FAILURE(clReleaseEvent(CopyEvent));
-  }
 
   return UR_RESULT_SUCCESS;
 }
@@ -369,20 +359,131 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
     return mapCLErrorToUR(CLErr);
   }
 
-  clEnqueueMemcpyINTEL_fn FuncPtr = nullptr;
-  ur_result_t RetVal = cl_ext::getExtFuncFromContext<clEnqueueMemcpyINTEL_fn>(
+  clGetMemAllocInfoINTEL_fn GetMemAllocInfo = nullptr;
+  UR_RETURN_ON_FAILURE(cl_ext::getExtFuncFromContext<clGetMemAllocInfoINTEL_fn>(
+      CLContext, cl_ext::ExtFuncPtrCache->clGetMemAllocInfoINTELCache,
+      cl_ext::GetMemAllocInfoName, &GetMemAllocInfo));
+
+  clEnqueueMemcpyINTEL_fn USMMemcpy = nullptr;
+  UR_RETURN_ON_FAILURE(cl_ext::getExtFuncFromContext<clEnqueueMemcpyINTEL_fn>(
       CLContext, cl_ext::ExtFuncPtrCache->clEnqueueMemcpyINTELCache,
-      cl_ext::EnqueueMemcpyName, &FuncPtr);
+      cl_ext::EnqueueMemcpyName, &USMMemcpy));
 
-  if (FuncPtr) {
-    RetVal = mapCLErrorToUR(
-        FuncPtr(cl_adapter::cast<cl_command_queue>(hQueue), blocking, pDst,
-                pSrc, size, numEventsInWaitList,
-                cl_adapter::cast<const cl_event *>(phEventWaitList),
-                cl_adapter::cast<cl_event *>(phEvent)));
+  clMemBlockingFreeINTEL_fn USMFree = nullptr;
+  UR_RETURN_ON_FAILURE(cl_ext::getExtFuncFromContext<clMemBlockingFreeINTEL_fn>(
+      CLContext, cl_ext::ExtFuncPtrCache->clMemBlockingFreeINTELCache,
+      cl_ext::MemBlockingFreeName, &USMFree));
+
+  // Check if the two allocations are DEVICE allocations from different
+  // devices, if they are we need to do the copy indirectly via a host
+  // allocation.
+  cl_device_id SrcDevice = 0, DstDevice = 0;
+  CL_RETURN_ON_FAILURE(
+      GetMemAllocInfo(CLContext, pSrc, CL_MEM_ALLOC_DEVICE_INTEL,
+                      sizeof(cl_device_id), &SrcDevice, nullptr));
+  CL_RETURN_ON_FAILURE(
+      GetMemAllocInfo(CLContext, pDst, CL_MEM_ALLOC_DEVICE_INTEL,
+                      sizeof(cl_device_id), &DstDevice, nullptr));
+
+  if ((SrcDevice && DstDevice) && SrcDevice != DstDevice) {
+    // We need a queue associated with each device, so first figure out which
+    // one we weren't given.
+    cl_device_id QueueDevice = nullptr;
+    CL_RETURN_ON_FAILURE(clGetCommandQueueInfo(
+        cl_adapter::cast<cl_command_queue>(hQueue), CL_QUEUE_DEVICE,
+        sizeof(QueueDevice), &QueueDevice, nullptr));
+
+    cl_command_queue MissingQueue = nullptr, SrcQueue = nullptr,
+                     DstQueue = nullptr;
+    if (QueueDevice == SrcDevice) {
+      MissingQueue = clCreateCommandQueue(CLContext, DstDevice, 0, &CLErr);
+      SrcQueue = cl_adapter::cast<cl_command_queue>(hQueue);
+      DstQueue = MissingQueue;
+    } else {
+      MissingQueue = clCreateCommandQueue(CLContext, SrcDevice, 0, &CLErr);
+      DstQueue = cl_adapter::cast<cl_command_queue>(hQueue);
+      SrcQueue = MissingQueue;
+    }
+    CL_RETURN_ON_FAILURE(CLErr);
+
+    cl_event HostCopyEvent = nullptr, FinalCopyEvent = nullptr;
+    clHostMemAllocINTEL_fn HostMemAlloc = nullptr;
+    UR_RETURN_ON_FAILURE(cl_ext::getExtFuncFromContext<clHostMemAllocINTEL_fn>(
+        CLContext, cl_ext::ExtFuncPtrCache->clHostMemAllocINTELCache,
+        cl_ext::HostMemAllocName, &HostMemAlloc));
+
+    auto HostAlloc = HostMemAlloc(CLContext, nullptr, size, 0, &CLErr);
+    CL_RETURN_ON_FAILURE(CLErr);
+
+    // Now that we've successfully allocated we should try to clean it up if we
+    // hit an error somewhere.
+    auto checkCLErr = [&](cl_int CLErr) -> ur_result_t {
+      if (CLErr != CL_SUCCESS) {
+        if (HostCopyEvent) {
+          clReleaseEvent(HostCopyEvent);
+        }
+        if (FinalCopyEvent) {
+          clReleaseEvent(FinalCopyEvent);
+        }
+        USMFree(CLContext, HostAlloc);
+        CL_RETURN_ON_FAILURE(CLErr);
+      }
+      return UR_RESULT_SUCCESS;
+    };
+
+    UR_RETURN_ON_FAILURE(checkCLErr(USMMemcpy(
+        SrcQueue, blocking, HostAlloc, pSrc, size, numEventsInWaitList,
+        cl_adapter::cast<const cl_event *>(phEventWaitList), &HostCopyEvent)));
+
+    UR_RETURN_ON_FAILURE(
+        checkCLErr(USMMemcpy(DstQueue, blocking, pDst, HostAlloc, size, 1,
+                             &HostCopyEvent, &FinalCopyEvent)));
+
+    // If this is a blocking operation we can do our cleanup immediately,
+    // otherwise we need to defer it to an event callback.
+    if (blocking) {
+      CL_RETURN_ON_FAILURE(USMFree(CLContext, HostAlloc));
+      CL_RETURN_ON_FAILURE(clReleaseEvent(HostCopyEvent));
+      CL_RETURN_ON_FAILURE(clReleaseCommandQueue(MissingQueue));
+      if (phEvent) {
+        *phEvent = cl_adapter::cast<ur_event_handle_t>(FinalCopyEvent);
+      } else {
+        CL_RETURN_ON_FAILURE(clReleaseEvent(FinalCopyEvent));
+      }
+    } else {
+      if (phEvent) {
+        *phEvent = cl_adapter::cast<ur_event_handle_t>(FinalCopyEvent);
+        // We are going to release this event in our callback so we need to
+        // retain if the user wants a copy.
+        CL_RETURN_ON_FAILURE(clRetainEvent(FinalCopyEvent));
+      }
+
+      // This self destructs taking the event and allocation with it.
+      auto DeleterInfo = new AllocDeleterCallbackInfoWithQueue(
+          USMFree, CLContext, HostAlloc, MissingQueue);
+
+      CLErr = clSetEventCallback(
+          HostCopyEvent, CL_COMPLETE,
+          AllocDeleterCallback<AllocDeleterCallbackInfoWithQueue>, DeleterInfo);
+
+      if (CLErr != CL_SUCCESS) {
+        // We can attempt to recover gracefully by attempting to wait for the
+        // copy to finish and deleting the info struct here.
+        clWaitForEvents(1, &HostCopyEvent);
+        delete DeleterInfo;
+        clReleaseEvent(HostCopyEvent);
+        CL_RETURN_ON_FAILURE(CLErr);
+      }
+    }
+  } else {
+    CL_RETURN_ON_FAILURE(
+        USMMemcpy(cl_adapter::cast<cl_command_queue>(hQueue), blocking, pDst,
+                  pSrc, size, numEventsInWaitList,
+                  cl_adapter::cast<const cl_event *>(phEventWaitList),
+                  cl_adapter::cast<cl_event *>(phEvent)));
   }
 
-  return RetVal;
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

source/adapters/opencl/usm.cpp

@@ -0,0 +1,60 @@
+//===--------------------- usm.hpp - OpenCL Adapter -----------------------===//
+//
+// Copyright (C) 2024 Intel Corporation
+//
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM
+// Exceptions. See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "CL/cl_ext.h"
+#include <CL/cl.h>
+
+// This struct is intended to be used in conjunction with the below callback via
+// clSetEventCallback to release temporary allocations created by the adapter to
+// implement certain USM operations.
+//
+// Example usage:
+//
+// auto Info = new AllocDeleterCallbackInfo(USMFreeFuncPtr, Context,
+// Allocation); clSetEventCallback(USMOpEvent, CL_COMPLETE,
+// AllocDeleterCallback, Info);
+struct AllocDeleterCallbackInfo {
+  AllocDeleterCallbackInfo(clMemBlockingFreeINTEL_fn USMFree,
+                           cl_context CLContext, void *Allocation)
+      : USMFree(USMFree), CLContext(CLContext), Allocation(Allocation) {
+    clRetainContext(CLContext);
+  }
+  ~AllocDeleterCallbackInfo() {
+    USMFree(CLContext, Allocation);
+    clReleaseContext(CLContext);
+  }
+  AllocDeleterCallbackInfo(const AllocDeleterCallbackInfo &) = delete;
+  AllocDeleterCallbackInfo &
+  operator=(const AllocDeleterCallbackInfo &) = delete;
+
+  clMemBlockingFreeINTEL_fn USMFree;
+  cl_context CLContext;
+  void *Allocation;
+};
+
+struct AllocDeleterCallbackInfoWithQueue : AllocDeleterCallbackInfo {
+  AllocDeleterCallbackInfoWithQueue(clMemBlockingFreeINTEL_fn USMFree,
+                                    cl_context CLContext, void *Allocation,
+                                    cl_command_queue CLQueue)
+      : AllocDeleterCallbackInfo(USMFree, CLContext, Allocation),
+        CLQueue(CLQueue) {
+    clRetainContext(CLContext);
+  }
+  ~AllocDeleterCallbackInfoWithQueue() { clReleaseCommandQueue(CLQueue); }
+  AllocDeleterCallbackInfoWithQueue(const AllocDeleterCallbackInfoWithQueue &) =
+      delete;
+  AllocDeleterCallbackInfoWithQueue &
+  operator=(const AllocDeleterCallbackInfoWithQueue &) = delete;
+
+  cl_command_queue CLQueue;
+};
+
+template <class T>
+void AllocDeleterCallback(cl_event event, cl_int, void *pUserData);