feat: add support for buffered mode send operations

src/KokkosComm_comm_mode.hpp

@@ -21,22 +21,25 @@ namespace KokkosComm {
 // Scoped enumeration to specify the communication mode of a sending operation.
 // See section 3.4 of the MPI standard for a complete specification.
 enum class CommMode {
-  // Default mode: lets the user override the send operations behavior at
-  // compile-time. E.g., this can be set to mode "Synchronous" for debug
-  // builds by defining KOKKOSCOMM_FORCE_SYNCHRONOUS_MODE.
+  // Default mode: lets the user override the send operations behavior at compile-time. E.g. this can be set to mode
+  // "Synchronous" for debug builds by defining KOKKOSCOMM_FORCE_SYNCHRONOUS_MODE.
   Default,
-  // Standard mode: MPI implementation decides whether outgoing messages will
-  // be buffered. Send operations can be started whether or not a matching
-  // receive has been started. They may complete before a matching receive is
-  // started. Standard mode is non-local: successful completion of the send
-  // operation may depend on the occurrence of a matching receive.
+  // Standard mode: MPI implementation decides whether outgoing messages will be buffered. Send operations can be
+  // started whether or not a matching receive has been started. They may complete before a matching receive is started.
+  // Standard mode is non-local: successful completion of the send operation may depend on the occurrence of a matching
+  // receive.
   Standard,
-  // Ready mode: Send operations may be started only if the matching receive is
-  // already started.
+  // Buffered mode: Send operation can be started whether or not a matching receive has been started. It may complete
+  // before a matching receive is started. However, unlike the standard send, this operation is local, and its
+  // completion does not depend on the occurrence of a matching receive. Thus, if a send is executed and no matching
+  // receive is started, then MPI must buffer the outgoing message, so as to allow the send call to complete. An error
+  // will occur if there is insufficient buffer space. The amount of available buffer space is controlled by the user
+  // (see Section 3.6). Buffer allocation by the user may be required for the buffered mode to be effective.
+  Buffered,
+  // Ready mode: Send operations may be started only if the matching receive is already started.
   Ready,
-  // Synchronous mode: Send operations complete successfully only if a matching
-  // receive is started, and the receive operation has started to receive the
-  // message sent.
+  // Synchronous mode: Send operations complete successfully only if a matching receive is started, and the receive
+  // operation has started to receive the message sent.
   Synchronous,
 };
 

src/impl/KokkosComm_isend.hpp

@@ -44,6 +44,8 @@ KokkosComm::Req isend(const ExecSpace &space, const SendView &sv, int dest, int
                          MPI_Comm mpi_comm, MPI_Request *mpi_req) {
     if constexpr (SendMode == CommMode::Standard) {
       MPI_Isend(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm, mpi_req);
+    } else if constexpr (SendMode == CommMode::Buffered) {
+      MPI_Ibsend(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm, mpi_req);
     } else if constexpr (SendMode == CommMode::Ready) {
       MPI_Irsend(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm, mpi_req);
     } else if constexpr (SendMode == CommMode::Synchronous) {

src/impl/KokkosComm_send.hpp

@@ -36,6 +36,8 @@ void send(const ExecSpace &space, const SendView &sv, int dest, int tag, MPI_Com
                         MPI_Comm mpi_comm) {
     if constexpr (SendMode == CommMode::Standard) {
       MPI_Send(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm);
+    } else if constexpr (SendMode == CommMode::Buffered) {
+      MPI_Bsend(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm);
     } else if constexpr (SendMode == CommMode::Ready) {
       MPI_Rsend(mpi_view, mpi_count, mpi_datatype, mpi_dest, mpi_tag, mpi_comm);
     } else if constexpr (SendMode == CommMode::Synchronous) {