GPU Cluster Pair Performance - #30
Merged
Merged
Conversation
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
…potentials Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
…_OPENMP=false Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
- neighbor.c: neighbor-list build always wrote via the AOS-only neighs_padded() macro while every read site used the layout-aware neighs() macro. CPU defaults to AOS (no mismatch) but GPU defaults to SOA, so force kernels read scrambled cross-cluster data instead of the real neighbor list, causing reneighbor-cycle crashes. Add the missing AOS->SOA compaction step (mirrors the existing CSR compaction) to both buildNeighborCPU and buildNeighborSuperclusters. - shell_methods.h/pbc.c/pbc.h: reverse() only folded ghost-cluster reaction forces back onto real atoms under _MPI, making it a no-op for non-MPI half-neighbor-list builds. Add reverseGhostForcesCPU() and wire it into the non-MPI path. - main.c: setup()/reneighbour() never refreshed ghost cluster positions via updatePbcCPU() after setupPbc() rebuilt the border map, leaving binJClusters() to bin ghosts by stale positions. Add the missing call. - neon_float.h/sve_float.h: add missing simd_real_h_reduce_sum and simd_real_masked_scatter_sub, needed for the verletlist half-neighbor SIMD kernel to link on NEON/SVE. - sve_double.h/sve_float.h: fix simd_real_masked_scatter_sub using svpnext_bXX to walk active predicate lanes indexed by loop counter i, mismatching values and indices once any mask lane is inactive (memory corruption on half-neighbor kernels). Use a plain bitmask + index loop instead, matching the already-correct AVX2 fallback. - force_lj_simd.c: replace raw operator subtraction with simd_real_sub() for gathered position deltas; SVE's ACLE vector types don't support arithmetic operators.
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
# Conflicts: # src/clusterpair/neighbor.c
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
…J_COMB_NONE references Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
Contributor
There was a problem hiding this comment.
Pull request overview
This PR expands MD-Bench’s benchmarking surface area around GPU Cluster Pair evaluation by adding new correctness/regression tests, introducing a tabulated Lennard‑Jones force-field option (lj_table), and extending runtime/build instrumentation (displacement-based reneighboring + verbose timing breakdown) across both optimization schemes.
Changes:
- Add tabulated/spline-interpolated LJ force field (
lj_table) with CPU+CUDA implementations (verletlist) and reference support scaffolding (clusterpair). - Add displacement-based neighbor-list rebuild option (
--displacement-reneigh) plus fine-grained timers and--verbosetiming breakdown. - Extend build/test infrastructure (binary tagging with
LJ_COMB_RULE, new OpenMP + displacement reneigh tests, CUDA arch override, CSR neighbor-list layout plumbing).
Reviewed changes
Copilot reviewed 82 out of 84 changed files in this pull request and generated 6 comments.
Show a summary per file
| File | Description |
|---|---|
| TODO.md | Updates project TODO items related to GPU CP + PBC on GPU. |
| tests/test_simd_vs_scalar.sh | Propagates LJ_COMB_RULE into builds and binary naming. |
| tests/test_openmp.sh | New OpenMP correctness regression across schemes/LJ rules/half-neigh. |
| tests/test_neighbor.c | Updates neighbor indexing call + adds displacement reneigh unit tests. |
| tests/test_mpi.sh | Adjusts MPI test binary tags to include LJ_COMB_RULE. |
| tests/test_lj_comb_rules.sh | Renames “none” rule references to “full” and updates cases. |
| tests/test_half_neigh.sh | Adds USE_SIMD_KERNEL control and updates binary tag naming. |
| tests/test_force.c | Adds unit tests validating ljtable spline accuracy vs analytic LJ. |
| tests/test_displacement_reneigh.sh | New end-to-end regression for displacement-based reneighboring. |
| tests/test_data_layout.sh | Includes LJ_COMB_RULE in binary tag derivation. |
| tests/regression_scheme_equiv.sh | Includes LJ_COMB_RULE in verletlist/clusterpair binary tags. |
| tests/regression_lj_table.sh | New regression: -f lj_table thermo equivalence vs -f lj. |
| src/verletlist/tracing.c | Updates neighs() access pattern + CSR/tracing guard. |
| src/verletlist/pbc.h | Splits setupPbc into CPU/CUDA variants via function pointer. |
| src/verletlist/pbc.c | Selects CPU/CUDA setupPbc implementation at init time. |
| src/verletlist/neighbor.h | Adds CSR layout support + numneigh_inner/neigh_start plumbing. |
| src/verletlist/main.c | Adds lj_table init, displacement reneigh, verbose timers, prune timing. |
| src/verletlist/main-stub.c | Aligns stub allocations with production + device init/copies. |
| src/verletlist/forceCuda.cu | Passes DeviceNeighbor into kernels + improves read-only loads. |
| src/verletlist/force.h | Adds FF_LJ_TABLE and SIMD kernel variant declarations. |
| src/verletlist/force.c | Selects among SIMD variants + adds lj_table dispatch (CPU/CUDA). |
| src/verletlist/force_lj.c | Uses inner neighbor counts + adds GEOM/FULL branches and new neighs() signature. |
| src/verletlist/force_lj_table.c | New CPU tabulated LJ implementation using ljtable coefficients. |
| src/verletlist/force_lj_table_cuda.cu | New CUDA tabulated LJ implementation + lazy table upload/free. |
| src/verletlist/force_eam.c | Updates neighs() signature usage. |
| src/verletlist/device_spec.c | Allocates/copies GEOM per-atom arrays + CSR neighbor device layout fields. |
| src/verletlist/atom.h | Adds reference-position storage + displacement reneigh API. |
| src/verletlist/atom.c | Preserves device buffers on grow + implements displacement reneigh logic. |
| src/common/util.h | Adds STR_EQ helper and related include changes. |
| src/common/util.c | Adds lj_table parsing and switches to whole-token force-field matching. |
| src/common/timers.h | Adds fine-grained timer buckets (neighbor sub-steps, integrate, thermo, checks). |
| src/common/thermo.c | Uses aligned allocator + treats FF_LJ_TABLE like LJ for thermo scaling. |
| src/common/simd/sve_float.h | Fixes SVE intrinsics forms + adds missing SIMD helpers (mask not, loads, sqrt, etc.). |
| src/common/simd/sve_double.h | Fixes SVE intrinsics forms + adds missing SIMD helpers and correct masked scatter fallback. |
| src/common/simd/scalar.h | Adds missing scalar SIMD helper APIs (loadu/sqrt/conversions/min). |
| src/common/simd/neon_float.h | Implements previously-stubbed mask conversion + adds missing SIMD helper APIs. |
| src/common/simd/neon_double.h | Implements previously-stubbed mask conversion + adds missing SIMD helper APIs. |
| src/common/simd/avx512_float.h | Adds missing helpers + implements horizontal sum for VL usage. |
| src/common/simd/avx512_double.h | Adds missing helpers (mask not, loadu, comparisons, sqrt, conversions, min). |
| src/common/simd/avx2_float.h | Adds missing helpers (mask not, loadu, eq compare, sqrt, conversions, min). |
| src/common/simd/avx2_double.h | Adds missing helpers (mask not, loadu, eq compare, sqrt, conversions, min). |
| src/common/simd/avx_float.h | Adds missing helpers (loadu, sqrt, conversions, min). |
| src/common/simd/avx_double.h | Adds missing helpers (loadu, sqrt, conversions, min). |
| src/common/shell_methods.h | Refines clusterpair reverse ghost force handling across CPU/CUDA/half-neigh. |
| src/common/parameter.h | Renames LJ “none” to “full” + adds lj_table_points, displacement reneigh, verbose flags. |
| src/common/parameter.c | Adds aligned allocs, STR_EQ parsing, lj_table_points, and parameter print enhancements. |
| src/common/ljtable.h | New API/struct for tabulated LJ coefficients and index-mode selection. |
| src/common/ljtable.c | New implementation building cubic Hermite spline coefficients for LJ force shapes. |
| src/common/grid.c | Switches several allocations to aligned allocator. |
| src/common/deviceGPU.cu | Adds reallocateGPUKeep() to preserve device buffers across growth. |
| src/common/device.h | Exposes reallocateGPUKeep() in the device API. |
| src/common/device.c | Adds stub for reallocateGPUKeep() in non-GPU builds. |
| src/common/comm.c | Formatting-only loop indentation adjustments. |
| src/clusterpair/vtk.c | Formatting-only loop indentation adjustments. |
| src/clusterpair/tracing.c | Updates neighs() access pattern + CSR/tracing guard. |
| src/clusterpair/pbc.h | Exposes reverseGhostForcesCPU() for non-MPI CPU path handling. |
| src/clusterpair/pbc.c | Implements reverseGhostForcesCPU() folding ghost reaction forces. |
| src/clusterpair/neighbor.h | Adds CSR support and changes neighs() signature to take Neighbor*. |
| src/clusterpair/main.c | Adds lj_table init, displacement reneigh, verbose timers, and timing breakdown. |
| src/clusterpair/main-stub.c | Aligns stub allocations with production + device init/copies + supercluster growth fix. |
| src/clusterpair/force.h | Adds FF_LJ_TABLE and declares table-kernel entrypoints. |
| src/clusterpair/force.c | Adds FF_LJ_TABLE dispatch (reference-only) and errors for SIMD/GPU for now. |
| src/clusterpair/force_lj_table.c | New reference tabulated LJ kernel for clusterpair. |
| src/clusterpair/atom.h | Adds reference-position storage and displacement reneigh API. |
| src/clusterpair/atom.c | Allocates reference positions + implements displacement reneigh logic. |
| README.md | Documents outer_skin, adds lj_table description, and updates LJ rule naming. |
| Makefile | Adds PREFIX_DIR, adds regression_lj_table + new tests to make test, links ljtable into test binary. |
| make/include_NVCC.mk | Makes CUDA SM arch configurable via CUDA_ARCH. |
| make/include_ICX.mk | Uses -qopenmp-simd when OpenMP disabled; adds -funroll-loops. |
| make/include_ICC.mk | Uses -qopenmp-simd when OpenMP disabled. |
| make/include_GCC.mk | Tweaks optimization flags and AVX512 flags. |
| make/include_CLANG.mk | Adds -funroll-loops. |
| config.mk | Adds CSR layout option, LJ table index config, SIMD neighbor build, SIMD compress, and tags binaries with LJ rule. |
| AGENTS.md | Updates build/testing guidance and documents new features/constraints. |
| .github/workflows/ci.yml | Updates CI binary tag construction to include LJ_COMB_RULE suffix. |
Suppressed comments (1)
src/verletlist/tracing.c:44
jis used in the tracing loop but is never assigned, so this block does not compile (and cannot trace neighbor atom accesses correctly). Assignjfrom the neighbor list entry before using it.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
38
to
42
| DIST_TRACE(neighbor->neighbors, i, numneighs, Nlocal, neighbor->maxneighs); | ||
|
|
||
| for (int k = 0; k < numneighs; k++) { | ||
| MEM_TRACE(neighs(neighbor->neighbors, i, k, Nlocal, neighbor->maxneighs), | ||
| 'R'); | ||
| MEM_TRACE(neighs(neighbor->neighbors, i, k, Nlocal, neighbor), 'R'); | ||
| MEM_TRACE(atom_x(j), 'R'); |
Comment on lines
+1059
to
+1061
| #ifdef _MPI | ||
| MPI_Allreduce(MPI_IN_PLACE, &max_sq, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD); | ||
| #endif |
Comment on lines
+32
to
+36
| #define NBLIST_DATA_LAYOUT "AoS" | ||
| #define neighs(l, i, j, M, nbr) (l)[(i) * (nbr)->maxneighs + (j)] | ||
| #elif defined(NBLIST_CSR) | ||
| #define NBLIST_DATA_LAYOUT "CSR" | ||
| #define neighs(l, i, j, M, nbr) (l)[(nbr)->neigh_start[(i)] + (j)] |
Comment on lines
+1304
to
+1306
| #ifdef _MPI | ||
| MPI_Allreduce(MPI_IN_PLACE, &max_sq, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD); | ||
| #endif |
Update README by stating lj_table is supported for CP cases Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
… memory tracer Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
It started as a performance evaluation of the GPU Cluster Pair kernels but it actually includes significant more stuff.
The changes, listed (AI generated):
GPU cluster-pair fixes & performance
SIMD/SVE work
New features
Build system
Tests
Docs
Note this branch also merged main at points (d1bbb36, 7218d75, etc.), so some of these commits reflect upstream work rather than this branch's own changes — the diff stat above is the net effect versus main's current tip either way.