synchronize cuda-aware mpi streams

[saareliad]

Background information

• v4.0.3

• installed from source (tar)

• cuda aware mpi

• cuda 10.2

• This is not a system problem, but suspected behavior/implementation issue in cuda-aware MPI. it will happen on all systems

---

Details of the problem

Inside cuda-aware MPI (here) you use async cuda streams to send messages. However, user's program run on other streams.
Therefore, your streams in cuda-aware implementation should be able to wait for work completion of user's streams,
otherwise, it would results in incorrect programs,
or it will force users to fully synchronize its streams before calling MPI.
See Pytorch discussion on the matter.

Possible solutions: expose the streams to the user, or (preferable) let the user allocate and manage them.

[bosilca]

MPI implementors, and OMPI developers in particular, have been talking for quite some time about similar concepts, but without converging to a suitable solution.

First of all, I think there is a misunderstanding on how OMPI uses CUDA streams internally. To be precise, we use streams in some cases, but not necessarily all the time, and these streams might be bound to one side on a send/recv communication pair (depending is we translate the point-to-point communication into a put or a get). Plus, with the new generations of hardware supporting GPU async and triggered operations we will be certainly moving further away from using these stream constructs.

Second, what looks from the user perspective as a single MPI call, might in reality be decomposed internally into a number of operations, some of them being able to be executed in parallel. In other words, in order to provide performance we need to use multiple streams, streams that would have no meaning to anybody outside the low-level internals of the MPI implementation.

Thus, in order to provide users with a stream to synchronize with, we have 2 major barriers.

• OMPI will need to expose unsavory (from a user perspective) details about the internal implementation of the CUDA support. Internally we divide the data movements across several stream (ipcStream for anything related to IPC, dtohStream and htodStream for transfers between host and GPU memory and finally memcpyStream for everything else). From a user perspective it is unknown what these streams are for, or how OMPI divides each data transfer into fragments and which streams hosts these fragments. So, even if we expose these streams to the user level, they will not be portable across MPI implementations (data transfer mechanisms differs) or even across different versions of the same implementations.
• performance in important. If OMPI has to limit the internal parallelism and sequentialize everything into a single, user exposed, stream, and ditch all internal pipelines, the performance will suffer. Not a good selling point for an MPI implementation.

Let me get back to the example posted on the pytorch discussion. If the goal is to be able to write code like

kernel1<<<..., mystream>>>(data, ...);
MPI_Isend(data);
kernel2<<<..., mystream>>>(data, ...);

and have the MPI_Isend communication send the data generated in kernel1, you should do it directly in pytorch, using futures and a .then/when relationship.

[sjeaugey]

@saareliad this is one of the key difference between MPI and NCCL. However, there are pros and cons on both sides.

To be able to guarantee the example George mentioned above would work, the MPI calls would need to become CUDA stream oriented, in both directions. You'd need MPI operations to be blocked from making progress until the stream is ready, and you would need all subsequent operations on the stream to block until the MPI operation is complete.

Achieving the first is relatively easy. You can use an event to know when the stream has reached that point and only start progressing the MPI operation after that.

Achieving the other direction is the main issue. It means that before the MPI operation can return, all data movement operations (cudaMemcpy or CUDA kernels) must be posted on the stream. Since they depend on external factors, e.g. data being ready to write or read, or even knowing which pointer to write to (for a send), this can severely limit the implementation as to which technique to use for copies. The only solution I see to achieve that, especially in the context of performing inter-node communication, is to do what NCCL does : run a CUDA kernel spinning and blocking, waiting for networks transfers to be complete. But this comes with side effects : multiple communication operations may block each other (even if run on different, asynchronous streams) and therefore cause deadlocks, which is the subject of numerous questions on the NCCL github project.

To conclude, at the moment you have an explicit choice which is to use NCCL which is running on the GPU, has CUDA stream semantics and CUDA constraints, or use MPI which is running on the CPU, with CPU semantics and CPU constraints.

cc @Akshay-Venkatesh

[leofang]

Hi, we also discussed this issue a while ago when attempting to revise the __cuda_array_interface__ protocol on the Python side, see numba/numba#5162. (It's a protocol that many Python GPU libraries, including mpi4py and PyTorch, support to interoperate with each other with zero copy of GPU data.) I thought in Open MPI cuStreamSynchronize or cuEventQuery is always called to sync the internal streams before an MPI call is returned? If my understanding on this is incorrect, this would open up a zoo of issues we worried about...

cc: @gmarkall @dalcinl

[Akshay-Venkatesh]

I thought in Open MPI cuStreamSynchronize or cuEventQuery is always called to sync the internal streams before an MPI call is returned? If my understanding on this is incorrect, this would open up a zoo of issues we worried about...

This isn't true. MPI implementations are expected to honor MPI semantics and this doesn't require synchronizing internal streams. The choice of how to treat internal streams and how MPI semantics are guaranteed are independent. As an example, for performance reasons a call to MPI_Isend would try not to synchronize on an internal stream that was, say, used to initiate inter-GPU data movement. If anything, the status of the transfer is queried during MPI_Wait*.

[saareliad]

@bosilca , I have 2 suggestions. (1) why not pass multiple (optional) streams to MPI?
if the user won't pass anything - do default (current) behavior. Otherwise, use user's streams.

