SYCL: Add support for Graphs #6912

masterleinad · 2024-04-02T22:40:51Z

This pull request adds support for Graphs based on https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc. The current status is sufficient to pass all related unit tests.
Note that:

immediate command lists are not supported in some implementations
SYCL_EXT_ONEAPI_GRAPH needs to be defined
we don't deal with scratch memory at all (same as for Cuda and HIP) so Team policy and parallel_reduce/parallel_scan are not properly supported.

masterleinad · 2024-05-08T20:52:54Z

@romintomasetti FYI

dalg24

Do the no_immediate_command_list changes stand on their own or are do they only make sense in the context of this PR?

dalg24 · 2024-05-09T12:53:36Z

core/unit_test/TestGraph.hpp

+                                                                   {0, 0}},
+            KOKKOS_LAMBDA(int, int, int&){}, count)
+        .then_parallel_reduce(
+            Kokkos::TeamPolicy<TEST_EXECSPACE>{0, 1},


Any reason not to use Kokkos::AUTO as 2nd argument here?

Just noting that this is increasing coverage of the current test and could be proposed on its own.

Any reason not to use Kokkos::AUTO as 2nd argument here?

No

Just noting that this is increasing coverage of the current test and could be proposed on its own.

Yes, we can do this independently. I only noticed that we don't test all the implementations (even though we can't do proper reductions anyway).

dalg24 · 2024-05-09T13:18:48Z

core/src/decl/Kokkos_Declare_SYCL.hpp

@@ -19,6 +19,9 @@

 #if defined(KOKKOS_ENABLE_SYCL)
 #include <SYCL/Kokkos_SYCL.hpp>
+#ifdef SYCL_EXT_ONEAPI_GRAPH
+#include <SYCL/Kokkos_SYCL_GraphNodeKernel.hpp>


I thought the current design was you get these when you explicitly include <Kokkos_Graph.hpp> meaning we should not add it here.
Am I missing something?

Cuda and HIP are abstracting the kernel launch away but we don't have such a mechanism for the SYCL backend (yet). This means that the parallel construct implementations are calling get_sycl_graph_node_from_kernel directly and need this header.

For a correct parallel_reduce/parallel_scan implementation, we will need to create multiple nodes anyway and can't get away with just attaching to the node in the kernel launch abstraction.

core/src/SYCL/Kokkos_SYCL_ParallelFor_Range.hpp

masterleinad · 2024-05-09T14:28:01Z

Do the no_immediate_command_list changes stand on their own or are do they only make sense in the context of this PR?

Whether or not to use immediate command lists is independent of this pull request (the difference is whether the runtime waits for a couple of kernels before executing them batched or immediately executes them). Thus, it's something users could play with depending on which mode gives better performance.
The reason to introduce the option here is that the initial implementation didn't support immediate command lists. This limitation was removed in intel/llvm#12279 and oneapi-src/unified-runtime#1218 but is not included in the default versions on the testbeds.

dalg24 · 2024-05-09T14:33:44Z

Then please open a separate PR for this change

romintomasetti · 2024-05-14T13:01:07Z

core/unit_test/TestGraph.hpp

