Copyright © 2024 Intel Corporation. All rights reserved.
Khronos® is a registered trademark and SYCL™ and SPIR™ are trademarks of The Khronos Group Inc. OpenCL™ is a trademark of Apple Inc. used by permission by Khronos.
To report problems with this extension, please open a new issue at:
This extension is written against the SYCL 2020 revision 8 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision.
This is an experimental extension specification, intended to provide early access to features and gather community feedback. Interfaces defined in this specification are implemented in DPC++, but they are not finalized and may change incompatibly in future versions of DPC++ without prior notice. Shipping software products should not rely on APIs defined in this specification.
The base extension is currently implemented in DPC++ only for GPU devices and only when using the CUDA or HIP backends. Attempting to use this extension in kernels that run on other devices or backends may result in undefined behavior. Be aware that the compiler is not able to issue a diagnostic to warn you if this happens.
The semantics in the SYCL-Graph Interaction section
however are implemented for all of Level-Zero, OpenCL, CUDA and HIP devices.
Where support is conditional on the device reporting the
aspect::ext_oneapi_limited_graph aspect.
This extension is derived from the experimental AdaptiveCpp extension,
enqueue_custom_operation which is documented
here.
The goal of ext_codeplay_enqueue_native_command is to integrate interop
work within the SYCL runtime’s creation of the asynchronous SYCL DAG. As such,
the user defined lambda must only enqueue asynchronous, as opposed to
synchronous, backend work within the user lambda. Asynchronous work must only
be submitted to the native queue obtained from
interop_handle::get_native_queue.
A callable submitted to ext_codeplay_enqueue_native_command won’t wait
on its dependent events to execute. The dependencies passed to an
ext_codeplay_enqueue_native_command submission will result in dependencies being
implicitly handled in the backend API, using the native queue object associated
with the SYCL queue that the sycl_ext_codeplay_enqueue_native_command is
submitted to. This gives different synchronization guarantees from normal SYCL
host_task s, which guarantee that the host_task callable will only begin
execution once all of its dependent events have completed.
In this example:
q.submit([&](sycl::handler &cgh) {
cgh.depends_on(dep_event);
cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle h) {
printf("This will print before dep_event has completed.\n");
// This stream has been synchronized with dep_event's underlying
// hipEvent_t
hipStream_t stream = h.get_native_queue<sycl::backend::ext_oneapi_hip>();
hipMemcpyAsync(target_ptr, native_mem, test_size * sizeof(int),
hipMemcpyDeviceToHost, stream);
});
});
q.wait();The print statement may print before dep_event has completed. However, the
asynchronous memcpy submitted to the native queue obtained by
interop_handle::get_native_queue is guaranteed to have the correct
dependencies, and therefore will only start once its dependent events have
completed.
By contrast, when using a host_task, it is guaranteed that the print statement
will only happen once the host task’s dependent events are observed to be
complete on the host.
A SYCL event returned by a submission of a
ext_codeplay_enqueue_native_command command is only complete once the
asynchronous work enqueued to the native queue obtained through
interop_handle::get_native_queue() has completed.
This extension provides a feature-test macro as described in the core SYCL
specification. An implementation supporting this extension must predefine the
macro SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND to one of the values defined
in the table below. Applications can test for the existence of this macro to
determine if the implementation supports this feature, or applications can test
the macro’s value to determine which of the extension’s features the
implementation supports.
| Value | Description |
|---|---|
1 |
The APIs of this experimental extension are not versioned, so the feature-test macro always has this value. |
This extension adds the following new member function to the SYCL handler
class:
class handler {
template <typename Func>
void ext_codeplay_enqueue_native_command(Func&& interopCallable);
};Constraints: The Func must a C++ callable object which takes a single
parameter of type interop_handle.
Effects: The interopCallable object is called exactly once, and this call
may be made asynchronously even after the calling thread returns from
ext_codeplay_enqueue_native_command.
The call to interopCallable may submit one or more asynchronous tasks to the
native backend object obtained from interop_handle::get_native_queue, and
these tasks become encapsulated in a SYCL command that is added to the queue.
If the enclosing command group has any dependencies, these dependencies are
propagated to the native asynchronous tasks. This happens, for example, if the
command group calls handler::depends_on or if it constructs an accessor. As a
result, there is typically no need to specify these dependencies through native
APIs. Note, however, that these dependencies are associated with the
asynchronous tasks submitted by interopCallable, not the call to
interopCallable. The call to interopCallable may happen even before the
dependencies are satisfied.
The SYCL command described above completes once all of the native asynchronous tasks it contains have completed.
The call to interopCallable must not add tasks to backend objects that underly
any other queue, aside from the queue that is associated with this handler,
otherwise, the behavior is undefined.
[Note: The function object interopCallable is invoked to enqueue commands to a
native queue or graph and therefore, APIs which block or synchronize could
prolong or interfere with other commands being enqueued to the backend.
{endnote}]
This section defines the interaction with the sycl_ext_oneapi_graph extension.
The interopCallable object will be invoked during command_graph::finalize()
when the backend object for the graph is available to give to the user as a
handle. The user may then add nodes using native APIs to the backend graph
object queried with interop_handle::ext_codeplay_get_native_graph(). The
runtime will schedule the dependencies of the user added nodes such
that they respect the graph node edges.
// Alias is for editorial brevity in the ext_codeplay_get_native_graph
// definition, and is non-normative.
using graph = ext::oneapi::experimental::command_graph<
ext::oneapi::experimental::graph_state::executable>;
class interop_handle {
bool ext_codeplay_has_graph() const;
template <backend Backend>
backend_return_t<Backend, graph> ext_codeplay_get_native_graph() const;
};Table 1. Additional member functions of the sycl::interop_handle class.
| Member function | Description |
|---|---|
bool interop_handle::ext_codeplay_has_graph() const; |
Returns: True if the [Note: that host-task nodes in a |
template <backend Backend>
backend_return_t<Backend, graph>
interop_handle::ext_codeplay_get_native_graph() const; |
Returns: The native graph object associated with the Throws: An exception with the |
When interop_handle::get_native_queue() is invoked in a native command
function object on graph finalize, the queue that is returned to the user is an
internal queue created by the SYCL runtime, as there is no user provided queue
at the point of graph finalization. This queue has the same device and context
as the graph was created with. The only valid usage of this queue is to perform
stream capture to a graph for backend APIs that provide this functionality.
Table 2. Native types for
template <backend Backend, class T> backend_return_t<Backend, T> where T is
instantiated as command_graph<graph_state::executable>.
| Backend | Native Graph Type |
|---|---|
|
|
|
|
|
|
|
|
This example demonstrates how to use this extension to enqueue asynchronous native tasks on the HIP backend.
sycl::queue q;
q.submit([&](sycl::handler &cgh) {
sycl::accessor acc{buf, cgh};
cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle h) {
// Can extract device pointers from accessors
void *native_mem = h.get_native_mem<sycl::backend::ext_oneapi_hip>(acc);
// Can extract stream
hipStream_t stream = h.get_native_queue<sycl::backend::ext_oneapi_hip>();
// Can enqueue arbitrary backend operations. This could also be a kernel
// launch or call to a library that enqueues operations on the stream etc
hipMemcpyAsync(target_ptr, native_mem, test_size * sizeof(int),
hipMemcpyDeviceToHost, stream);
});
});
q.wait();This example demonstrates how to use this extension to add a native command to a SYCL-Graph object on the Level-Zero backend. The command is doing a memory copy between two USM pointers.
Graph.add([&](sycl::handler &CGH) {
CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
ze_command_list_handle_t NativeGraph =
IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_level_zero>();
zeCommandListAppendMemoryCopy(
NativeGraph, PtrY, PtrX, Size * sizeof(int), nullptr, 0, nullptr);
});
});This example demonstrates how to use this extension to add a native command to a SYCL-Graph object on the OpenCL backend. The command is doing a copy between two buffer objects.
sycl::queue Queue;
auto Platform = get_native<sycl::backend::opencl>(Queue.get_context().get_platform());
clCommandCopyBufferKHR_fn clCommandCopyBufferKHR =
reinterpret_cast<clCommandCopyBufferKHR_fn>(
clGetExtensionFunctionAddressForPlatform(Platform, "clCommandCopyBufferKHR"));
Graph.add([&](sycl::handler &CGH) {
auto AccX = BufX.get_access(CGH);
auto AccY = BufY.get_access(CGH);
CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
cl_command_buffer_khr NativeGraph =
IH.ext_codeplay_get_native_graph<sycl::backend::opencl>();
auto SrcBuffer = IH.get_native_mem<sycl::backend::opencl>(AccX);
auto DstBuffer = IH.get_native_mem<sycl::backend::opencl>(AccY);
clCommandCopyBufferKHR(
NativeGraph, nullptr, nullptr, SrcBuffer[0], DstBuffer[0], 0, 0,
Size * sizeof(int), 0, nullptr, nullptr, nullptr);
});
});This example demonstrates how to use this extension to add a native command to a SYCL-Graph object on the CUDA backend. The command is doing a memory copy between two device USM pointers.
Graph.add([&](sycl::handler &CGH) {
CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
CUgraph NativeGraph =
IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_cuda>();
CUDA_MEMCPY3D Params;
std::memset(&Params, 0, sizeof(CUDA_MEMCPY3D));
Params.srcMemoryType = CU_MEMORYTYPE_DEVICE;
Params.srcDevice = (CUdeviceptr)PtrX;
Params.srcHost = nullptr;
Params.dstMemoryType = CU_MEMORYTYPE_DEVICE;
Params.dstDevice = (CUdeviceptr)PtrY;
Params.dstHost = nullptr;
Params.WidthInBytes = Size * sizeof(int);
Params.Height = 1;
Params.Depth = 1;
CUgraphNode Node;
CUcontext Context = IH.get_native_context<sycl::backend::ext_oneapi_cuda>();
cuGraphAddMemcpyNode(&Node, NativeGraph, nullptr, 0, &Params, Context);
});
});This example demonstrates how to use this extension to add a native command to a SYCL-Graph object on the HIP backend. The command is doing a memory copy between two device USM pointers.
Graph.add([&](sycl::handler &CGH) {
CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
HIPGraph NativeGraph =
IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_hip>();
HIPGraphNode Node;
hipGraphAddMemcpyNode1D(&Node, NativeGraph, nullptr, 0, PtrY, PtrX,
Size * sizeof(int), hipMemcpyDefault);
});
});This example demonstrates how to use this extension to add stream recorded native nodes to a SYCL-Graph object on the CUDA backend.
q.submit([&](sycl::handler &CGH) {
CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
auto NativeStream = h.get_native_queue<cuda>();
if (IH.ext_codeplay_has_graph()) {
auto NativeGraph =
IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_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 {
myNativeLibraryCall(NativeStream);
}
});
});