[SYCL][Graph] Add External Events APIs #355

mfrancepillois · 2024-02-13T17:41:26Z

Updates specification with new APIs for handling external events in Record&Replay and Explicit mode.
Motivations:

Allow synchronization between separate graphs to support workloads using similar features in other APIs.
Allowing synchronizing on host part-way through a graph execution, for example to launch some other host-side work once part of the graph has executed.

Major changes:

Introduce new concept of "external" graph events - events which can be use for synchronization outside of their associated graph, both in other graphs and on the host.
Extend normal SYCL events rather than creating new types
Codify types of SYCL events, including what is now known as "graph-limited" events, which give a name to events returned from queue graph recordings which may not be used outside of their associated graph.
queue::ext_oneapi_external_event() used to obtain an external event in the Record & Replay API.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

EwanC · 2024-02-14T09:08:59Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+If a graph containing one or more external events is run multiple times, 
+the external events are automatically reset when the graph is resubmitted to 
+the queue.


After reading through this, I'm think about whether we want to overload SYCL events with "external events" at all, as resetting an event isn't a concept in SYCL events and i'm not sure if we'll run into issues in the DPC++ runtime code because it assumes this can't happen. It's also not very intuitive to a user that they now need to maintain the knowledge of if a sycl::event variable is an external event, a graph-limited event, or a normal event.

An example of an alternative would be to define our own synchronization primitive that can be reset, let's call it
a "semaphore" for the sake of argument. This is what OpenCL have done in cl_khr_semaphore, then if we defined an API for it like:

handler.wait_semaphore() handler.signal_semaphore() handler.reset_semaphore() // might not need this, not in cl_khr_semaphore

It could be used in CGF functions of both graph nodes and queue submission and would provide a simpler API for achieving the same goals. Need to check if it does indeed solve the same usecases, be implementable across backends, and it could probably be an extension in itself that we are compatible with (which is more work).

So currently SYCL events don't have a concept of being reset, but I'm not sure this is necessarily a concern as prior to SYCL graphs a DAG could only be traversed once, so it only makes sense for events to be set once. But if a DAG (in the form of a command_graph) can be traversed multiple times I think it's a reasonable extension of the SYCL execution model that events within that graph are reset each time the graph is executed.

I do agree that it's a bit odd from the perspective of non-graph SYCL code for an event to be reset, though I would be wary of creating a new type as this could create duplicate mechanisms for doing the same thing in non-graph code. Perhaps we need a new concept to represent this to the existing SYCL execution model, such as a multi-shot event which can be complete and reset, and a way to query for this.

I think it could also be useful to have an API which can be used to query whether an event is external or graph-restricted, so that code which is simply passed an event object can create an "external event" if necessary.

That's for the comment Gordon, it's good to know there's not an immediate objection to adding reset semantics to an event.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

Bensuo · 2024-02-21T18:58:36Z

The use-case of sharing events between graphs seems pretty easy to deal with and fairly problem free, but I feel like the use-case of waiting on graph events externally from normal user code presents a lot of issues that this PR has to solve. This includes things like:

If we're building on top of sycl::event which have no concept of being reset, how do we handle this?
If we do allow resetting events how does the user reason about when this may happen if there are multiple executions in flight?

The use case for this kind of event seems to be signalling to user code that some portion of the graph has finished so that the application can do something with this information. However, I feel like rather than supporting this directly in the API this could simply be accomplished with a host_task that depends on some other nodes which can be used to signal in user code when those nodes have completed. For example:

// Callback function which could signal some other waiting part of the application
void GraphCallback();

// Add a kernel node
auto KernelNode = Graph.add(...);

// Add a host_task node which depends on that node which signals a callback for example
Graph.add([&](sycl::handler& CGH){
   CGH.host_task([](){ 
      // Callback on the host now that KernelNode has finished executing
      GraphCallback();
   });
}, property::depends_on{KernelNode});

mfrancepillois · 2024-02-22T09:47:45Z

The use-case of sharing events between graphs seems pretty easy to deal with and fairly problem free, but I feel like the use-case of waiting on graph events externally from normal user code presents a lot of issues that this PR has to solve. This includes things like:
* If we're building on top of `sycl::event` which have no concept of being reset, how do we handle this?