It would also allow additional level of control we currently miss. For example, for intra-node communication, its sometimes useful using streams with different priorities.

Let's say we implement the Isend example in pytorch with futures.
Without an exposed cuda-stream, we would need a CPU thread blocking (synchronize()) on some cuda Event, which is "previous work Event" of user's streams.
The CPU thread would call MPI only when this "previous work" ends.
This means a thread per send, and probably waiting the GPU-CPU latency twice just to make it work.

This brings me to additional option (2) - user gives MPI a "previous work" cuda Event to wait on. MPI cuda internals wait for than event to know the data is ready.
This could work without exposing any internals.
Moreover when you don't use cuda streams, wait on the event with the CPU.

(By the way, I spotted this issue with intra-node communication actually, so I assumed cuda-streams doing the send)

[dalcinl]

@saareliad Have you guys ever though about using MPI_Send_init() and then MPI_Start() the (persistent) MPI request on an event callback [using cudaLaunchHostFunc()]? This approach should at least fix the issues on the sending side.

PS: Do not forget to explicitly free the request after waiting for it.

Edit: It may not be so straightforward. IIRC, you cannot wait on a persistent request that has not been started. Maybe wrapping things on a generalized MPI request would work, but at this point I'm not sure about the all details on the CUDA side.

[saareliad]

@saareliad Have you guys ever though about using MPI_Send_init() and then MPI_Start() the (persistent) MPI request on an event callback [using cudaLaunchHostFunc()]? This approach should at least fix the issues on the sending side.

PS: Do not forget to explicitly free the request after waiting for it.

Edit: It may not be so straightforward. IIRC, you cannot wait on a persistent request that has not been started. Maybe wrapping things on a generalized MPI request would work, but at this point I'm not sure about the all details on the CUDA side.

It seems like a good idea at first, but I briefly checked now, it seems problematic (see 1, 2):
"Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback."

[sjeaugey]

@saareliad tying an MPI communicator to a CUDA stream is something we explored and as I mentioned in my comment above, the main problem is on the receive side. Blocking the progress of an MPI_Isend based on a CUDA stream wouldn't be too hard. But blocking a CUDA stream on an MPI_Irecv is a much harder thing to do. And it would be weird to provide a CUDA-oriented mechanism that only works in a single direction.

Do you really only need to link MPI and CUDA on the send side ?

Besides, having a CPU thread waiting until the CUDA stream is ready and launching the send isn't as bad as you seem to think. The extra latency is needed anyway unless you can completely offload the MPI operation to CUDA, and having an active CPU thread ensures quick progress and completion of the MPI operation if multiple steps are needed.

[leofang]

Hi @bosilca @sjeaugey @Akshay-Venkatesh I am revisiting this issue to ensure my understanding on what CUDA-awareness can and cannot do, as it's little documented AFAIK...So far we've put this in production without noticing issues, so I am still surprised how complicated the problem is.

Let's revisit @bosilca's earlier (C/C++) example:

kernel1<<<..., mystream>>>(data, ...);
MPI_Isend(data);  // or the sync version MPI_send; personally I'm more interested in MPI_Allreduce/MPI_Iallreduce
kernel2<<<..., mystream>>>(data, ...);

My question is, for such common CUDA/MPI mix-up operations, what do we need to do to guarantee safety for the 4 combinations (default/non-default stream x sync/async)? Could you provide a rule of thumb as general guidance? Thanks.

(ps. I am mainly concerned with the C/C++ semantics; we will figure out how to make it safe in Python once we have a solid understanding.)

[sjeaugey]

CUDA-aware MPI only means you can use pointers to CUDA memory. It does not mean anything w.r.t. streams (at least today).

In the example, the only thing you need is to add a cudaStreamSynchronize before MPI_Isend to make it correct (the stream being default/sync/async doesn't matter):

kernel1<<<..., mystream>>>(data, ...);
cudaStreamSynchronize(mystream);
MPI_Isend(data);

[leofang]

Thanks, @sjeaugey!

It does not mean anything w.r.t. streams (at least today).

Right, which is why this issue was opened and why I'm seeking guidance 🙂

In the example, the only thing you need is to add a cudaStreamSynchronize before MPI_Isend to make it correct (the stream being default/sync/async doesn't matter):

I should have been clearer. By "sync/async" I am referring to the MPI call (ex: MPI_Send vs MPI_Isend, see the comment I left in the code snippet), not the cuda stream.

Next, how about the subsequent kernel launch following the MPI call? What do we need to ensure the correctness of its operation?

[sjeaugey]

In the CUDA -> MPI direction you need a call to cudaStreamSynchronize to make sure CUDA operations are complete.

Similarly, in the MPI -> CUDA direction you need to wait for MPI operations to be complete, i.e. call MPI_Wait[all/any], or use synchronous operations, before starting new CUDA kernels.

Now I didn't reuse the example above, as kernel2 could still be launched after MPI_Isend (without waiting for the send operation to be complete) as long as it doesn't modify data (only reads it), which wasn't specified. But in the case where the MPI operation is receiving data (e.g. MPI_Irecv) we need to wait until the receive is complete before we start a CUDA kernel on that same data.
But that's all standard MPI semantics, those are not CUDA specific.