Stream and Graph Based MPI Operations

[jdinan]

What

Support for enqueueing MPI operations into accelerator work queues (streams) and compute graphs.

Why

Integration of communication with the computation scheduling model for accelerators improves the programming model, can improve communication/computation overlap, and reduces overheads.

Implementations

• MPI ACX
• MPICH MPIX_Stream Extensions

Slides

• 2020-12-16 -- Hybrid-Accelerator WG (Jim Dinan)
• 2021-03-10 -- Lazy GPU Programming (Stephen Jones)
• 2021-05-18 -- Task Graphs vs MPI (Stephen Jones)
• 2022-09-28 -- Hybrid WG MPIX_Queue (Jim Dinan)
• 2022-09-28 -- EuroMPI-MPIX-Stream (Hui Zhou)

Papers

• MPIX Stream: An Explicit Solution to Hybrid MPI+X Programming. Zhou et al.

[patrickb314]

Looking at the first set of slides, when MPI_Start_enqueue is called on a two-sided operation, when does the match occur, particularly with respect to other issued two-sided operations? When it's enqueued or when it's executed by the assocuated stream?

E.G. if a sender calls:

MPI_Send_init(&req[0]);
MPI_Start_enqueue(&req[0], stream);
MPI_Isend(&req[1]);

Are the sends guaranteed to be matched in the order 0 then 1 (since start_enqueue happened before the isend), or could they be matched in the other order if the match doen't happen until the stream actually executes the enqueued start?

[jdinan]

Matching would occur based on when the operation gets executed by the stream.

I think of streams as being similar to threads. In this case, you have a thread stream that will start the persistent send and the main process that starts the Isend. Since there's no synchronization between the two, the matching order is nondeterministic.

[patrickb314]

Let me ask a slightly different question, then. Is MPI allowed to start matching when Start_enqueue is called, or does it have to wait until the stream triggers the progress engine to start matching? Is there any way the application could observe if MPI started the match early? Obviously MPI has to wait to move data in stream order, but it could be useful for implementations if the match could happen sooner.

The main challenge in doing this would be that Start_enqueue is a local operation, and matching is non-local operation, so you couldn't block Start_enqueue to do the match.

[jdinan]

Taking the thread analogy, the start_enqueue operation is simply putting a work descriptor into a queue, not performing an MPI operation in the traditional sense. Therefore, I would expect that MPI shouldn't start matching an enqueued operation until it's actually executed by whatever is taking work out of the queue and executing it. An exception to this would be partitioned operations, which match as soon as the persistent operation is initialized.

[patrickb314]

The reason I'm asking is that if MPI can optimistically match when Start_enqueue is called, it can potentially separate the non-local matching portion of the enqueued operation from on-stream one-sided data movement.

This really comes down to how much ambiguity is in the standard. If the application cannot observe if MPI started matching when Start_enqueue is called, then MPI could match optimistically and enqueue the appropriate one-sided get or put on the stream, taking matching out of the stream's critical path. If the application can observe if MPI started the match early, however, then this API implicitly requires the stream to either perform matching itself or to synchronize with something else which performs the matching.

EDIT: spelling corrections, removed incorrect thoughts on psend matching

[jdinan]

I think this comes back to the discussion on "logically concurrent" ambiguity in the standard. I made an analogy between enqueued operations and threads. If we strengthened that to a semantic (i.e. each queue has the MPI semantics of a thread), then it would be subject to the "logically concurrent" discussion that @Wee-Free-Scot has been leading. Should we land on the side that "logically concurrent" means that the application can't enforce a specific ordering (e.g. by synchronization between queued operations using as an example CUDA events between streams or the host CPU), then I think we could make optimizations like you described.

To answer your last question -- with CUDA streams you can create an ordering across streams and with the host CPU using CUDA events. So, it would be possible for an application to observe that operations didn't match in the order that they attempted to create using such synchronization. I would expect the MPI communicator to still be the serialization point for matching when the same communicator is used across streams (ignoring relaxations possible with info keys). We could introduce new info assertions to optimize matching for queued operations. An interesting difference between queues and threads is that the MPI library can actually see queues because they're explicit in the API.

[hzhou]

Looking at the first set of slides, when MPI_Start_enqueue is called on a two-sided operation, when does the match occur, particularly with respect to other issued two-sided operations? When it's enqueued or when it's executed by the assocuated stream?

E.G. if a sender calls:

MPI_Send_init(&req[0]);
MPI_Start_enqueue(&req[0], stream);
MPI_Isend(&req[1]);

Are the sends guaranteed to be matched in the order 0 then 1 (since start_enqueue happened before the isend), or could they be matched in the other order if the match doen't happen until the stream actually executes the enqueued start?

Assume all above calls are issued on the same stream (a same communicator with serial context), then mixing immediate operations between an "enqueue" operation and the next stream "synchronization" is undefined or illegal. The usage essentially breaks the "serial" context of a stream.

An "immediate" operation essentially is an "enqueue" followed with an immediate "synchronization". This interpretation allows the "undefined" scenario. That is, it is possible with a reasonable but likely counter-intuitive outcome.

[Wee-Free-Scot]