* If we do allow resetting events how does the user reason about when this may happen if there are multiple executions in flight?
The use case for this kind of event seems to be signalling to user code that some portion of the graph has finished so that the application can do something with this information. However, I feel like rather than supporting this directly in the API this could simply be accomplished with a host_task that depends on some other nodes which can be used to signal in user code when those nodes have completed. For example:
// Callback function which could signal some other waiting part of the application
void GraphCallback();

// Add a kernel node
auto KernelNode = Graph.add(...);

// Add a host_task node which depends on that node which signals a callback for example
Graph.add([&](sycl::handler& CGH){
   CGH.host_task([](){ 
      // Callback on the host now that KernelNode has finished executing
      GraphCallback();
   });
}, property::depends_on{KernelNode});

Although this solution fulfills the requested feature in a relatively simple way (from our point of view at least), it involves resynchronizing the device flow with the host to execute this callback function which: 1) degrades the performance, 2) increases the workload of the host.

The use of events makes it possible to directly synchronize two device flows without host intervention.
As for the reset problem, we can solve it by requesting from users to update the event between two graph executions.

AerialMantis

This sounds like a great feature that will be a really powerful addition to SYCL Graphs and I think the proposed design looks good, though I think we need to think about how best to represent the capability for a "regular" event to be complete and then reset again.

AerialMantis · 2024-02-29T00:05:26Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+
+* `updatable` - Flag indicating that users will update the external event 
+ after the graph finalization. Please, note that setting this flag may degrade 
+ graph execution performance for level-zero backend, as updatable events 


Does enabling the updatable flag only cause a potential performance degradation on the Level Zero backend, or could this also impact other backends as well or can those support external events within their native graph APIs?

I'm assuming if external events are natively supported, these events can be signalled to with little or no interuption to the execution of the graph, but if this would require splitting the graph up into sub-graphs around the external event?

If this is the case, perhaps it would be useful to have an spect which can be queried for whether device supports native external events.

External events are natively supported by CUDA and HIP backend (I don't know for OpenCL @EwanC ?), and can be updated.
As for level-zero, the current API does not allow us to update events once the command-buffer is created. So, in this case, we have to split the graph and handle these updatable external events on the host. (Hopefully, mutable command-list will help to improve that.)

Thanks for clarifying, I think it would be useful to have an aspect for querying whether external events are supported natively.

Synchronization outside a command-buffer isn't supported in OpenCL at all, so that backend won't be able to support any of this proposal not just update.

We've got pushback before when trying to add extra aspects as it makes the code less portable, i'd be inclined to leave it out and throw an exception for things that aren't supported natively (documented in spec as a limitation)

Thanks for clarifying, so the OpenCL backend as of now wouldn't be able to support this extension at all, is there interest in adding this capability to the OpenCL command buffer extension as well?

I would probably argue that if we make this an optional feature and add an aspect for it this would improve portability because then users can query for the capability to avoid it failing on devices which don't support it. Saying that I think it would be okay to leave it out for the extension, but if we were to progress this to a KHR I think we would need it then.

I would probably argue that if we make this an optional feature and add an aspect for it this would improve portability because then users can query for the capability to avoid it failing on devices which don't support it. Saying that I think it would be okay to leave it out for the extension, but if we were to progress this to a KHR I think we would need it then.

See intel#12486 (comment) (and intel#12486 (comment)) for more general discussion about the role of aspects in SYCL-Graph. It might be worth commenting there if you have strong feelings about aspects for optional features.

Thanks for clarifying, so the OpenCL backend as of now wouldn't be able to support this extension at all, is there interest in adding this capability to the OpenCL command buffer extension as well?

There's definitely interest, it was an early feature request KhronosGroup/OpenCL-Docs#793
The only thing is that the timelines for adding things to the command-buffer spec aren't as quick as a vendor extension we can control, but we could start nudging this in the OpenCL side.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

AerialMantis · 2024-02-29T01:18:18Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+If a graph containing one or more external events is run multiple times, 
+the external events are automatically reset when the graph is resubmitted to 
+the queue.


So currently SYCL events don't have a concept of being reset, but I'm not sure this is necessarily a concern as prior to SYCL graphs a DAG could only be traversed once, so it only makes sense for events to be set once. But if a DAG (in the form of a command_graph) can be traversed multiple times I think it's a reasonable extension of the SYCL execution model that events within that graph are reset each time the graph is executed.

I do agree that it's a bit odd from the perspective of non-graph SYCL code for an event to be reset, though I would be wary of creating a new type as this could create duplicate mechanisms for doing the same thing in non-graph code. Perhaps we need a new concept to represent this to the existing SYCL execution model, such as a multi-shot event which can be complete and reset, and a way to query for this.

I think it could also be useful to have an API which can be used to query whether an event is external or graph-restricted, so that code which is simply passed an event object can create an "external event" if necessary.

AerialMantis · 2024-02-29T01:43:49Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+graphA.end_recording();
+
+event externalEventA;
+queueA.ext_oneapi_external_event(EventA, externalEventA, flase /* non-updatable */);


I imagine another way to do this would be to pass a property to the submit call which requires the event to always be external, even when produced by a recording queue. Though I guess the reason is to avoid this kind of design is that it would require modifying existing SYCL code to be used in record/replay?

This is also because some backends don't record events by default (typically CUDA and HIP). So when users request an external event, we have to add a new graph node to record this event.
In the current implementation we have a 1:1 map between events and nodes, so if we only transform the previous graph-limited event into an external event we will lose this.
Moreover, the complexity of implementation would be considerably higher.

Bensuo · 2024-04-04T16:15:17Z

I've pushed some changes as a separate commit which can be viewed here: 3d3350e

To summarise the changes:

As mentioned in some comments I've removed updating and profiling capabilities from external events for now. These can be revisited if we get a use-case/need for them.
Simplified the interfaces for adding external events to remove having to pass in a sycl::event which is transformed into an external event, and instead the user gets the external event from our APIs.
Added to and modified some wording around events to try and explain the types of events better.

Open questions/discussions:

I'm not sure the event::is_graph_limited() query is sufficient, since there are still some restrictions on external events compared to regular sycl events. I considered adding an event type and associated query instead, but this is a more invasive solution that I am less keen on for that reason.
I've disallowed using external events inside the graph they are associated with. This simplifies the recording interface (don't need to return the graph-limited event and so can remove passing in an event to transform as mentioned above) and I don't think there's a compelling reason why a user would want or need to do this, but it might be unintuitive.

