[SYCL][Doc] Graph fusion extension proposal

[sommerlukas]

Experimental SYCL extension proposal for kernel fusion on top of the SYCL graphs API.

Constructing the sequence of kernels to fuse is completely left to the graphs proposal, which provides two APIs to this end. One recording API similar to the fusion mode for queues in the initial kernel fusion proposal, and an explicit graph construction APIs. Both APIs are supported for kernel fusion.

This proposal mainly introduces a number of properties to trigger fusion of the graph and internalization of dataflow in the fused kernel.

This proposal continues some of the ideas of the experimental SYCL extension for kernel fusion. In contrast to the original kernel fusion proposal, this proposal now also allows internalization of USM pointers.

[EwanC]

We are considering removing this add_malloc_device API (due to issue outlined here), and instead advising users to use core sycl::malloc_device outside the graph. i.e tending towards option 0 from the call.

Would using the property_list parameter to sycl::malloc_device to pass the internalization information, and then using that device pointer in a graph be satisfactory for fusion?

[sommerlukas]

As long as it is possible for the implementation to determine that a pointer was created with an internalization property, using sycl::malloc_device to allocate USM pointers for internalization should be fine for fusion.

The downside of using sycl::malloc_device with the internalization properties outside of the graph is mainly granularity: When passing the properties to add_malloc_device, the additional semantics of the properties are limited to one graph. If, on the other hand, the property is passed sycl::malloc_device, it will apply to all graphs that use the pointer.

[EwanC]

I agree the granularity isn't ideal, but good to know that the finer granularity is a nice-to-have rather must-have for you. Right now I don't think we have confidence in any of our potential graph memory allocation designs, and am reluctant to specify something that user applications start using, and we then have to depreciate (even through the extension is experimental). So if core SYCL USM is good enough, it may be the best way forward for now so we that can focus on the critical path to get something into the hands of users. Then revisit this again once we have real world feedback.

[sommerlukas]

Yes, core SYCL USM should be fine as long as there is a way to match a USM pointer with its properties specified during allocation. I'm no expert on the implementation of USM, so I can't predict how complicated that matching is. If device allocation from multiple devices are in play, the same address could actually be used for two completely different pointers from two different devices, making the matching more complicated. But this is probably an implementation detail.

Regarding granularity, a while back we had also considered properties for USM internalization, which would simply take a pointer as value of the property, and that you could pass to complete_fusion (or finalize for this proposal).
Something along the lines of:
sycl::ext::oneapi::experimental::property::promote_private(void* ptr).

This would allow to limit the internalization semantic for a specific pointer to just this graph. However, IMHO it is less elegant than a data-less property and would also mean you need two separate properties for USM and buffer internalization (unless you allow nullptr to be passed).

[sommerlukas]

Another advantage of having add_malloc_device and passing the properties to that API: You can completely eliminate the allocation of the device pointer if internalization is successful. After the JIT compilation as part of finalize, it is known whether the allocation is needed at all. If it is not needed, it can be removed from the graph, saving some time for the device allocation and improving performance.

This optimization is not possible if the USM allocations are not part of the graph and happen outside of the graph with the core SYCL USM API.

[EwanC]

We have a draft PR for this change now, feel free to comment - reble#79 Can hold off on merging this until we have our sync-up meeting if you'd like to discuss it over a call

[sommerlukas]

Thanks! I've put a comment on reble#79, might be worth to discuss this further in the sync-up meeting.

[EwanC]

With this method compared to event complete_fusion(const property_list &propList = {}); doing fusion will become a synchronous operation compared to an asynchronous one. From my naive perspective I think that's fine, but do you have any concerns about how this change will affect any applications you're targetting?

[sommerlukas]

The JIT compilation part of complete_fusion has always been synchronous. The only asynchronous part about complete_fusion was the execution of the fused kernel, similar to how kernel execution in queue::submit is asynchronous. So I think this is fine. Having the JIT compilation (part of finalize) and the submission of the fused graph as two separate APIs even gives the user more control over when the JIT compilation overhead is encountered, potentially even giving him the freedom to perform finalization in a separate thread in parallel to some computation.

[EwanC]

I agree the granularity isn't ideal, but good to know that the finer granularity is a nice-to-have rather must-have for you. Right now I don't think we have confidence in any of our potential graph memory allocation designs, and am reluctant to specify something that user applications start using, and we then have to depreciate (even through the extension is experimental). So if core SYCL USM is good enough, it may be the best way forward for now so we that can focus on the critical path to get something into the hands of users. Then revisit this again once we have real world feedback.

[EwanC]

Do we need to restrict the use of this graphs entry-point which updates an executable graph with the inputs/outputs of another topologically identical graph?

void command_graph<graph_state::executable> update(const command_graph<graph_state::modifiable>& graph)

Guessing if an executable graph has already been fusued you won't be able to swap out an input with another using different internalization properties.

There's also a pending PR - reble#77 - that proposes another entry-point with a more direct update.

[sommerlukas]

Interesting point, thanks for bringing this up.

The question here (next to some technical/implementation details) probably is, how computationally expensive the update operation is allowed to be. The graph proposal does not mention computational complexity or whether the update is finished when returning from update.

If update is allowed to be a computationally expensive operation, updating a fused graph with new inputs with different internalization properties could simply trigger re-compilation (=fusion) of the graph, yielding a new fused graph. In that case, no additional restrictions would need to be applied to the update API.

In the current implementation of kernel fusion, we also use a cache to avoid re-compilation in case the same sequence of kernels is fused repeatedly and the internalization properties (and some other context) have not changed. The update scenario would be somewhat similar, only triggering re-compilation on update if relevant parameters, including internalization properties, have changed.

[EwanC]

I envisage that the computational complexity ultimately will depend on backend support, e.g cl_khr_command_buffer_mutable_dispatch, Vulkan, or CUDA Graph backends lets you update the arguments to an existing command-buffer easily, but Level Zero or vanilla OpenCL backends do not provide a mechanism so the command-lists would have to be recreated in a slower path.

The current PI interface also doesn't provide a way to expose this, so our extensions to PI for graphs will need to expose kernel arguments differently so that cl_khr_command_buffer_mutable_dispatch/Vulkan/CUDA-Graphs could take advantage it.

So the update() entry-point is definitely allowed to be expensive, equivalent to finalizing again. In the immediate future of our graphs implementation it will be the slower emulated path until we get round to adding the relevant PI, which isn't of my top of our priority list as it's a change for performance rather than correctness.

[sommerlukas]

So the update() entry-point is definitely allowed to be expensive, equivalent to finalizing again

In that case, re-compiling on an update with different internalization properties (or otherwise different context) seems like the best route to go, which means that no additional restrictions would need to be applied to the update API.

[sommerlukas]

Closing this PR, as it is now superseded by intel#8678.

The updated version in that PR removes reference to the add_malloc_device API that was removed from the SYCL graph proposal for now.