[SYCL][Graphs] Squashed Graphs UR changes

- Squashed commit of command-buffer UR changes
- Level Zero implementation
- Stubs for all other adapters

Author: Bensuo

sycl/include/sycl/detail/pi.def

@@ -141,6 +141,16 @@ _PI_API(piextKernelSetArgSampler)
 
 _PI_API(piextPluginGetOpaqueData)
 
+/// command-buffer Extension
+_PI_API(piextCommandBufferCreate)
+_PI_API(piextCommandBufferRetain)
+_PI_API(piextCommandBufferRelease)
+_PI_API(piextCommandBufferFinalize)
+_PI_API(piextCommandBufferNDRangeKernel)
+_PI_API(piextCommandBufferMemcpyUSM)
+_PI_API(piextCommandBufferMemBufferCopy)
+_PI_API(piextCommandBufferMemBufferCopyRect)
+_PI_API(piextEnqueueCommandBuffer)
 _PI_API(piPluginGetLastError)
 
 _PI_API(piTearDown)

sycl/include/sycl/detail/pi.h

@@ -94,9 +94,9 @@
 // info query.
 // 12.32 Removed backwards compatibility of piextQueueCreateWithNativeHandle and
 // piextQueueGetNativeHandle
-
+// 12.33 Added command-buffer extension methods
 #define _PI_H_VERSION_MAJOR 13
-#define _PI_H_VERSION_MINOR 32
+#define _PI_H_VERSION_MINOR 33
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -474,6 +474,7 @@ typedef enum {
   PI_COMMAND_TYPE_SVM_MEMFILL = 0x120B,
   PI_COMMAND_TYPE_SVM_MAP = 0x120C,
   PI_COMMAND_TYPE_SVM_UNMAP = 0x120D,
+  PI_COMMAND_TYPE_EXT_COMMAND_BUFFER = 0x12A8,
   PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_READ = 0x418E,
   PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_WRITE = 0x418F
 } _pi_command_type;
@@ -2100,6 +2101,144 @@ __SYCL_EXPORT pi_result piGetDeviceAndHostTimer(pi_device Device,
                                                 uint64_t *DeviceTime,
                                                 uint64_t *HostTime);
 