EwanC

Couple of minor points from a first read over the new changes. It's probably worth updating the PR description to describe the problem we are trying to solve rather than enumerating the API that's changing frequently.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

EwanC · 2024-04-09T09:37:44Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+graph to finish.
+
+External events are considered complete when all their dependent graph nodes
+have finished execution, and their execution status is automatically reset when


Noting the point we discussed verbally about where the reset should be done by the user rather than automatically

Since there is some other discussion going on around update which may affect these I'll leave this unresolved for now until we have more clarity on this aspect.

I also thought this language about being automatically reset was too vague. If you read this literally, then I think your example using the explicit API is not well defined:

// Add an external event node to Graph A auto ExtEventNode = GraphA.add_external_event(); auto ExternalEventA = ExtEventNode.get_external_event(); // Add a wait external event node in GraphB auto ExtWaitNode = GraphB.add_wait_external_event({ExternalEventA}); queueA.submit([&](handler& cgh) { cgh.ext_oneapi_graph(execGraphA); }); // Is "ExternalEventA" really reset at this point? // If so, then is the use of "ExternalEventA" in GraphB below ever signalled? queueB.submit([&](handler& cgh) { cgh.ext_oneapi_graph(execGraphB); });

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

guoyejun · 2024-04-10T07:46:09Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+event get_external_event();
+----
+|Gets the external event associated with this node. Can only be called for nodes
+with type `node_type::external_event`.


how about other member functions if the node type is node_type::external_event

shall we add something elsewhere for spec completeness?

Other member functions are valid for node_type::external_event unless otherwise noted, for example the update_range/nd_range() member functions do say this under exceptions:

Throws with error code `invalid` if the type of the node is not a kernel execution.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

guoyejun · 2024-04-11T07:41:34Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+
+ queueB.submit([&](handler& cgh) {
+ cgh.ext_oneapi_graph(execGraphB);
+ })


what will happen if we first submit execGraphB and then submit execGraphA, endless wait within execGraphB?

