Skip to content

Commit 8c4cad4

Browse files
authored
[SYCL][Doc] Graph fusion extension proposal (#8678)
Experimental SYCL extension proposal for kernel fusion on top of the [SYCL graphs API](https://github.com/reble/llvm/blob/sycl-graph-update/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc). 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](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc). In contrast to the original kernel fusion proposal, this proposal now also allows internalization of USM pointers. --------- Signed-off-by: Lukas Sommer <[email protected]>
1 parent e423b5a commit 8c4cad4

File tree

3 files changed

+792
-1
lines changed

3 files changed

+792
-1
lines changed

sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,18 @@ or may be in a preliminary state. The specification itself may also change in
4848
incompatible ways before it is finalized. *Shipping software products should
4949
not rely on APIs defined in this specification.*
5050

51+
[NOTE]
52+
====
53+
There is a link:../proposed/sycl_ext_oneapi_graph_fusion.asciidoc[follow-up
54+
proposal] for fusion based on the
55+
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc[SYCL graph API].
56+
That proposal continues some of the ideas presented in this proposal, but uses
57+
the more versatile SYCL graphs API to define the sequence of kernels to
58+
execute.
59+
60+
Once accepted and implemented, the new proposal will supersede this proposal.
61+
====
62+
5163
[NOTE]
5264
====
5365
This is an experimental extension for the SYCL specification.

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1207,7 +1207,7 @@ The `sycl::ext::codeplay::experimental::property::queue::enable_fusion` property
12071207
defined by the extension is ignored by queue recording.
12081208

12091209
To enable kernel fusion in a `command_graph` see the
1210-
https://github.com/sommerlukas/llvm/blob/proposal/graph-fusion/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph_fusion.asciidoc[sycl_ext_oneapi_graph_fusion extension proposal]
1210+
link:../proposed/sycl_ext_oneapi_graph_fusion.asciidoc[sycl_ext_oneapi_graph_fusion extension proposal]
12111211
which is layered ontop of `sycl_ext_oneapi_graph`.
12121212

12131213
==== sycl_ext_oneapi_kernel_properties

0 commit comments

Comments
 (0)