diff --git a/.github/workflows/ci-short.yml b/.github/workflows/ci-short.yml index ecb4052411eea..ec56441a564f4 100644 --- a/.github/workflows/ci-short.yml +++ b/.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: diff --git a/CHANGELOG.md b/CHANGELOG.md index 17f8a009242d1..c29941247e438 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/src/interface/swarm_comms.cpp b/src/interface/swarm_comms.cpp index 7b236b6e4e1f2..054ded34fabb1 100644 --- a/src/interface/swarm_comms.cpp +++ b/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; + } } } diff --git a/src/utils/sort.hpp b/src/utils/sort.hpp index 97e9c77a88e43..3c91b1ec85ca6 100644 --- a/src/utils/sort.hpp +++ b/src/utils/sort.hpp @@ -59,9 +59,11 @@ KOKKOS_INLINE_FUNCTION int upper_bound(const T &arr, Real val) { template void sort(ParArray1D 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 data, KeyComparator comparator, size_t min_idx, thrust::device_ptr 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::value) { std::sort(data.data() + min_idx, data.data() + max_idx + 1, comparator); @@ -83,13 +92,15 @@ void sort(ParArray1D 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 void sort(ParArray1D 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 data, size_t min_idx, size_t max_idx) { thrust::device_ptr 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::value) { std::sort(data.data() + min_idx, data.data() + max_idx + 1);