Since they are separate graphs with no implicit dependencies from accessors between them, then no I believe that execGraphA could still execute after execGraphB. It may endlessly wait if there were implicit dependencies though I think.

not quite understand that 'It may endlessly wait if there were implicit dependencies'.

I agree that execGraphA could run successfully since it does not depend on others.

While for execGraphB, there's a clear explicit dependency that it depends on externalEventA, not "if there were implicit dependencies", so my point is that it always endlessly waits if launch first, not "may endlessly wait".

So the external events should be implemented as a device-side dependency, so execGraphB will not block on the host waiting for the external event because there is no host synchronization there. Since execGraphB does not block on host that allows the application to proceed and execute execGraphA, and once it reaches the point where the external event is signaled then execGraphB (on the device) can also proceed to finish execution. Does that make sense?

My point about implicit dependencies is more of a runtime implementation detail I suppose, but in SYCL using accessors creates implicit dependencies between commands (and we also extend this to graphs), so if you were using accessors which created a dependency this may cause a blocking host wait on attempting to execute execGraphA which would deadlock in this case.

This could potentially be mitigated by specifying exactly when/how the external event status changes. For example, if we consider the external event to begin in the complete state and be reset upon the owning graph being executed then that would prevent this scenario from being a deadlock, at least.

So the external events should be implemented as a device-side dependency, so execGraphB will not block on the host waiting for the external event because there is no host synchronization there. Since execGraphB does not block on host that allows the application to proceed and execute execGraphA, and once it reaches the point where the external event is signaled then execGraphB (on the device) can also proceed to finish execution. Does that make sense?

yes, for this case, execGraphB waits on the device side until there is a host function changes the event status to complete.

My point about implicit dependencies is more of a runtime implementation detail I suppose, but in SYCL using accessors creates implicit dependencies between commands (and we also extend this to graphs), so if you were using accessors which created a dependency this may cause a blocking host wait on attempting to execute execGraphA which would deadlock in this case.

just wonder, for most popular cases, the buffer is always accessed by the device, and there is no internal D2H or H2D copy, so, my understanding is that the wait also happens at the device side, why the host side?

This could potentially be mitigated by specifying exactly when/how the external event status changes. For example, if we consider the external event to begin in the complete state and be reset upon the owning graph being executed then that would prevent this scenario from being a deadlock, at least.

yes, good point.

just wonder, for most popular cases, the buffer is always accessed by the device, and there is no internal D2H or H2D copy, so, my understanding is that the wait also happens at the device side, why the host side?

You are right actually this should be a device side wait in all cases, I was remembering old workarounds in the backend.

guoyejun · 2024-04-11T07:50:56Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+To overcome the limitations of "graph-limited" events, users can add external
+event nodes to a graph. These represent a point of execution within the graph
+and provide a SYCL event object which can be used to both manage synchronization
+between different graphs as well as between the graph and regular operations


regarding "between the graph and regular operations",

For the case that,
a) warmup stage: some regular operations generate an event (name it as eventA).
b) recording stage: it depends on eventA (it is not allowed)

and at the time of warmup stage, the implementation does not know if there will be recording stage.

Any suggestion on how external event can resolve the above issue?

I think we could allow this behaviour you suggest, and I have gone back and forth on it a few times. Probably allowing normal SYCL events to be passed to command_graph::add_wait_external_event() and allowing them as dependencies for recorded queue submissions would be fine.

There may be a potential for some undesirable overhead though: the event dependency in the backend could not be removed from subsequent graph executions once it has completed, so there could be a small overhead there particularly if a lot of these types of events were used (assuming there is still some small overhead from checking completed events in the backend/driver).

The user could avoid this by waiting on those one time operations before executing the graph, but this is problematic in the library record case.

I've removed the restrictions on this in the spec for now, and added some wording explicitly allowing it.

Bensuo · 2024-04-22T15:28:35Z

@gmlueck Can't seem to request a review on this one but would appreciate some feedback on this proposal.

Updates specification with new APIs for handling external events in Record&Replay and Explicit mode. Adds 3 functions to `graph`: - add_barrier() - make_external_event() - update_external_event() Adds 1 function to `queue`: - ext_oneapi_external_event() Adds examples to show how to use these functions. update spec Fix issue in the second code example.

