Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/workflows/sycl_nightly.yml

@@ -16,7 +16,8 @@ jobs:
     with:
       build_cache_root: "/__w/"
       build_artifact_suffix: default
-      lts_config: "ocl_gen9;ocl_x64;hip_amdgpu"
+      build_configure_extra_args: ''
+      lts_config: "ocl_gen9;ocl_x64"
 
   windows_default:
     name: Windows

devops/actions/llvm_test_suite/action.yml

@@ -93,7 +93,6 @@ runs:
       echo "::endgroup::"
       echo "::group::SYCL_PI_TRACE=-1 sycl-ls"
       echo $LD_LIBRARY_PATH
-      ldd $PWD/toolchain/lib/libpi_hip.so
       SYCL_PI_TRACE=-1 sycl-ls
       echo "::endgroup::"
       cd build

sycl/doc/design/DeviceGlobal.md

@@ -563,15 +563,37 @@ instance of a device global variable in a `pi_program`.  This functionality is
 exposed as two new PI interfaces:
 
 ```
-pi_result piextCopyToDeviceVariable(pi_device Device, pi_program Program,
-  const char *name, const void *src, size_t count, size_t offset);
-
-pi_result piextCopyFromDeviceVariable(pi_device Device, pi_program Program,
-  const char *name, void *dst, size_t count, size_t offset);
+pi_result piextEnqueueDeviceVariableRead(pi_queue Queue, pi_program Program,
+                                         const char *Name, pi_bool BlockingRead,
+                                         size_t Count, size_t Offset, void *Dst,
+                                         pi_uint32 NumEventsInWaitList,
+                                         const pi_event *EventsWaitList,
+                                         pi_event *Event);
+
+pi_result piextEnqueueDeviceVariableWrite(pi_queue Queue, pi_program Program,
+                                          const char *Name,
+                                          pi_bool BlockingWrite, size_t Count,
+                                          size_t Offset, const void *Src,
+                                          pi_uint32 NumEventsInWaitList,
+                                          const pi_event *EventsWaitList,
+                                          pi_event *Event);
 ```
 
-In both cases the `name` parameter is the same as the `sycl-unique-id` string
-that is associated with the device global variable.
+The `piextEnqueueDeviceVariableRead` function reads `Count` bytes at byte-offset
+`Offset` from a device global variable in `Program` identified by the name
+`Name`. The read data is stored in `Dst`. Likewise, the
+`piextEnqueueDeviceVariableWrite` function reads `Count` bytes from `Dst` and
+stores them at byte-offset `Offset` in the device global variable in `Program`
+identified by the name `Name`.
+
+Both functions will enqueue the associated memory command on `Queue` where it
+will first wait for `NumEventsInWaitList` events in `EventsWaitList` to finish.
+`Event` will be populated with the event associated with resulting enqueued
+command. If either `BlockingRead` or `BlockingWrite` is `true` the call will
+block on the host until the enqueued command finishes execution.
+
+For `device_global` variables the `Name` parameter in calls to these functions
+is the same as the associated `sycl-unique-id` string.
 
 The Level Zero backend has existing APIs that can implement these PI
 interfaces.  The plugin first calls [`zeModuleGetGlobalPointer()`][8] to get a
@@ -616,8 +638,8 @@ depends upon implementation of that OpenCL extension.
 
 [10]: <opencl-extensions/cl_intel_global_variable_access.asciidoc>
 
-The CUDA backend has existing APIs `cudaMemcpyToSymbol()` and
-`cudaMemcpyFromSymbol()` which can be used to implement these PI interfaces.
+The CUDA backend has existing APIs `cuModuleGetGlobal()` and `cuMemcpyAsync()`
+which can be used to implement these PI interfaces.
 
 
 ## Design choices

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -976,8 +976,8 @@ struct _pi_buffer final : _pi_mem {
   _pi_buffer(pi_context Ctx, char *Mem, char *HostPtr, bool OwnZeMemHandle,
              _pi_mem *Parent = nullptr, size_t Origin = 0, size_t Size = 0,
              bool MemOnHost = false, bool ImportedHostPtr = false)
-      : _pi_mem(Ctx, HostPtr, OwnZeMemHandle, MemOnHost), ZeMem{Mem},
-        SubBuffer{Parent, Origin, Size} {}
+      : _pi_mem(Ctx, HostPtr, OwnZeMemHandle, MemOnHost, ImportedHostPtr),
+        ZeMem{Mem}, SubBuffer{Parent, Origin, Size} {}
 
   void *getZeHandle() override { return ZeMem; }
 

sycl/source/detail/error_handling/enqueue_kernel.cpp

@@ -268,6 +268,31 @@ bool handleInvalidWorkItemSize(const device_impl &DeviceImpl,
   return 0;
 }
 
+bool handleInvalidValue(const device_impl &DeviceImpl,
+                        const NDRDescT &NDRDesc) {
+  const plugin &Plugin = DeviceImpl.getPlugin();
+  RT::PiDevice Device = DeviceImpl.getHandleRef();
+
+  size_t MaxNWGs[] = {0, 0, 0};
+  Plugin.call<PiApiKind::piDeviceGetInfo>(
+      Device, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D, sizeof(MaxNWGs),
+      &MaxNWGs, nullptr);
+  for (unsigned int I = 0; I < NDRDesc.Dims; I++) {
+    size_t NWgs = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
+    if (NWgs > MaxNWGs[I])
+      throw sycl::nd_range_error(
+          "Number of work-groups exceed limit for dimension " +
+              std::to_string(I) + " : " + std::to_string(NWgs) + " > " +
+              std::to_string(MaxNWGs[I]),
+          PI_INVALID_VALUE);
+  }
+
+  // fallback
+  constexpr pi_result Error = PI_INVALID_VALUE;
+  throw runtime_error(
+      "Native API failed. Native API returns: " + codeToString(Error), Error);
+}
+
 bool handleError(pi_result Error, const device_impl &DeviceImpl,
                  pi_kernel Kernel, const NDRDescT &NDRDesc) {
   assert(Error != PI_SUCCESS &&
@@ -315,6 +340,9 @@ bool handleError(pi_result Error, const device_impl &DeviceImpl,
         "slice pitch) are not supported by device associated with queue",
         PI_INVALID_IMAGE_SIZE);
 
+  case PI_INVALID_VALUE:
+    return handleInvalidValue(DeviceImpl, NDRDesc);
+
     // TODO: Handle other error codes
 
   default:

sycl/source/detail/sycl_mem_obj_t.cpp

@@ -23,7 +23,7 @@ SYCLMemObjT::SYCLMemObjT(cl_mem MemObject, const context &SyclContext,
                   SizeInBytes, AvailableEvent, std::move(Allocator)) {}
 
 SYCLMemObjT::SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext,
-                         const size_t SizeInBytes, event AvailableEvent,
+                         const size_t, event AvailableEvent,
                          std::unique_ptr<SYCLMemObjAllocator> Allocator)
     : SYCLMemObjT(MemObject, SyclContext, true, AvailableEvent,
                   std::move(Allocator)) {}