Device-wide synchronization in SYCL on NVIDIA GPUs

[Luigi-Crisci]

Context
I'm porting a complex CUDA application to SYCL which uses multiple cudaStream to launch the kernels. In addition, it also uses the default Stream in some cases, forcing a device-wide synchronization.

Problem
Cuda Streams can be mapped quite easily to in order SYCL Queues, however when encountering a device-wide synchronization point (i.e. cudaDeviceSyncronize()), I must explicitly wait on all the queues as queue::wait() waits just on the commands submitted to that queue.

Question
Is there a way to wait on all the commands for a specific device without having to call wait() on every queue explicitly?

[bader]

@YuriPlyakhin do you have thoughts on this?

[YuriPlyakhin]

I would recommend to look at the API implemented here: https://github.com/oneapi-src/SYCLomatic/blob/SYCLomatic/clang/test/dpct/helper_files_ref/include/device.hpp

There is a method: dpct::device_ext::queues_wait_and_throw that might be what you are looking for.
You would need to use device_ext to manage all queues for the device to get benefit of this method though.

@zhimingwang36 to comment as well

[zhimingwang36]

Example code: https://github.com/oneapi-src/SYCLomatic/blob/SYCLomatic/clang/test/dpct/device002.cu#L211
we use dpct::queues_wait_and_throw() to migrate the cudaDeviceSynchronize()

[Luigi-Crisci]

Thank you for providing the code. So basically, it is handled as I was doing: a for-loop over all the queues is the only solution.
Anyway, shouldn't a feature like this be in the standard? It seems meaningful and useful. Does it go against the SYCL philosophy in some way?