From 9a85832096611e809b9dd46d00c62fa681d8e5d7 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 24 Jun 2026 17:09:41 +0000 Subject: [PATCH 1/2] KOKKOS: add GPU-aware Pair forward/reverse comm to CommTiledKokkos CommTiledKokkos::forward_comm(Pair*) and reverse_comm(Pair*) previously fell back to the CPU CommTiled implementation, which calls the host pack/unpack_forward_comm() routines. For Kokkos pair styles running on the GPU this dereferences device pointers on the host and segfaults during dynamics (issue #4942). Mirror the approach already used in CommKokkos (comm_style brick): route device-resident pairs through new templated forward_comm_device<>(Pair*) and reverse_comm_device<>(Pair*) methods that pack/unpack on the device via pack_forward_comm_kokkos()/unpack_forward_comm_kokkos() (and the reverse equivalents), while preserving the tiled multi-proc-per-swap communication pattern with per-proc recv offsets. Dedicated device-resident k_buf_send_pair/k_buf_recv_pair buffers are sized through grow_buf_pair(), which also fixes the buffer-undersizing issue for pair communication. Host execution and the *_pair_comm_legacy paths continue to use the CPU base-class implementation. Non-GPU-aware MPI is supported by staging through the host mirror of the dual-view buffers, matching CommKokkos. Co-Authored-By: Stan Moore Co-Authored-By: Claude Opus 4.8 Claude-Session: https://claude.ai/code/session_01KDG2tvUWemKJvvW2HCsTxT --- src/KOKKOS/comm_tiled_kokkos.cpp | 218 ++++++++++++++++++++++++++++++- src/KOKKOS/comm_tiled_kokkos.h | 6 + 2 files changed, 220 insertions(+), 4 deletions(-) diff --git a/src/KOKKOS/comm_tiled_kokkos.cpp b/src/KOKKOS/comm_tiled_kokkos.cpp index 32d151b65cd..584f1391a18 100644 --- a/src/KOKKOS/comm_tiled_kokkos.cpp +++ b/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(pair, size); + } +} + +/* ---------------------------------------------------------------------- */ + +template +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(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().data(); + buf_recv_pair = k_buf_recv_pair.view().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(); + 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(); + } + for (i = 0; i < nrecv; i++) { + DAT::tdual_double_1d k_buf_recv_offset = + Kokkos::subview(k_buf_recv_pair,std::pair(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(pair, size); + } +} + +/* ---------------------------------------------------------------------- */ + +template +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(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().data(); + buf_recv_pair = k_buf_recv_pair.view().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(); + 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(); + } + 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(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); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/comm_tiled_kokkos.h b/src/KOKKOS/comm_tiled_kokkos.h index 2c396c98f82..dd21a6937d5 100644 --- a/src/KOKKOS/comm_tiled_kokkos.h +++ b/src/KOKKOS/comm_tiled_kokkos.h @@ -62,6 +62,8 @@ class CommTiledKokkos : public CommTiled { template void forward_comm_device(); template void reverse_comm_device(); + template void forward_comm_device(Pair *pair, int size=0); + template void reverse_comm_device(Pair *pair, int size=0); protected: int nprocmaxtot; @@ -69,6 +71,10 @@ class CommTiledKokkos : public CommTiled { 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); From fd37efb05de04b1b19ffa9da0618df12461cabdf Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Fri, 3 Jul 2026 11:07:18 -0400 Subject: [PATCH 2/2] Save unit test results as JUnit XML artifacts on post-merge runs Add --output-junit to the ctest invocations of the six unit test workflows (Linux BIGBIG, Linux FFT_SINGLE, macOS, Windows MSVC, Linux ARM64, KOKKOS host backends) and upload the resulting XML file as a workflow artifact. The upload step is gated on the push event (i.e. post-merge runs on the develop or maintenance branches) so that results from pull request test runs are not published; those are for the submitters. The step uses 'if: always()' so that results are uploaded especially when tests fail. The XML files are meant for downstream reporting, e.g. aggregating test results across the whole platform matrix on a test status website. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01RZmFEmpx4SmV2vaUWBF4Yk --- .github/workflows/compile-msvc.yml | 11 ++++++++++- .github/workflows/unittest-arm64.yml | 11 ++++++++++- .github/workflows/unittest-kokkos.yml | 11 ++++++++++- .github/workflows/unittest-linux.yml | 11 ++++++++++- .github/workflows/unittest-macos.yml | 11 ++++++++++- .github/workflows/unittest-single.yml | 11 ++++++++++- 6 files changed, 60 insertions(+), 6 deletions(-) diff --git a/.github/workflows/compile-msvc.yml b/.github/workflows/compile-msvc.yml index e6635984d9e..5ff3d6f985e 100644 --- a/.github/workflows/compile-msvc.yml +++ b/.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 diff --git a/.github/workflows/unittest-arm64.yml b/.github/workflows/unittest-arm64.yml index 27f0e22126c..4743325eae7 100644 --- a/.github/workflows/unittest-arm64.yml +++ b/.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 diff --git a/.github/workflows/unittest-kokkos.yml b/.github/workflows/unittest-kokkos.yml index 0001df261f1..3125c8113d0 100644 --- a/.github/workflows/unittest-kokkos.yml +++ b/.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 diff --git a/.github/workflows/unittest-linux.yml b/.github/workflows/unittest-linux.yml index d5877e2925e..e818bf614cc 100644 --- a/.github/workflows/unittest-linux.yml +++ b/.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 diff --git a/.github/workflows/unittest-macos.yml b/.github/workflows/unittest-macos.yml index 1a11e08b367..0b09cd00963 100644 --- a/.github/workflows/unittest-macos.yml +++ b/.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 diff --git a/.github/workflows/unittest-single.yml b/.github/workflows/unittest-single.yml index 53c5d168f7e..03c4866968d 100644 --- a/.github/workflows/unittest-single.yml +++ b/.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