MPIX_Win_post_stream

Starts an RMA exposure epoch with GPU stream awareness.

Definitions

C/C++ Synopsis

int MPIX_Win_post_stream (   MPI_Group group,
                                 MPI_Win win,
                                 void * stream       )

Arguments

group                group of origin processes (handle) passed as input
win                  window object (handle) passed as input
stream               GPU stream handle passed as input

Description

Enqueues operations to start an RMA exposure epoch for the local window associated with win. Only the set of processes belonging to group can access the window with one-sided RMA operations on win during the exposure epoch.

MPIX_Win_post_stream provides support for GPU stream awareness through the GPU stream handle stream. Operations required to start the RMA exposure epoch are enqueued to the GPU stream specified by the stream handle.

MPIX_Win_post_stream is non-blocking with respect to the application process running on the CPU. Once the required operations to start the exposure epoch are enqueued, the application process returns immediately. In-order execution guarantees are provided for the enqueued operations.

Trying to close an exposure epoch created with GPU stream awareness using MPIX_Win_post_stream with calls to MPI_Win_wait is undefined. An exposure epoch created by a call to MPIX_Win_post_stream must be closed with a call to MPIX_Win_wait_stream.

The standard function MPI_Win_start starts an RMA access epoch for an MPI RMA window, with no change in semantics and syntax from the existing MPI specification. Using the new constant MPIX_MODE_STREAM as the assert argument to this function allows the MPI implementation to provide GPU stream awareness to the access epoch opened.

The standard function MPI_Win_start starts an RMA access epoch for an MPI RMA window, with no change in semantics and syntax from the existing MPI specification. Using the new constant MPIX_MODE_STREAM as the assert argument to this function allows the MPI implementation to provide GPU stream awareness to the access epoch opened.

Return Values

None.

Examples

C/C++ Example

Example code snippet showing the usage of basic stream triggered communication operations and comparing about regular non-GPU stream aware MPI active RMA operation.

Non-GPU Stream-aware Example

for (int i = 0; i < iterations; i++ ) {
    MPI_Win_post(group,0,win);

    increment<<<nb,nt,0,stream>>>(n,source);
    CHECK(hipStreamSynchronize(stream));

    MPI_Win_start(group,0,win);
    MPI_Put(source,n,MPI_INT,left,fromL-base,n,MPI_INT,win);
    MPI_Put(source,n,MPI_INT,right,fromR-base,n,MPI_INT,win);
    MPI_Win_complete(win);

    MPI_Win_wait(win);

    compare<<<nb,nt,0,stream>>>(i,n,fromL);
    compare<<<nb,nt,0,stream>>>(i,n,fromR);
    CHECK(hipStreamSynchronize(stream));
}

GPU Stream-aware Example

for (int i = 0; i < iterations; i++ ) {
    MPIX_Win_post_stream(group,win,stream);
    increment<<<nb,nt,0,stream>>>(n,source);

    MPI_Win_start(group,MPI_MODE_STREAM,win);
    MPI_Put(source,n,MPI_INT,left,fromL-base,n,MPI_INT,win);
    MPI_Put(source,n,MPI_INT,right,fromR-base,n,MPI_INT,win);
    MPIX_Win_complete_stream(win,stream);

    MPIX_Win_wait_stream(win,stream);

    compare<<<nb,nt,0,stream>>>(i,n,fromL);
    compare<<<nb,nt,0,stream>>>(i,n,fromR);
}
CHECK(hipStreamSynchronize(stream));