Optimization: Reduce the number of pi_event-related PI calls

[al42and]

Is your feature request related to a problem? Please describe

When in-order queues are used, some events returned from queue::submit can be discarded. This leads to a repeated "createEvent/recordEvent/destroyEvent" around any submitted operation.

Here is an example from our application running with the CUDA backend, where we see the following sequence of operations:

• Submission of BondedKernel:
  • cuEventCreate
  • Enqueueing the kernel
  • cuEventRecord
  • cuEventDestroy, when the return sycl::event is discarded
• Submission of NbnxmKernel:
  • cuEventCreate
  • Enqueueing the kernel
  • cuEventRecord
  • cuEventDestroy, when the return sycl::event is again discarded.

This can be relevant in the cases when multiple small kernels are submitted to the device

I know of two existing mitigations:

• sycl_ext_oneapi_discard_queue_events: Does not allow any events, which is overly limiting. E.g., one might wish to submit a chain of kernels, then get an event for their completion.
• SYCL_PI_LEVEL_ZERO_REUSE_DISCARDED_EVENTS: Level-Zero specific and requires the use of the previous one.

Describe the solution you would like

• HIPSYCL_EXT_COARSE_GRAINED_EVENTS seems to offer a good balance by letting the developers annotate when the recording of an event can be skipped.
  • For GROMACS, returning either an OpenSYCL-style coarse-grained event (which does a queue::wait when waited upon) or a DPC++-style invalid event (which throws when waited upon) is fine.
• Caching and reusing events instead of creating and destroying them. This seems to be done in [SYCL] Reuse discarded L0 events in scope of command list #7256, but only for LevelZero and only when the whole queue is in "discard_events" mode.
  • CUDA events can be re-used just fine. I strongly suspect the same for HIP.

Describe alternatives you have considered

• Using sycl_ext_oneapi_discard_queue_events and a host-task in a separate, non-discarding queue doing queue::wait whenever an event is needed. It seems overly complicated and highly likely inefficient when the event is needed for synchronizing between two GPU queues.

[al42and]

We got the annotations working :) So here's an example, NVIDIA backend, RTX3060, CUDA 11.8, IntelLLVM afebb25. In-order queues only.

In CUDA, we're doing cudaStreamWaitEvent to wait for another stream/queue, launch three kernels back-to-back, record an event after the last one to use for synchronization later. ~8µs for all:

With DPC++, this is interspersed with event manipulations, and takes ~21µs:

Another example, not so wide. Launching two memsets:

Native CUDA, 3.4µs. Ask for two memsets and receive two memsets:

DPC++, 6.8µs. Extra cuStreamWaitEvent, 2*cuEventCreate, 2*cuEventRecord, 2*cuEventDestroy:

Overall, in this short run, we have 35k cuEventRecord with DPC++, and only 21k actual GPU operations (kernel launch, memcpy, memset). That seems excessive.

Num Calls	Name
35610	cuEventRecord
25406	cuEventCreate
25397	cuEventDestroy_v2
15493	cuLaunchKernel
15182	cuStreamWaitEvent
3105	cuMemsetD8Async
2592	cuMemcpyAsync
1282	cuEventSynchronize
86	cuMemAllocHost_v2
86	cuMemFreeHost
47	cuStreamSynchronize
46	cuMemAlloc_v2
34	cuMemFree_v2
23	cuModuleLoadDataEx
23	cuModuleUnload
8	cuStreamCreate
8	cuStreamDestroy_v2
1	cuMemcpyDtoHAsync_v2
1	cuCtxSetCurrent

For long-running kernels, the impact of such things is minuscule. But when we're talking about shorter kernels, which are common in MD, that becomes a problem.

[MartinWehking]

@al42and The issue seems very interesting. I'm having some trouble reproducing the issue locally. Could you perhaps share your code (for the first example in your second comment)?

[al42and]

@MartinWehking: it's part of a larger application trace, obtained with ONEAPI_DEVICE_SELECTOR=cuda:gpu nsys profile --trace=cuda,nvtx gmx mdrun -nb gpu -pme gpu -update gpu -ntmpi 1 -ntomp 8 -nsteps 1000 -noconfout.

I can look into making a reduced example. Are you more interested in timings or just the number of extraneous API calls?

[MartinWehking]

That would be great!
I'm more interested in the number of API calls and the order (timelines) in which they occur in the sycl and cuda example.

[al42and]

Here's a toy example. Submit the same two kernels five times, then copy the data. We use a single in-order queue, so no synchronization is strictly necessary.

If we discard the first (JIT / warm-up) and the last (extra data copy) iterations, with oneAPI submitting two kernels takes ~7.4µs; with CUDA: ~3.4µs. Kernel performance is approximately the same (no surprises here). The gaps between kernels in GPU utilization are ~1.5µs with oneAPI, and ~0.6µs with CUDA.

$ clang++ -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend=nvptx64-nvidia-cuda --offload-arch=sm_86 8390.cpp -L/usr/local/cuda/lib64/ -lnvToolsExt -o 8390 && ONEAPI_DEVICE_SELECTOR=cuda:gpu nsys profile --force-overwrite true --output=sycl --trace=cuda,nvtx ./8390

The small yellow bits around each kernel are cuEventCreate, cuEventRecord, cuEventDestroy.

$ nvcc -O3 8390.cu -lnvToolsExt -o 8390 && nsys profile --force-overwrite true --output=cuda --trace=cuda,nvtx ./8390

With CUDA, we can much more quickly schedule all the GPU work and, potentially, do some CPU work while GPU is busy. With oneAPI, however, the CPU can barely schedule things for the GPU.

Code: https://gist.github.com/al42and/1c9e166abd02c76772214d50b76150a0

NVIDIA Nsight Systems traces (gzipped):

• cuda.nsys-rep.gz
• sycl.nsys-rep.gz

[G-071]

I would like second the suggestion of having some functionality that caches and reuses events when possible.

When running octotiger with oneAPI (via the Kokkos SYCL execution space), we see a non-negligible overhead for the event creations/destructions. When using plain CUDA we mostly get around this by using the existing event caching within HPX, but that does not seem to be an option when using SYCL (as we can't create and use our own events with the API).

[al42and]

Related discussion about making the return of events optional: KhronosGroup/SYCL-Docs#404

[al42and]

May I inquire whether there are any planned work here?

I ran some tests on AMD platform with ROCm 5.3.3 and MI250X, and the things are worse there. Submitting two kernels typically takes ~5µs with native HIP:

And >15µs with IntelLLVM 2023-09-15:

So, we're talking about >3x slower submission.

[al42and]

According to a question on StackOverflow, this problem could be exacerbated when running on WSL.