+/// Command buffer extension
+struct _pi_ext_command_buffer;
+struct _pi_ext_sync_point;
+using pi_ext_command_buffer = _pi_ext_command_buffer *;
+using pi_ext_sync_point = pi_uint32;
+
+typedef enum {
+  PI_EXT_STRUCTURE_TYPE_COMMAND_BUFFER_DESC = 0
+} pi_ext_structure_type;
+
+struct pi_ext_command_buffer_desc final {
+  pi_ext_structure_type stype;
+  const void *pNext;
+  pi_queue_properties *properties;
+};
+
+/// API to create a command-buffer.
+/// \param context The context to associate the command-buffer with.
+/// \param device The device to associate the command-buffer with.
+/// \param desc Descriptor for the new command-buffer.
+/// \param ret_command_buffer Pointer to fill with the address of the new
+/// command-buffer.
+__SYCL_EXPORT pi_result
+piextCommandBufferCreate(pi_context context, pi_device device,
+                         const pi_ext_command_buffer_desc *desc,
+                         pi_ext_command_buffer *ret_command_buffer);
+
+/// API to increment the reference count of the command-buffer
+/// \param command_buffer The command_buffer to retain.
+__SYCL_EXPORT pi_result
+piextCommandBufferRetain(pi_ext_command_buffer command_buffer);
+
+/// API to decrement the reference count of the command-buffer. After the
+/// command_buffer reference count becomes zero and has finished execution, the
+/// command-buffer is deleted. \param command_buffer The command_buffer to
+/// release.
+__SYCL_EXPORT pi_result
+piextCommandBufferRelease(pi_ext_command_buffer command_buffer);
+
+/// API to stop command-buffer recording such that no more commands can be
+/// appended, and makes the command-buffer ready to enqueue on a command-queue.
+/// \param command_buffer The command_buffer to finalize.
+__SYCL_EXPORT pi_result
+piextCommandBufferFinalize(pi_ext_command_buffer command_buffer);
+
+/// API to append a kernel execution command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param kernel The kernel to append.
+/// \param work_dim Dimension of the kernel execution.
+/// \param global_work_offset Offset to use when executing kernel.
+/// \param global_work_size Global work size to use when executing kernel.
+/// \param local_work_size Local work size to use when executing kernel.
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this kernel execution.
+__SYCL_EXPORT pi_result piextCommandBufferNDRangeKernel(
+    pi_ext_command_buffer command_buffer, pi_kernel kernel, pi_uint32 work_dim,
+    const size_t *global_work_offset, const size_t *global_work_size,
+    const size_t *local_work_size, pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to append a USM memcpy command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param dst_ptr is the location the data will be copied
+/// \param src_ptr is the data to be copied
+/// \param size is number of bytes to copy
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this memory operation.
+__SYCL_EXPORT pi_result piextCommandBufferMemcpyUSM(
+    pi_ext_command_buffer command_buffer, void *dst_ptr, const void *src_ptr,
+    size_t size, pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to append a mem buffer copy command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param src_buffer is the data to be copied
+/// \param dst_buffer is the location the data will be copied
+/// \param src_offset offset into \p src_buffer
+/// \param dst_offset offset into \p dst_buffer
+/// \param size is number of bytes to copy
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this memory operation.
+__SYCL_EXPORT pi_result piextCommandBufferMemBufferCopy(
+    pi_ext_command_buffer command_buffer, pi_mem src_buffer, pi_mem dst_buffer,
+    size_t src_offset, size_t dst_offset, size_t size,
+    pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to append a rectangular mem buffer copy command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param src_buffer is the data to be copied
+/// \param dst_buffer is the location the data will be copied
+/// \param src_origin offset for the start of the region to copy in src_buffer
+/// \param dst_origin offset for the start of the region to copy in dst_buffer
+/// \param region The size of the region to be copied
+/// \param src_row_pitch Row pitch for the src data
+/// \param src_slice_pitch Slice pitch for the src data
+/// \param dst_row_pitch Row pitch for the dst data
+/// \param dst_slice_pitch Slice pitch for the dst data
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this memory operation.
+__SYCL_EXPORT pi_result piextCommandBufferMemBufferCopyRect(
+    pi_ext_command_buffer command_buffer, pi_mem src_buffer, pi_mem dst_buffer,
+    pi_buff_rect_offset src_origin, pi_buff_rect_offset dst_origin,
+    pi_buff_rect_region region, size_t src_row_pitch, size_t src_slice_pitch,
+    size_t dst_row_pitch, size_t dst_slice_pitch,
+    pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to submit the command-buffer to queue for execution, returns an error if
+/// command-buffer not finalized or another instance of same command-buffer
+/// currently executing.
+/// \param command_buffer The command-buffer to be submitted.
+/// \param queue The PI queue to submit on.
+/// \param num_events_in_wait_list The number of events that this execution
+/// depends on.
+/// \param event_wait_list List of pi_events to wait on.
+/// \param event The pi_event associated with this enqueue.
+__SYCL_EXPORT pi_result
+piextEnqueueCommandBuffer(pi_ext_command_buffer command_buffer, pi_queue queue,
+                          pi_uint32 num_events_in_wait_list,
+                          const pi_event *event_wait_list, pi_event *event);
+
 struct _pi_plugin {
   // PI version supported by host passed to the plugin. The Plugin
   // checks and writes the appropriate Function Pointers in

sycl/include/sycl/detail/pi.hpp

@@ -148,6 +148,9 @@ using PiMemObjectType = ::pi_mem_type;
 using PiMemImageChannelOrder = ::pi_image_channel_order;
 using PiMemImageChannelType = ::pi_image_channel_type;
 using PiKernelCacheConfig = ::pi_kernel_cache_config;
+using PiExtSyncPoint = ::pi_ext_sync_point;
+using PiExtCommandBuffer = ::pi_ext_command_buffer;
+using PiExtCommandBufferDesc = ::pi_ext_command_buffer_desc;
 
 __SYCL_EXPORT void contextSetExtendedDeleter(const sycl::context &constext,
                                              pi_context_extended_deleter func,

sycl/plugins/cuda/CMakeLists.txt

@@ -79,6 +79,8 @@ add_sycl_plugin(cuda
     "../unified_runtime/ur/adapters/cuda/tracing.cpp"
     "../unified_runtime/ur/adapters/cuda/ur_interface_loader.cpp"
     "../unified_runtime/ur/adapters/cuda/usm.cpp"
+    "../unified_runtime/ur/adapters/cuda/command_buffer.hpp"
+    "../unified_runtime/ur/adapters/cuda/command_buffer.cpp"
     # --- 
     "${sycl_inc_dir}/sycl/detail/pi.h"
     "${sycl_inc_dir}/sycl/detail/pi.hpp"

sycl/plugins/cuda/pi_cuda.cpp

@@ -186,6 +186,19 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piGetDeviceAndHostTimer, pi2ur::piGetDeviceAndHostTimer)
   _PI_CL(piPluginGetBackendOption, pi2ur::piPluginGetBackendOption)
 
+  // command-buffer
+  _PI_CL(piextCommandBufferCreate, pi2ur::piextCommandBufferCreate)
+  _PI_CL(piextCommandBufferRetain, pi2ur::piextCommandBufferRetain)
+  _PI_CL(piextCommandBufferRelease, pi2ur::piextCommandBufferRelease)
+  _PI_CL(piextCommandBufferNDRangeKernel,
+         pi2ur::piextCommandBufferNDRangeKernel)
+  _PI_CL(piextCommandBufferMemcpyUSM, pi2ur::piextCommandBufferMemcpyUSM)
+  _PI_CL(piextCommandBufferMemBufferCopy,
+         pi2ur::piextCommandBufferMemBufferCopy)
+  _PI_CL(piextCommandBufferMemBufferCopyRect,
+         pi2ur::piextCommandBufferMemBufferCopyRect)
+  _PI_CL(piextEnqueueCommandBuffer, pi2ur::piextEnqueueCommandBuffer)
+
 #undef _PI_CL
 
   return PI_SUCCESS;

sycl/plugins/cuda/pi_cuda.hpp

@@ -25,6 +25,7 @@
 #define _PI_CUDA_PLUGIN_VERSION_STRING                                         \
   _PI_PLUGIN_VERSION_STRING(_PI_CUDA_PLUGIN_VERSION)
 
+#include <ur/adapters/cuda/command_buffer.hpp>
 #include <ur/adapters/cuda/context.hpp>
 #include <ur/adapters/cuda/device.hpp>
 #include <ur/adapters/cuda/event.hpp>
@@ -76,4 +77,8 @@ struct _pi_sampler : ur_sampler_handle_t_ {
   using ur_sampler_handle_t_::ur_sampler_handle_t_;
 };
 
+struct _pi_ext_command_buffer : ur_exp_command_buffer_handle_t_ {
+  using ur_exp_command_buffer_handle_t_::ur_exp_command_buffer_handle_t_;
+};
+
 #endif // PI_CUDA_HPP

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -2098,6 +2098,70 @@ pi_result piextEnqueueDeviceGlobalVariableRead(pi_queue, pi_program,
   DIE_NO_IMPLEMENTATION;
 }
 
+pi_result piextCommandBufferCreate(pi_context context, pi_device device,
+                                   const pi_ext_command_buffer_desc *desc,
+                                   pi_ext_command_buffer *ret_command_buffer) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferRetain(pi_ext_command_buffer command_buffer) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferRelease(pi_ext_command_buffer command_buffer) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferFinalize(pi_ext_command_buffer command_buffer) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferNDRangeKernel(
+    pi_ext_command_buffer command_buffer, pi_kernel kernel, pi_uint32 work_dim,
+    const size_t *global_work_offset, const size_t *global_work_size,
+    const size_t *local_work_size, pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result
+piextCommandBufferMemcpyUSM(pi_ext_command_buffer command_buffer, void *dst_ptr,
+                            const void *src_ptr, size_t size,
+                            pi_uint32 num_sync_points_in_wait_list,
+                            const pi_ext_sync_point *sync_point_wait_list,
+                            pi_ext_sync_point *sync_point) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferMemBufferCopy(
+    pi_ext_command_buffer command_buffer, pi_mem src_buffer, pi_mem dst_buffer,
+    size_t src_offset, size_t dst_offset, size_t size,
+    pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextCommandBufferMemBufferCopyRect(
+    pi_ext_command_buffer command_buffer, pi_mem src_buffer, pi_mem dst_buffer,
+    pi_buff_rect_offset src_origin, pi_buff_rect_offset dst_origin,
+    pi_buff_rect_region region, size_t src_row_pitch, size_t src_slice_pitch,
+    size_t dst_row_pitch, size_t dst_slice_pitch,
+    pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer command_buffer,
+                                    pi_queue queue,
+                                    pi_uint32 num_events_in_wait_list,
+                                    const pi_event *event_wait_list,
+                                    pi_event *event) {
+  DIE_NO_IMPLEMENTATION;
+}
+
 pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) {
   *OpaqueDataReturn = reinterpret_cast<void *>(PiESimdDeviceAccess);
   return PI_SUCCESS;

sycl/plugins/esimd_emulator/pi_esimd_emulator.hpp

@@ -218,4 +218,6 @@ struct _pi_kernel : _pi_object {
   _pi_kernel() {}
 };
 
+struct _pi_ext_command_buffer {};
+
 #include <sycl/ext/intel/esimd/emu/detail/esimd_emulator_device_interface.hpp>