- Remove updating and profiling capabilities - Simplify and align interfaces for getting external events - Restrict external event use within the graph (prevents needing graph-limited event from queue shortcut)

- Restore forbidding outside graph deps that are not external events

- Rename add_barrier to add_wait_external_event - Clarify wording about edges and ext events - Minor formatting and wording changes

- Useful for warmups etc.

guoyejun · 2024-04-28T02:18:25Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+create an edge, it must be of a specific type (see <<event-class-modifications,
+the "Event Class Modifications" section>> for more information).
+`handler::depends_on()` can be used to express edges when a user is working with
+USM memory rather than SYCL buffers. For a graph recorded with an in-order


Suggested change

USM memory rather than SYCL buffers. For a graph recorded with an in-order

USM memory rather than SYCL buffers. Thirdly, for a graph recorded with an in-order

since there's "three ways" at line 278.

guoyejun · 2024-04-28T02:40:02Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+- Using the event outside of the recording scope will throw synchronously with
+error code `invalid`.
+
+==== External Events


Suggested change

==== External Events

==== Graph-External Events

gmlueck

Here are some initial comments after a first read through.

gmlueck · 2024-04-27T14:22:33Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -632,6 +654,10 @@ public:

 void make_edge(node& src, node& dest);

+ node add_wait_external_event(const std::vector<event> eventList = {}, const property_list& propList = {});


Suggested change

node add_wait_external_event(const std::vector<event> eventList = {}, const property_list& propList = {});

node add_wait_external_event(const std::vector<event>& eventList = {}, const property_list& propList = {});

gmlueck · 2024-04-27T14:23:02Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+|
+[source,c++]
+----
+node add_wait_external_event(const std::vector<event> eventList, const property_list& propList = {});


Suggested change

node add_wait_external_event(const std::vector<event> eventList, const property_list& propList = {});

node add_wait_external_event(const std::vector<event>& eventList, const property_list& propList = {});

