From 0981355d634f7af2f0042d5aba007c0a40f0b9ac Mon Sep 17 00:00:00 2001
From: Carl Pearson <cwpears@sandia.gov>
Date: Mon, 13 May 2024 15:44:18 -0600
Subject: [PATCH 1/2] mpi: contiguous isend wrapper

---
 src/impl/KokkosComm_isend.hpp | 12 ++++++++++++
 1 file changed, 12 insertions(+)

diff --git a/src/impl/KokkosComm_isend.hpp b/src/impl/KokkosComm_isend.hpp
index 1e2a8ba4..97105afa 100644
--- a/src/impl/KokkosComm_isend.hpp
+++ b/src/impl/KokkosComm_isend.hpp
@@ -31,6 +31,18 @@
 
 namespace KokkosComm::Impl {
 
+template <KokkosView SendView>
+void isend(const SendView &sv, int dest, int tag, MPI_Comm comm, MPI_Request &req) {
+  using KCT = typename KokkosComm::Traits<SendView>;
+
+  if (KCT::is_contiguous(sv)) {
+    using SendScalar = typename SendView::non_const_value_type;
+    MPI_Isend(KCT::data_handle(sv), KCT::span(sv), mpi_type_v<SendScalar>, dest, tag, comm, &req);
+  } else {
+    throw std::runtime_error("only contiguous views supported for low-level isend");
+  }
+}
+
 template <CommMode SendMode = CommMode::Default, KokkosExecutionSpace ExecSpace, KokkosView SendView>
 KokkosComm::Req isend(const ExecSpace &space, const SendView &sv, int dest, int tag, MPI_Comm comm) {
   Kokkos::Tools::pushRegion("KokkosComm::Impl::isend");

From 0017caa11c255a217ff47a6685c085595c5daa13 Mon Sep 17 00:00:00 2001
From: Carl Pearson <cwpears@sandia.gov>
Date: Wed, 15 May 2024 12:39:23 -0600
Subject: [PATCH 2/2] mpi: isend profiling region, space handling

---
 src/impl/KokkosComm_isend.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/src/impl/KokkosComm_isend.hpp b/src/impl/KokkosComm_isend.hpp
index 97105afa..3243c021 100644
--- a/src/impl/KokkosComm_isend.hpp
+++ b/src/impl/KokkosComm_isend.hpp
@@ -33,6 +33,7 @@ namespace KokkosComm::Impl {
 
 template <KokkosView SendView>
 void isend(const SendView &sv, int dest, int tag, MPI_Comm comm, MPI_Request &req) {
+  Kokkos::Tools::pushRegion("KokkosComm::Impl::isend");
   using KCT = typename KokkosComm::Traits<SendView>;
 
   if (KCT::is_contiguous(sv)) {
@@ -41,6 +42,7 @@ void isend(const SendView &sv, int dest, int tag, MPI_Comm comm, MPI_Request &re
   } else {
     throw std::runtime_error("only contiguous views supported for low-level isend");
   }
+  Kokkos::Tools::popRegion();
 }
 
 template <CommMode SendMode = CommMode::Default, KokkosExecutionSpace ExecSpace, KokkosView SendView>
@@ -79,6 +81,7 @@ KokkosComm::Req isend(const ExecSpace &space, const SendView &sv, int dest, int
     req.keep_until_wait(args.view);
   } else {
     using SendScalar = typename SendView::value_type;
+    space.fence();  // can't issue isend until work in space is complete
     mpi_isend_fn(KCT::data_handle(sv), KCT::span(sv), mpi_type_v<SendScalar>, dest, tag, comm, &req.mpi_req());
     if (KCT::is_reference_counted()) {
       req.keep_until_wait(sv);