[pull] develop from lammps:develop

.github/workflows/compile-msvc.yml

@@ -66,4 +66,13 @@ jobs:
 
     - name: Run Unit Tests
       working-directory: build
-      run: ctest -V -E FixTimestep:python_move_nve
+      run: ctest -V -E FixTimestep:python_move_nve --output-junit junit-windows-msvc.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-windows-msvc
+        path: build/junit-windows-msvc.xml

.github/workflows/unittest-arm64.yml

@@ -78,4 +78,13 @@ jobs:
     - name: Run Tests
       working-directory: build
       shell: bash
-      run: ctest -V -LE unstable
+      run: ctest -V -LE unstable --output-junit junit-linux-arm64.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-linux-arm64
+        path: build/junit-linux-arm64.xml

.github/workflows/unittest-kokkos.yml

@@ -77,4 +77,13 @@ jobs:
     - name: Run Tests
       working-directory: build
       shell: bash
-      run: ctest --output-on-failure -R '^(MolPairStyle|AtomicPairStyle|ManybodyPairStyle|EllipsoidPairStyle|KSpaceStyle|BondStyle|AngleStyle|DihedralStyle|ImproperStyle|FixTimestep|Library)'
+      run: ctest --output-on-failure -R '^(MolPairStyle|AtomicPairStyle|ManybodyPairStyle|EllipsoidPairStyle|KSpaceStyle|BondStyle|AngleStyle|DihedralStyle|ImproperStyle|FixTimestep|Library)' --output-junit junit-kokkos-${{ matrix.backend }}-${{ matrix.precision }}.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-kokkos-${{ matrix.backend }}-${{ matrix.precision }}
+        path: build/junit-kokkos-${{ matrix.backend }}-${{ matrix.precision }}.xml

.github/workflows/unittest-linux.yml

@@ -84,4 +84,13 @@ jobs:
     - name: Run Tests
       working-directory: build
       shell: bash
-      run: ctest -V
+      run: ctest -V --output-junit junit-linux-x86_64-bigbig.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-linux-x86_64-bigbig
+        path: build/junit-linux-x86_64-bigbig.xml

.github/workflows/unittest-macos.yml

@@ -67,4 +67,13 @@ jobs:
     - name: Run Tests
       working-directory: build
       shell: bash
-      run: ctest -V
+      run: ctest -V --output-junit junit-macos-x86_64.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-macos-x86_64
+        path: build/junit-macos-x86_64.xml

.github/workflows/unittest-single.yml

@@ -77,4 +77,13 @@ jobs:
     - name: Run Tests
       working-directory: build
       shell: bash
-      run: ctest -V
+      run: ctest -V --output-junit junit-linux-x86_64-fftsingle.xml
+
+    # test results from pull request branches are only for the submitters,
+    # only post-merge results are published for downstream reporting
+    - name: Upload JUnit test results for post-merge runs
+      if: ${{ always() && (github.event_name == 'push') }}
+      uses: actions/upload-artifact@v7
+      with:
+        name: junit-linux-x86_64-fftsingle
+        path: build/junit-linux-x86_64-fftsingle.xml

src/KOKKOS/comm_tiled_kokkos.cpp

@@ -22,9 +22,11 @@
 #include "fix.h"
 #include "force.h"
 #include "kokkos.h"
+#include "kokkos_base.h"
 #include "memory_kokkos.h"
 #include "modify.h"
 #include "output.h"
+#include "pair.h"
 
 using namespace LAMMPS_NS;
 
@@ -38,6 +40,7 @@ CommTiledKokkos::CommTiledKokkos(LAMMPS *_lmp) : CommTiled(_lmp)
   sendlist = nullptr;
   maxsendlist = nullptr;
   nprocmaxtot = 0;
+  max_buf_pair = 0;
 }
 
 /* ---------------------------------------------------------------------- */