gmlueck · 2024-04-27T14:35:07Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+
+[source, c++]
+----
+namespace sycl::info {


New types should be added to the namespace sycl::ext::oneapi::experimental::info.

gmlueck · 2024-04-27T14:38:08Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+[source, c++]
+----
+namespace sycl::info {
+ enum class ext_oneapi_event_t{


Do not prefix new types with "ext_oneapi" because they are already in a descriptive namespace. Do not indent declarations within a namespace.

Why not name the enum event_type? That seems more descriptive and more in line with existing enums (e.g. sycl::info::device_type).

gmlueck · 2024-04-27T14:45:11Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+|Return Type
+|Description
+
+|`info::event::ext_oneapi_event_type`


I think this should be named sycl::ext::oneapi::experimental::event::event_type. This is consistent with the way sycl::info::device::device_type and sycl::info::device_type are named.

gmlueck · 2024-04-29T13:01:04Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

-graph.
+`handler::depends_on()` will also create a graph edge when passed certain kinds
+of events, for more information see <<event-class-modifications, the "Event
+Class Modifications" section>>.


This link looks a little awkward in the HTML render. I think it would look better by rewording to:

... see the section <<event-class-modifications, Event Class Modifications>>.

gmlueck · 2024-04-29T15:07:03Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+Exceptions:
+
+* Throws synchronously with error code `invalid` if the queue is not 
+ in recording mode. 


I think these two member functions are intended to be used when recording a graph. How realistic is it that the caller will know the depEvents? I thought the usage model is that the caller starts recording a queue, and then calls some non-graph-aware function, which puts commands on the queue. Since the called function is not aware it is being recorded, it can't call this new graph-specific API. However, the depEvents are probably internal details of the function, so the caller isn't aware of them.

I wonder if it would be better to have a new API that returns an external event which represent a dependency on all the commands that have been recorded to the queue so far. This would allow a caller to record a dependency point inside a recorded graph without knowing the specific depEvents.

I also think the ext_oneapi_external_event does not work well when recording an in-order queue. In this case, there typically is not any depEvents. Will the external event returned by this function have an implicit dependency on all previous commands in this case?

gmlueck · 2024-04-29T15:20:08Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+graph to finish.
+
+External events are considered complete when all their dependent graph nodes
+have finished execution, and their execution status is automatically reset when


I also thought this language about being automatically reset was too vague. If you read this literally, then I think your example using the explicit API is not well defined:

// Add an external event node to Graph A auto ExtEventNode = GraphA.add_external_event(); auto ExternalEventA = ExtEventNode.get_external_event(); // Add a wait external event node in GraphB auto ExtWaitNode = GraphB.add_wait_external_event({ExternalEventA}); queueA.submit([&](handler& cgh) { cgh.ext_oneapi_graph(execGraphA); }); // Is "ExternalEventA" really reset at this point? // If so, then is the use of "ExternalEventA" in GraphB below ever signalled? queueB.submit([&](handler& cgh) { cgh.ext_oneapi_graph(execGraphB); });

Bensuo · 2024-05-22T09:26:09Z

Closing this in favour of #372 which presents a simpler API for accomplishing the same general goals as this.

mfrancepillois added the Graph Specification Extension Specification related label Feb 13, 2024

mfrancepillois requested review from EwanC, ori-sky, reble, Bensuo and julianmi February 13, 2024 17:41

EwanC reviewed Feb 14, 2024

View reviewed changes

AerialMantis reviewed Feb 29, 2024

View reviewed changes

Bensuo force-pushed the maxime/external-events-spec branch from e07eb71 to 3d3350e Compare April 4, 2024 16:03

EwanC reviewed Apr 5, 2024

View reviewed changes

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Outdated Show resolved Hide resolved

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Outdated Show resolved Hide resolved

EwanC requested a review from fabiomestre April 8, 2024 10:38

EwanC reviewed Apr 9, 2024

View reviewed changes

ori-sky reviewed Apr 9, 2024

View reviewed changes

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Show resolved Hide resolved

guoyejun reviewed Apr 10, 2024

View reviewed changes

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Outdated Show resolved Hide resolved

guoyejun reviewed Apr 10, 2024

View reviewed changes

guoyejun reviewed Apr 11, 2024

View reviewed changes

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Outdated Show resolved Hide resolved

guoyejun reviewed Apr 11, 2024

View reviewed changes

mfrancepillois and others added 8 commits April 22, 2024 17:18

[SYCL][Graph] External event changes

00fae5b

- Remove updating and profiling capabilities - Simplify and align interfaces for getting external events - Restrict external event use within the graph (prevents needing graph-limited event from queue shortcut)

Fix typos

e0dc8d1

Restore event limitations

69b1d04

- Restore forbidding outside graph deps that are not external events

Replace is_graph_limited with event info query

898e58d

Addressing review comments

3714f8f

- Rename add_barrier to add_wait_external_event - Clarify wording about edges and ext events - Minor formatting and wording changes

Remove duplicate wording

9c27142

Minor wording update

df48b8a

Bensuo added 2 commits April 22, 2024 17:20

Restore using events from outside graph

75f8c14

- Useful for warmups etc.

Add some more external event examples

35788b5

Bensuo force-pushed the maxime/external-events-spec branch from 6b908b5 to 35788b5 Compare April 22, 2024 16:21

guoyejun reviewed Apr 28, 2024

View reviewed changes

gmlueck reviewed Apr 29, 2024

View reviewed changes

Bensuo closed this May 22, 2024

	USM memory rather than SYCL buffers. For a graph recorded with an in-order
	USM memory rather than SYCL buffers. Thirdly, for a graph recorded with an in-order

		@@ -632,6 +654,10 @@ public:

		void make_edge(node& src, node& dest);

		node add_wait_external_event(const std::vector<event> eventList = {}, const property_list& propList = {});

[SYCL][Graph] Add External Events APIs #355

[SYCL][Graph] Add External Events APIs #355

Conversation

mfrancepillois commented Feb 13, 2024 • edited by EwanC Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Bensuo commented Feb 21, 2024

mfrancepillois commented Feb 22, 2024

AerialMantis left a comment • edited Loading

Choose a reason for hiding this comment

AerialMantis Feb 29, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

AerialMantis Feb 29, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Bensuo commented Apr 4, 2024

EwanC left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Bensuo commented Apr 22, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

gmlueck left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Bensuo commented May 22, 2024

mfrancepillois commented Feb 13, 2024 •

edited by EwanC

Loading

AerialMantis left a comment •

edited

Loading

AerialMantis Feb 29, 2024 •

edited

Loading

AerialMantis Feb 29, 2024 •

edited

Loading