Merge 42a9356 into ddc6500

.github/workflows/ci-short.yml

@@ -154,17 +154,19 @@ jobs:
             -DMACHINE_VARIANT=hip-mpi \
             -DCMAKE_CXX_COMPILER=hipcc
       # Test example with "variables" and output
-      - name: advection
-        run: |
-          cmake --build build -t advection-example
-          cd build
-          ctest -R regression_mpi_test:output_hdf5
+      #- name: advection
+      #  run: |
+      #    cmake --build build -t advection-example
+      #    cd build
+      #    ctest -R regression_mpi_test:output_hdf5
       # Test example with swarms
       - name: particle-leapfrog
         run: |
           cmake --build build -t particle-leapfrog
           cd build
-          ctest -R regression_mpi_test:particle_leapfrog
+          /usr/bin/mpiexec -n 2 --allow-run-as-root /__w/parthenon/parthenon/build/example/particle_leapfrog/particle-leapfrog -i /__w/parthenon/parthenon/tst/regression/test_suites/particle_leapfrog/parthinput.particle_leapfrog parthenon/job/problem_id=gold --kokkos-map-device-id-by=mpi_rank
+          ls
+          ctest -R regression_mpi_test:particle_leapfrog -V
 
       - uses: actions/upload-artifact@v3
         with:

CHANGELOG.md

@@ -11,6 +11,7 @@
 - [[PR 1161]](https://github.com/parthenon-hpc-lab/parthenon/pull/1161) Make flux field Metadata accessible, add Metadata::CellMemAligned flag, small perfomance upgrades
 
 ### Changed (changing behavior/API/variables/...)
+- [[PR 1206]](https://github.com/parthenon-hpc-lab/parthenon/pull/1206) Leapfrog fix
 - [[PR1203]](https://github.com/parthenon-hpc-lab/parthenon/pull/1203) Pin Ubuntu CI image
 - [[PR1177]](https://github.com/parthenon-hpc-lab/parthenon/pull/1177) Make mesh-level boundary conditions usable without the "user" flag
 - [[PR 1187]](https://github.com/parthenon-hpc-lab/parthenon/pull/1187) Make DataCollection::Add safer and generalize MeshBlockData::Initialize

src/interface/swarm_comms.cpp

@@ -269,6 +269,11 @@ void Swarm::LoadBuffers_() {
 
     // Remove particles that were loaded to send to another block from this block
     RemoveMarkedParticles();
+  } else {
+    for (int n = 0; n < pmb->neighbors.size(); n++) {
+      const int bufid = pmb->neighbors[n].bufid;
+      vbswarm->send_size[bufid] = 0;
+    }
   }
 }
 

src/utils/sort.hpp

@@ -59,9 +59,11 @@ KOKKOS_INLINE_FUNCTION int upper_bound(const T &arr, Real val) {
 template <class Key, class KeyComparator>
 void sort(ParArray1D<Key> data, KeyComparator comparator, size_t min_idx,
           size_t max_idx) {
+  Kokkos::fence();
+  printf("%s:%i\n", __FILE__, __LINE__);
   PARTHENON_DEBUG_REQUIRE(min_idx < data.extent(0), "Invalid minimum sort index!");
   PARTHENON_DEBUG_REQUIRE(max_idx < data.extent(0), "Invalid maximum sort index!");
-#ifdef KOKKOS_ENABLE_CUDA
+#if defined(KOKKOS_ENABLE_CUDA)
 #ifdef __clang__
   PARTHENON_FAIL("sort is using thrust and there exists an incompatibility with clang, "
                  "see https://github.com/lanl/parthenon/issues/647 for more details. We "
@@ -74,6 +76,13 @@ void sort(ParArray1D<Key> data, KeyComparator comparator, size_t min_idx,
   thrust::device_ptr<Key> last_d = thrust::device_pointer_cast(data.data()) + max_idx + 1;
   thrust::sort(first_d, last_d, comparator);
 #endif
+#elif defined(KOKKOS_ENABLE_HIP)
+  auto data_h = Kokkos::create_mirror_view_and_copy(HostMemSpace(), data);
+  std::sort(data_h.data() + min_idx, data_h.data() + max_idx + 1, comparator);
+  Kokkos::deep_copy(data, data_h);
+  // TODO(BRR) With Kokkos 4.4, switch to Kokkos::sort
+  // auto sub_data = Kokkos::subview(data, std::make_pair(min_idx, max_idx + 1));
+  // Kokkos::sort(sub_data, comparator);
 #else
   if (std::is_same<DevExecSpace, HostExecSpace>::value) {
     std::sort(data.data() + min_idx, data.data() + max_idx + 1, comparator);
@@ -83,13 +92,15 @@ void sort(ParArray1D<Key> data, KeyComparator comparator, size_t min_idx,
                    "touch by opening an issue on the Parthenon GitHub.");
   }
 #endif // KOKKOS_ENABLE_CUDA
+  Kokkos::fence();
+  printf("%s:%i\n", __FILE__, __LINE__);
 }
 
 template <class Key>
 void sort(ParArray1D<Key> data, size_t min_idx, size_t max_idx) {
   PARTHENON_DEBUG_REQUIRE(min_idx < data.extent(0), "Invalid minimum sort index!");
   PARTHENON_DEBUG_REQUIRE(max_idx < data.extent(0), "Invalid maximum sort index!");
-#ifdef KOKKOS_ENABLE_CUDA
+#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
 #ifdef __clang__
   PARTHENON_FAIL("sort is using thrust and there exists an incompatibility with clang, "
                  "see https://github.com/lanl/parthenon/issues/647 for more details. We "
@@ -102,6 +113,12 @@ void sort(ParArray1D<Key> data, size_t min_idx, size_t max_idx) {
   thrust::device_ptr<Key> last_d = thrust::device_pointer_cast(data.data()) + max_idx + 1;
   thrust::sort(first_d, last_d);
 #endif
+  auto data_h = Kokkos::create_mirror_view_and_copy(HostMemSpace(), data);
+  std::sort(data_h.data() + min_idx, data_h.data() + max_idx + 1);
+  Kokkos::deep_copy(data, data_h);
+  // TODO(BRR) With Kokkos 4.4, switch to Kokkos::sort
+  // auto sub_data = Kokkos::subview(data, std::make_pair(min_idx, max_idx + 1));
+  // Kokkos::sort(sub_data);
 #else
   if (std::is_same<DevExecSpace, HostExecSpace>::value) {
     std::sort(data.data() + min_idx, data.data() + max_idx + 1);