@@ -52,6 +55,7 @@ CommTiledKokkos::CommTiledKokkos(LAMMPS *_lmp, Comm *oldcomm) : CommTiled(_lmp,o
   sendlist = nullptr;
   maxsendlist = nullptr;
   nprocmaxtot = 0;
+  max_buf_pair = 0;
 }
 
 /* ---------------------------------------------------------------------- */
@@ -414,8 +418,106 @@ void CommTiledKokkos::borders()
 
 void CommTiledKokkos::forward_comm(Pair *pair, int size)
 {
-  k_sendlist.sync_host();
-  CommTiled::forward_comm(pair, size);
+  if (pair->execution_space == Host || pair->execution_space == HostKK || forward_pair_comm_legacy) {
+    k_sendlist.sync_host();
+    CommTiled::forward_comm(pair, size);
+  } else {
+    k_sendlist.sync_device();
+    forward_comm_device<LMPDeviceType>(pair, size);
+  }
+}
+
+/* ---------------------------------------------------------------------- */
+
+template<class DeviceType>
+void CommTiledKokkos::forward_comm_device(Pair *pair, int size)
+{
+  int i,n,nsize,nsend,nrecv;
+
+  if (size) nsize = size;
+  else nsize = pair->comm_forward;
+
+  KokkosBase* pairKKBase = dynamic_cast<KokkosBase*>(pair);
+
+  // ensure send/recv buffers are large enough for all swaps/procs
+  // send buf holds one forward send to a single proc
+  // recv buf holds all forward recvs in one swap (packed at per-proc offsets)
+
+  int nmax = max_buf_pair;
+  for (int iswap = 0; iswap < nswap; iswap++) {
+    for (i = 0; i < nsendproc[iswap]; i++)
+      nmax = MAX(nmax,nsize*sendnum[iswap][i]);
+    n = nrecvproc[iswap];
+    if (n)
+      nmax = MAX(nmax,nsize*(forward_recv_offset[iswap][n-1] + recvnum[iswap][n-1]));
+  }
+  if (nmax > max_buf_pair)
+    grow_buf_pair(nmax);
+
+  // exchange data with another set of procs in each swap
+  // post recvs from all procs except self
+  // send data to all procs except self
+  // copy data to self if sendself is set
+  // wait on all procs except self and unpack received data
+
+  double* buf_send_pair;
+  double* buf_recv_pair;
+  if (lmp->kokkos->gpu_aware_flag) {
+    buf_send_pair = k_buf_send_pair.view<DeviceType>().data();
+    buf_recv_pair = k_buf_recv_pair.view<DeviceType>().data();
+  } else {
+    buf_send_pair = k_buf_send_pair.view_host().data();
+    buf_recv_pair = k_buf_recv_pair.view_host().data();
+  }
+
+  for (int iswap = 0; iswap < nswap; iswap++) {
+    nsend = nsendproc[iswap] - sendself[iswap];
+    nrecv = nrecvproc[iswap] - sendself[iswap];
+
+    if (recvother[iswap]) {
+      for (i = 0; i < nrecv; i++)
+        MPI_Irecv(&buf_recv_pair[nsize*forward_recv_offset[iswap][i]],
+                  nsize*recvnum[iswap][i],MPI_DOUBLE,recvproc[iswap][i],0,world,&requests[i]);
+    }
+
+    if (sendother[iswap]) {
+      for (i = 0; i < nsend; i++) {
+        auto k_sendlist_small = Kokkos::subview(k_sendlist,iswap,i,Kokkos::ALL);
+        n = pairKKBase->pack_forward_comm_kokkos(sendnum[iswap][i],k_sendlist_small,
+                                    k_buf_send_pair,pbc_flag[iswap][i],pbc[iswap][i]);
+        if (!lmp->kokkos->gpu_aware_flag) {
+          k_buf_send_pair.modify<DeviceType>();
+          k_buf_send_pair.sync_host();
+        }
+        DeviceType().fence();
+        MPI_Send(buf_send_pair,n,MPI_DOUBLE,sendproc[iswap][i],0,world);
+      }
+    }
+
+    if (sendself[iswap]) {
+      auto k_sendlist_small = Kokkos::subview(k_sendlist,iswap,nsend,Kokkos::ALL);
+      pairKKBase->pack_forward_comm_kokkos(sendnum[iswap][nsend],k_sendlist_small,
+                              k_buf_send_pair,pbc_flag[iswap][nsend],pbc[iswap][nsend]);
+      pairKKBase->unpack_forward_comm_kokkos(recvnum[iswap][nrecv],firstrecv[iswap][nrecv],
+                              k_buf_send_pair);
+    }
+
+    if (recvother[iswap]) {
+      MPI_Waitall(nrecv,requests,MPI_STATUS_IGNORE);
+      DeviceType().fence();
+      if (!lmp->kokkos->gpu_aware_flag) {
+        k_buf_recv_pair.modify_host();
+        k_buf_recv_pair.sync<DeviceType>();
+      }
+      for (i = 0; i < nrecv; i++) {
+        DAT::tdual_double_1d k_buf_recv_offset =
+          Kokkos::subview(k_buf_recv_pair,std::pair<int,int>(nsize*forward_recv_offset[iswap][i],
+                          (int)k_buf_recv_pair.extent(0)));
+        pairKKBase->unpack_forward_comm_kokkos(recvnum[iswap][i],firstrecv[iswap][i],
+                                  k_buf_recv_offset);
+      }
+    }
+  }
 }
 
 /* ----------------------------------------------------------------------
@@ -429,8 +531,116 @@ void CommTiledKokkos::forward_comm(Pair *pair, int size)
 
 void CommTiledKokkos::reverse_comm(Pair *pair, int size)
 {
-  k_sendlist.sync_host();
-  CommTiled::reverse_comm(pair, size);
+  if (pair->execution_space == Host || pair->execution_space == HostKK ||
+      !pair->reverse_comm_device || reverse_pair_comm_legacy) {
+    k_sendlist.sync_host();
+    CommTiled::reverse_comm(pair, size);
+  } else {
+    k_sendlist.sync_device();
+    reverse_comm_device<LMPDeviceType>(pair, size);
+  }
+}
+
+/* ---------------------------------------------------------------------- */
+
+template<class DeviceType>
+void CommTiledKokkos::reverse_comm_device(Pair *pair, int size)
+{
+  int i,n,nsize,nsend,nrecv;
+
+  if (size) nsize = MAX(pair->comm_reverse, pair->comm_reverse_off);
+  else nsize = pair->comm_reverse;
+
+  KokkosBase* pairKKBase = dynamic_cast<KokkosBase*>(pair);
+
+  // ensure send/recv buffers are large enough for all swaps/procs
+  // send buf holds one reverse send to a single proc
+  // recv buf holds all reverse recvs in one swap (packed at per-proc offsets)
+
+  int nmax = max_buf_pair;
+  for (int iswap = 0; iswap < nswap; iswap++) {
+    for (i = 0; i < nrecvproc[iswap]; i++)
+      nmax = MAX(nmax,nsize*recvnum[iswap][i]);
+    n = nsendproc[iswap];
+    if (n)
+      nmax = MAX(nmax,nsize*(reverse_recv_offset[iswap][n-1] + sendnum[iswap][n-1]));
+  }
+  if (nmax > max_buf_pair)
+    grow_buf_pair(nmax);
+
+  // exchange data with another set of procs in each swap
+  // post recvs from all procs except self
+  // send data to all procs except self
+  // copy data to self if sendself is set
+  // wait on all procs except self and unpack received data
+
+  double* buf_send_pair;
+  double* buf_recv_pair;
+  if (lmp->kokkos->gpu_aware_flag) {
+    buf_send_pair = k_buf_send_pair.view<DeviceType>().data();
+    buf_recv_pair = k_buf_recv_pair.view<DeviceType>().data();
+  } else {
+    buf_send_pair = k_buf_send_pair.view_host().data();
+    buf_recv_pair = k_buf_recv_pair.view_host().data();
+  }
+
+  for (int iswap = nswap-1; iswap >= 0; iswap--) {
+    nsend = nsendproc[iswap] - sendself[iswap];
+    nrecv = nrecvproc[iswap] - sendself[iswap];
+
+    if (sendother[iswap]) {
+      for (i = 0; i < nsend; i++)
+        MPI_Irecv(&buf_recv_pair[nsize*reverse_recv_offset[iswap][i]],
+                  nsize*sendnum[iswap][i],MPI_DOUBLE,sendproc[iswap][i],0,world,&requests[i]);
+    }
+
+    if (recvother[iswap]) {
+      for (i = 0; i < nrecv; i++) {
+        n = pairKKBase->pack_reverse_comm_kokkos(recvnum[iswap][i],firstrecv[iswap][i],
+                                    k_buf_send_pair);
+        if (!lmp->kokkos->gpu_aware_flag) {
+          k_buf_send_pair.modify<DeviceType>();
+          k_buf_send_pair.sync_host();
+        }
+        DeviceType().fence();
+        MPI_Send(buf_send_pair,n,MPI_DOUBLE,recvproc[iswap][i],0,world);
+      }
+    }
+
+    if (sendself[iswap]) {
+      auto k_sendlist_small = Kokkos::subview(k_sendlist,iswap,nsend,Kokkos::ALL);
+      pairKKBase->pack_reverse_comm_kokkos(recvnum[iswap][nrecv],firstrecv[iswap][nrecv],
+                              k_buf_send_pair);
+      pairKKBase->unpack_reverse_comm_kokkos(sendnum[iswap][nsend],k_sendlist_small,
+                              k_buf_send_pair);
+    }
+
+    if (sendother[iswap]) {
+      MPI_Waitall(nsend,requests,MPI_STATUS_IGNORE);
+      DeviceType().fence();
+      if (!lmp->kokkos->gpu_aware_flag) {
+        k_buf_recv_pair.modify_host();
+        k_buf_recv_pair.sync<DeviceType>();
+      }
+      for (i = 0; i < nsend; i++) {
+        auto k_sendlist_small = Kokkos::subview(k_sendlist,iswap,i,Kokkos::ALL);
+        DAT::tdual_double_1d k_buf_recv_offset =
+          Kokkos::subview(k_buf_recv_pair,std::pair<int,int>(nsize*reverse_recv_offset[iswap][i],
+                          (int)k_buf_recv_pair.extent(0)));
+        pairKKBase->unpack_reverse_comm_kokkos(sendnum[iswap][i],k_sendlist_small,
+                                  k_buf_recv_offset);
+      }
+    }
+  }
+}
+
+/* ---------------------------------------------------------------------- */
+
+void CommTiledKokkos::grow_buf_pair(int n)
+{
+  max_buf_pair = n * BUFFACTOR;
+  k_buf_send_pair.resize(max_buf_pair);
+  k_buf_recv_pair.resize(max_buf_pair);
 }
 
 /* ----------------------------------------------------------------------

src/KOKKOS/comm_tiled_kokkos.h

@@ -62,13 +62,19 @@ class CommTiledKokkos : public CommTiled {
 
   template<class DeviceType> void forward_comm_device();
   template<class DeviceType> void reverse_comm_device();
+  template<class DeviceType> void forward_comm_device(Pair *pair, int size=0);
+  template<class DeviceType> void reverse_comm_device(Pair *pair, int size=0);
 
  protected:
   int nprocmaxtot;
 
   DAT::tdual_int_3d_lr k_sendlist;
   DAT::tdual_double_2d_lr k_buf_send,k_buf_recv;
 
+  int max_buf_pair;
+  DAT::tdual_double_1d k_buf_send_pair, k_buf_recv_pair;
+  void grow_buf_pair(int);
+
   void grow_send(int, int) override;             // reallocate send buffer
   void grow_recv(int, int flag = 0) override;    // free/allocate recv buffer
   void grow_send_kokkos(int, int, ExecutionSpace space = Host);