@@ -21,6 +21,7 @@

 namespace Test {

+#if !defined(KOKKOS_ENABLE_SYCL) || defined(SYCL_EXT_ONEAPI_GRAPH)


I'm not sure to understand why you want to disable the whole code in this file if KOKKOS_ENABLE_SYCL is defined...

I'm also wondering who is defining SYCL_EXT_ONEAPI_GRAPH ?

Shouldn't the logic be in the CMake itself ? Something like:

if(KOKKOS_ENABLE_SYCL AND SYCL_EXT_ONEAPI_GRAPH) .... enable the SYCL graph test ... endif()

The condition reads: only enable the test if the SYCL backend is not enabled, or if the SYCL backend is enabled and the SYCL_EXT_ONEAPI_GRAPH feature macro is defined.

Shouldn't the logic be in the CMake itself ?

The feature macro is not a CMake variable.

OK. I would suggest to move the line

#if !defined(KOKKOS_ENABLE_SYCL) || defined(SYCL_EXT_ONEAPI_GRAPH)

to TestSYCL_Graph.hpp then.

We typically auto-generate the *.cpp test files (such as TestSYCL_Graph.cpp) which we might want to do in a follow-up. Moving the guard to the *.cpp file would prevent that.

OK. I would strongly advise a "CMake experimental feature flag" like KOKKOS_ENABLE_SYCL_GRAPH or so for several reasons 🔢

I feel strongly against introducing backend-specific preprocessing guard that deactivate the whole bunch of generic tests in TestGraph.hpp.
This goes against the idea of a portable test file and is a backdoor for future maintenance headaches.

It feels strange that if a user has the SYCL backend enabled, but hasn't support for SYCL_EXT_ONEAPI_GRAPH, still there will be a test target generated
for TestSYCL_Graph.cpp that actually does nothing. To me, such a test deactivation should be done in the CMakeLists.txt for sure.

Am I right to say that if KOKKOS_ENABLE_SYCL is defined, but SYCL_EXT_ONEAPI_GRAPH is not, then Kokkos::Graph<SYCL> will fallback on the
default GraphImpl? If so, I see that as a very big warning sign ⚠️. Indeed, as a user that wants the SYCL specialization, I would not be happy to see that
Kokkos::Graph<SYCL> silently falls back on the default Graph implementation...

For these reasons, I strongly suggest that you introduce KOKKOS_ENABLE_SYCL_GRAPH or so. It would be an "experimental feature" flag that, if set by the user,
enforces the SYCL Graph specialization (and therefore ensures that compilation fails if SYCL_EXT_ONEAPI_GRAPH is not supported).

You could also check for the SYCL_EXT_ONEAPI_GRAPH support with some kind of try_compile.

The logic could be align these lines:

if(NOT DEFINED KOKKOS_ENABLE_SYCL_GRAPH) try_compile(...) if(COMPILATION_SUCCEEDED) set(KOKKOS_ENABLE_SYCL_GRAPH ON) endif() endif() ... if(KOKKOS_ENABLE_SYCL_GRAPH) add_test(TestSYCL_Graph.cpp) endif()

What do you think ?

I feel strongly against introducing backend-specific preprocessing guard that deactivate the whole bunch of generic tests in TestGraph.hpp.

It feels strange that if a user has the SYCL backend enabled, but hasn't support for SYCL_EXT_ONEAPI_GRAPH, still there will be a test target generated
for TestSYCL_Graph.cpp that actually does nothing. To me, such a test deactivation should be done in the CMakeLists.txt for sure.

We shouldn't actually need that guard and can just always execute the test even if no specialization is available, see #7011.

Am I right to say that if KOKKOS_ENABLE_SYCL is defined, but SYCL_EXT_ONEAPI_GRAPH is not, then Kokkos::Graph will fallback on the
default GraphImpl? If so, I see that as a very big warning sign ⚠️. Indeed, as a user that wants the SYCL specialization, I would not be happy to see that
Kokkos::Graph silently falls back on the default Graph implementation...

That's how the Graph implementation is designed and we should not start breaking Graphs for users if there is a specialization but it's not active in their configuration. We always reserve the right to improve performance based on the backend capability without the need for the user to pitch in. I added information about SYCL_EXT_ONEAPI_GRAPH to print_configuration, though.

We shouldn't actually need that guard and can just always execute the test even if no specialization is available, see #7011.

Nice 👍

That's how the Graph implementation is designed and we should not start breaking Graphs for users if there is a specialization but it's not active in their configuration. We always reserve the right to improve performance based on the backend capability without the need for the user to pitch in. I added information about SYCL_EXT_ONEAPI_GRAPH to print_configuration, though.

My suggestion is about allowing the user to "fore opt-in into the specialization or stop compilation" 😉 So if I define KOKKOS_ENABLE_SYCL_GRAPH_IMPL, and for some reason it will not be possible to compile the SYCL specialization, I would prefer that Kokkos does not compile, rather than falling back to the default implementation. However, if I don't define KOKKOS_ENABLE_SYCL_GRAPH_IMPL, then I agree with you that Kokkos can do whatever it want.

I don't understand what your use case is. Why would you rather fail compilation than use a fallback implementation? Since KOKKOS_ENABLE_SYCL_GRAPH_IMPL is not portable, you can just as well check for SYCL_EXT_ONEAPI_GRAPH in user code.

Note that we already have

#if !((HIP_VERSION_MAJOR == 5) && (HIP_VERSION_MINOR == 2)) #define KOKKOS_IMPL_HIP_GRAPH_ENABLED #endif

which silently enables or disables the Graph specialization based on the compiler version.

core/unit_test/CMakeLists.txt

dalg24 · 2024-05-23T20:52:05Z

core/src/SYCL/Kokkos_SYCL_ParallelFor_Range.hpp

-  static sycl::event sycl_direct_launch(const Policy& policy,
-                                        const Functor& functor,
-                                        const sycl::event& memcpy_event) {
+  sycl::event sycl_direct_launch(const Policy& policy, const Functor& functor,


Just observing that this PR makes all the Parallel<Construct>::sycl_direct_launch have a somewhat misleading name because they now either launch a kernel or append a node to some graph.
Any thought on whether/how we can improve?

Yeah, we could inline sycl_direct_launch into execute but that's just as misleading (but there is no good reason anymore not to do it other than that there is little gain from it).
Just as for all other Graph specializations, the implementation calls execute since that's where the functor is set up. In the case of HIP and Cuda, the functor is then passed to a launch mechanism outside, and only there treating it as a regular kernel or graph node happens. This won't work once a parallel construct requires more than one node though.

We could think about renaming execute everywhere and split it into a prepare_kernel and a launch step and then introduce a setup_graph_node function (which might just call launch for Cuda and HIP but would call sycl_attach_kernel_to_node for SYCL).

dalg24 · 2024-05-23T20:54:32Z

core/src/SYCL/Kokkos_SYCL_ParallelReduce_Range.hpp

+      if constexpr (Policy::is_graph_kernel::value)
+        Kokkos::abort(
+            "parallel_reduce not implemented for graph kernels if result is "
+            "not device-accessible!");


Is that something we could turn into a compile error?
(Not necessarily blocking, just wondering if you look into this and if it is doable)

It's a little tricky since we don't know at compile-time at this point whether parallel_reduce was called with a result parameter or not.

core/unit_test/TestGraph.hpp

core/src/SYCL/Kokkos_SYCL_Graph_Impl.hpp

dalg24 · 2024-05-28T22:48:19Z

core/src/SYCL/Kokkos_SYCL_GraphNodeKernel.hpp

+  }
+
+  void set_sycl_graph_node_ptr(
+      std::optional<sycl::ext::oneapi::experimental::node>* arg_node) {


So you would prefer pointer to optional to resemble the CUDA/HIP code?
I suppose that's fine.

core/src/SYCL/Kokkos_SYCL_Graph_Impl.hpp

core/src/SYCL/Kokkos_SYCL_GraphNodeKernel.hpp

core/unit_test/TestGraph.hpp

masterleinad · 2024-05-31T17:58:57Z

I can confirm that this works correctly on the Intel testbeds with two different compiler drops (one of them requires SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0).

masterleinad added 8 commits April 2, 2024 18:13

Clean up some static_asserts

ea68e86

Introduce KOKKOS_ENABLE_IMPL_SYCL_NO_IMMEDIATE_COMMAND_LISTS

6fb2002

Make SYCL parallel constructs copy-constructible

924d41c

SYCL: Add support for Graphs

6a47e6b

Guard with SYCL_EXT_ONEAPI_GRAPH

afbc380

Test parallel_reduce for all policies

58b275a

Test all reduce policies

ebe0501

Fix TeamPolicy

7bb46ab

masterleinad force-pushed the sycl_command_graph branch from 32b3243 to 7bb46ab Compare April 9, 2024 15:15

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

921b7d6

masterleinad mentioned this pull request May 3, 2024

SYCL: Prepare Parallel* for Graphs #6988

Merged

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

3ef4b14

masterleinad force-pushed the sycl_command_graph branch from f07fc23 to 3ef4b14 Compare May 8, 2024 20:51

masterleinad marked this pull request as ready for review May 8, 2024 20:51

dalg24 reviewed May 9, 2024

View reviewed changes

This was referenced May 9, 2024

SYCL: Add Kokkos_ENABLE_SYCL_NO_IMMEDIATE_COMMAND_LISTS option #7003

Closed

SYCL: Print submission command queue property #7004

Merged

Remove KOKKOS_ENABLE_IMPL_SYCL_NO_IMMEDIATE_COMMAND_LISTS

67597df

masterleinad force-pushed the sycl_command_graph branch from 277b470 to 67597df Compare May 10, 2024 13:39

romintomasetti reviewed May 14, 2024

View reviewed changes

masterleinad mentioned this pull request May 14, 2024

Always run Graph tests #7011

Merged

masterleinad added 4 commits May 14, 2024 19:28

Remove guard in TestGraph.hpp

0f9f6c0

Print value of SYCL_EXT_ONEAPI_GRAPH

69a1f3c

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

938a18e

Avoid generic lambda in TestGraph.hpp to fix Windows+Cuda

8da765f

masterleinad force-pushed the sycl_command_graph branch from 16ed361 to 8da765f Compare May 15, 2024 16:27

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

ebed93b

masterleinad added 4 commits May 22, 2024 16:55

Revert unintentional changes

40b70b5

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

a7921b3

fix

d89e3a7

Fix indentation

3a5bbda

dalg24 reviewed May 22, 2024

View reviewed changes

core/unit_test/CMakeLists.txt Outdated Show resolved Hide resolved

masterleinad added 6 commits May 22, 2024 18:30

Add sycl_attach_kernel_to_node

fb67aa7

Fix indentation in CMakeLists.txt

3a6a55d

Guard sycl_attach_kernel_to_node

0f2b7c9

Merge remote-tracking branch 'upstream/develop' into sycl_command_graph

067cac7

Update TestGraph.hpp

ffb6fa3

Indentation

a951dbf

masterleinad force-pushed the sycl_command_graph branch from 64a959a to 96209c1 Compare May 23, 2024 16:06

Fix Cuda

358416e

masterleinad force-pushed the sycl_command_graph branch from 96209c1 to 358416e Compare May 23, 2024 18:49

dalg24 reviewed May 23, 2024

View reviewed changes

masterleinad added 2 commits May 24, 2024 14:03

Fix HIP

0f6ca2a

Fix OpenMPTarget

a38f4b3

masterleinad mentioned this pull request May 28, 2024

Improve Kokkos Graphs #7039

Merged

dalg24 reviewed May 28, 2024

View reviewed changes

masterleinad added 4 commits May 29, 2024 12:53

Fix typos

5b76cca

Add comments

caaf6cc

Remove unnecessary KOKKOS_ENSURES

0897f0c

Remove m_driver_storage

548d4dc

masterleinad force-pushed the sycl_command_graph branch from 5e530ad to 548d4dc Compare May 29, 2024 12:59

Merge branch 'develop' into sycl_command_graph

5f12b9b

masterleinad commented May 31, 2024

View reviewed changes

core/unit_test/TestGraph.hpp Outdated Show resolved Hide resolved

Fix TestGraph.hpp after merging develop

7ecf1a0

dalg24 approved these changes May 31, 2024

View reviewed changes

masterleinad requested a review from nliber May 31, 2024 17:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

SYCL: Add support for Graphs #6912

SYCL: Add support for Graphs #6912

masterleinad commented Apr 2, 2024 •

edited

masterleinad commented May 8, 2024

dalg24 left a comment

dalg24 May 9, 2024

dalg24 May 9, 2024

masterleinad May 9, 2024

dalg24 May 9, 2024

masterleinad May 9, 2024

masterleinad commented May 9, 2024

dalg24 commented May 9, 2024

romintomasetti May 14, 2024

romintomasetti May 14, 2024

masterleinad May 14, 2024

masterleinad May 14, 2024

romintomasetti May 14, 2024

masterleinad May 14, 2024

romintomasetti May 14, 2024

masterleinad May 14, 2024

romintomasetti May 15, 2024

masterleinad May 15, 2024

dalg24 May 23, 2024

masterleinad May 23, 2024

dalg24 May 23, 2024

masterleinad May 23, 2024

dalg24 May 28, 2024

masterleinad commented May 31, 2024

		@@ -21,6 +21,7 @@

		namespace Test {

		#if !defined(KOKKOS_ENABLE_SYCL) \|\| defined(SYCL_EXT_ONEAPI_GRAPH)

SYCL: Add support for Graphs #6912

Are you sure you want to change the base?

SYCL: Add support for Graphs #6912

Conversation

masterleinad commented Apr 2, 2024 • edited

masterleinad commented May 8, 2024

dalg24 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

masterleinad commented May 9, 2024

dalg24 commented May 9, 2024

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

masterleinad commented May 31, 2024

masterleinad commented Apr 2, 2024 •

edited