Skip to content

Commit

Permalink
[SYCL][Graph] Implement dynamic command-groups
Browse files Browse the repository at this point in the history
Implement Dynamic Command-Group feature specified in
PR [[SYCL][Graph] Add specification for kernel binary updates](intel#14896)

This feature enables updating `ur_kernel_handle_t` objects in graph nodes
between executions as well as parameters and execution range of nodes.

This functionality is currently supported on CUDA & HIP which are used
for testing in the new E2E tests. Level Zero support will follow
shortly, resulting in the removal of the `XFAIL` labels from the E2E
tests.

The code for adding nodes to a graph has been refactored to split out
verification of edges, and marking memory objects used in a node, as
separate helper functions. This allows path for adding a command-group
node to do this functions over each CG in the list before creating the
node itself.

The `dynamic_parameter_impl` code has also been refactored so the code
is shared for updating a dynamic parameter used in both a regular kernel
node and a dynamic command-group node.

See the addition to the design doc for further details on the
implementation.
  • Loading branch information
EwanC committed Oct 17, 2024
1 parent e2075c7 commit 6f38259
Show file tree
Hide file tree
Showing 25 changed files with 1,723 additions and 203 deletions.
17 changes: 17 additions & 0 deletions sycl/doc/design/CommandGraph.md
Original file line number Diff line number Diff line change
Expand Up @@ -282,6 +282,23 @@ 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.

### Dynamic Command-Group

To implement the `dynamic_command_group` class for updating the command-groups (CG)
associated with nodes, the CG member of the node implementation class changes
from a `std::unique_ptr` to a `std::shared_ptr` so that multiple nodes and the
`dynamic_command_group_impl` object can share the same CG object. This avoids
the overhead of having to allocate and free copies of the CG when a new active
CG is selected.

The `dynamic_command_group_impl` class contains weak pointers to the nodes which
have been created with it, so that when a new active CG is selected it can
propagate the change to those nodes. The `node_impl` class also contains a
reference to the dynamic command-group that created it, so that when the graph
is finalized each node can use the list of kernels in its dynamic command-group
as part of the `urCommandBufferAppendKernelLaunchExp` call to pass the possible
alternative kernels.

## Optimizations
### Interactions with Profiling

Expand Down
46 changes: 46 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/graph.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -93,6 +93,7 @@ class node_impl;
class graph_impl;
class exec_graph_impl;
class dynamic_parameter_impl;
class dynamic_command_group_impl;
} // namespace detail

enum class node_type {
Expand Down Expand Up @@ -213,6 +214,23 @@ class depends_on_all_leaves : public ::sycl::detail::DataLessProperty<
} // namespace node
} // namespace property

class __SYCL_EXPORT dynamic_command_group {
public:
dynamic_command_group(
const command_graph<graph_state::modifiable> &Graph,
const std::vector<std::function<void(handler &)>> &CGFList);

size_t get_active_cgf() const;
void set_active_cgf(size_t Index);

private:
template <class Obj>
friend const decltype(Obj::impl) &
sycl::detail::getSyclObjImpl(const Obj &SyclObject);

std::shared_ptr<detail::dynamic_command_group_impl> impl;
};

namespace detail {
// Templateless modifiable command-graph base class.
class __SYCL_EXPORT modifiable_command_graph {
Expand Down Expand Up @@ -269,6 +287,28 @@ class __SYCL_EXPORT modifiable_command_graph {
return Node;
}

/// Add a Dynamic command-group node to the graph.
/// @param DynamicCG Dynamic command-group function to create node with.
/// @param PropList Property list used to pass [0..n] predecessor nodes.
/// @return Constructed node which has been added to the graph.
node add(dynamic_command_group &DynamicCG,
const property_list &PropList = {}) {
if (PropList.has_property<property::node::depends_on>()) {
auto Deps = PropList.get_property<property::node::depends_on>();
node Node = addImpl(DynamicCG, Deps.get_dependencies());
if (PropList.has_property<property::node::depends_on_all_leaves>()) {
addGraphLeafDependencies(Node);
}
return Node;
}

node Node = addImpl(DynamicCG, {});
if (PropList.has_property<property::node::depends_on_all_leaves>()) {
addGraphLeafDependencies(Node);
}
return Node;
}

/// Add a dependency between two nodes.
/// @param Src Node which will be a dependency of \p Dest.
/// @param Dest Node which will be dependent on \p Src.
Expand Down Expand Up @@ -328,6 +368,12 @@ class __SYCL_EXPORT modifiable_command_graph {
modifiable_command_graph(const std::shared_ptr<detail::graph_impl> &Impl)
: impl(Impl) {}

/// Template-less implementation of add() for dynamic command-group nodes.
/// @param DynCGF Dynamic Command-group function object to add.
/// @param Dep List of predecessor nodes.
/// @return Node added to the graph.
node addImpl(dynamic_command_group &DynCGF, const std::vector<node> &Dep);

/// Template-less implementation of add() for CGF nodes.
/// @param CGF Command-group function to add.
/// @param Dep List of predecessor nodes.
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/handler.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -3376,6 +3376,7 @@ class __SYCL_EXPORT handler {
size_t Size, bool Block = false);
friend class ext::oneapi::experimental::detail::graph_impl;
friend class ext::oneapi::experimental::detail::dynamic_parameter_impl;
friend class ext::oneapi::experimental::detail::dynamic_command_group_impl;

bool DisableRangeRounding();

Expand Down
Loading

0 comments on commit 6f38259

Please sign in to comment.