From 84adf7054a03f4a9fbb1e7b0fa983f79f9156aba Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 11:01:49 +0200 Subject: [PATCH 01/32] Add std cout to exchange boundaries --- src/Field/HaloCells.hpp | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 542e067fc..2ccfd91e3 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -3,6 +3,7 @@ // The guard / ghost cells of BareField. // +#include #include #include @@ -39,11 +40,14 @@ namespace ippl { auto& comm = layout->comm; + std::cout << "Step 1" << std::endl; const neighbor_list& neighbors = layout->getNeighbors(); + std::cout << "Step 2" << std::endl; const range_list &sendRanges = layout->getNeighborsSendRange(), &recvRanges = layout->getNeighborsRecvRange(); - + std::cout << "Step 3" << std::endl; auto ldom = layout->getLocalNDIndex(); + std::cout << "Step 4" << std::endl; for (const auto& axis : ldom) { if ((axis.length() == 1) && (Dim != 1)) { throw std::runtime_error( @@ -57,11 +61,12 @@ namespace ippl { // exchange when we set HALO_TO_INTERNAL_NOGHOST const auto domain = layout->getDomain(); const auto& ldomains = layout->getHostLocalDomains(); - + std::cout << "Step 5" << std::endl; size_t totalRequests = 0; for (const auto& componentNeighbors : neighbors) { totalRequests += componentNeighbors.size(); } + std::cout << "Step 6" << std::endl; int me = Comm->rank(); @@ -71,6 +76,7 @@ namespace ippl { // sending loop constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; size_t requestIndex = 0; + std::cout << "Step 7" << std::endl; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + index; const auto& componentNeighbors = neighbors[index]; @@ -117,6 +123,7 @@ namespace ippl { } // receiving loop + std::cout << "Step 8" << std::endl; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + Layout_t::getMatchingIndex(index); const auto& componentNeighbors = neighbors[index]; @@ -157,12 +164,15 @@ namespace ippl { unpack(range, view, haloData_m); } } + std::cout << "Step 9" << std::endl; if (totalRequests > 0) { MPI_Waitall(totalRequests, requests.data(), MPI_STATUSES_IGNORE); } + std::cout << "Step 10" << std::endl; comm.freeAllBuffers(); + std::cout << "Step 11" << std::endl; } template From 6adf2d91862f2294db114de3e7315c4e162154da Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 11:10:57 +0200 Subject: [PATCH 02/32] Seg fault happens in the loop, add cout inside --- src/Field/HaloCells.hpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 2ccfd91e3..151b0c3ed 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -84,6 +84,7 @@ namespace ippl { int targetRank = componentNeighbors[i]; bound_type range; + std::cout << "Step 7.1" << std::endl; if (order == INTERNAL_TO_HALO) { /*We store only the sending and receiving ranges * of INTERNAL_TO_HALO and use the fact that the @@ -113,12 +114,16 @@ namespace ippl { } size_type nsends; + std::cout << "Step 7.2" << std::endl; pack(range, view, haloData_m, nsends); + std::cout << "Step 7.3" << std::endl; buffer_type buf = comm.template getBuffer(nsends); - + std::cout << "Step 7.4" << std::endl; comm.isend(targetRank, tag, haloData_m, *buf, requests[requestIndex++], nsends); + std::cout << "Step 7.5" << std::endl; buf->resetWritePos(); + std::cout << "Step 7.6" << std::endl; } } From 7e65139c3ee7c1a0c82922d21185f64c549a2d73 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 11:38:41 +0200 Subject: [PATCH 03/32] Also add cout inside pack for the HaloCells --- src/Field/HaloCells.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 151b0c3ed..4cda8e5aa 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -184,8 +184,10 @@ namespace ippl { void HaloCells::pack(const bound_type& range, const view_type& view, databuffer_type& fd, size_type& nsends) { auto subview = makeSubview(view, range); + std::cout << "Step 7.2.1" << std::endl; auto& buffer = fd.buffer; + std::cout << "Step 7.2.2" << std::endl; size_t size = subview.size(); nsends = size; @@ -193,9 +195,11 @@ namespace ippl { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); } + std::cout << "Step 7.2.3" << std::endl; using index_array_type = typename RangePolicy::index_array_type; + std::cout << "Step 7.2.4" << std::endl; ippl::parallel_for( "HaloCells::pack()", getRangePolicy(subview), KOKKOS_LAMBDA(const index_array_type& args) { @@ -211,6 +215,7 @@ namespace ippl { buffer(l) = apply(subview, args); }); + std::cout << "Step 7.2.5" << std::endl; Kokkos::fence(); } From 01e2e645c6d3cab2a7a2141872c4f7ada805f685 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 11:53:58 +0200 Subject: [PATCH 04/32] Remove std::cout and don't use the subview wrapper: use direct view access instead; similar to how FieldBufferOps already does it --- src/Field/HaloCells.hpp | 66 +++++++++++++++++++---------------------- 1 file changed, 30 insertions(+), 36 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 4cda8e5aa..45d155651 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -3,7 +3,6 @@ // The guard / ghost cells of BareField. // -#include #include #include @@ -40,14 +39,10 @@ namespace ippl { auto& comm = layout->comm; - std::cout << "Step 1" << std::endl; const neighbor_list& neighbors = layout->getNeighbors(); - std::cout << "Step 2" << std::endl; const range_list &sendRanges = layout->getNeighborsSendRange(), &recvRanges = layout->getNeighborsRecvRange(); - std::cout << "Step 3" << std::endl; auto ldom = layout->getLocalNDIndex(); - std::cout << "Step 4" << std::endl; for (const auto& axis : ldom) { if ((axis.length() == 1) && (Dim != 1)) { throw std::runtime_error( @@ -61,12 +56,10 @@ namespace ippl { // exchange when we set HALO_TO_INTERNAL_NOGHOST const auto domain = layout->getDomain(); const auto& ldomains = layout->getHostLocalDomains(); - std::cout << "Step 5" << std::endl; size_t totalRequests = 0; for (const auto& componentNeighbors : neighbors) { totalRequests += componentNeighbors.size(); } - std::cout << "Step 6" << std::endl; int me = Comm->rank(); @@ -76,7 +69,6 @@ namespace ippl { // sending loop constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; size_t requestIndex = 0; - std::cout << "Step 7" << std::endl; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + index; const auto& componentNeighbors = neighbors[index]; @@ -84,7 +76,6 @@ namespace ippl { int targetRank = componentNeighbors[i]; bound_type range; - std::cout << "Step 7.1" << std::endl; if (order == INTERNAL_TO_HALO) { /*We store only the sending and receiving ranges * of INTERNAL_TO_HALO and use the fact that the @@ -114,21 +105,15 @@ namespace ippl { } size_type nsends; - std::cout << "Step 7.2" << std::endl; pack(range, view, haloData_m, nsends); - std::cout << "Step 7.3" << std::endl; buffer_type buf = comm.template getBuffer(nsends); - std::cout << "Step 7.4" << std::endl; comm.isend(targetRank, tag, haloData_m, *buf, requests[requestIndex++], nsends); - std::cout << "Step 7.5" << std::endl; buf->resetWritePos(); - std::cout << "Step 7.6" << std::endl; } } // receiving loop - std::cout << "Step 8" << std::endl; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + Layout_t::getMatchingIndex(index); const auto& componentNeighbors = neighbors[index]; @@ -169,53 +154,62 @@ namespace ippl { unpack(range, view, haloData_m); } } - std::cout << "Step 9" << std::endl; if (totalRequests > 0) { MPI_Waitall(totalRequests, requests.data(), MPI_STATUSES_IGNORE); } - std::cout << "Step 10" << std::endl; comm.freeAllBuffers(); - std::cout << "Step 11" << std::endl; } template void HaloCells::pack(const bound_type& range, const view_type& view, databuffer_type& fd, size_type& nsends) { - auto subview = makeSubview(view, range); - std::cout << "Step 7.2.1" << std::endl; - auto& buffer = fd.buffer; - std::cout << "Step 7.2.2" << std::endl; - size_t size = subview.size(); + size_type size = 1; + for (unsigned d = 0; d < Dim; ++d) { + const long extent = range.hi[d] - range.lo[d]; + if (extent <= 0) { + nsends = 0; + return; + } + size *= static_cast(extent); + } + nsends = size; if (buffer.size() < size) { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); } - std::cout << "Step 7.2.3" << std::endl; + + auto packedBuffer = buffer; + auto fullView = view; + + using exec_space = typename view_type::execution_space; + using index_type = typename RangePolicy::index_type; + Kokkos::Array begin, end, extent; + for (unsigned d = 0; d < Dim; ++d) { + begin[d] = static_cast(range.lo[d]); + end[d] = static_cast(range.hi[d]); + extent[d] = end[d] - begin[d]; + } using index_array_type = - typename RangePolicy::index_array_type; - std::cout << "Step 7.2.4" << std::endl; + typename RangePolicy::index_array_type; ippl::parallel_for( - "HaloCells::pack()", getRangePolicy(subview), + "HaloCells::pack()", createRangePolicy(begin, end), KOKKOS_LAMBDA(const index_array_type& args) { - int l = 0; + size_type l = 0; + size_type stride = 1; - for (unsigned d1 = 0; d1 < Dim; d1++) { - int next = args[d1]; - for (unsigned d2 = 0; d2 < d1; d2++) { - next *= subview.extent(d2); - } - l += next; + for (unsigned d = 0; d < Dim; ++d) { + l += static_cast(args[d] - begin[d]) * stride; + stride *= static_cast(extent[d]); } - buffer(l) = apply(subview, args); + packedBuffer(l) = apply(fullView, args); }); - std::cout << "Step 7.2.5" << std::endl; Kokkos::fence(); } From 0c3b0ce423ae84991cc5c696c7a28a0ba4e8081b Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 13:38:55 +0200 Subject: [PATCH 05/32] Lot's of AI generated debugging statements inside scatter --- src/Particle/ParticleAttrib.hpp | 83 +++++++++++++++++++++++++++++++++ 1 file changed, 83 insertions(+) diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index af484d2bb..55501415c 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -15,6 +15,8 @@ // #include "Ippl.h" +#include + #include "Communicate/DataTypes.h" #include "Utility/IpplTimings.h" @@ -154,6 +156,78 @@ namespace ippl { } auto dview = dview_m; auto ppview = pp.getView(); + + const auto policyBegin = iteration_policy.begin(); + const auto policyEnd = iteration_policy.end(); + const size_type dviewExtent = dview.extent(0); + const size_type ppExtent = ppview.extent(0); + Kokkos::Array fieldExtent; + for (unsigned d = 0; d < Dim; ++d) { + fieldExtent[d] = view.extent(d); + } + + std::cout << "[rank " << Comm->rank() << "] ParticleAttrib::scatter begin" + << " policy=[" << policyBegin << ", " << policyEnd << ")" + << " useHashView=" << useHashView + << " hashExtent=" << hash_array.extent(0) + << " attribExtent=" << dviewExtent + << " posExtent=" << ppExtent + << " fieldExtent=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cout << fieldExtent[d] << (d + 1 == Dim ? ")" : ", "); + } + std::cout << " nghost=" << nghost << " lDomFirst=("; + const auto lDomFirst = lDom.first(); + for (unsigned d = 0; d < Dim; ++d) { + std::cout << lDomFirst[d] << (d + 1 == Dim ? ")" : ", "); + } + std::cout << std::endl; + + size_type invalidMappedCount = 0; + Kokkos::parallel_reduce( + "ParticleAttrib::scatter debug mapped index", iteration_policy, + KOKKOS_LAMBDA(const size_t idx, size_type& localInvalid) { + const size_t mapped_idx = useHashView ? hash_array(idx) : idx; + if (mapped_idx >= dviewExtent || mapped_idx >= ppExtent) { + localInvalid++; + } + }, + invalidMappedCount); + Kokkos::fence(); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter mapped-index debug invalidCount=" + << invalidMappedCount << std::endl; + + size_type invalidStencilCount = 0; + Kokkos::parallel_reduce( + "ParticleAttrib::scatter debug CIC stencil bounds", iteration_policy, + KOKKOS_LAMBDA(const size_t idx, size_type& localInvalid) { + const size_t mapped_idx = useHashView ? hash_array(idx) : idx; + if (mapped_idx >= dviewExtent || mapped_idx >= ppExtent) { + return; + } + + vector_type l = (ppview(mapped_idx) - origin) * invdx + 0.5; + Vector index = l; + Vector args = index - lDomFirst + nghost; + + bool inBounds = true; + for (unsigned d = 0; d < Dim; ++d) { + inBounds = inBounds && args[d] > 0 + && args[d] < static_cast(fieldExtent[d]); + } + if (!inBounds) { + localInvalid++; + } + }, + invalidStencilCount); + Kokkos::fence(); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter CIC-stencil debug invalidCount=" + << invalidStencilCount << std::endl; + + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter launching CIC kernel" << std::endl; Kokkos::parallel_for( "ParticleAttrib::scatter", iteration_policy, KOKKOS_LAMBDA(const size_t idx) { // map index to possible hash_map @@ -172,12 +246,21 @@ namespace ippl { detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi, args, val); }); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter launched CIC kernel, entering fence" << std::endl; + Kokkos::fence(); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter completed CIC fence" << std::endl; IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter entering accumulateHalo" << std::endl; IpplTimings::startTimer(accumulateHaloTimer); f.accumulateHalo(); IpplTimings::stopTimer(accumulateHaloTimer); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter completed accumulateHalo" << std::endl; } template From 2c897d15f2c900a09bd0afed8ddcfe251dcdc7e1 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:01:03 +0200 Subject: [PATCH 06/32] More debugging output for and a explicit 3D path for testing --- src/Field/HaloCells.hpp | 89 ++++++++++++++++++++++++++++++++++------- 1 file changed, 74 insertions(+), 15 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 45d155651..bcdfd5f13 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -3,6 +3,7 @@ // The guard / ghost cells of BareField. // +#include #include #include @@ -177,11 +178,14 @@ namespace ippl { size *= static_cast(extent); } - nsends = size; + const auto bufferSizeBefore = buffer.size(); + + nsends = size; if (buffer.size() < size) { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); } + const auto bufferSizeAfter = buffer.size(); auto packedBuffer = buffer; auto fullView = view; @@ -195,22 +199,77 @@ namespace ippl { extent[d] = end[d] - begin[d]; } - using index_array_type = - typename RangePolicy::index_array_type; - ippl::parallel_for( - "HaloCells::pack()", createRangePolicy(begin, end), - KOKKOS_LAMBDA(const index_array_type& args) { - size_type l = 0; - size_type stride = 1; - - for (unsigned d = 0; d < Dim; ++d) { - l += static_cast(args[d] - begin[d]) * stride; - stride *= static_cast(extent[d]); - } + const int rank = Comm->rank(); + std::cout << "[rank " << rank << "] HaloCells::pack begin range.lo=("; + for (unsigned d = 0; d < Dim; ++d) { + if (d != 0) { + std::cout << ", "; + } + std::cout << range.lo[d]; + } + std::cout << ") range.hi=("; + for (unsigned d = 0; d < Dim; ++d) { + if (d != 0) { + std::cout << ", "; + } + std::cout << range.hi[d]; + } + std::cout << ") view.extent=("; + for (unsigned d = 0; d < Dim; ++d) { + if (d != 0) { + std::cout << ", "; + } + std::cout << view.extent(d); + } + std::cout << ") nsends=" << nsends << " bufferBefore=" << bufferSizeBefore + << " bufferAfter=" << bufferSizeAfter << std::endl; + + if constexpr (Dim == 3) { + const index_type first0 = begin[0]; + const index_type first1 = begin[1]; + const index_type first2 = begin[2]; + const index_type last0 = end[0]; + const index_type last1 = end[1]; + const index_type last2 = end[2]; + const size_type extent0 = static_cast(extent[0]); + const size_type extent1 = static_cast(extent[1]); + + using mdrange_type = Kokkos::MDRangePolicy>; + std::cout << "[rank " << rank << "] HaloCells::pack launching raw 3D kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(raw3d)", + mdrange_type({first0, first1, first2}, {last0, last1, last2}), + KOKKOS_LAMBDA(const index_type i, const index_type j, const index_type k) { + const size_type local0 = static_cast(i - first0); + const size_type local1 = static_cast(j - first1); + const size_type local2 = static_cast(k - first2); + const size_type l = local0 + local1 * extent0 + + local2 * extent0 * extent1; + + packedBuffer(l) = fullView(i, j, k); + }); + } else { + using index_array_type = + typename RangePolicy::index_array_type; + ippl::parallel_for( + "HaloCells::pack()", createRangePolicy(begin, end), + KOKKOS_LAMBDA(const index_array_type& args) { + size_type l = 0; + size_type stride = 1; + + for (unsigned d = 0; d < Dim; ++d) { + l += static_cast(args[d] - begin[d]) * stride; + stride *= static_cast(extent[d]); + } - packedBuffer(l) = apply(fullView, args); - }); + packedBuffer(l) = apply(fullView, args); + }); + } + std::cout << "[rank " << rank << "] HaloCells::pack launched kernel, entering fence" + << std::endl; Kokkos::fence(); + std::cout << "[rank " << rank << "] HaloCells::pack completed fence" << std::endl; } template From 6b3e414e7f18dc3514e779a1226f64a008c29d1d Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:15:38 +0200 Subject: [PATCH 07/32] More checks and debugging in the buffer system --- src/Field/HaloCells.hpp | 66 ++++++++++++++++++++++++++++++++++++++++- 1 file changed, 65 insertions(+), 1 deletion(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index bcdfd5f13..398db7fd5 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -184,6 +184,13 @@ namespace ippl { if (buffer.size() < size) { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); + std::cout << "[rank " << Comm->rank() + << "] HaloCells::pack reallocated typed buffer, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << Comm->rank() + << "] HaloCells::pack completed typed-buffer realloc fence" + << std::endl; } const auto bufferSizeAfter = buffer.size(); @@ -222,7 +229,9 @@ namespace ippl { std::cout << view.extent(d); } std::cout << ") nsends=" << nsends << " bufferBefore=" << bufferSizeBefore - << " bufferAfter=" << bufferSizeAfter << std::endl; + << " bufferAfter=" << bufferSizeAfter + << " viewData=" << static_cast(view.data()) + << " bufferData=" << static_cast(buffer.data()) << std::endl; if constexpr (Dim == 3) { const index_type first0 = begin[0]; @@ -235,6 +244,61 @@ namespace ippl { const size_type extent1 = static_cast(extent[1]); using mdrange_type = Kokkos::MDRangePolicy>; + using range_policy_type = Kokkos::RangePolicy; + + std::cout << "[rank " << rank + << "] HaloCells::pack launching debug buffer-touch kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug buffer touch)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) { + packedBuffer(0) = T(); + }); + std::cout << "[rank " << rank + << "] HaloCells::pack launched debug buffer-touch kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed debug buffer-touch fence" << std::endl; + + std::cout << "[rank " << rank + << "] HaloCells::pack launching debug one-cell copy kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug one-cell copy)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) { + packedBuffer(0) = fullView(first0, first1, first2); + }); + std::cout << "[rank " << rank + << "] HaloCells::pack launched debug one-cell copy kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed debug one-cell copy fence" << std::endl; + + std::cout << "[rank " << rank + << "] HaloCells::pack launching debug raw 3D buffer-fill kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug raw3d buffer fill)", + mdrange_type({first0, first1, first2}, {last0, last1, last2}), + KOKKOS_LAMBDA(const index_type i, const index_type j, const index_type k) { + const size_type local0 = static_cast(i - first0); + const size_type local1 = static_cast(j - first1); + const size_type local2 = static_cast(k - first2); + const size_type l = local0 + local1 * extent0 + + local2 * extent0 * extent1; + + packedBuffer(l) = T(); + }); + std::cout << "[rank " << rank + << "] HaloCells::pack launched debug raw 3D buffer-fill kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed debug raw 3D buffer-fill fence" + << std::endl; + std::cout << "[rank " << rank << "] HaloCells::pack launching raw 3D kernel" << std::endl; Kokkos::parallel_for( From 5b4cc249b4fcfbf026e39c0c9f040a3ff3d2f1b2 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:25:17 +0200 Subject: [PATCH 08/32] Also add raw kernel launch just to be sure --- src/Field/HaloCells.hpp | 82 +++++++++++++++++++++++++++++++++++++++-- 1 file changed, 79 insertions(+), 3 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 398db7fd5..d200c3681 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -168,6 +168,14 @@ namespace ippl { databuffer_type& fd, size_type& nsends) { auto& buffer = fd.buffer; + using exec_space = typename view_type::execution_space; + using view_memory_space = typename view_type::memory_space; + using buffer_view_type = typename databuffer_type::view_type; + using buffer_memory_space = typename buffer_view_type::memory_space; + using range_policy_type = Kokkos::RangePolicy; + + const int rank = Comm->rank(); + size_type size = 1; for (unsigned d = 0; d < Dim; ++d) { const long extent = range.hi[d] - range.lo[d]; @@ -178,6 +186,21 @@ namespace ippl { size *= static_cast(extent); } + std::cout << "[rank " << rank + << "] HaloCells::pack launching pre-realloc noop kernel sizeofT=" + << sizeof(T) << " execSpace=" << exec_space::name() + << " viewMemorySpace=" << view_memory_space::name() + << " bufferMemorySpace=" << buffer_memory_space::name() << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug pre-realloc noop)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + std::cout << "[rank " << rank + << "] HaloCells::pack launched pre-realloc noop kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed pre-realloc noop fence" << std::endl; + const auto bufferSizeBefore = buffer.size(); nsends = size; @@ -197,7 +220,6 @@ namespace ippl { auto packedBuffer = buffer; auto fullView = view; - using exec_space = typename view_type::execution_space; using index_type = typename RangePolicy::index_type; Kokkos::Array begin, end, extent; for (unsigned d = 0; d < Dim; ++d) { @@ -206,7 +228,6 @@ namespace ippl { extent[d] = end[d] - begin[d]; } - const int rank = Comm->rank(); std::cout << "[rank " << rank << "] HaloCells::pack begin range.lo=("; for (unsigned d = 0; d < Dim; ++d) { if (d != 0) { @@ -244,7 +265,62 @@ namespace ippl { const size_type extent1 = static_cast(extent[1]); using mdrange_type = Kokkos::MDRangePolicy>; - using range_policy_type = Kokkos::RangePolicy; + + std::cout << "[rank " << rank + << "] HaloCells::pack launching post-realloc noop kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug post-realloc noop)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + std::cout << "[rank " << rank + << "] HaloCells::pack launched post-realloc noop kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed post-realloc noop fence" << std::endl; + + buffer_view_type freshBuffer("HaloCells::pack(debug fresh typed buffer)", 1); + std::cout << "[rank " << rank + << "] HaloCells::pack allocated fresh typed buffer data=" + << static_cast(freshBuffer.data()) << ", entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed fresh typed-buffer alloc fence" + << std::endl; + + std::cout << "[rank " << rank + << "] HaloCells::pack launching debug fresh-buffer touch kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug fresh buffer touch)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) { + freshBuffer(0) = T(); + }); + std::cout << "[rank " << rank + << "] HaloCells::pack launched debug fresh-buffer touch kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed debug fresh-buffer touch fence" + << std::endl; + + T* packedBufferPtr = buffer.data(); + std::cout << "[rank " << rank + << "] HaloCells::pack launching debug raw-pointer buffer-touch kernel" + << std::endl; + Kokkos::parallel_for( + "HaloCells::pack(debug raw pointer buffer touch)", range_policy_type(0, 1), + KOKKOS_LAMBDA(const int) { + packedBufferPtr[0] = T(); + }); + std::cout << "[rank " << rank + << "] HaloCells::pack launched debug raw-pointer buffer-touch kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::pack completed debug raw-pointer buffer-touch fence" + << std::endl; std::cout << "[rank " << rank << "] HaloCells::pack launching debug buffer-touch kernel" From d5eb029d091f678ee2fcdb6a41413b6890f4f3f9 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:34:10 +0200 Subject: [PATCH 09/32] A lot more debugging calls. It looks like the kernel launch itself is already failing????? --- src/Field/BareField.hpp | 16 ++++++++++++++++ src/Field/HaloCells.hpp | 15 +++++++++++++++ src/Particle/ParticleAttrib.hpp | 13 +++++++++++++ 3 files changed, 44 insertions(+) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 6cd9aade9..6759e9a38 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -9,6 +9,7 @@ #include #include +#include #include #include @@ -157,6 +158,21 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { + using exec_space = typename view_type::execution_space; + using policy_type = Kokkos::RangePolicy; + const int rank = Comm->rank(); + std::cout << "[rank " << rank + << "] BareField::accumulateHalo launching entry noop kernel" << std::endl; + Kokkos::parallel_for( + "BareField::accumulateHalo debug entry noop", policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + std::cout << "[rank " << rank + << "] BareField::accumulateHalo launched entry noop kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] BareField::accumulateHalo completed entry noop fence" << std::endl; + halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index d200c3681..ca5264243 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -18,6 +18,21 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { + using exec_space = typename view_type::execution_space; + using policy_type = Kokkos::RangePolicy; + const int rank = Comm->rank(); + std::cout << "[rank " << rank + << "] HaloCells::accumulateHalo launching entry noop kernel" << std::endl; + Kokkos::parallel_for( + "HaloCells::accumulateHalo debug entry noop", policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + std::cout << "[rank " << rank + << "] HaloCells::accumulateHalo launched entry noop kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << rank + << "] HaloCells::accumulateHalo completed entry noop fence" << std::endl; + exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index 55501415c..3cbec02c4 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -253,6 +253,19 @@ namespace ippl { << "] ParticleAttrib::scatter completed CIC fence" << std::endl; IpplTimings::stopTimer(scatterTimer); + using post_scatter_policy_type = Kokkos::RangePolicy; + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter launching post-CIC noop kernel" << std::endl; + Kokkos::parallel_for( + "ParticleAttrib::scatter debug post-CIC noop", post_scatter_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter launched post-CIC noop kernel, entering fence" + << std::endl; + Kokkos::fence(); + std::cout << "[rank " << Comm->rank() + << "] ParticleAttrib::scatter completed post-CIC noop fence" << std::endl; + static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); std::cout << "[rank " << Comm->rank() << "] ParticleAttrib::scatter entering accumulateHalo" << std::endl; From 7ed6a1a0e2d3ddcba54437f924c49c22559edf44 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:47:30 +0200 Subject: [PATCH 10/32] Remove debugging output and only keep the empty kernel launch that fixed the seg-fault before --- src/Field/BareField.hpp | 16 -- src/Field/HaloCells.hpp | 270 +++----------------------------- src/Particle/ParticleAttrib.hpp | 96 +----------- 3 files changed, 25 insertions(+), 357 deletions(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 6759e9a38..6cd9aade9 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -9,7 +9,6 @@ #include #include -#include #include #include @@ -158,21 +157,6 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { - using exec_space = typename view_type::execution_space; - using policy_type = Kokkos::RangePolicy; - const int rank = Comm->rank(); - std::cout << "[rank " << rank - << "] BareField::accumulateHalo launching entry noop kernel" << std::endl; - Kokkos::parallel_for( - "BareField::accumulateHalo debug entry noop", policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - std::cout << "[rank " << rank - << "] BareField::accumulateHalo launched entry noop kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] BareField::accumulateHalo completed entry noop fence" << std::endl; - halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index ca5264243..542e067fc 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -3,7 +3,6 @@ // The guard / ghost cells of BareField. // -#include #include #include @@ -18,21 +17,6 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - using exec_space = typename view_type::execution_space; - using policy_type = Kokkos::RangePolicy; - const int rank = Comm->rank(); - std::cout << "[rank " << rank - << "] HaloCells::accumulateHalo launching entry noop kernel" << std::endl; - Kokkos::parallel_for( - "HaloCells::accumulateHalo debug entry noop", policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - std::cout << "[rank " << rank - << "] HaloCells::accumulateHalo launched entry noop kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::accumulateHalo completed entry noop fence" << std::endl; - exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } @@ -58,6 +42,7 @@ namespace ippl { const neighbor_list& neighbors = layout->getNeighbors(); const range_list &sendRanges = layout->getNeighborsSendRange(), &recvRanges = layout->getNeighborsRecvRange(); + auto ldom = layout->getLocalNDIndex(); for (const auto& axis : ldom) { if ((axis.length() == 1) && (Dim != 1)) { @@ -72,6 +57,7 @@ namespace ippl { // exchange when we set HALO_TO_INTERNAL_NOGHOST const auto domain = layout->getDomain(); const auto& ldomains = layout->getHostLocalDomains(); + size_t totalRequests = 0; for (const auto& componentNeighbors : neighbors) { totalRequests += componentNeighbors.size(); @@ -124,6 +110,7 @@ namespace ippl { pack(range, view, haloData_m, nsends); buffer_type buf = comm.template getBuffer(nsends); + comm.isend(targetRank, tag, haloData_m, *buf, requests[requestIndex++], nsends); buf->resetWritePos(); } @@ -181,250 +168,35 @@ namespace ippl { template void HaloCells::pack(const bound_type& range, const view_type& view, databuffer_type& fd, size_type& nsends) { - auto& buffer = fd.buffer; - - using exec_space = typename view_type::execution_space; - using view_memory_space = typename view_type::memory_space; - using buffer_view_type = typename databuffer_type::view_type; - using buffer_memory_space = typename buffer_view_type::memory_space; - using range_policy_type = Kokkos::RangePolicy; - - const int rank = Comm->rank(); - - size_type size = 1; - for (unsigned d = 0; d < Dim; ++d) { - const long extent = range.hi[d] - range.lo[d]; - if (extent <= 0) { - nsends = 0; - return; - } - size *= static_cast(extent); - } - - std::cout << "[rank " << rank - << "] HaloCells::pack launching pre-realloc noop kernel sizeofT=" - << sizeof(T) << " execSpace=" << exec_space::name() - << " viewMemorySpace=" << view_memory_space::name() - << " bufferMemorySpace=" << buffer_memory_space::name() << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug pre-realloc noop)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - std::cout << "[rank " << rank - << "] HaloCells::pack launched pre-realloc noop kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed pre-realloc noop fence" << std::endl; + auto subview = makeSubview(view, range); - const auto bufferSizeBefore = buffer.size(); + auto& buffer = fd.buffer; - nsends = size; + size_t size = subview.size(); + nsends = size; if (buffer.size() < size) { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); - std::cout << "[rank " << Comm->rank() - << "] HaloCells::pack reallocated typed buffer, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << Comm->rank() - << "] HaloCells::pack completed typed-buffer realloc fence" - << std::endl; } - const auto bufferSizeAfter = buffer.size(); - auto packedBuffer = buffer; - auto fullView = view; - - using index_type = typename RangePolicy::index_type; - Kokkos::Array begin, end, extent; - for (unsigned d = 0; d < Dim; ++d) { - begin[d] = static_cast(range.lo[d]); - end[d] = static_cast(range.hi[d]); - extent[d] = end[d] - begin[d]; - } + using index_array_type = + typename RangePolicy::index_array_type; + ippl::parallel_for( + "HaloCells::pack()", getRangePolicy(subview), + KOKKOS_LAMBDA(const index_array_type& args) { + int l = 0; - std::cout << "[rank " << rank << "] HaloCells::pack begin range.lo=("; - for (unsigned d = 0; d < Dim; ++d) { - if (d != 0) { - std::cout << ", "; - } - std::cout << range.lo[d]; - } - std::cout << ") range.hi=("; - for (unsigned d = 0; d < Dim; ++d) { - if (d != 0) { - std::cout << ", "; - } - std::cout << range.hi[d]; - } - std::cout << ") view.extent=("; - for (unsigned d = 0; d < Dim; ++d) { - if (d != 0) { - std::cout << ", "; - } - std::cout << view.extent(d); - } - std::cout << ") nsends=" << nsends << " bufferBefore=" << bufferSizeBefore - << " bufferAfter=" << bufferSizeAfter - << " viewData=" << static_cast(view.data()) - << " bufferData=" << static_cast(buffer.data()) << std::endl; - - if constexpr (Dim == 3) { - const index_type first0 = begin[0]; - const index_type first1 = begin[1]; - const index_type first2 = begin[2]; - const index_type last0 = end[0]; - const index_type last1 = end[1]; - const index_type last2 = end[2]; - const size_type extent0 = static_cast(extent[0]); - const size_type extent1 = static_cast(extent[1]); - - using mdrange_type = Kokkos::MDRangePolicy>; - - std::cout << "[rank " << rank - << "] HaloCells::pack launching post-realloc noop kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug post-realloc noop)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - std::cout << "[rank " << rank - << "] HaloCells::pack launched post-realloc noop kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed post-realloc noop fence" << std::endl; - - buffer_view_type freshBuffer("HaloCells::pack(debug fresh typed buffer)", 1); - std::cout << "[rank " << rank - << "] HaloCells::pack allocated fresh typed buffer data=" - << static_cast(freshBuffer.data()) << ", entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed fresh typed-buffer alloc fence" - << std::endl; - - std::cout << "[rank " << rank - << "] HaloCells::pack launching debug fresh-buffer touch kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug fresh buffer touch)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) { - freshBuffer(0) = T(); - }); - std::cout << "[rank " << rank - << "] HaloCells::pack launched debug fresh-buffer touch kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed debug fresh-buffer touch fence" - << std::endl; - - T* packedBufferPtr = buffer.data(); - std::cout << "[rank " << rank - << "] HaloCells::pack launching debug raw-pointer buffer-touch kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug raw pointer buffer touch)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) { - packedBufferPtr[0] = T(); - }); - std::cout << "[rank " << rank - << "] HaloCells::pack launched debug raw-pointer buffer-touch kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed debug raw-pointer buffer-touch fence" - << std::endl; - - std::cout << "[rank " << rank - << "] HaloCells::pack launching debug buffer-touch kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug buffer touch)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) { - packedBuffer(0) = T(); - }); - std::cout << "[rank " << rank - << "] HaloCells::pack launched debug buffer-touch kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed debug buffer-touch fence" << std::endl; - - std::cout << "[rank " << rank - << "] HaloCells::pack launching debug one-cell copy kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug one-cell copy)", range_policy_type(0, 1), - KOKKOS_LAMBDA(const int) { - packedBuffer(0) = fullView(first0, first1, first2); - }); - std::cout << "[rank " << rank - << "] HaloCells::pack launched debug one-cell copy kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed debug one-cell copy fence" << std::endl; - - std::cout << "[rank " << rank - << "] HaloCells::pack launching debug raw 3D buffer-fill kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(debug raw3d buffer fill)", - mdrange_type({first0, first1, first2}, {last0, last1, last2}), - KOKKOS_LAMBDA(const index_type i, const index_type j, const index_type k) { - const size_type local0 = static_cast(i - first0); - const size_type local1 = static_cast(j - first1); - const size_type local2 = static_cast(k - first2); - const size_type l = local0 + local1 * extent0 - + local2 * extent0 * extent1; - - packedBuffer(l) = T(); - }); - std::cout << "[rank " << rank - << "] HaloCells::pack launched debug raw 3D buffer-fill kernel, entering fence" - << std::endl; - Kokkos::fence(); - std::cout << "[rank " << rank - << "] HaloCells::pack completed debug raw 3D buffer-fill fence" - << std::endl; - - std::cout << "[rank " << rank << "] HaloCells::pack launching raw 3D kernel" - << std::endl; - Kokkos::parallel_for( - "HaloCells::pack(raw3d)", - mdrange_type({first0, first1, first2}, {last0, last1, last2}), - KOKKOS_LAMBDA(const index_type i, const index_type j, const index_type k) { - const size_type local0 = static_cast(i - first0); - const size_type local1 = static_cast(j - first1); - const size_type local2 = static_cast(k - first2); - const size_type l = local0 + local1 * extent0 - + local2 * extent0 * extent1; - - packedBuffer(l) = fullView(i, j, k); - }); - } else { - using index_array_type = - typename RangePolicy::index_array_type; - ippl::parallel_for( - "HaloCells::pack()", createRangePolicy(begin, end), - KOKKOS_LAMBDA(const index_array_type& args) { - size_type l = 0; - size_type stride = 1; - - for (unsigned d = 0; d < Dim; ++d) { - l += static_cast(args[d] - begin[d]) * stride; - stride *= static_cast(extent[d]); + for (unsigned d1 = 0; d1 < Dim; d1++) { + int next = args[d1]; + for (unsigned d2 = 0; d2 < d1; d2++) { + next *= subview.extent(d2); } + l += next; + } - packedBuffer(l) = apply(fullView, args); - }); - } - std::cout << "[rank " << rank << "] HaloCells::pack launched kernel, entering fence" - << std::endl; + buffer(l) = apply(subview, args); + }); Kokkos::fence(); - std::cout << "[rank " << rank << "] HaloCells::pack completed fence" << std::endl; } template diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index 3cbec02c4..c9cc72926 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -15,8 +15,6 @@ // #include "Ippl.h" -#include - #include "Communicate/DataTypes.h" #include "Utility/IpplTimings.h" @@ -156,78 +154,6 @@ namespace ippl { } auto dview = dview_m; auto ppview = pp.getView(); - - const auto policyBegin = iteration_policy.begin(); - const auto policyEnd = iteration_policy.end(); - const size_type dviewExtent = dview.extent(0); - const size_type ppExtent = ppview.extent(0); - Kokkos::Array fieldExtent; - for (unsigned d = 0; d < Dim; ++d) { - fieldExtent[d] = view.extent(d); - } - - std::cout << "[rank " << Comm->rank() << "] ParticleAttrib::scatter begin" - << " policy=[" << policyBegin << ", " << policyEnd << ")" - << " useHashView=" << useHashView - << " hashExtent=" << hash_array.extent(0) - << " attribExtent=" << dviewExtent - << " posExtent=" << ppExtent - << " fieldExtent=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cout << fieldExtent[d] << (d + 1 == Dim ? ")" : ", "); - } - std::cout << " nghost=" << nghost << " lDomFirst=("; - const auto lDomFirst = lDom.first(); - for (unsigned d = 0; d < Dim; ++d) { - std::cout << lDomFirst[d] << (d + 1 == Dim ? ")" : ", "); - } - std::cout << std::endl; - - size_type invalidMappedCount = 0; - Kokkos::parallel_reduce( - "ParticleAttrib::scatter debug mapped index", iteration_policy, - KOKKOS_LAMBDA(const size_t idx, size_type& localInvalid) { - const size_t mapped_idx = useHashView ? hash_array(idx) : idx; - if (mapped_idx >= dviewExtent || mapped_idx >= ppExtent) { - localInvalid++; - } - }, - invalidMappedCount); - Kokkos::fence(); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter mapped-index debug invalidCount=" - << invalidMappedCount << std::endl; - - size_type invalidStencilCount = 0; - Kokkos::parallel_reduce( - "ParticleAttrib::scatter debug CIC stencil bounds", iteration_policy, - KOKKOS_LAMBDA(const size_t idx, size_type& localInvalid) { - const size_t mapped_idx = useHashView ? hash_array(idx) : idx; - if (mapped_idx >= dviewExtent || mapped_idx >= ppExtent) { - return; - } - - vector_type l = (ppview(mapped_idx) - origin) * invdx + 0.5; - Vector index = l; - Vector args = index - lDomFirst + nghost; - - bool inBounds = true; - for (unsigned d = 0; d < Dim; ++d) { - inBounds = inBounds && args[d] > 0 - && args[d] < static_cast(fieldExtent[d]); - } - if (!inBounds) { - localInvalid++; - } - }, - invalidStencilCount); - Kokkos::fence(); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter CIC-stencil debug invalidCount=" - << invalidStencilCount << std::endl; - - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter launching CIC kernel" << std::endl; Kokkos::parallel_for( "ParticleAttrib::scatter", iteration_policy, KOKKOS_LAMBDA(const size_t idx) { // map index to possible hash_map @@ -246,34 +172,20 @@ namespace ippl { detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi, args, val); }); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter launched CIC kernel, entering fence" << std::endl; - Kokkos::fence(); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter completed CIC fence" << std::endl; - IpplTimings::stopTimer(scatterTimer); using post_scatter_policy_type = Kokkos::RangePolicy; - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter launching post-CIC noop kernel" << std::endl; + // A separate launch before halo accumulation avoids a GH200 CUDA release-mode crash + // seen when the next launch happens inside HaloCells::pack(). Kokkos::parallel_for( - "ParticleAttrib::scatter debug post-CIC noop", post_scatter_policy_type(0, 1), + "ParticleAttrib::scatter post-CIC launch guard", post_scatter_policy_type(0, 1), KOKKOS_LAMBDA(const int) {}); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter launched post-CIC noop kernel, entering fence" - << std::endl; Kokkos::fence(); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter completed post-CIC noop fence" << std::endl; + IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter entering accumulateHalo" << std::endl; IpplTimings::startTimer(accumulateHaloTimer); f.accumulateHalo(); IpplTimings::stopTimer(accumulateHaloTimer); - std::cout << "[rank " << Comm->rank() - << "] ParticleAttrib::scatter completed accumulateHalo" << std::endl; } template From f6174c85d2857fabd015d18bb81940b38deac2bc Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Tue, 7 Jul 2026 14:58:22 +0200 Subject: [PATCH 11/32] Add other stuff back in --- src/Field/BareField.hpp | 6 ++++++ src/Field/HaloCells.hpp | 6 ++++++ 2 files changed, 12 insertions(+) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 6cd9aade9..b9c7df142 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -157,6 +157,12 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { + using guard_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for( + "BareField::accumulateHalo launch guard", guard_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + Kokkos::fence(); + halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 542e067fc..4ae2e311a 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,6 +17,12 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { + using guard_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for( + "HaloCells::accumulateHalo launch guard", guard_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + Kokkos::fence(); + exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } From e12129a94dcc64ccce3978ce0f3550b75a5cec41 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Wed, 8 Jul 2026 09:32:26 +0200 Subject: [PATCH 12/32] Add env var's to isolate different kernel launches to investigate - saves the recompilationg --- src/Field/BareField.hpp | 9 +++--- src/Field/HaloCells.hpp | 9 +++--- src/Particle/ParticleAttrib.hpp | 9 +++--- src/Utility/LaunchGuard.h | 53 +++++++++++++++++++++++++++++++++ 4 files changed, 65 insertions(+), 15 deletions(-) create mode 100644 src/Utility/LaunchGuard.h diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index b9c7df142..55ccb38df 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -16,6 +16,7 @@ #include "Utility/Inform.h" #include "Utility/IpplInfo.h" +#include "Utility/LaunchGuard.h" #include "BareField.h" namespace Kokkos { @@ -157,11 +158,9 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { - using guard_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for( - "BareField::accumulateHalo launch guard", guard_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - Kokkos::fence(); + detail::launchGuard( + "IPPL_GH200_GUARD_BAREFIELD_ACCUMULATE", + "BareField::accumulateHalo launch guard"); halo_m.accumulateHalo(dview_m, layout_m); } diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 4ae2e311a..4e30aea0d 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -7,6 +7,7 @@ #include #include "Utility/IpplException.h" +#include "Utility/LaunchGuard.h" #include "Communicate/Communicator.h" @@ -17,11 +18,9 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - using guard_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for( - "HaloCells::accumulateHalo launch guard", guard_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - Kokkos::fence(); + detail::launchGuard( + "IPPL_GH200_GUARD_HALOCELLS_ACCUMULATE", + "HaloCells::accumulateHalo launch guard"); exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index c9cc72926..d8937e599 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -18,6 +18,7 @@ #include "Communicate/DataTypes.h" #include "Utility/IpplTimings.h" +#include "Utility/LaunchGuard.h" namespace ippl { @@ -173,13 +174,11 @@ namespace ippl { args, val); }); - using post_scatter_policy_type = Kokkos::RangePolicy; // A separate launch before halo accumulation avoids a GH200 CUDA release-mode crash // seen when the next launch happens inside HaloCells::pack(). - Kokkos::parallel_for( - "ParticleAttrib::scatter post-CIC launch guard", post_scatter_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - Kokkos::fence(); + detail::launchGuard( + "IPPL_GH200_GUARD_SCATTER_POST_CIC", + "ParticleAttrib::scatter post-CIC launch guard"); IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); diff --git a/src/Utility/LaunchGuard.h b/src/Utility/LaunchGuard.h new file mode 100644 index 000000000..30bb30e0e --- /dev/null +++ b/src/Utility/LaunchGuard.h @@ -0,0 +1,53 @@ +// +// Small host-side launch guard helper for isolating backend launch-order issues. +// +#ifndef IPPL_LAUNCH_GUARD_H +#define IPPL_LAUNCH_GUARD_H + +#include +#include + +#include + +namespace ippl::detail { + enum class LaunchGuardMode { Disabled, Fence, Launch, LaunchAndFence }; + + inline LaunchGuardMode launchGuardMode(const char* envName) { + const char* value = std::getenv(envName); + if (value == nullptr) { + return LaunchGuardMode::LaunchAndFence; + } + + if (value[0] == '0' || value[0] == 'n' || value[0] == 'N' + || std::strcmp(value, "off") == 0 || std::strcmp(value, "OFF") == 0 + || std::strcmp(value, "false") == 0 || std::strcmp(value, "FALSE") == 0) { + return LaunchGuardMode::Disabled; + } + if (std::strcmp(value, "fence") == 0 || std::strcmp(value, "FENCE") == 0) { + return LaunchGuardMode::Fence; + } + if (std::strcmp(value, "launch") == 0 || std::strcmp(value, "LAUNCH") == 0) { + return LaunchGuardMode::Launch; + } + + return LaunchGuardMode::LaunchAndFence; + } + + template + inline void launchGuard(const char* envName, const char* label) { + const LaunchGuardMode mode = launchGuardMode(envName); + if (mode == LaunchGuardMode::Disabled) { + return; + } + + if (mode == LaunchGuardMode::Launch || mode == LaunchGuardMode::LaunchAndFence) { + using guard_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for(label, guard_policy_type(0, 1), KOKKOS_LAMBDA(const int) {}); + } + if (mode == LaunchGuardMode::Fence || mode == LaunchGuardMode::LaunchAndFence) { + Kokkos::fence(); + } + } +} // namespace ippl::detail + +#endif From 04ff0fdc1c8d5624f59431e7a844cf7a0e04450d Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Wed, 8 Jul 2026 10:19:41 +0200 Subject: [PATCH 13/32] Somehow the empty launches have to be inlined? Also add some debugging output --- src/Field/BareField.hpp | 13 +++++- src/Field/HaloCells.hpp | 13 +++++- src/Particle/ParticleAttrib.hpp | 13 +++++- src/Utility/LaunchGuard.h | 77 +++++++++++++++++++++------------ 4 files changed, 86 insertions(+), 30 deletions(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 55ccb38df..5aeabc571 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -158,9 +158,20 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { - detail::launchGuard( + const auto accumulateGuardMode = detail::launchGuardMode( "IPPL_GH200_GUARD_BAREFIELD_ACCUMULATE", "BareField::accumulateHalo launch guard"); + if (accumulateGuardMode == detail::LaunchGuardMode::Launch + || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + using guard_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for( + "BareField::accumulateHalo launch guard", guard_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + } + if (accumulateGuardMode == detail::LaunchGuardMode::Fence + || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + Kokkos::fence(); + } halo_m.accumulateHalo(dview_m, layout_m); } diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 4e30aea0d..f05b94c04 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -18,9 +18,20 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - detail::launchGuard( + const auto accumulateGuardMode = detail::launchGuardMode( "IPPL_GH200_GUARD_HALOCELLS_ACCUMULATE", "HaloCells::accumulateHalo launch guard"); + if (accumulateGuardMode == detail::LaunchGuardMode::Launch + || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + using guard_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for( + "HaloCells::accumulateHalo launch guard", guard_policy_type(0, 1), + KOKKOS_LAMBDA(const int) {}); + } + if (accumulateGuardMode == detail::LaunchGuardMode::Fence + || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + Kokkos::fence(); + } exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index d8937e599..5a383cf06 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -176,9 +176,20 @@ namespace ippl { // A separate launch before halo accumulation avoids a GH200 CUDA release-mode crash // seen when the next launch happens inside HaloCells::pack(). - detail::launchGuard( + const auto scatterGuardMode = detail::launchGuardMode( "IPPL_GH200_GUARD_SCATTER_POST_CIC", "ParticleAttrib::scatter post-CIC launch guard"); + if (scatterGuardMode == detail::LaunchGuardMode::Launch + || scatterGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + using post_scatter_policy_type = Kokkos::RangePolicy; + Kokkos::parallel_for( + "ParticleAttrib::scatter post-CIC launch guard", + post_scatter_policy_type(0, 1), KOKKOS_LAMBDA(const int) {}); + } + if (scatterGuardMode == detail::LaunchGuardMode::Fence + || scatterGuardMode == detail::LaunchGuardMode::LaunchAndFence) { + Kokkos::fence(); + } IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); diff --git a/src/Utility/LaunchGuard.h b/src/Utility/LaunchGuard.h index 30bb30e0e..490e90763 100644 --- a/src/Utility/LaunchGuard.h +++ b/src/Utility/LaunchGuard.h @@ -1,52 +1,75 @@ // -// Small host-side launch guard helper for isolating backend launch-order issues. +// Host-side mode parser for isolating backend launch-order issues. // #ifndef IPPL_LAUNCH_GUARD_H #define IPPL_LAUNCH_GUARD_H #include #include - -#include +#include namespace ippl::detail { enum class LaunchGuardMode { Disabled, Fence, Launch, LaunchAndFence }; + inline const char* launchGuardModeName(LaunchGuardMode mode) { + switch (mode) { + case LaunchGuardMode::Disabled: + return "disabled"; + case LaunchGuardMode::Fence: + return "fence"; + case LaunchGuardMode::Launch: + return "launch"; + case LaunchGuardMode::LaunchAndFence: + return "launch+fence"; + } + return "unknown"; + } + inline LaunchGuardMode launchGuardMode(const char* envName) { const char* value = std::getenv(envName); - if (value == nullptr) { - return LaunchGuardMode::LaunchAndFence; - } + LaunchGuardMode mode = LaunchGuardMode::LaunchAndFence; - if (value[0] == '0' || value[0] == 'n' || value[0] == 'N' - || std::strcmp(value, "off") == 0 || std::strcmp(value, "OFF") == 0 - || std::strcmp(value, "false") == 0 || std::strcmp(value, "FALSE") == 0) { - return LaunchGuardMode::Disabled; - } - if (std::strcmp(value, "fence") == 0 || std::strcmp(value, "FENCE") == 0) { - return LaunchGuardMode::Fence; - } - if (std::strcmp(value, "launch") == 0 || std::strcmp(value, "LAUNCH") == 0) { - return LaunchGuardMode::Launch; + if (value != nullptr) { + if (value[0] == '0' || value[0] == 'n' || value[0] == 'N' + || std::strcmp(value, "off") == 0 || std::strcmp(value, "OFF") == 0 + || std::strcmp(value, "false") == 0 || std::strcmp(value, "FALSE") == 0) { + mode = LaunchGuardMode::Disabled; + } else if (std::strcmp(value, "fence") == 0 || std::strcmp(value, "FENCE") == 0) { + mode = LaunchGuardMode::Fence; + } else if (std::strcmp(value, "launch") == 0 || std::strcmp(value, "LAUNCH") == 0) { + mode = LaunchGuardMode::Launch; + } } - return LaunchGuardMode::LaunchAndFence; + return mode; } - template - inline void launchGuard(const char* envName, const char* label) { + inline LaunchGuardMode launchGuardMode(const char* envName, const char* label) { const LaunchGuardMode mode = launchGuardMode(envName); - if (mode == LaunchGuardMode::Disabled) { - return; - } - if (mode == LaunchGuardMode::Launch || mode == LaunchGuardMode::LaunchAndFence) { - using guard_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for(label, guard_policy_type(0, 1), KOKKOS_LAMBDA(const int) {}); + static bool scatterPrinted = false; + static bool bareFieldPrinted = false; + static bool haloCellsPrinted = false; + static bool fallbackPrinted = false; + bool* printed = &fallbackPrinted; + + if (std::strcmp(envName, "IPPL_GH200_GUARD_SCATTER_POST_CIC") == 0) { + printed = &scatterPrinted; + } else if (std::strcmp(envName, "IPPL_GH200_GUARD_BAREFIELD_ACCUMULATE") == 0) { + printed = &bareFieldPrinted; + } else if (std::strcmp(envName, "IPPL_GH200_GUARD_HALOCELLS_ACCUMULATE") == 0) { + printed = &haloCellsPrinted; } - if (mode == LaunchGuardMode::Fence || mode == LaunchGuardMode::LaunchAndFence) { - Kokkos::fence(); + + if (!*printed) { + const char* value = std::getenv(envName); + std::cout << "[IPPL launch guard] " << label << ": " << envName << "=" + << (value != nullptr ? value : "") + << " -> " << launchGuardModeName(mode) << std::endl; + *printed = true; } + + return mode; } } // namespace ippl::detail From 8febce771d39ed96553187e7b3faeb5c5b350efb Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Wed, 8 Jul 2026 10:47:16 +0200 Subject: [PATCH 14/32] Remove again the debug kernel launches and only add make indices explicitly long instead of int/long conversion --- src/Field/BareField.hpp | 16 ------- src/Field/HaloCells.hpp | 16 ------- src/Particle/ParticleAttrib.hpp | 63 ++++++++++++++++++--------- src/Utility/LaunchGuard.h | 76 --------------------------------- 4 files changed, 43 insertions(+), 128 deletions(-) delete mode 100644 src/Utility/LaunchGuard.h diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 5aeabc571..6cd9aade9 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -16,7 +16,6 @@ #include "Utility/Inform.h" #include "Utility/IpplInfo.h" -#include "Utility/LaunchGuard.h" #include "BareField.h" namespace Kokkos { @@ -158,21 +157,6 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { - const auto accumulateGuardMode = detail::launchGuardMode( - "IPPL_GH200_GUARD_BAREFIELD_ACCUMULATE", - "BareField::accumulateHalo launch guard"); - if (accumulateGuardMode == detail::LaunchGuardMode::Launch - || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - using guard_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for( - "BareField::accumulateHalo launch guard", guard_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - } - if (accumulateGuardMode == detail::LaunchGuardMode::Fence - || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - Kokkos::fence(); - } - halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index f05b94c04..542e067fc 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -7,7 +7,6 @@ #include #include "Utility/IpplException.h" -#include "Utility/LaunchGuard.h" #include "Communicate/Communicator.h" @@ -18,21 +17,6 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - const auto accumulateGuardMode = detail::launchGuardMode( - "IPPL_GH200_GUARD_HALOCELLS_ACCUMULATE", - "HaloCells::accumulateHalo launch guard"); - if (accumulateGuardMode == detail::LaunchGuardMode::Launch - || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - using guard_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for( - "HaloCells::accumulateHalo launch guard", guard_policy_type(0, 1), - KOKKOS_LAMBDA(const int) {}); - } - if (accumulateGuardMode == detail::LaunchGuardMode::Fence - || accumulateGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - Kokkos::fence(); - } - exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index 5a383cf06..b7e9d22a6 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -15,10 +15,11 @@ // #include "Ippl.h" +#include + #include "Communicate/DataTypes.h" #include "Utility/IpplTimings.h" -#include "Utility/LaunchGuard.h" namespace ippl { @@ -155,10 +156,25 @@ namespace ippl { } auto dview = dview_m; auto ppview = pp.getView(); + const int rank = Comm->rank(); Kokkos::parallel_for( "ParticleAttrib::scatter", iteration_policy, KOKKOS_LAMBDA(const size_t idx) { // map index to possible hash_map - size_t mapped_idx = useHashView ? hash_array(idx) : idx; + const long long mapped_idx_signed = + useHashView ? static_cast(hash_array(idx)) + : static_cast(idx); + if (mapped_idx_signed < 0 + || mapped_idx_signed >= static_cast(dview.extent(0)) + || mapped_idx_signed >= static_cast(ppview.extent(0))) { + printf("[rank %d] ParticleAttrib::scatter invalid mapped index: idx=%llu " + "mapped=%lld attribExtent=%lld posExtent=%lld useHashView=%d\n", + rank, static_cast(idx), mapped_idx_signed, + static_cast(dview.extent(0)), + static_cast(ppview.extent(0)), static_cast(useHashView)); + Kokkos::abort("ParticleAttrib::scatter invalid mapped index"); + } + + size_t mapped_idx = static_cast(mapped_idx_signed); // find nearest grid point vector_type l = (ppview(mapped_idx) - origin) * invdx + 0.5; @@ -166,30 +182,37 @@ namespace ippl { Vector whi = l - index; Vector wlo = 1.0 - whi; - Vector args = index - lDom.first() + nghost; + Vector args = index - lDom.first() + nghost; + + for (unsigned point = 0; point < (1u << Field::dim); ++point) { + for (unsigned d = 0; d < Field::dim; ++d) { + const int target = args[d] - ((point & (1u << d)) ? 1 : 0); + const int extent = static_cast(view.extent(d)); + if (target < 0 || target >= extent) { + printf("[rank %d] ParticleAttrib::scatter invalid CIC target: idx=%llu " + "mapped=%llu point=%u dim=%u target=%d extent=%d base=(", + rank, static_cast(idx), + static_cast(mapped_idx), point, d, target, + extent); + for (unsigned p = 0; p < Field::dim; ++p) { + printf("%d%s", args[p], (p + 1 == Field::dim) ? "" : ", "); + } + printf(") viewExtent=("); + for (unsigned p = 0; p < Field::dim; ++p) { + printf("%lld%s", static_cast(view.extent(p)), + (p + 1 == Field::dim) ? "" : ", "); + } + printf(")\n"); + Kokkos::abort("ParticleAttrib::scatter invalid CIC target"); + } + } + } // scatter const value_type& val = dview(mapped_idx); detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi, args, val); }); - - // A separate launch before halo accumulation avoids a GH200 CUDA release-mode crash - // seen when the next launch happens inside HaloCells::pack(). - const auto scatterGuardMode = detail::launchGuardMode( - "IPPL_GH200_GUARD_SCATTER_POST_CIC", - "ParticleAttrib::scatter post-CIC launch guard"); - if (scatterGuardMode == detail::LaunchGuardMode::Launch - || scatterGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - using post_scatter_policy_type = Kokkos::RangePolicy; - Kokkos::parallel_for( - "ParticleAttrib::scatter post-CIC launch guard", - post_scatter_policy_type(0, 1), KOKKOS_LAMBDA(const int) {}); - } - if (scatterGuardMode == detail::LaunchGuardMode::Fence - || scatterGuardMode == detail::LaunchGuardMode::LaunchAndFence) { - Kokkos::fence(); - } IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); diff --git a/src/Utility/LaunchGuard.h b/src/Utility/LaunchGuard.h deleted file mode 100644 index 490e90763..000000000 --- a/src/Utility/LaunchGuard.h +++ /dev/null @@ -1,76 +0,0 @@ -// -// Host-side mode parser for isolating backend launch-order issues. -// -#ifndef IPPL_LAUNCH_GUARD_H -#define IPPL_LAUNCH_GUARD_H - -#include -#include -#include - -namespace ippl::detail { - enum class LaunchGuardMode { Disabled, Fence, Launch, LaunchAndFence }; - - inline const char* launchGuardModeName(LaunchGuardMode mode) { - switch (mode) { - case LaunchGuardMode::Disabled: - return "disabled"; - case LaunchGuardMode::Fence: - return "fence"; - case LaunchGuardMode::Launch: - return "launch"; - case LaunchGuardMode::LaunchAndFence: - return "launch+fence"; - } - return "unknown"; - } - - inline LaunchGuardMode launchGuardMode(const char* envName) { - const char* value = std::getenv(envName); - LaunchGuardMode mode = LaunchGuardMode::LaunchAndFence; - - if (value != nullptr) { - if (value[0] == '0' || value[0] == 'n' || value[0] == 'N' - || std::strcmp(value, "off") == 0 || std::strcmp(value, "OFF") == 0 - || std::strcmp(value, "false") == 0 || std::strcmp(value, "FALSE") == 0) { - mode = LaunchGuardMode::Disabled; - } else if (std::strcmp(value, "fence") == 0 || std::strcmp(value, "FENCE") == 0) { - mode = LaunchGuardMode::Fence; - } else if (std::strcmp(value, "launch") == 0 || std::strcmp(value, "LAUNCH") == 0) { - mode = LaunchGuardMode::Launch; - } - } - - return mode; - } - - inline LaunchGuardMode launchGuardMode(const char* envName, const char* label) { - const LaunchGuardMode mode = launchGuardMode(envName); - - static bool scatterPrinted = false; - static bool bareFieldPrinted = false; - static bool haloCellsPrinted = false; - static bool fallbackPrinted = false; - bool* printed = &fallbackPrinted; - - if (std::strcmp(envName, "IPPL_GH200_GUARD_SCATTER_POST_CIC") == 0) { - printed = &scatterPrinted; - } else if (std::strcmp(envName, "IPPL_GH200_GUARD_BAREFIELD_ACCUMULATE") == 0) { - printed = &bareFieldPrinted; - } else if (std::strcmp(envName, "IPPL_GH200_GUARD_HALOCELLS_ACCUMULATE") == 0) { - printed = &haloCellsPrinted; - } - - if (!*printed) { - const char* value = std::getenv(envName); - std::cout << "[IPPL launch guard] " << label << ": " << envName << "=" - << (value != nullptr ? value : "") - << " -> " << launchGuardModeName(mode) << std::endl; - *printed = true; - } - - return mode; - } -} // namespace ippl::detail - -#endif From 958b76f1b6d5fc006a2d90fcbeb4c084e849adc4 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Wed, 8 Jul 2026 11:32:57 +0200 Subject: [PATCH 15/32] Initialize BufferHandler correctly and add logging to communicator --- src/Communicate/BufferHandler.h | 4 +- src/Communicate/Communicator.h | 88 +++++++++++++++++++++++++++++++-- 2 files changed, 87 insertions(+), 5 deletions(-) diff --git a/src/Communicate/BufferHandler.h b/src/Communicate/BufferHandler.h index 9975a3631..720e24f0c 100644 --- a/src/Communicate/BufferHandler.h +++ b/src/Communicate/BufferHandler.h @@ -154,8 +154,8 @@ namespace ippl { buffer_type reallocateLargestFreeBuffer(size_type requiredSize); buffer_type allocateNewBuffer(size_type requiredSize); - size_type usedSize_m; ///< Total size of all allocated buffers - size_type freeSize_m; ///< Total size of all free buffers + size_type usedSize_m = 0; ///< Total size of all allocated buffers + size_type freeSize_m = 0; ///< Total size of all free buffers protected: buffer_set_type used_buffers{ diff --git a/src/Communicate/Communicator.h b/src/Communicate/Communicator.h index 9819ab3bd..400dbd23c 100644 --- a/src/Communicate/Communicator.h +++ b/src/Communicate/Communicator.h @@ -5,8 +5,10 @@ #ifndef IPPL_MPI_COMMUNICATOR_H #define IPPL_MPI_COMMUNICATOR_H +#include #include #include +#include #include "Communicate/BufferHandler.h" #include "Communicate/LoggingBufferHandler.h" @@ -25,6 +27,10 @@ #include "Communicate/Tags.h" //////////////////////////////////////////////////// +#if defined(KOKKOS_ENABLE_CUDA) +#include +#endif + namespace ippl { namespace mpi { @@ -140,6 +146,78 @@ namespace ippl { using buffer_handler_type = typename detail::ContainerForAllSpaces::type; + static bool debugArchiveMpiEnabled() { + static const bool enabled = [] { + const char* value = std::getenv("IPPL_DEBUG_ARCHIVE_MPI"); + return value != nullptr && value[0] != '\0' + && !(value[0] == '0' && value[1] == '\0'); + }(); + return enabled; + } + + static int debugArchiveMpiLimit() { + static const int limit = [] { + const char* value = std::getenv("IPPL_DEBUG_ARCHIVE_MPI_LIMIT"); + if (value == nullptr || value[0] == '\0') { + return 64; + } + long parsed = std::strtol(value, nullptr, 10); + return parsed < 0 ? 0 : static_cast(parsed); + }(); + return limit; + } + + template + void debugArchiveMpiBuffer(const char* op, int peer, int tag, Archive& ar, + detail::size_type mpiBytes) { + if (!debugArchiveMpiEnabled()) { + return; + } + + static int printed = 0; + auto* ptr = ar.getBuffer(); + const auto cap = ar.getBufferSize(); + const auto written = ar.getSize(); + const bool invalid = (mpiBytes > 0 && ptr == nullptr) || mpiBytes > cap; + const bool emit = invalid || printed < debugArchiveMpiLimit(); + + if (emit) { + ++printed; + std::cerr << "[rank " << rank_m << "] " << op << " archive buffer" + << " peer=" << peer << " tag=" << tag << " ptr=" + << static_cast(ptr) << " mpiBytes=" << mpiBytes + << " capacity=" << cap << " writeSize=" << written << std::endl; + +#if defined(KOKKOS_ENABLE_CUDA) + using memory_space = typename Archive::buffer_type::memory_space; + if constexpr (std::is_same_v) { + if (ptr != nullptr) { + cudaPointerAttributes attrs; + cudaError_t err = cudaPointerGetAttributes(&attrs, ptr); + if (err == cudaSuccess) { + std::cerr << "[rank " << rank_m << "] " << op + << " cudaPointerGetAttributes type=" + << static_cast(attrs.type) + << " device=" << attrs.device << std::endl; + } else { + std::cerr << "[rank " << rank_m << "] " << op + << " cudaPointerGetAttributes failed: " + << cudaGetErrorString(err) << std::endl; + cudaGetLastError(); + } + } + } +#endif + } + + if (invalid) { + std::cerr << "[rank " << rank_m << "] invalid archive MPI buffer for " << op + << ": mpiBytes=" << mpiBytes << " capacity=" << cap + << " ptr=" << static_cast(ptr) << std::endl; + this->abort(); + } + } + public: using size_type = detail::size_type; double getDefaultOverallocation() const { return defaultOveralloc_m; } @@ -168,6 +246,7 @@ namespace ippl { this->abort(); } MPI_Status status; + debugArchiveMpiBuffer("MPI_Recv", src, tag, ar, msize); MPI_Recv(ar.getBuffer(), msize, MPI_BYTE, src, tag, *comm_m, &status); buffer.deserialize(ar, nrecvs); @@ -176,12 +255,14 @@ namespace ippl { template void isend(int dest, int tag, Buffer& buffer, Archive& ar, MPI_Request& request, size_type nsends) { - if (ar.getSize() > INT_MAX) { + buffer.serialize(ar, nsends); + const size_type sendSize = ar.getSize(); + if (sendSize > INT_MAX) { std::cerr << "Message size exceeds range of int" << std::endl; this->abort(); } - buffer.serialize(ar, nsends); - MPI_Isend(ar.getBuffer(), ar.getSize(), MPI_BYTE, dest, tag, *comm_m, &request); + debugArchiveMpiBuffer("MPI_Isend", dest, tag, ar, sendSize); + MPI_Isend(ar.getBuffer(), sendSize, MPI_BYTE, dest, tag, *comm_m, &request); } template @@ -190,6 +271,7 @@ namespace ippl { std::cerr << "Message size exceeds range of int" << std::endl; this->abort(); } + debugArchiveMpiBuffer("MPI_Irecv", src, tag, ar, msize); MPI_Irecv(ar.getBuffer(), msize, MPI_BYTE, src, tag, *comm_m, &request); } From e4fadc033c305f2e00c04f10613ecaf72758dde9 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Wed, 8 Jul 2026 11:48:29 +0200 Subject: [PATCH 16/32] more debugging output --- src/Field/BareField.hpp | 19 +++++++++ src/Field/HaloCells.hpp | 68 ++++++++++++++++++++++++++++++++- src/Particle/ParticleAttrib.hpp | 37 ++++++++++++++++++ src/Utility/Debug.h | 19 +++++++++ 4 files changed, 142 insertions(+), 1 deletion(-) create mode 100644 src/Utility/Debug.h diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 6cd9aade9..f92c279c0 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -9,12 +9,14 @@ #include #include +#include #include #include #include "Communicate/DataTypes.h" #include "Utility/Inform.h" +#include "Utility/Debug.h" #include "Utility/IpplInfo.h" #include "BareField.h" @@ -156,13 +158,30 @@ namespace ippl { template void BareField::accumulateHalo() { + const bool debugScatterHalo = detail::debugScatterHaloEnabled(); + if (debugScatterHalo) { + std::cerr << "[rank " << Comm->rank() << "] BareField::accumulateHalo begin" + << " commSize=" << layout_m->comm.size() << " viewExtent=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << dview_m.extent(d) << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ")" << std::endl; + } if (layout_m->comm.size() > 1) { halo_m.accumulateHalo(dview_m, layout_m); + if (debugScatterHalo) { + std::cerr << "[rank " << Comm->rank() + << "] BareField::accumulateHalo completed halo exchange" << std::endl; + } } if (layout_m->isAllPeriodic_m) { using Op = typename detail::HaloCells::rhs_plus_assign; halo_m.template applyPeriodicSerialDim(dview_m, layout_m, nghost_m); } + if (debugScatterHalo) { + std::cerr << "[rank " << Comm->rank() << "] BareField::accumulateHalo end" + << std::endl; + } } template diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 542e067fc..90867c6db 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -4,11 +4,13 @@ // #include +#include #include #include "Utility/IpplException.h" #include "Communicate/Communicator.h" +#include "Utility/Debug.h" namespace ippl { namespace detail { @@ -64,12 +66,23 @@ namespace ippl { } int me = Comm->rank(); + constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; + const bool debugScatterHalo = detail::debugScatterHaloEnabled(); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries begin" + << " order=" << static_cast(order) + << " totalRequests=" << totalRequests + << " cubeCount=" << cubeCount << " viewExtent=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << view.extent(d) << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ")" << std::endl; + } using memory_space = typename view_type::memory_space; using buffer_type = mpi::Communicator::buffer_type; std::vector requests(totalRequests); // sending loop - constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; size_t requestIndex = 0; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + index; @@ -107,10 +120,33 @@ namespace ippl { } size_type nsends; + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries pack" + << " target=" << targetRank << " tag=" << tag + << " range.lo=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << range.lo[d] << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ") range.hi=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << range.hi[d] << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ")" << std::endl; + } pack(range, view, haloData_m, nsends); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries packed" + << " target=" << targetRank << " tag=" << tag + << " nsends=" << nsends << std::endl; + } buffer_type buf = comm.template getBuffer(nsends); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries isend" + << " target=" << targetRank << " tag=" << tag + << " nsends=" << nsends << std::endl; + } comm.isend(targetRank, tag, haloData_m, *buf, requests[requestIndex++], nsends); buf->resetWritePos(); } @@ -151,18 +187,35 @@ namespace ippl { buffer_type buf = comm.template getBuffer(nrecvs); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries recv" + << " source=" << sourceRank << " tag=" << tag + << " nrecvs=" << nrecvs << std::endl; + } comm.recv(sourceRank, tag, haloData_m, *buf, nrecvs * sizeof(T), nrecvs); buf->resetReadPos(); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries unpack" + << " source=" << sourceRank << " tag=" << tag + << " nrecvs=" << nrecvs << std::endl; + } unpack(range, view, haloData_m); } } if (totalRequests > 0) { + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries waitall" + << " totalRequests=" << totalRequests << std::endl; + } MPI_Waitall(totalRequests, requests.data(), MPI_STATUSES_IGNORE); } comm.freeAllBuffers(); + if (debugScatterHalo) { + std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries end" << std::endl; + } } template @@ -179,6 +232,15 @@ namespace ippl { Kokkos::realloc(buffer, size * overalloc); } + if (detail::debugScatterHaloEnabled()) { + std::cerr << "[rank " << Comm->rank() << "] HaloCells::pack launch" + << " nsends=" << nsends << " bufferExtent=" << buffer.extent(0) + << " subviewExtent=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << subview.extent(d) << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ")" << std::endl; + } using index_array_type = typename RangePolicy::index_array_type; ippl::parallel_for( @@ -197,6 +259,10 @@ namespace ippl { buffer(l) = apply(subview, args); }); Kokkos::fence(); + if (detail::debugScatterHaloEnabled()) { + std::cerr << "[rank " << Comm->rank() << "] HaloCells::pack completed" + << " nsends=" << nsends << std::endl; + } } template diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index b7e9d22a6..a454f3e2d 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -16,9 +16,11 @@ #include "Ippl.h" #include +#include #include "Communicate/DataTypes.h" +#include "Utility/Debug.h" #include "Utility/IpplTimings.h" namespace ippl { @@ -157,6 +159,25 @@ namespace ippl { auto dview = dview_m; auto ppview = pp.getView(); const int rank = Comm->rank(); + const bool debugScatterHalo = detail::debugScatterHaloEnabled(); + if (debugScatterHalo) { + std::cerr << "[rank " << rank << "] ParticleAttrib::scatter begin policy=[" + << iteration_policy.begin() << ", " << iteration_policy.end() + << ") useHashView=" << useHashView + << " hashExtent=" << hash_array.extent(0) + << " attribExtent=" << dview.extent(0) + << " posExtent=" << ppview.extent(0) + << " fieldExtent=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << view.extent(d) << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ") nghost=" << nghost << " lDomFirst=("; + for (unsigned d = 0; d < Dim; ++d) { + std::cerr << lDom[d].first() << (d + 1 == Dim ? "" : ", "); + } + std::cerr << ")" << std::endl; + } + Kokkos::parallel_for( "ParticleAttrib::scatter", iteration_policy, KOKKOS_LAMBDA(const size_t idx) { // map index to possible hash_map @@ -213,11 +234,27 @@ namespace ippl { detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi, args, val); }); + if (debugScatterHalo) { + std::cerr << "[rank " << rank + << "] ParticleAttrib::scatter launched kernel, entering debug fence" + << std::endl; + Kokkos::fence("ParticleAttrib::scatter debug fence"); + std::cerr << "[rank " << rank << "] ParticleAttrib::scatter completed debug fence" + << std::endl; + } IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); IpplTimings::startTimer(accumulateHaloTimer); + if (debugScatterHalo) { + std::cerr << "[rank " << rank << "] ParticleAttrib::scatter entering accumulateHalo" + << std::endl; + } f.accumulateHalo(); + if (debugScatterHalo) { + std::cerr << "[rank " << rank << "] ParticleAttrib::scatter completed accumulateHalo" + << std::endl; + } IpplTimings::stopTimer(accumulateHaloTimer); } diff --git a/src/Utility/Debug.h b/src/Utility/Debug.h new file mode 100644 index 000000000..2e1a6a177 --- /dev/null +++ b/src/Utility/Debug.h @@ -0,0 +1,19 @@ +#ifndef IPPL_UTILITY_DEBUG_H +#define IPPL_UTILITY_DEBUG_H + +#include + +namespace ippl { + namespace detail { + inline bool debugScatterHaloEnabled() { + static const bool enabled = [] { + const char* value = std::getenv("IPPL_DEBUG_SCATTER_HALO"); + return value != nullptr && value[0] != '\0' + && !(value[0] == '0' && value[1] == '\0'); + }(); + return enabled; + } + } // namespace detail +} // namespace ippl + +#endif From 263eed10998dbfa1c31fa1003de89de9a61785f3 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 11:05:57 +0200 Subject: [PATCH 17/32] Prepare BareField accumulateHalo env probe --- src/Field/BareField.hpp | 19 ++----------------- 1 file changed, 2 insertions(+), 17 deletions(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index f92c279c0..e3237f21c 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -9,7 +9,6 @@ #include #include -#include #include #include @@ -159,29 +158,15 @@ namespace ippl { template void BareField::accumulateHalo() { const bool debugScatterHalo = detail::debugScatterHaloEnabled(); - if (debugScatterHalo) { - std::cerr << "[rank " << Comm->rank() << "] BareField::accumulateHalo begin" - << " commSize=" << layout_m->comm.size() << " viewExtent=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << dview_m.extent(d) << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ")" << std::endl; - } + (void)debugScatterHalo; + if (layout_m->comm.size() > 1) { halo_m.accumulateHalo(dview_m, layout_m); - if (debugScatterHalo) { - std::cerr << "[rank " << Comm->rank() - << "] BareField::accumulateHalo completed halo exchange" << std::endl; - } } if (layout_m->isAllPeriodic_m) { using Op = typename detail::HaloCells::rhs_plus_assign; halo_m.template applyPeriodicSerialDim(dview_m, layout_m, nghost_m); } - if (debugScatterHalo) { - std::cerr << "[rank " << Comm->rank() << "] BareField::accumulateHalo end" - << std::endl; - } } template From 0ef0f30fe33dda2f4a392f217f96743f149dc488 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 12:19:35 +0200 Subject: [PATCH 18/32] Revert GH200 halo debug probes --- src/Communicate/BufferHandler.h | 4 +- src/Communicate/Communicator.h | 88 ++------------------------------- src/Field/BareField.hpp | 4 -- src/Field/HaloCells.hpp | 68 +------------------------ src/Particle/ParticleAttrib.hpp | 82 +----------------------------- src/Utility/Debug.h | 19 ------- 6 files changed, 8 insertions(+), 257 deletions(-) delete mode 100644 src/Utility/Debug.h diff --git a/src/Communicate/BufferHandler.h b/src/Communicate/BufferHandler.h index 720e24f0c..9975a3631 100644 --- a/src/Communicate/BufferHandler.h +++ b/src/Communicate/BufferHandler.h @@ -154,8 +154,8 @@ namespace ippl { buffer_type reallocateLargestFreeBuffer(size_type requiredSize); buffer_type allocateNewBuffer(size_type requiredSize); - size_type usedSize_m = 0; ///< Total size of all allocated buffers - size_type freeSize_m = 0; ///< Total size of all free buffers + size_type usedSize_m; ///< Total size of all allocated buffers + size_type freeSize_m; ///< Total size of all free buffers protected: buffer_set_type used_buffers{ diff --git a/src/Communicate/Communicator.h b/src/Communicate/Communicator.h index 400dbd23c..9819ab3bd 100644 --- a/src/Communicate/Communicator.h +++ b/src/Communicate/Communicator.h @@ -5,10 +5,8 @@ #ifndef IPPL_MPI_COMMUNICATOR_H #define IPPL_MPI_COMMUNICATOR_H -#include #include #include -#include #include "Communicate/BufferHandler.h" #include "Communicate/LoggingBufferHandler.h" @@ -27,10 +25,6 @@ #include "Communicate/Tags.h" //////////////////////////////////////////////////// -#if defined(KOKKOS_ENABLE_CUDA) -#include -#endif - namespace ippl { namespace mpi { @@ -146,78 +140,6 @@ namespace ippl { using buffer_handler_type = typename detail::ContainerForAllSpaces::type; - static bool debugArchiveMpiEnabled() { - static const bool enabled = [] { - const char* value = std::getenv("IPPL_DEBUG_ARCHIVE_MPI"); - return value != nullptr && value[0] != '\0' - && !(value[0] == '0' && value[1] == '\0'); - }(); - return enabled; - } - - static int debugArchiveMpiLimit() { - static const int limit = [] { - const char* value = std::getenv("IPPL_DEBUG_ARCHIVE_MPI_LIMIT"); - if (value == nullptr || value[0] == '\0') { - return 64; - } - long parsed = std::strtol(value, nullptr, 10); - return parsed < 0 ? 0 : static_cast(parsed); - }(); - return limit; - } - - template - void debugArchiveMpiBuffer(const char* op, int peer, int tag, Archive& ar, - detail::size_type mpiBytes) { - if (!debugArchiveMpiEnabled()) { - return; - } - - static int printed = 0; - auto* ptr = ar.getBuffer(); - const auto cap = ar.getBufferSize(); - const auto written = ar.getSize(); - const bool invalid = (mpiBytes > 0 && ptr == nullptr) || mpiBytes > cap; - const bool emit = invalid || printed < debugArchiveMpiLimit(); - - if (emit) { - ++printed; - std::cerr << "[rank " << rank_m << "] " << op << " archive buffer" - << " peer=" << peer << " tag=" << tag << " ptr=" - << static_cast(ptr) << " mpiBytes=" << mpiBytes - << " capacity=" << cap << " writeSize=" << written << std::endl; - -#if defined(KOKKOS_ENABLE_CUDA) - using memory_space = typename Archive::buffer_type::memory_space; - if constexpr (std::is_same_v) { - if (ptr != nullptr) { - cudaPointerAttributes attrs; - cudaError_t err = cudaPointerGetAttributes(&attrs, ptr); - if (err == cudaSuccess) { - std::cerr << "[rank " << rank_m << "] " << op - << " cudaPointerGetAttributes type=" - << static_cast(attrs.type) - << " device=" << attrs.device << std::endl; - } else { - std::cerr << "[rank " << rank_m << "] " << op - << " cudaPointerGetAttributes failed: " - << cudaGetErrorString(err) << std::endl; - cudaGetLastError(); - } - } - } -#endif - } - - if (invalid) { - std::cerr << "[rank " << rank_m << "] invalid archive MPI buffer for " << op - << ": mpiBytes=" << mpiBytes << " capacity=" << cap - << " ptr=" << static_cast(ptr) << std::endl; - this->abort(); - } - } - public: using size_type = detail::size_type; double getDefaultOverallocation() const { return defaultOveralloc_m; } @@ -246,7 +168,6 @@ namespace ippl { this->abort(); } MPI_Status status; - debugArchiveMpiBuffer("MPI_Recv", src, tag, ar, msize); MPI_Recv(ar.getBuffer(), msize, MPI_BYTE, src, tag, *comm_m, &status); buffer.deserialize(ar, nrecvs); @@ -255,14 +176,12 @@ namespace ippl { template void isend(int dest, int tag, Buffer& buffer, Archive& ar, MPI_Request& request, size_type nsends) { - buffer.serialize(ar, nsends); - const size_type sendSize = ar.getSize(); - if (sendSize > INT_MAX) { + if (ar.getSize() > INT_MAX) { std::cerr << "Message size exceeds range of int" << std::endl; this->abort(); } - debugArchiveMpiBuffer("MPI_Isend", dest, tag, ar, sendSize); - MPI_Isend(ar.getBuffer(), sendSize, MPI_BYTE, dest, tag, *comm_m, &request); + buffer.serialize(ar, nsends); + MPI_Isend(ar.getBuffer(), ar.getSize(), MPI_BYTE, dest, tag, *comm_m, &request); } template @@ -271,7 +190,6 @@ namespace ippl { std::cerr << "Message size exceeds range of int" << std::endl; this->abort(); } - debugArchiveMpiBuffer("MPI_Irecv", src, tag, ar, msize); MPI_Irecv(ar.getBuffer(), msize, MPI_BYTE, src, tag, *comm_m, &request); } diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index e3237f21c..6cd9aade9 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -15,7 +15,6 @@ #include "Communicate/DataTypes.h" #include "Utility/Inform.h" -#include "Utility/Debug.h" #include "Utility/IpplInfo.h" #include "BareField.h" @@ -157,9 +156,6 @@ namespace ippl { template void BareField::accumulateHalo() { - const bool debugScatterHalo = detail::debugScatterHaloEnabled(); - (void)debugScatterHalo; - if (layout_m->comm.size() > 1) { halo_m.accumulateHalo(dview_m, layout_m); } diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 90867c6db..542e067fc 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -4,13 +4,11 @@ // #include -#include #include #include "Utility/IpplException.h" #include "Communicate/Communicator.h" -#include "Utility/Debug.h" namespace ippl { namespace detail { @@ -66,23 +64,12 @@ namespace ippl { } int me = Comm->rank(); - constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; - const bool debugScatterHalo = detail::debugScatterHaloEnabled(); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries begin" - << " order=" << static_cast(order) - << " totalRequests=" << totalRequests - << " cubeCount=" << cubeCount << " viewExtent=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << view.extent(d) << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ")" << std::endl; - } using memory_space = typename view_type::memory_space; using buffer_type = mpi::Communicator::buffer_type; std::vector requests(totalRequests); // sending loop + constexpr size_t cubeCount = detail::countHypercubes(Dim) - 1; size_t requestIndex = 0; for (size_t index = 0; index < cubeCount; index++) { int tag = mpi::tag::HALO + index; @@ -120,33 +107,10 @@ namespace ippl { } size_type nsends; - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries pack" - << " target=" << targetRank << " tag=" << tag - << " range.lo=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << range.lo[d] << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ") range.hi=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << range.hi[d] << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ")" << std::endl; - } pack(range, view, haloData_m, nsends); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries packed" - << " target=" << targetRank << " tag=" << tag - << " nsends=" << nsends << std::endl; - } buffer_type buf = comm.template getBuffer(nsends); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries isend" - << " target=" << targetRank << " tag=" << tag - << " nsends=" << nsends << std::endl; - } comm.isend(targetRank, tag, haloData_m, *buf, requests[requestIndex++], nsends); buf->resetWritePos(); } @@ -187,35 +151,18 @@ namespace ippl { buffer_type buf = comm.template getBuffer(nrecvs); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries recv" - << " source=" << sourceRank << " tag=" << tag - << " nrecvs=" << nrecvs << std::endl; - } comm.recv(sourceRank, tag, haloData_m, *buf, nrecvs * sizeof(T), nrecvs); buf->resetReadPos(); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries unpack" - << " source=" << sourceRank << " tag=" << tag - << " nrecvs=" << nrecvs << std::endl; - } unpack(range, view, haloData_m); } } if (totalRequests > 0) { - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries waitall" - << " totalRequests=" << totalRequests << std::endl; - } MPI_Waitall(totalRequests, requests.data(), MPI_STATUSES_IGNORE); } comm.freeAllBuffers(); - if (debugScatterHalo) { - std::cerr << "[rank " << me << "] HaloCells::exchangeBoundaries end" << std::endl; - } } template @@ -232,15 +179,6 @@ namespace ippl { Kokkos::realloc(buffer, size * overalloc); } - if (detail::debugScatterHaloEnabled()) { - std::cerr << "[rank " << Comm->rank() << "] HaloCells::pack launch" - << " nsends=" << nsends << " bufferExtent=" << buffer.extent(0) - << " subviewExtent=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << subview.extent(d) << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ")" << std::endl; - } using index_array_type = typename RangePolicy::index_array_type; ippl::parallel_for( @@ -259,10 +197,6 @@ namespace ippl { buffer(l) = apply(subview, args); }); Kokkos::fence(); - if (detail::debugScatterHaloEnabled()) { - std::cerr << "[rank " << Comm->rank() << "] HaloCells::pack completed" - << " nsends=" << nsends << std::endl; - } } template diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index a454f3e2d..af484d2bb 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -15,12 +15,8 @@ // #include "Ippl.h" -#include -#include - #include "Communicate/DataTypes.h" -#include "Utility/Debug.h" #include "Utility/IpplTimings.h" namespace ippl { @@ -158,44 +154,10 @@ namespace ippl { } auto dview = dview_m; auto ppview = pp.getView(); - const int rank = Comm->rank(); - const bool debugScatterHalo = detail::debugScatterHaloEnabled(); - if (debugScatterHalo) { - std::cerr << "[rank " << rank << "] ParticleAttrib::scatter begin policy=[" - << iteration_policy.begin() << ", " << iteration_policy.end() - << ") useHashView=" << useHashView - << " hashExtent=" << hash_array.extent(0) - << " attribExtent=" << dview.extent(0) - << " posExtent=" << ppview.extent(0) - << " fieldExtent=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << view.extent(d) << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ") nghost=" << nghost << " lDomFirst=("; - for (unsigned d = 0; d < Dim; ++d) { - std::cerr << lDom[d].first() << (d + 1 == Dim ? "" : ", "); - } - std::cerr << ")" << std::endl; - } - Kokkos::parallel_for( "ParticleAttrib::scatter", iteration_policy, KOKKOS_LAMBDA(const size_t idx) { // map index to possible hash_map - const long long mapped_idx_signed = - useHashView ? static_cast(hash_array(idx)) - : static_cast(idx); - if (mapped_idx_signed < 0 - || mapped_idx_signed >= static_cast(dview.extent(0)) - || mapped_idx_signed >= static_cast(ppview.extent(0))) { - printf("[rank %d] ParticleAttrib::scatter invalid mapped index: idx=%llu " - "mapped=%lld attribExtent=%lld posExtent=%lld useHashView=%d\n", - rank, static_cast(idx), mapped_idx_signed, - static_cast(dview.extent(0)), - static_cast(ppview.extent(0)), static_cast(useHashView)); - Kokkos::abort("ParticleAttrib::scatter invalid mapped index"); - } - - size_t mapped_idx = static_cast(mapped_idx_signed); + size_t mapped_idx = useHashView ? hash_array(idx) : idx; // find nearest grid point vector_type l = (ppview(mapped_idx) - origin) * invdx + 0.5; @@ -203,58 +165,18 @@ namespace ippl { Vector whi = l - index; Vector wlo = 1.0 - whi; - Vector args = index - lDom.first() + nghost; - - for (unsigned point = 0; point < (1u << Field::dim); ++point) { - for (unsigned d = 0; d < Field::dim; ++d) { - const int target = args[d] - ((point & (1u << d)) ? 1 : 0); - const int extent = static_cast(view.extent(d)); - if (target < 0 || target >= extent) { - printf("[rank %d] ParticleAttrib::scatter invalid CIC target: idx=%llu " - "mapped=%llu point=%u dim=%u target=%d extent=%d base=(", - rank, static_cast(idx), - static_cast(mapped_idx), point, d, target, - extent); - for (unsigned p = 0; p < Field::dim; ++p) { - printf("%d%s", args[p], (p + 1 == Field::dim) ? "" : ", "); - } - printf(") viewExtent=("); - for (unsigned p = 0; p < Field::dim; ++p) { - printf("%lld%s", static_cast(view.extent(p)), - (p + 1 == Field::dim) ? "" : ", "); - } - printf(")\n"); - Kokkos::abort("ParticleAttrib::scatter invalid CIC target"); - } - } - } + Vector args = index - lDom.first() + nghost; // scatter const value_type& val = dview(mapped_idx); detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi, args, val); }); - if (debugScatterHalo) { - std::cerr << "[rank " << rank - << "] ParticleAttrib::scatter launched kernel, entering debug fence" - << std::endl; - Kokkos::fence("ParticleAttrib::scatter debug fence"); - std::cerr << "[rank " << rank << "] ParticleAttrib::scatter completed debug fence" - << std::endl; - } IpplTimings::stopTimer(scatterTimer); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); IpplTimings::startTimer(accumulateHaloTimer); - if (debugScatterHalo) { - std::cerr << "[rank " << rank << "] ParticleAttrib::scatter entering accumulateHalo" - << std::endl; - } f.accumulateHalo(); - if (debugScatterHalo) { - std::cerr << "[rank " << rank << "] ParticleAttrib::scatter completed accumulateHalo" - << std::endl; - } IpplTimings::stopTimer(accumulateHaloTimer); } diff --git a/src/Utility/Debug.h b/src/Utility/Debug.h deleted file mode 100644 index 2e1a6a177..000000000 --- a/src/Utility/Debug.h +++ /dev/null @@ -1,19 +0,0 @@ -#ifndef IPPL_UTILITY_DEBUG_H -#define IPPL_UTILITY_DEBUG_H - -#include - -namespace ippl { - namespace detail { - inline bool debugScatterHaloEnabled() { - static const bool enabled = [] { - const char* value = std::getenv("IPPL_DEBUG_SCATTER_HALO"); - return value != nullptr && value[0] != '\0' - && !(value[0] == '0' && value[1] == '\0'); - }(); - return enabled; - } - } // namespace detail -} // namespace ippl - -#endif From a0817c612cc23447f0913b272d86db0f3cf50564 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 12:20:49 +0200 Subject: [PATCH 19/32] Test BareField pre-halo fence --- src/Field/BareField.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 6cd9aade9..e66531b7b 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -156,6 +156,8 @@ namespace ippl { template void BareField::accumulateHalo() { + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); + if (layout_m->comm.size() > 1) { halo_m.accumulateHalo(dview_m, layout_m); } From f35e7d93eee8f8c6757cb7bc5e2222acc1c181d9 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 12:46:27 +0200 Subject: [PATCH 20/32] Try putting the fence only inside scatte --- src/Field/BareField.hpp | 2 +- src/Particle/ParticleAttrib.hpp | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index e66531b7b..dab10cf65 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -156,7 +156,7 @@ namespace ippl { template void BareField::accumulateHalo() { - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); + //Kokkos::fence("BareField::accumulateHalo pre-halo fence"); if (layout_m->comm.size() > 1) { halo_m.accumulateHalo(dview_m, layout_m); diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index af484d2bb..cf115078e 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -173,6 +173,8 @@ namespace ippl { args, val); }); IpplTimings::stopTimer(scatterTimer); + + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); IpplTimings::startTimer(accumulateHaloTimer); From 02505060fa7fce876bb3b53f24081dcd2af429c5 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 12:58:29 +0200 Subject: [PATCH 21/32] Fence before the accumulateHalo call, but after the timer? --- src/Particle/ParticleAttrib.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index cf115078e..8729e7070 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -173,11 +173,10 @@ namespace ippl { args, val); }); IpplTimings::stopTimer(scatterTimer); - - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); IpplTimings::startTimer(accumulateHaloTimer); + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); f.accumulateHalo(); IpplTimings::stopTimer(accumulateHaloTimer); } From 249e57dc601e28ac0423f1a6674eaa71bf14b827 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 13:14:04 +0200 Subject: [PATCH 22/32] Fence only inside HaloCells? --- src/Field/HaloCells.hpp | 2 +- src/Particle/ParticleAttrib.hpp | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 542e067fc..9c58d3e6d 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -171,7 +171,7 @@ namespace ippl { auto subview = makeSubview(view, range); auto& buffer = fd.buffer; - + Kokkos::fence("HaloCells::pack pre-fence"); size_t size = subview.size(); nsends = size; if (buffer.size() < size) { diff --git a/src/Particle/ParticleAttrib.hpp b/src/Particle/ParticleAttrib.hpp index 8729e7070..af484d2bb 100644 --- a/src/Particle/ParticleAttrib.hpp +++ b/src/Particle/ParticleAttrib.hpp @@ -176,7 +176,6 @@ namespace ippl { static IpplTimings::TimerRef accumulateHaloTimer = IpplTimings::getTimer("accumulateHalo"); IpplTimings::startTimer(accumulateHaloTimer); - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); f.accumulateHalo(); IpplTimings::stopTimer(accumulateHaloTimer); } From f93cd19afbad958798748b6acfec5d2b5c75fc8e Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 13:26:44 +0200 Subject: [PATCH 23/32] Put fence inside accumulateHalo() but only in the multi rank path --- src/Field/BareField.hpp | 3 +-- src/Field/HaloCells.hpp | 1 - 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index dab10cf65..0b03b47ee 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -156,9 +156,8 @@ namespace ippl { template void BareField::accumulateHalo() { - //Kokkos::fence("BareField::accumulateHalo pre-halo fence"); - if (layout_m->comm.size() > 1) { + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 9c58d3e6d..0cfb5eadb 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -171,7 +171,6 @@ namespace ippl { auto subview = makeSubview(view, range); auto& buffer = fd.buffer; - Kokkos::fence("HaloCells::pack pre-fence"); size_t size = subview.size(); nsends = size; if (buffer.size() < size) { From c48afc6532fc56e9fa31c4109c0dad7a3af77d45 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 13:39:55 +0200 Subject: [PATCH 24/32] Put the fence only inside exchangeBoundaries --- src/Field/BareField.hpp | 1 - src/Field/HaloCells.hpp | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/src/Field/BareField.hpp b/src/Field/BareField.hpp index 0b03b47ee..6cd9aade9 100644 --- a/src/Field/BareField.hpp +++ b/src/Field/BareField.hpp @@ -157,7 +157,6 @@ namespace ippl { template void BareField::accumulateHalo() { if (layout_m->comm.size() > 1) { - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); halo_m.accumulateHalo(dview_m, layout_m); } if (layout_m->isAllPeriodic_m) { diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 0cfb5eadb..31f98c22d 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -36,7 +36,7 @@ namespace ippl { SendOrder order, int nghost) { using neighbor_list = typename Layout_t::neighbor_list; using range_list = typename Layout_t::neighbor_range_list; - + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); auto& comm = layout->comm; const neighbor_list& neighbors = layout->getNeighbors(); From 899ba7f208080a96317f2e719a66da60f9430063 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 13:51:37 +0200 Subject: [PATCH 25/32] And what if we put the fence ONLY before exchangeBoundary? --- src/Field/HaloCells.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 31f98c22d..404e70594 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,6 +17,7 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } @@ -36,7 +37,7 @@ namespace ippl { SendOrder order, int nghost) { using neighbor_list = typename Layout_t::neighbor_list; using range_list = typename Layout_t::neighbor_range_list; - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); + auto& comm = layout->comm; const neighbor_list& neighbors = layout->getNeighbors(); From 0b626fbeae426bed6046f8907b0bae01560fae84 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 14:35:36 +0200 Subject: [PATCH 26/32] Initialize buffer handler accounting --- src/Communicate/BufferHandler.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/Communicate/BufferHandler.h b/src/Communicate/BufferHandler.h index 9975a3631..720e24f0c 100644 --- a/src/Communicate/BufferHandler.h +++ b/src/Communicate/BufferHandler.h @@ -154,8 +154,8 @@ namespace ippl { buffer_type reallocateLargestFreeBuffer(size_type requiredSize); buffer_type allocateNewBuffer(size_type requiredSize); - size_type usedSize_m; ///< Total size of all allocated buffers - size_type freeSize_m; ///< Total size of all free buffers + size_type usedSize_m = 0; ///< Total size of all allocated buffers + size_type freeSize_m = 0; ///< Total size of all free buffers protected: buffer_set_type used_buffers{ From 979f6f251ee786589cdda13a2c68f9a7501ec1c3 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 14:37:30 +0200 Subject: [PATCH 27/32] Test non-template halo exchange boundaries --- src/Field/HaloCells.h | 3 --- src/Field/HaloCells.hpp | 14 ++++++++------ 2 files changed, 8 insertions(+), 9 deletions(-) diff --git a/src/Field/HaloCells.h b/src/Field/HaloCells.h index c4d87b9bd..6e19594f6 100644 --- a/src/Field/HaloCells.h +++ b/src/Field/HaloCells.h @@ -135,10 +135,7 @@ namespace ippl { * @param view is the original field data * @param layout the field layout storing the domain decomposition * @param order the data send orientation - * @tparam Op the data assigment operator of the - * unpack function call */ - template void exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost = 1); /*! diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 404e70594..472bd16d0 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,22 +17,20 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); - exchangeBoundaries(view, layout, HALO_TO_INTERNAL); + exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } template void HaloCells::accumulateHalo_noghost(view_type& view, Layout_t* layout, int nghost) { - exchangeBoundaries(view, layout, HALO_TO_INTERNAL_NOGHOST, nghost); + exchangeBoundaries(view, layout, HALO_TO_INTERNAL_NOGHOST, nghost); } template void HaloCells::fillHalo(view_type& view, Layout_t* layout) { - exchangeBoundaries(view, layout, INTERNAL_TO_HALO); + exchangeBoundaries(view, layout, INTERNAL_TO_HALO); } template - template void HaloCells::exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost) { using neighbor_list = typename Layout_t::neighbor_list; @@ -155,7 +153,11 @@ namespace ippl { comm.recv(sourceRank, tag, haloData_m, *buf, nrecvs * sizeof(T), nrecvs); buf->resetReadPos(); - unpack(range, view, haloData_m); + if (order == INTERNAL_TO_HALO) { + unpack(range, view, haloData_m); + } else { + unpack(range, view, haloData_m); + } } } From 900681abf56823ff36ed67c17c51ab3ef731cc7d Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Thu, 9 Jul 2026 14:53:46 +0200 Subject: [PATCH 28/32] Revert non-template halo exchange test --- src/Field/HaloCells.h | 3 +++ src/Field/HaloCells.hpp | 14 ++++++-------- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/src/Field/HaloCells.h b/src/Field/HaloCells.h index 6e19594f6..c4d87b9bd 100644 --- a/src/Field/HaloCells.h +++ b/src/Field/HaloCells.h @@ -135,7 +135,10 @@ namespace ippl { * @param view is the original field data * @param layout the field layout storing the domain decomposition * @param order the data send orientation + * @tparam Op the data assigment operator of the + * unpack function call */ + template void exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost = 1); /*! diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 472bd16d0..404e70594 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,20 +17,22 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - exchangeBoundaries(view, layout, HALO_TO_INTERNAL); + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); + exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } template void HaloCells::accumulateHalo_noghost(view_type& view, Layout_t* layout, int nghost) { - exchangeBoundaries(view, layout, HALO_TO_INTERNAL_NOGHOST, nghost); + exchangeBoundaries(view, layout, HALO_TO_INTERNAL_NOGHOST, nghost); } template void HaloCells::fillHalo(view_type& view, Layout_t* layout) { - exchangeBoundaries(view, layout, INTERNAL_TO_HALO); + exchangeBoundaries(view, layout, INTERNAL_TO_HALO); } template + template void HaloCells::exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost) { using neighbor_list = typename Layout_t::neighbor_list; @@ -153,11 +155,7 @@ namespace ippl { comm.recv(sourceRank, tag, haloData_m, *buf, nrecvs * sizeof(T), nrecvs); buf->resetReadPos(); - if (order == INTERNAL_TO_HALO) { - unpack(range, view, haloData_m); - } else { - unpack(range, view, haloData_m); - } + unpack(range, view, haloData_m); } } From b1dbe43194d17edf4947848f4266798bc5e9119b Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Fri, 10 Jul 2026 10:14:25 +0200 Subject: [PATCH 29/32] Force no-inline templated exchangeBoundaries. Perhaps the fence only worked because it prevented inlining --- src/Field/HaloCells.h | 3 +++ src/Field/HaloCells.hpp | 1 - 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/src/Field/HaloCells.h b/src/Field/HaloCells.h index c4d87b9bd..511940675 100644 --- a/src/Field/HaloCells.h +++ b/src/Field/HaloCells.h @@ -139,6 +139,9 @@ namespace ippl { * unpack function call */ template +#if defined(__GNUC__) || defined(__clang__) + __attribute__((noinline)) +#endif void exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost = 1); /*! diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 404e70594..0cfb5eadb 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,7 +17,6 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } From 00583172fa8ca0f90176fad008e165e7c50fcaa5 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Fri, 10 Jul 2026 10:46:37 +0200 Subject: [PATCH 30/32] Revert "Force no-inline templated exchangeBoundaries. Perhaps the fence only worked because it prevented inlining" This reverts commit b1dbe43194d17edf4947848f4266798bc5e9119b. --- src/Field/HaloCells.h | 3 --- src/Field/HaloCells.hpp | 1 + 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/src/Field/HaloCells.h b/src/Field/HaloCells.h index 511940675..c4d87b9bd 100644 --- a/src/Field/HaloCells.h +++ b/src/Field/HaloCells.h @@ -139,9 +139,6 @@ namespace ippl { * unpack function call */ template -#if defined(__GNUC__) || defined(__clang__) - __attribute__((noinline)) -#endif void exchangeBoundaries(view_type& view, Layout_t* layout, SendOrder order, int nghost = 1); /*! diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 0cfb5eadb..404e70594 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -17,6 +17,7 @@ namespace ippl { template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { + Kokkos::fence("BareField::accumulateHalo pre-halo fence"); exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } From 565f73ccff5d1435e32a6442e69d860ba34ec632 Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Fri, 10 Jul 2026 10:47:44 +0200 Subject: [PATCH 31/32] Avoid CUDA extended-lambda helper mismatch in halo kernels --- src/Field/HaloCells.hpp | 138 +++++++++++++++++++++++----------------- 1 file changed, 81 insertions(+), 57 deletions(-) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 404e70594..88a654faa 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -12,12 +12,77 @@ namespace ippl { namespace detail { + // These kernels are instantiated in multiple translation units. Named functors avoid + // NVCC's translation-unit-local extended-lambda copy/delete helper state. + template + struct HaloPackFunctor { + Subview subview; + Buffer buffer; + + KOKKOS_INLINE_FUNCTION void operator()(const IndexArray& args) const { + int l = 0; + + for (unsigned d1 = 0; d1 < Dim; ++d1) { + int next = args[d1]; + for (unsigned d2 = 0; d2 < d1; ++d2) { + next *= subview.extent(d2); + } + l += next; + } + + buffer(l) = apply(subview, args); + } + }; + + template + struct HaloUnpackFunctor { + Subview subview; + Buffer buffer; + Op op; + + KOKKOS_INLINE_FUNCTION void operator()(const IndexArray& args) const { + int l = 0; + + for (unsigned d1 = 0; d1 < Dim; ++d1) { + int next = args[d1]; + for (unsigned d2 = 0; d2 < d1; ++d2) { + next *= subview.extent(d2); + } + l += next; + } + + op(apply(subview, args), buffer(l)); + } + }; + + template + struct HaloPeriodicFunctor { + View view; + Op op; + unsigned dimension; + int extent; + int nghost; + + KOKKOS_INLINE_FUNCTION void operator()(IndexArray coords) const { + coords[dimension] += nghost; + auto&& left = apply(view, coords); + + coords[dimension] = extent - coords[dimension]; + auto&& right = apply(view, coords); + + coords[dimension] += 2 * nghost - 1 - extent; + op(apply(view, coords), right); + + coords[dimension] = extent - coords[dimension]; + op(apply(view, coords), left); + } + }; + template HaloCells::HaloCells() {} template void HaloCells::accumulateHalo(view_type& view, Layout_t* layout) { - Kokkos::fence("BareField::accumulateHalo pre-halo fence"); exchangeBoundaries(view, layout, HALO_TO_INTERNAL); } @@ -172,8 +237,8 @@ namespace ippl { auto subview = makeSubview(view, range); auto& buffer = fd.buffer; - size_t size = subview.size(); - nsends = size; + size_t size = subview.size(); + nsends = size; if (buffer.size() < size) { int overalloc = Comm->getDefaultOverallocation(); Kokkos::realloc(buffer, size * overalloc); @@ -181,21 +246,11 @@ namespace ippl { using index_array_type = typename RangePolicy::index_array_type; - ippl::parallel_for( - "HaloCells::pack()", getRangePolicy(subview), - KOKKOS_LAMBDA(const index_array_type& args) { - int l = 0; - - for (unsigned d1 = 0; d1 < Dim; d1++) { - int next = args[d1]; - for (unsigned d2 = 0; d2 < d1; d2++) { - next *= subview.extent(d2); - } - l += next; - } - - buffer(l) = apply(subview, args); - }); + using buffer_view_type = typename databuffer_type::view_type; + using functor_type = + HaloPackFunctor; + ippl::parallel_for("HaloCells::pack()", getRangePolicy(subview), + functor_type{subview, buffer}); Kokkos::fence(); } @@ -212,21 +267,10 @@ namespace ippl { using index_array_type = typename RangePolicy::index_array_type; - ippl::parallel_for( - "HaloCells::unpack()", getRangePolicy(subview), - KOKKOS_LAMBDA(const index_array_type& args) { - int l = 0; - - for (unsigned d1 = 0; d1 < Dim; d1++) { - int next = args[d1]; - for (unsigned d2 = 0; d2 < d1; d2++) { - next *= subview.extent(d2); - } - l += next; - } - - op(apply(subview, args), buffer(l)); - }); + using functor_type = + HaloUnpackFunctor; + ippl::parallel_for("HaloCells::unpack()", getRangePolicy(subview), + functor_type{subview, buffer, op}); Kokkos::fence(); } @@ -271,30 +315,10 @@ namespace ippl { using index_array_type = typename RangePolicy::index_array_type; - ippl::parallel_for( - "applyPeriodicSerialDim", createRangePolicy(begin, end), - KOKKOS_LAMBDA(index_array_type & coords) { - // The ghosts are filled starting from the inside - // of the domain proceeding outwards for both lower - // and upper faces. The extra brackets and explicit - // mention - - // nghost + i - coords[d] += nghost; - auto&& left = apply(view, coords); - - // N - nghost - i - coords[d] = N - coords[d]; - auto&& right = apply(view, coords); - - // nghost - 1 - i - coords[d] += 2 * nghost - 1 - N; - op(apply(view, coords), right); - - // N - (nghost - 1 - i) = N - (nghost - 1) + i - coords[d] = N - coords[d]; - op(apply(view, coords), left); - }); + using functor_type = HaloPeriodicFunctor; + ippl::parallel_for("applyPeriodicSerialDim", + createRangePolicy(begin, end), + functor_type{view, op, d, N, nghost}); Kokkos::fence(); } From f7a3db26bda5a6c0e8e40d1fc363037597d7b81f Mon Sep 17 00:00:00 2001 From: Alexander Liemen <47395086+aliemen@users.noreply.github.com> Date: Fri, 10 Jul 2026 23:11:31 +0200 Subject: [PATCH 32/32] Forgot to add also copy the comments into the functor --- src/Field/HaloCells.hpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/src/Field/HaloCells.hpp b/src/Field/HaloCells.hpp index 88a654faa..0f15339af 100644 --- a/src/Field/HaloCells.hpp +++ b/src/Field/HaloCells.hpp @@ -64,15 +64,22 @@ namespace ippl { int nghost; KOKKOS_INLINE_FUNCTION void operator()(IndexArray coords) const { + // The ghosts are filled starting from the inside of the domain proceeding outwards + // for both lower and upper faces. The extra brackets and explicit mention + + // nghost + i coords[dimension] += nghost; auto&& left = apply(view, coords); + // N - nghost - i coords[dimension] = extent - coords[dimension]; auto&& right = apply(view, coords); + // nghost - 1 - i coords[dimension] += 2 * nghost - 1 - extent; op(apply(view, coords), right); + // N - (nghost - 1 - i) = N - (nghost - 1) + i coords[dimension] = extent - coords[dimension]; op(apply(view, coords), left); }