[SYCL][Graph] Implementation of explicit update with indices

sycl/doc/design/CommandGraph.md

@@ -46,6 +46,7 @@ with the following entry-points:
 | `urCommandBufferAppendMemBufferReadRectExp`  | Append a rectangular memory read command to a command-buffer object. |
 | `urCommandBufferAppendMemBufferFillExp`      | Append a memory fill command to a command-buffer object. |
 | `urCommandBufferEnqueueExp`                  | Submit command-buffer to a command-queue for execution. |
+| `urCommandBufferUpdateKernelLaunchExp`       | Updates the parameters of a previous kernel launch command. |
 
 See the [UR EXP-COMMAND-BUFFER](https://oneapi-src.github.io/unified-runtime/core/EXP-COMMAND-BUFFER.html)
 specification for more details.
@@ -230,6 +231,47 @@ on buffer usage in a graph so that their lifetime semantics are compatible with
 a lazy work execution model. However these changes to storage lifetimes have not
 yet been implemented.
 
+## Graph Update
+
+### Design Challenges
+
+Graph update faces significant design challenges in SYCL:
+
+* Lambda capture order is explicitly undefined in C++, so the user cannot reason
+  about the indices of arguments captured by kernel lambdas.
+* Once arguments have been captured the actual type information is lost in the
+  transition through the integration header and extracting arguments in the SYCL
+  runtime, therefore we cannot automatically match new argument values by
+  querying the captured arguments without significant possibility for
+  collisions. For example, if a kernel captures two USM pointers and the user
+  wishes to update one, we cannot reason about which pointer they actually want
+  to update when we only know that: they are pointer args of a certain size.
+
+The current approach is to limit graph update to the explicit APIs and where the
+user is using `handler::set_arg()` or some equivalent to manually set kernel
+arguments using indices. Therefore when updating we can use indices to avoid
+collisions. In practice there are only a few current scenarios where `set_arg()`
+can be used:
+
+* The proposed ["Free Function Kernel"
+  extension](../extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc)
+* OpenCL interop kernels created from SPIR-V source at runtime.
+
+A possible future workaround lambda capture issues could be "Whole-Graph Update"
+where if we can guarantee that lambda capture order is the same across two
+different recordings we can then match parameter order when updating.
+
+### Scheduler Integration
+
+Graph updates in the runtime are sychronous calls however they can optionally be
+done through the scheduler using a new command,
+`sycl::detail::UpdateCommandBufferCommand`. This is needed when dealing with
+accessor updates. Since a new buffer which the user creates for updating may not
+yet have been lazily initialized on device we schedule a new command which has
+requirements for these new accessors to correctly trigger allocations before
+updating. This is similar to how individual graph commands are enqueued when
+accessors are used in a graph node.
+
 ## Backend Implementation
 
 Implementation of UR command-buffers

sycl/include/sycl/detail/pi.def

@@ -183,6 +183,7 @@ _PI_API(piextCommandBufferFillUSM)
 _PI_API(piextCommandBufferPrefetchUSM)
 _PI_API(piextCommandBufferAdviseUSM)
 _PI_API(piextEnqueueCommandBuffer)
+_PI_API(piextCommandBufferUpdateKernelLaunch)
 
 _PI_API(piextUSMPitchedAlloc)
 

sycl/include/sycl/detail/pi.h

@@ -154,9 +154,10 @@
 // 15.44 Add coarse-grain memory advice flag for HIP.
 // 15.45 Added piextKernelSuggestMaxCooperativeGroupCount and
 //       piextEnqueueCooperativeKernelLaunch.
+// 15.46 Add CommandBuffer update definitions
 
 #define _PI_H_VERSION_MAJOR 15
-#define _PI_H_VERSION_MINOR 45
+#define _PI_H_VERSION_MINOR 46
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -443,6 +444,9 @@ typedef enum {
   // Composite device
   PI_EXT_ONEAPI_DEVICE_INFO_COMPONENT_DEVICES = 0x20111,
   PI_EXT_ONEAPI_DEVICE_INFO_COMPOSITE_DEVICE = 0x20112,
+
+  // Command Buffers
+  PI_EXT_ONEAPI_DEVICE_INFO_COMMAND_BUFFER_UPDATE_SUPPORT = 0x20113,
 } _pi_device_info;
 
 typedef enum {
@@ -2314,7 +2318,10 @@ __SYCL_EXPORT pi_result piGetDeviceAndHostTimer(pi_device Device,
 /// Command buffer extension
 struct _pi_ext_command_buffer;
 struct _pi_ext_sync_point;
+struct _pi_ext_command_buffer_command;
+
 using pi_ext_command_buffer = _pi_ext_command_buffer *;
+using pi_ext_command_buffer_command = _pi_ext_command_buffer_command *;
 using pi_ext_sync_point = pi_uint32;
 
 typedef enum {
@@ -2324,7 +2331,40 @@ typedef enum {
 struct pi_ext_command_buffer_desc final {
   pi_ext_structure_type stype;
   const void *pNext;
-  pi_queue_properties *properties;
+  pi_bool is_updateable;
+};
+
+// Command Buffer Update types
+struct pi_ext_command_buffer_update_memobj_arg_desc_t final {
+  uint32_t arg_index;
+  const pi_mem_obj_property *properties;
+  pi_mem new_mem_obj;
+};
+
+struct pi_ext_command_buffer_update_pointer_arg_desc_t final {
+  uint32_t arg_index;
+  void *new_ptr;
+};
+
+struct pi_ext_command_buffer_update_value_arg_desc_t final {
+  uint32_t arg_index;
+  uint32_t arg_size;
+  void *new_value;
+};
+
+struct pi_ext_command_buffer_update_kernel_launch_desc final {
+  uint32_t num_mem_obj_args;
+  uint32_t num_ptr_args;
+  uint32_t num_value_args;
+  uint32_t num_work_dim;
+
+  pi_ext_command_buffer_update_memobj_arg_desc_t *mem_obj_arg_list;
+  pi_ext_command_buffer_update_pointer_arg_desc_t *ptr_arg_list;
+  pi_ext_command_buffer_update_value_arg_desc_t *value_arg_list;
+
+  size_t *global_work_offset;
+  size_t *global_work_size;
+  size_t *local_work_size;
 };
 
 /// API to create a command-buffer.
@@ -2368,12 +2408,14 @@ piextCommandBufferFinalize(pi_ext_command_buffer command_buffer);
 /// \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.
+/// \param command Return pointer to the command representing 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);
+    pi_ext_sync_point *sync_point, pi_ext_command_buffer_command *command);
 
 /// API to append a USM memcpy command to the command-buffer.
 /// \param command_buffer The command-buffer to append onto.
@@ -2601,6 +2643,14 @@ 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);
 
+/// API to update a kernel launch command inside of a command-buffer.
+/// @param command The command to be updated.
+/// @param desc Descriptor which describes the updated parameters of the kernel
+/// launch.
+__SYCL_EXPORT pi_result piextCommandBufferUpdateKernelLaunch(
+    pi_ext_command_buffer_command command,
+    pi_ext_command_buffer_update_kernel_launch_desc *desc);
+
 /// API to destroy bindless unsampled image handles.
 ///
 /// \param context is the pi_context

sycl/include/sycl/detail/pi.hpp

@@ -156,6 +156,7 @@ 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;
+using PiExtCommandBufferCommand = ::pi_ext_command_buffer_command;
 using PiPeerAttr = ::pi_peer_attr;
 using PiImageHandle = ::pi_image_handle;
 using PiImageMemHandle = ::pi_image_mem_handle;

sycl/include/sycl/detail/property_helper.hpp

@@ -47,8 +47,9 @@ enum DataLessPropKind {
   GraphAssumeDataOutlivesBuffer = 22,
   GraphAssumeBufferOutlivesGraph = 23,
   GraphDependOnAllLeaves = 24,
+  GraphUpdateable = 25,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 24,
+  LastKnownDataLessPropKind = 25,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -8,10 +8,13 @@
 
 #pragma once
 
-#include <sycl/context.hpp>                // for context
-#include <sycl/detail/export.hpp>          // for __SYCL_EXPORT
+#include <sycl/accessor.hpp>
+#include <sycl/context.hpp>       // for context
+#include <sycl/detail/export.hpp> // for __SYCL_EXPORT
+#include <sycl/detail/kernel_desc.hpp>
 #include <sycl/detail/property_helper.hpp> // for DataLessPropKind, PropWith...
 #include <sycl/device.hpp>                 // for device
+#include <sycl/nd_range.hpp>
 #include <sycl/properties/property_traits.hpp> // for is_property, is_property_of
 #include <sycl/property_list.hpp>              // for property_list
 
@@ -30,6 +33,15 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 
+/// State to template the command_graph class on.
+enum class graph_state {
+  modifiable, ///< In modifiable state, commands can be added to graph.
+  executable, ///< In executable state, the graph is ready to execute.
+};
+
+// Forward declare Graph class
+template <graph_state State> class command_graph;
+
 namespace detail {
 // List of sycl features and extensions which are not supported by graphs. Used
 // for throwing errors when these features are used with graphs.
@@ -73,15 +85,9 @@ UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) {
 class node_impl;
 class graph_impl;
 class exec_graph_impl;
-
+class dynamic_parameter_impl;
 } // namespace detail
 
-/// State to template the command_graph class on.
-enum class graph_state {
-  modifiable, ///< In modifiable state, commands can be added to graph.
-  executable, ///< In executable state, the graph is ready to execute.
-};
-
 enum class node_type {
   empty = 0,
   subgraph = 1,
@@ -113,6 +119,13 @@ class __SYCL_EXPORT node {
   /// submission.
   static node get_node_from_event(event nodeEvent);
 
+  /// Update the ND-Range of this node if it is a kernel execution node
+  template <int Dimensions>
+  void update_nd_range(nd_range<Dimensions> executionRange);
+
+  /// Update the Range of this node if it is a kernel execution node
+  template <int Dimensions> void update_range(range<Dimensions> executionRange);
+
 private:
   node(const std::shared_ptr<detail::node_impl> &Impl) : impl(Impl) {}
 
@@ -146,6 +159,14 @@ class assume_buffer_outlives_graph
 public:
   assume_buffer_outlives_graph() = default;
 };
+
+/// Property passed to command_graph<graph_state::modifiable>::finalize() to
+/// mark the resulting executable command_graph as able to be updated.
+class updateable
+    : public ::sycl::detail::DataLessProperty<::sycl::detail::GraphUpdateable> {
+public:
+  updateable() = default;
+};
 } // namespace graph
 
 namespace node {
@@ -336,12 +357,24 @@ class __SYCL_EXPORT executable_command_graph {
   /// @param Graph Graph to use the inputs and outputs of.
   void update(const command_graph<graph_state::modifiable> &Graph);
 
+  /// Updates a single node in this graph based on the contents of the provided
+  /// node.
+  /// @param Node The node to use for updating the graph.
+  void update(const node &Node);
+
+  /// Updates a number of nodes in this graph based on the contents of the
+  /// provided nodes.
+  /// @param Nodes The nodes to use for updating the graph.
+  void update(const std::vector<node> &Nodes);
+
 protected:
   /// Constructor used by internal runtime.
   /// @param Graph Detail implementation class to construct with.
   /// @param Ctx Context to use for graph.
+  /// @param PropList Optional list of properties to pass.
   executable_command_graph(const std::shared_ptr<detail::graph_impl> &Graph,
-                           const sycl::context &Ctx);
+                           const sycl::context &Ctx,
+                           const property_list &PropList = {});
 
   template <class Obj>
   friend decltype(Obj::impl)
@@ -385,13 +418,60 @@ class command_graph : public detail::modifiable_command_graph {
 template <>
 class command_graph<graph_state::executable>
     : public detail::executable_command_graph {
-
 protected:
   friend command_graph<graph_state::executable>
   detail::modifiable_command_graph::finalize(const sycl::property_list &) const;
   using detail::executable_command_graph::executable_command_graph;
 };
 
+namespace detail {
+class __SYCL_EXPORT dynamic_parameter_base {
+public:
+  dynamic_parameter_base(
+      sycl::ext::oneapi::experimental::command_graph<graph_state::modifiable>
+          Graph);
+
+  void register_with_node(handler &CGH, int ArgIndex);
+
+protected:
+  void updateValue(void *NewValue, size_t Size);
+
+  void updateAccessor(sycl::detail::AccessorBaseHost *Acc);
+  std::shared_ptr<dynamic_parameter_impl> impl;
+};
+} // namespace detail
+
+template <typename ValueT>
+class dynamic_parameter : public detail::dynamic_parameter_base {
+  static constexpr bool IsAccessor =
+      std::is_base_of_v<sycl::detail::AccessorBaseHost, ValueT>;
+  static constexpr sycl::detail::kernel_param_kind_t ParamType =
+      IsAccessor ? sycl::detail::kernel_param_kind_t::kind_accessor
+                 : std::is_pointer_v<ValueT>
+                       ? sycl::detail::kernel_param_kind_t::kind_pointer
+                       : sycl::detail::kernel_param_kind_t::kind_std_layout;
+
+public:
+  dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph)
+      : detail::dynamic_parameter_base(Graph), MValue() {}
+
+  dynamic_parameter(ValueT InitialValue,
+                    experimental::command_graph<graph_state::modifiable> Graph)
+      : detail::dynamic_parameter_base(Graph), MValue(InitialValue) {}
+
+  void update(const ValueT &NewValue) {
+    MValue = NewValue;
+    if constexpr (IsAccessor) {
+      detail::dynamic_parameter_base::updateAccessor(&MValue);
+    } else {
+      detail::dynamic_parameter_base::updateValue(&MValue, sizeof(ValueT));
+    }
+  }
+
+private:
+  ValueT MValue;
+};
+
 /// Additional CTAD deduction guide.
 template <graph_state State = graph_state::modifiable>
 command_graph(const context &SyclContext, const device &SyclDevice,