Looking at the first set of slides, when MPI_Start_enqueue is called on a two-sided operation, when does the match occur, particularly with respect to other issued two-sided operations? When it's enqueued or when it's executed by the assocuated stream?
E.G. if a sender calls:

MPI_Send_init(&req[0]);
MPI_Start_enqueue(&req[0], stream);
MPI_Isend(&req[1]);

Are the sends guaranteed to be matched in the order 0 then 1 (since start_enqueue happened before the isend), or could they be matched in the other order if the match doen't happen until the stream actually executes the enqueued start?

Assume all above calls are issued on the same stream (a same communicator with serial context), then mixing immediate operations between an "enqueue" operation and the next stream "synchronization" is undefined or illegal. The usage essentially breaks the "serial" context of a stream.

An "immediate" operation essentially is an "enqueue" followed with an immediate "synchronization". This interpretation allows the "undefined" scenario. That is, it is possible with a reasonable but likely counter-intuitive outcome.

Personally, I see this as sufficiently similar to forking a separate CPU thread to be able to reason about that analogous situation and draw conclusions that are valid for the enqueue situation.

What I mean is -- calling MPI_Start_enqueue is sufficiently similar to "start a pThread with MPI_Start as the function pointer (or, equivalently, a simple wrapper user function around MPI_Start)" that we substitute that code and expect the same behaviour. This substitution replaces a proposed function, with unproven/not-yet-well-defined semantics, with an extant function with well-known behaviour and outcome.

MPI_Send_init(&req[0]);
MPI_Request *ptrReq = &req[0]; pthread_create(&thread, NULL, &MPI_Start, &ptrReq);
MPI_Isend(&req[1]);

Comments:

• This new code requires MPI_THREAD_MULTIPLE because there is no guard preventing concurrent execution of the MPI_Start at the time time as (or with an overlapping time interval with) the MPI_Isend.

• The threading model does not define a relative order between code executing on the main thread and code executing on the pThread, so the logically concurrent rule applies.

• Both camps of thought on the interpretation of the logically concurrent rule agree that this situation is one where the user cannot know or predict the ordering of the two send operations -- even if MPI sees and obeys an ordering (or imposes and enforces an ordering), the user cannot code the receive side to assume either one ordering or the other.

• This is not undefined or illegal (or erroneous, to use the MPI terminology) but it does run head-on into the ambiguity around the phrase "logically concurrent", which is sub-optimal, IMHO.

• According to Camp A (in the logically concurrent debate), the user could guarantee one ordering or the other by adding synchronisation code, e.g. some kind of pThread barrier in the pThread case or some kind of stream synchronisation in the stream case.

• According to Camp B, the user has no remedy for the ambiguous ordering and must accept that either message could be matched before the other.

• A Camp B implementation might have better performance than a Camp A implementation (tbc - open question) but a Camp A implementation is likely to be seen as easier to understand and use than a Camp B implementation.

• We should strive to avoid this reliance on the interpretation of "logically concurrent" (as far as possible) in the definition of new interfaces.

[hzhou]

• We should strive to avoid this reliance on the interpretation of "logically concurrent" (as far as possible) in the definition of new interfaces.

Agreed! We started to play around with this new idea called "MPI Stream" -- pmodels/mpich#5908. @Wee-Free-Scot It certainly can use some of your early input.

[patrickb314]

Clarification - the match order for persistent sends is defined by when Send_init is called, not Start (or start enqueue) is called so perhaps there is no ambiguity in the matching order in the example I gave. This was a misunderstanding on my part. Get Outlook for iOS<https://aka.ms/o0ukef>

…

