[SYCL][Graph] Support for native-command #16871
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
EwanC
force-pushed
the
graph_native_enqueue
branch
from
4fa2a2f to
2ead142
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
2ead142 to
80269a8
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
80269a8 to
d40e0d1
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
d40e0d1 to
4aeba98
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
2 times, most recently
from
e4498ce to
7d12878
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
7d12878 to
764aebc
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
764aebc to
4e62db9
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
3 times, most recently
from
047eb0e to
c07c67a
Compare
EwanC
changed the title
EwanC
force-pushed
the
graph_native_enqueue
branch
2 times, most recently
from
f499268 to
19bb3dc
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
ac21cad to
3661bfe
Compare
EwanC
force-pushed
the
graph_native_enqueue
branch
from
3661bfe to
41687a3
Compare
Support [sycl_ext_codeplay_enqueue_native_command](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc) with SYCL-Graph. Introduces `interop_handle::ext_codeplay_get_native_graph<backend>()` to give the user access to the native graph object which native commands can be appended to. To use CUDA as an example, code using `ext_codeplay_enqueue_native_command` eagerly can be updated from: ```cpp CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) { auto NativeStream = IH.get_native_queue<cuda>(); myNativeLibraryCall(NativeStream); } ``` To ```cpp CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) { if (IH.ext_codeplay_has_graph()) { auto NativeGraph = IH.ext_codeplay_get_native_graph<cuda>(); auto NativeStream = IH.get_native_queue<cuda>(); // Start capture stream calls into graph cuStreamBeginCaptureToGraph(NativeStream, NativeGraph, nullptr, nullptr, 0, CU_STREAM_CAPTURE_MODE_GLOBAL); myNativeLibraryCall(NativeStream); // Stop capturing stream calls into graph cuStreamEndCapture(NativeStream, &NativeGraph); } else { auto NativeStream = IH.get_native_queue<cuda>(); myNativeLibraryCall(NativeStream ); } } ``` Example of how this integration could work in GROMACS https://gitlab.com/gromacs/gromacs/-/merge_requests/4954
EwanC
force-pushed
the
graph_native_enqueue
branch
from
41687a3 to
d82bb30
Compare
|
@intel/llvm-gatekeepers This is ready to merge, thanks |
sommerlukas
pushed a commit
that referenced
this pull request
Mar 21, 2025
The merged PR #16871 introduced new UR entry-points `urCommandBufferAppendNativeCommandExp` and `urCommandBufferGetNativeHandleExp`, however there was no UR CTS coverage for `urCommandBufferAppendNativeCommandExp` included, as verification was done via SYCL E2E testing. This PR adds UR CTS testing for `urCommandBufferAppendNativeCommandExp` for all of the L0, CUDA, HIP, and OpenCL adapters. Resolves #17448
kbenzie
pushed a commit
to oneapi-src/unified-runtime
that referenced
this pull request
Mar 21, 2025
The merged PR intel/llvm#16871 introduced new UR entry-points `urCommandBufferAppendNativeCommandExp` and `urCommandBufferGetNativeHandleExp`, however there was no UR CTS coverage for `urCommandBufferAppendNativeCommandExp` included, as verification was done via SYCL E2E testing. This PR adds UR CTS testing for `urCommandBufferAppendNativeCommandExp` for all of the L0, CUDA, HIP, and OpenCL adapters. Resolves intel/llvm#17448
EwanC
added a commit
to reble/llvm
that referenced
this pull request
May 5, 2025
In intel#16871 the `sycl_ext_codeplay_enqueue_native_command` extensions was extended to add new `interop_handler` APIs for working with SYCL-Graph. However, the extension macro was not bumped. This is problematic for users that want to use the extension with graph support, but also use older oneAPI releases. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ``` By bumping the feature test macro, users can write code that supports old DPC++ versions. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ```
EwanC
added a commit
to reble/llvm
that referenced
this pull request
May 13, 2025
In intel#16871 the `sycl_ext_codeplay_enqueue_native_command` extensions was extended to add new `interop_handler` APIs for working with SYCL-Graph. However, the extension macro was not bumped. This is problematic for users that want to use the extension with graph support, but also use older oneAPI releases. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ``` By bumping the feature test macro, users can write code that supports old DPC++ versions. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ```
EwanC
added a commit
to reble/llvm
that referenced
this pull request
May 13, 2025
In intel#16871 the `sycl_ext_codeplay_enqueue_native_command` extensions was extended to add new `interop_handler` APIs for working with SYCL-Graph. However, the extension macro was not bumped. This is problematic for users that want to use the extension with graph support, but also use older oneAPI releases. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ``` By bumping the feature test macro, users can write code that supports old DPC++ versions. ```cpp CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph() /* is this defined?*/) { // Graph path } else { // Eager path } CGH.host_task(...) ```
martygrant
pushed a commit
that referenced
this pull request
May 14, 2025
In #16871 the `sycl_ext_codeplay_enqueue_native_command` extensions was extended to add new `interop_handler` APIs for working with SYCL-Graph. However the extension macro was not bumped, which I think was an oversight. This is problematic for users that want to use the extension with graph support, but also use older oneAPI releases. ```cpp #ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { if (IH.ext_codeplay_has_graph()) { // Is this API defined? // Graph path } else { // Eager path } #else CGH.host_task(...) #endif ``` By bumping the feature test macro users can write code that supports old DPC++ versions. ```cpp #ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) { #if SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND > 1 if (IH.ext_codeplay_has_graph()) { // Graph path } else #endif { // Eager path } #else CGH.host_task(...) #endif ```
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Support sycl_ext_codeplay_enqueue_native_command with SYCL-Graph for all of L0, CUDA, HIP, and OpenCL backends.
Introduces
interop_handle::ext_oneapi_get_native_graph<backend>()to give the user access to the native graph object which native commands can be appended to. Implemented using new UR command-buffer entry-pointsurCommandBufferAppendNativeCommandExpandurCommandBufferGetNativeHandleExp.To use CUDA as an example, code using
ext_codeplay_enqueue_native_commandeagerly can be updated from:CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) { auto NativeStream = IH.get_native_queue<cuda>(); myNativeLibraryCall(NativeStream); }To
CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) { if (IH.ext_oneapi_has_graph()) { auto NativeGraph = IH.ext_oneapi_get_native_graph<cuda>(); auto NativeStream = IH.get_native_queue<cuda>(); // Start capture stream calls into graph cuStreamBeginCaptureToGraph(NativeStream, NativeGraph, nullptr, nullptr, 0, CU_STREAM_CAPTURE_MODE_GLOBAL); myNativeLibraryCall(NativeStream); // Stop capturing stream calls into graph cuStreamEndCapture(NativeStream, &NativeGraph); } else { auto NativeStream = IH.get_native_queue<cuda>(); myNativeLibraryCall(NativeStream ); } }Example of how this integration could work in GROMACS https://gitlab.com/gromacs/gromacs/-/merge_requests/4954