[SYCL] Add device_global initialization and PI APIs (#7797)

This commit adds initialization of device_global without
device_image_scope. This is done as follows:
1. Allocate USM memory for the device_global.
2. Fill the USM memory with zeros. Store this event as a dependency for
related kernels.
3. When enqueuing a kernel using one or more device_globals, get the
initialization events for each of them (unless they have finished) and
add them as dependencies for the kernel launch.

Additionally, to write the USM pointer to the device_global, PI has two
new functions: piextEnqueueDeviceVariableWrite and
piextEnqueueDeviceVariableRead. These allow for writing to and reading
from device_globals in a given program. These are implemented for
OpenCL, L0, and CUDA. For CUDA sycl-post-link will generate a mapping
from the device_global's unique ID and the name in the binary.

Design doc:
[sycl/doc/design/DeviceGlobal.md](https://github.com/intel/llvm/blob/sycl/sycl/doc/design/DeviceGlobal.md)

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/detail/pi.def

@@ -147,4 +147,9 @@ _PI_API(piextUSMEnqueueMemset2D)
 _PI_API(piextUSMEnqueueMemcpy2D)
 
 _PI_API(piGetDeviceAndHostTimer)
+
+// Device global variable
+_PI_API(piextEnqueueDeviceGlobalVariableWrite)
+_PI_API(piextEnqueueDeviceGlobalVariableRead)
+
 #undef _PI_API

sycl/include/sycl/detail/pi.h

@@ -75,9 +75,11 @@
 // PI_EXT_ONEAPI_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT context info query
 // descriptors.
 // 12.22 Add piGetDeviceAndHostTimer to query device wall-clock timestamp
+// 12.23 Added new piextEnqueueDeviceGlobalVariableWrite and
+// piextEnqueueDeviceGlobalVariableRead functions.
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 22
+#define _PI_H_VERSION_MINOR 23
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -435,7 +437,9 @@ typedef enum {
   PI_COMMAND_TYPE_SVM_MEMCPY = 0x120A,
   PI_COMMAND_TYPE_SVM_MEMFILL = 0x120B,
   PI_COMMAND_TYPE_SVM_MAP = 0x120C,
-  PI_COMMAND_TYPE_SVM_UNMAP = 0x120D
+  PI_COMMAND_TYPE_SVM_UNMAP = 0x120D,
+  PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_READ = 0x418E,
+  PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_WRITE = 0x418F
 } _pi_command_type;
 
 typedef enum {
@@ -808,6 +812,7 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 /// must appear after the kernel name.
 #define __SYCL_PI_PROGRAM_METADATA_TAG_REQD_WORK_GROUP_SIZE                    \
   "@reqd_work_group_size"
+#define __SYCL_PI_PROGRAM_METADATA_GLOBAL_ID_MAPPING "@global_id_mapping"
 
 /// This struct is a record of the device binary information. If the Kind field
 /// denotes a portable binary type (SPIR-V or LLVM IR), the DeviceTargetSpec
@@ -1876,6 +1881,51 @@ __SYCL_EXPORT pi_result piextUSMEnqueueMemcpy2D(
     pi_uint32 num_events_in_waitlist, const pi_event *events_waitlist,
     pi_event *event);
 
+///
+/// Device global variable
+///
+
+/// API for writing data from host to a device global variable.
+///
+/// \param queue is the queue
+/// \param program is the program containing the device global variable
+/// \param blocking_write is true if the write should block
+/// \param name is the unique identifier for the device global variable
+/// \param count is the number of bytes to copy
+/// \param offset is the byte offset into the device global variable to start
+/// copying
+/// \param src is a pointer to where the data must be copied from
+/// \param num_events_in_wait_list is a number of events in the wait list
+/// \param event_wait_list is the wait list
+/// \param event is the resulting event
+pi_result piextEnqueueDeviceGlobalVariableWrite(
+    pi_queue queue, pi_program program, const char *name,
+    pi_bool blocking_write, size_t count, size_t offset, const void *src,
+    pi_uint32 num_events_in_wait_list, const pi_event *event_wait_list,
+    pi_event *event);
+
+/// API reading data from a device global variable to host.
+///
+/// \param queue is the queue
+/// \param program is the program containing the device global variable
+/// \param blocking_read is true if the read should block
+/// \param name is the unique identifier for the device global variable
+/// \param count is the number of bytes to copy
+/// \param offset is the byte offset into the device global variable to start
+/// copying
+/// \param dst is a pointer to where the data must be copied to
+/// \param num_events_in_wait_list is a number of events in the wait list
+/// \param event_wait_list is the wait list
+/// \param event is the resulting event
+pi_result piextEnqueueDeviceGlobalVariableRead(
+    pi_queue queue, pi_program program, const char *name, pi_bool blocking_read,
+    size_t count, size_t offset, void *dst, pi_uint32 num_events_in_wait_list,
+    const pi_event *event_wait_list, pi_event *event);
+
+///
+/// Plugin
+///
+
 /// API to get Plugin internal data, opaque to SYCL RT. Some devices whose
 /// device code is compiled by the host compiler (e.g. CPU emulators) may use it
 /// to access some device code functionality implemented in/behind the plugin.

sycl/include/sycl/ext/oneapi/device_global/device_global.hpp

@@ -135,13 +135,6 @@ class
   static_assert(is_property_list<property_list_t>::value,
                 "Property list is invalid.");
 
-  // TODO: Remove when support has been added for device_global without the
-  // device_image_scope property.
-  static_assert(
-      property_list_t::template has_property<device_image_scope_key>(),
-      "device_global without the device_image_scope property is currently "
-      "unavailable.");
-
   device_global() = default;
 
   device_global(const device_global &) = delete;

sycl/plugins/cuda/pi_cuda.cpp

@@ -670,31 +670,26 @@ _pi_program::_pi_program(pi_context ctxt)
 
 _pi_program::~_pi_program() { cuda_piContextRelease(context_); }
 
-bool get_kernel_metadata(std::string metadataName, const char *tag,
-                         std::string &kernelName) {
-  const size_t tagLength = strlen(tag);
-  const size_t metadataNameLength = metadataName.length();
-  if (metadataNameLength >= tagLength &&
-      metadataName.compare(metadataNameLength - tagLength, tagLength, tag) ==
-          0) {
-    kernelName = metadataName.substr(0, metadataNameLength - tagLength);
-    return true;
-  }
-  return false;
+std::pair<std::string, std::string>
+splitMetadataName(const std::string &metadataName) {
+  size_t splitPos = metadataName.rfind('@');
+  if (splitPos == std::string::npos)
+    return std::make_pair(metadataName, std::string{});
+  return std::make_pair(metadataName.substr(0, splitPos),
+                        metadataName.substr(splitPos, metadataName.length()));
 }
 
 pi_result _pi_program::set_metadata(const pi_device_binary_property *metadata,
                                     size_t length) {
   for (size_t i = 0; i < length; ++i) {
     const pi_device_binary_property metadataElement = metadata[i];
     std::string metadataElementName{metadataElement->Name};
-    std::string kernelName;
 
-    // If metadata is reqd_work_group_size record it for the corresponding
-    // kernel name.
-    if (get_kernel_metadata(metadataElementName,
-                            __SYCL_PI_PROGRAM_METADATA_TAG_REQD_WORK_GROUP_SIZE,
-                            kernelName)) {
+    auto [prefix, tag] = splitMetadataName(metadataElementName);
+
+    if (tag == __SYCL_PI_PROGRAM_METADATA_TAG_REQD_WORK_GROUP_SIZE) {
+      // If metadata is reqd_work_group_size, record it for the corresponding
+      // kernel name.
       size_t MDElemsSize = metadataElement->ValSize - sizeof(std::uint64_t);
 
       // Expect between 1 and 3 32-bit integer values.
@@ -709,9 +704,16 @@ pi_result _pi_program::set_metadata(const pi_device_binary_property *metadata,
       // Read values and pad with 1's for values not present.
       std::uint32_t reqdWorkGroupElements[] = {1, 1, 1};
       std::memcpy(reqdWorkGroupElements, ValuePtr, MDElemsSize);
-      kernelReqdWorkGroupSizeMD_[kernelName] =
+      kernelReqdWorkGroupSizeMD_[prefix] =
           std::make_tuple(reqdWorkGroupElements[0], reqdWorkGroupElements[1],
                           reqdWorkGroupElements[2]);
+    } else if (tag == __SYCL_PI_PROGRAM_METADATA_GLOBAL_ID_MAPPING) {
+      const char *metadataValPtr =
+          reinterpret_cast<const char *>(metadataElement->ValAddr) +
+          sizeof(std::uint64_t);
+      const char *metadataValPtrEnd =
+          metadataValPtr + metadataElement->ValSize - sizeof(std::uint64_t);
+      globalIDMD_[prefix] = std::string{metadataValPtr, metadataValPtrEnd};
     }
   }
   return PI_SUCCESS;
@@ -5505,6 +5507,82 @@ pi_result cuda_piextUSMGetMemAllocInfo(pi_context context, const void *ptr,
   return result;
 }
 
+pi_result cuda_piextEnqueueDeviceGlobalVariableWrite(
+    pi_queue queue, pi_program program, const char *name,
+    pi_bool blocking_write, size_t count, size_t offset, const void *src,
+    pi_uint32 num_events_in_wait_list, const pi_event *event_wait_list,
+    pi_event *event) {
+  assert(queue != nullptr);
+  assert(program != nullptr);
+
+  if (name == nullptr || src == nullptr)
+    return PI_ERROR_INVALID_VALUE;
+
+  // Since CUDA requires a the global variable to be referenced by name, we use
+  // metadata to find the correct name to access it by.
+  auto device_global_name_it = program->globalIDMD_.find(name);
+  if (device_global_name_it == program->globalIDMD_.end())
+    return PI_ERROR_INVALID_VALUE;
+  std::string device_global_name = device_global_name_it->second;
+
+  pi_result result = PI_SUCCESS;
+  try {
+    CUdeviceptr device_global = 0;
+    size_t device_global_size = 0;
+    result = PI_CHECK_ERROR(
+        cuModuleGetGlobal(&device_global, &device_global_size, program->get(),
+                          device_global_name.c_str()));
+
+    if (offset + count > device_global_size)
+      return PI_ERROR_INVALID_VALUE;
+
+    return cuda_piextUSMEnqueueMemcpy(
+        queue, blocking_write, reinterpret_cast<void *>(device_global + offset),
+        src, count, num_events_in_wait_list, event_wait_list, event);
+  } catch (pi_result error) {
+    result = error;
+  }
+  return result;
+}
+
+pi_result cuda_piextEnqueueDeviceGlobalVariableRead(
+    pi_queue queue, pi_program program, const char *name, pi_bool blocking_read,
+    size_t count, size_t offset, void *dst, pi_uint32 num_events_in_wait_list,
+    const pi_event *event_wait_list, pi_event *event) {
+  assert(queue != nullptr);
+  assert(program != nullptr);
+
+  if (name == nullptr || dst == nullptr)
+    return PI_ERROR_INVALID_VALUE;
+
+  // Since CUDA requires a the global variable to be referenced by name, we use
+  // metadata to find the correct name to access it by.
+  auto device_global_name_it = program->globalIDMD_.find(name);
+  if (device_global_name_it == program->globalIDMD_.end())
+    return PI_ERROR_INVALID_VALUE;
+  std::string device_global_name = device_global_name_it->second;
+
+  pi_result result = PI_SUCCESS;
+  try {
+    CUdeviceptr device_global = 0;
+    size_t device_global_size = 0;
+    result = PI_CHECK_ERROR(
+        cuModuleGetGlobal(&device_global, &device_global_size, program->get(),
+                          device_global_name.c_str()));
+
+    if (offset + count > device_global_size)
+      return PI_ERROR_INVALID_VALUE;
+
+    return cuda_piextUSMEnqueueMemcpy(
+        queue, blocking_read, dst,
+        reinterpret_cast<const void *>(device_global + offset), count,
+        num_events_in_wait_list, event_wait_list, event);
+  } catch (pi_result error) {
+    result = error;
+  }
+  return result;
+}
+
 // This API is called by Sycl RT to notify the end of the plugin lifetime.
 // TODO: add a global variable lifetime management code here (see
 // pi_level_zero.cpp for reference) Currently this is just a NOOP.
@@ -5685,6 +5763,11 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMEnqueueMemset2D, cuda_piextUSMEnqueueMemset2D)
   _PI_CL(piextUSMEnqueueMemcpy2D, cuda_piextUSMEnqueueMemcpy2D)
   _PI_CL(piextUSMGetMemAllocInfo, cuda_piextUSMGetMemAllocInfo)
+  // Device global variable
+  _PI_CL(piextEnqueueDeviceGlobalVariableWrite,
+         cuda_piextEnqueueDeviceGlobalVariableWrite)
+  _PI_CL(piextEnqueueDeviceGlobalVariableRead,
+         cuda_piextEnqueueDeviceGlobalVariableRead)
 
   _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj)
   _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler)

sycl/plugins/cuda/pi_cuda.hpp

@@ -773,6 +773,7 @@ struct _pi_program {
   // Metadata
   std::unordered_map<std::string, std::tuple<uint32_t, uint32_t, uint32_t>>
       kernelReqdWorkGroupSizeMD_;
+  std::unordered_map<std::string, std::string> globalIDMD_;
 
   constexpr static size_t MAX_LOG_SIZE = 8192u;
 

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -2029,6 +2029,20 @@ pi_result piextUSMEnqueuePrefetch(pi_queue, const void *, size_t,
   DIE_NO_IMPLEMENTATION;
 }
 
+pi_result piextEnqueueDeviceGlobalVariableWrite(pi_queue, pi_program,
+                                                const char *, pi_bool, size_t,
+                                                size_t, const void *, pi_uint32,
+                                                const pi_event *, pi_event *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextEnqueueDeviceGlobalVariableRead(pi_queue, pi_program,
+                                               const char *, pi_bool, size_t,
+                                               size_t, void *, pi_uint32,
+                                               const pi_event *, pi_event *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
 pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) {
   *OpaqueDataReturn = reinterpret_cast<void *>(PiESimdDeviceAccess);
   return PI_SUCCESS;

sycl/plugins/hip/pi_hip.cpp

@@ -5228,6 +5228,47 @@ pi_result hip_piextUSMGetMemAllocInfo(pi_context context, const void *ptr,
   return result;
 }
 
+pi_result hip_piextEnqueueDeviceGlobalVariableWrite(
+    pi_queue queue, pi_program program, const char *name,
+    pi_bool blocking_write, size_t count, size_t offset, const void *src,
+    pi_uint32 num_events_in_wait_list, const pi_event *event_wait_list,
+    pi_event *event) {
+  (void)queue;
+  (void)program;
+  (void)name;
+  (void)blocking_write;
+  (void)count;
+  (void)offset;
+  (void)src;
+  (void)num_events_in_wait_list;
+  (void)event_wait_list;
+  (void)event;
+
+  sycl::detail::pi::die(
+      "hip_piextEnqueueDeviceGlobalVariableWrite not implemented");
+  return {};
+}
+
+pi_result hip_piextEnqueueDeviceGlobalVariableRead(
+    pi_queue queue, pi_program program, const char *name, pi_bool blocking_read,
+    size_t count, size_t offset, void *dst, pi_uint32 num_events_in_wait_list,
+    const pi_event *event_wait_list, pi_event *event) {
+  (void)queue;
+  (void)program;
+  (void)name;
+  (void)blocking_read;
+  (void)count;
+  (void)offset;
+  (void)dst;
+  (void)num_events_in_wait_list;
+  (void)event_wait_list;
+  (void)event;
+
+  sycl::detail::pi::die(
+      "hip_piextEnqueueDeviceGlobalVariableRead not implemented");
+  return {};
+}
+
 // This API is called by Sycl RT to notify the end of the plugin lifetime.
 // TODO: add a global variable lifetime management code here (see
 // pi_level_zero.cpp for reference) Currently this is just a NOOP.
@@ -5401,6 +5442,11 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMEnqueueFill2D, hip_piextUSMEnqueueFill2D)
   _PI_CL(piextUSMEnqueueMemset2D, hip_piextUSMEnqueueMemset2D)
   _PI_CL(piextUSMGetMemAllocInfo, hip_piextUSMGetMemAllocInfo)
+  // Device global variable
+  _PI_CL(piextEnqueueDeviceGlobalVariableWrite,
+         hip_piextEnqueueDeviceGlobalVariableWrite)
+  _PI_CL(piextEnqueueDeviceGlobalVariableRead,
+         hip_piextEnqueueDeviceGlobalVariableRead)
 
   _PI_CL(piextKernelSetArgMemObj, hip_piextKernelSetArgMemObj)
   _PI_CL(piextKernelSetArgSampler, hip_piextKernelSetArgSampler)