________________________________ From: Hui Zhou ***@***.***> Sent: Sunday, March 27, 2022 12:49:56 PM To: mpiwg-hybrid/hybrid-issues ***@***.***> Cc: Patrick Bridges ***@***.***>; Comment ***@***.***> Subject: Re: [mpiwg-hybrid/hybrid-issues] Stream and Graph Based MPI Operations (#5) * We should strive to avoid this reliance on the interpretation of "logically concurrent" (as far as possible) in the definition of new interfaces. Agreed! We started to play around with this new idea called "MPI Stream" -- pmodels/mpich#5908<pmodels/mpich#5908>. @Wee-Free-Scot<https://github.com/Wee-Free-Scot> It certainly can use some of your early input. — Reply to this email directly, view it on GitHub<#5 (comment)>, or unsubscribe<https://github.com/notifications/unsubscribe-auth/ACQTKTQ34POWWHZGCZOHLNDVCCUVJANCNFSM5F5QHLFQ>. You are receiving this because you commented.Message ID: ***@***.***>

[Wee-Free-Scot]

the match order for persistent sends is defined by when Send_init is called, not Start (or start enqueue) is called

This is only true for partitioned operations. Regular persistent point-to-point has always matched anew each time it is active (conceptually during the starting stage -- MPI_Start or MPI_Start_all -- although in practice the protocol messages might happen any time before the associated completion stage) whereas partitioned point-to-point matches once (conceptually during the initialisation stage -- MPI_Psend_init -- although in practice the protocol messages might happen any time before the first completion stage).

MPI-4.0 p107 lines 21-23 (comparison of partitioned and regular persistent point-to-point match order):

Advice to implementors. Unlike MPI_SEND_INIT, MPI_PSEND_INIT can be matched as early as the initialization call.

MPI-4.0 p107 lines 13-15 (partitioned point-to-point match order defined by initialisation procedure order):

In the event that the communicator, tag, and source do not uniquely identify a message, the order in which partitioned communication initialization calls are made is the order in which they will eventually match.

MPI-4.0 p101 line 21 (regular persistent point-to-point must be started to permit matching):

A send operation started with MPI_START can be matched ...

MPI-4.0 p94 lines 24-27 (regular persistent point-to-point forms a half-channel):

In the case of point-to-point communication, the persistent communication request thus created can be thought of as a communication port or a “half-channel.” It does not provide the full functionality of a conventional channel, since there is no binding of the send port to the receive port.

Interestingly, I can find no slam-dunk quote from MPI-4.0 stating that regular persistent point-to-point matching order is determined by the starting procedure order. This is implied by the "half-channel" statement and by the "started with MPI_START" statement (quoted above) but there is no equivalent to the ordering statement made for nonblocking point-to-point (see below).

MPI-4.0 p74 lines 40-42:

Nonblocking communication operations are ordered according to the execution order of the calls that initiate the communication.

Side-note: we should fix this omission -- we should add a new subsubsection "3.9.1 Semantics of Persistent Communications" and state explicitly the tribal knowledge of the semantic rules pertaining to these operations.

[Wee-Free-Scot]

• We should strive to avoid this reliance on the interpretation of "logically concurrent" (as far as possible) in the definition of new interfaces.

Agreed! We started to play around with this new idea called "MPI Stream" -- pmodels/mpich#5908. @Wee-Free-Scot It certainly can use some of your early input.

I have responded with some initial thoughts on the linked issue. Thanks for taking the time to write up your idea clearly.

[patrickb314]

MPI-4.0 p74 lines 40-42:

Nonblocking communication operations are ordered according to the execution order of the calls that initiate the communication.

Thanks, Dan, this is the source of my confusion. I had initially thought matching was in Start order, read the partitioned spec which explicitly states that they match in order, wanted to make sure about regular persistent requests, read carefully about the state of MPI requests, and then saw the various quotes you provided.

From a standards terminology perspective, my uncertainty stems from whether:

1. matching of point to point operations happens in the order of point to point request initialization or starting. (p15 terminology)
2. If inactive requests made by persistent calls are initializing in the page 16 sense.
3. If partitioned init requests are actually initializing or initiating

This quote from p.71 seems relevant to point 2:

A persistent communication request and the handle to it are inactive if the request is not associated with any ongoing communication (see Section 3.9).

But this could just be that I’m still not fully familiar with all of the relevant terminology and latest abstractions in the standard yet. If so, please excuse me while I get up to speed with standard. It’s been many years since I waded into it in depth and lots has changed.

[Wee-Free-Scot]

I think some confusion is arising from trying to determine whether MPI requests are (or exhibit) operation stages, which they are (do) not.

The initialisation stage of an MPI operation creates an MPI request that represents that operation. The MPI operation is inactive (because the starting stage has not been done yet). We call the request an inactive request because it is a request that represents an inactive operation.

The starting stage of an MPI operation changes the state of the operation from inactive to active. Any request that represents this operation is now called an active request because it is a request that represents an active operation.

The completion stage of an MPI operation changes the state of the operation from active to inactive. Any request that represents this operation is now called an inactive request because it is a request that represents an inactive operation.

The freeing stage of an MPI operation deallocates/destroys the request that represents the operation.

This can be gleaned (we hope) from the state transition diagrams provided in the Terms chapter (see MPI-4.0 §2.4.1).

There is no API to discover-without-permitting-change the stage of the operation represented by a request.

Thus,

1. Matching for regular point-to-point operations (blocking, nonblocking, or persistent) happens once per starting stage of the operation. For blocking operations, this means a match occurs for every call to the single procedure that expresses that operation, e.g. MPI_Send. For nonblocking operations, this means a match occurs for every call to the single initiation procedure for that operation, e.g. MPI_Isend. For persistent operations, this means a match occurs for every call to the starting procedure for that operation, e.g. MPI_Start.

• Matching for partitioned point-to-point operations (persistent only) happens once per initialisation stage of the operation. This means that a match occurs for every call to the initialisation procedure for that operation, e.g. MPI_Psend_init.
• I am being very careful with the wording here -- a match occurs, permitting us to count how many receive operations are needed without bringing the ordering rules(s) into our discussion. The matching order (if any) is determined by (when it is defined at all) the order of the procedure calls that give rise to the match. This is somewhat orthogonal, although very important for programmability/predictability/performance/etc.

2. Persistent requests (of all types) are created inactive by the initialisation procedure, which performs the initialisation stage of the operation. No requests are initialising (category error).

3. Partitioned communication operations are never initiated, they are only ever initialised and, perhaps, later/separately started (the two component parts of initiated) because they are persistent operations, rather than nonblocking operations. No requests are initialising or initiating (category error).