(Question) Does `SequentialOrderManager` track externally-submitted work on a shared `sycl::queue`

[abagusetty]

@ndgrigorian @vlad-perevezentsev @antonwolfy @alvarovm

Just wanted to reach out related to a setup we have been working with an application using cupy and dpnp (porting).

App: Heavily uses cupy and hand-written CUDA kernels. These were ported to DPNP and SYCL kernels for Intel PVC (on Aurora).

Intent: The app relies on the use of stream = cupy.cuda.get_current_stream() , stream.synchronize(), and passing the stream obtained by the above to pass it down to user CUDA kernels. A pipeline methodology where a given instance/process is tied to a single stream through the upstream cupy to downstream user CUDA kernels.

With the port to DPNP & SYCL (currently), we have relied on using a single queue written for user kernels and dpnp/dpctl ops wasnt aware of this queue. Since they were all created with the same default context (using ZE_SERIALIZE=2, OverrideImmediateCmdListSynchronousMode=1) helped with the initial correctness testing. Since everything got synchronized and default-context shared by any number of queues wasnt a problem.

Goal: The intent now is to move from the above design and explore if upstream dpctl provides us a single master (in-order) queue per-device q = dpctl.SyclQueue(devs[device_id], property="in_order") such that we can use it to the downstream user SYCL kernels as well. Since everything is in-order we expect similar FIFO behavior as cudaStream

Few questions if you can provide any pointers:
0. Would the above mentioned design be feasible.

1. Since there is no such equivalent as stream = cupy.cuda.get_current_stream(), we have tried something like from ._sycl_queue_manager import get_device_cached_queue but the obtained queue from get_device_cached_queue is inherently an OoO queue which doesnt really suite our broader needs in the user-app. Since we are trying to mimic the infrastructure based on cudaStream FIFO. So we had to start using q = dpctl.SyclQueue(devs[device_id], property="in_order").
2. Does SequentialOrderManager track externally-submitted work on a dpctl obtained sycl::queue

More details:
Nothing in SequentialOrder observes the queue itself. So if a user extension does:

pythonq = dpctl.SyclQueue()
a = dpt.arange(..., sycl_queue=q)          # dpctl registers events
my_ext.launch_kernel(q, a.usm_data, ...)   # user submits directly, NOT registered
b = dpt.sum(a, ...)                        # dpctl uses its tracked deps only

then the my_ext.launch_kernel event is invisible to the manager. A subsequent dpt.sum submission's depends_on will not include it, and SequentialOrderManager[q].wait() will return before the user kernel completes.
Is the following characterization correct?

In-order queue: user-submitted kernels are still correctly ordered against later dpctl submissions because the queue itself serializes, so the data dependency is safe. But SequentialOrderManager[q].wait() may return early because it only waits on events it knows about.
Out-of-order queue: nothing orders the user's kernel against later dpctl submissions, so without the user explicitly passing som.submitted_events as deps and registering their own event back, this is a genuine race on shared USM.

I understand this is a lot and would appreciate any feedback. There are a rich set of DPCPP extensions too that can help in making this robust and lightweight without explicit book-keeping and would be more than welcome to contribute if there is sufficient interest

[ndgrigorian]

@abagusetty
It's fine that it's a lot, it's nice in general to get some questions about the order manager -- good to see it can be useful to an extension other than dpnp. Unfortunately, this response will be quite a lot too.

For point 1, about a master queue, we don't have the same idea. Like you pointed out, OoO is default behavior. One possibility is that we can implement some env var which defaults the constructed cached queue to given properties, it sounds kind of interesting, at the very least in-order does seem to have value. Then the cached queue may fit your use a bit better.

For point 2, the intended work-flow for using order manager can be seen in this example:
https://github.com/IntelPython/dpctl/blob/feature/enable-free-threaded-python/examples/python/using_order_manager.py
this is not yet in the master branch, and it may not even really work without building that branch (some bugs in GIL/no-GIL needed to be ironed out) but you can see the idea:
SequentialOrderManager does not attach itself to a queue, nor are queues constructed knowing about it. A user is expected to construct the SequentialOrderManager via SequentialOrderManager[q] for some queue q that they want to track events on. Then they have to provide it events: a host_task which manages the lifetime of the Python args (to prevent GC), and usually this uses the keep_args_alive construct, and the actual computation event.

SequentialOrderManager itself does give you back the same events when later constructed again with SequentialOrderManager[q], with the important caveat that, the queue must be the same queue as before. Then once again, it is expected the user will pass the submitted events as dependency to the next function that offloads (assuming the events are needed), and later add the returned events.

The expected workflow would be something like, for your example

pythonq = dpctl.SyclQueue()
a = dpt.arange(..., sycl_queue=q)          # dpctl registers events
som = SequentialOrderManager[q]
deps = som.submitted_events
ht_ev, comp_ev = my_ext.launch_kernel(q, a.usm_data, ..., deps=deps)   # user submits directly via argument, now registered
som.add_event_pair(ht_ev, comp_ev)
b = dpt.sum(a, ...)                        # dpctl uses the tracked deps ht_ev, comp_ev

Instead of returning the host_task from the extension, the user can also use q._submit_keep_args_alive(...) with their Python arguments, and use the event returned from there. That's what that example is doing. It would be good for us to document all of these things, as well, and maybe to add a pybind11 example.

And in terms of contribution, dpctl is trying to support more compiler extensions. I'd be curious to know which in particular are of interest, but also, you are more than welcome to contribute to implement them, and it would be greatly appreciated.

[ndgrigorian]

It may also be possible to implement some kind of construct like, set_cached_device_queue where user provides a device and a queue, and it explicitly forces the cache to use them, thus later SyclQueue() constructors would default to that queue for that device... the only issue is if one does this between two calls, it could cause strange behavior.

I will continue to think on what we could do here to ease user experience with the default queues.