GitLab

Introduce OPENMM_HIP_PME_FAST_GRID_LEVEL as a HIP-only experimental
control for reducing automatically selected PME grid dimensions.  A
positive level repeatedly steps each grid dimension down to the previous
legal FFT size, while level 0 preserves the existing OpenMM behavior.

Keep explicit user-provided PME grid parameters unchanged so the feature
only affects automatic grid selection.  This is intended for performance
and accuracy exploration, not as a default tuning policy.

Use explicit 128-thread launches for the three LangevinMiddle integration kernels to improve HIP throughput while
preserving the existing PME launch heuristics.

Use explicit 128-thread block launches for selected HIP PME kernels that
benefit from larger blocks.  Keep the platform default block size unchanged,
and leave small-system grid indexing and charge spreading on the existing
default launch configuration.

The heuristic applies 128-thread launches to finishSpreadCharge on HIP, and
uses 128-thread launches for findAtomGridIndex and gridSpreadCharge only for
larger systems.  Coulomb PME and LJPME dispersion paths are handled in
parallel, while interpolation and energy evaluation remain unchanged.

Apply heuristics for HIP neighbor-list construction:
use fewer nonbonded force blocks for small neighbor-list systems, use two
tiles per batch for larger atom-block counts, and increase the
findBlocksWithInteractions thread block size for small atom-block counts.

Standard concurrent validation shows no clear per-case regression and a
small geomean throughput improvement over the current blocksPerCU baseline.

- Use bitwise prefix accounting when storing sparse interactions as single pairs in the HIP pair-list kernel. This reduces the number of ballot operations needed to compute per-lane single-pair offsets.
- For HIP CDNA single precision, raise MAX_BITS_FOR_PAIRS to 8 so more sparse interactions are emitted as single pairs instead of full tiles. Keep the existing double precision and RDNA thresholds unchanged.
- Also simplify the HIP LJPME direct correction by computing alpha^2*r2

Avoid forcing Coulomb PME to re-sort whenever LJ-PME is enabled, and give dispersion PME its own atom-grid index and sort state so the performance impact can be measured independently.

* PythonForce can be restricted to a subset of particles

* Fix exception with CUDA

* optimization, switch to unordered map

* Remove emplace, reduce reserve allocation

* Cache coefficients for long range correction

* updateParametersInContext() clears cache

* Fixed issue that caused inefficient sorting when a block contained only one atom

* Add the fix to OpenCL and HIP

* Make reference/CPU minimizer into a kernel

* Add per-platform support for GPU minimization

* Initial implementation of GPU minimization

* Fixes

* Increase robustness when initial gradient is huge

* Handle overflow leading to non-finite values gracefully

* Handle large forces in single precision more robustly

* Optimize kernels

* Fix kernel launch size

* Update banner years

* Don't create MinimizeKernel until first minimization requested

* Make some compile-time constants into kernel arguments

* Consolidate scale calculation kernel

* Condense alpha/beta reduction kernels using atomics

* Condense line search dot kernels with reductions

* Remove a download, and download grad norm separately

* Asynchronously check lbfgs convergence condition

* Restructure line search to avoid download waiting

* Start line search preemptively in case CPU evaluation is not needed

* In rare cases, constraint error might not decrease after one optimization round

* Better handling of unsupported 64-bit atomics, use FLT_MAX

* Pick gradient mode based on GPU vs. CPU evaluation

* Rework getDiff/getScale reduction, remove reduceBuffer

* Older CUDA might not like float hex literals

* Fix error in a comment

* API for querying devices

* CUDA and HIP implementations of getDevices()

* Fix test failures

* Fix test failures

* CUDA returns correct devices even if no context has been created

* Return a single device for Reference and CPU

* Fix CI failure

* Remove std::enable_if, warpRotateLeft is always used with TILE_SIZE

* Do not use built-in warpSize in constexpr contexts

Starting from ROCm 7 warpSize is no longer constexpr.
findInteractingBlocks.hip uses it for sizes of __shared__ arrays.

* Check if hipHostMallocNumaUser is allowed before using it

adaptFrictionPart1 is executed first, it reads randomForce and
segmentVelocity buffers before other kernels write them so they can
contain garbage after allocation (cuMemAlloc/hipMalloc/clCreateBuffer
do not clear allocated memory).

* Basic LCPO support

* Add basic test for LCPO from a prmtop file

* API for LCPOForce

* Started LCPO reference implementation

* Finished reference forces & test cases

* Use other test for finite difference since grid might have discontinuous forces

* Reference platform formatting

* Initial implementation of CPU platform

* Bugfixes

* More vectorization and improve neighbor list query speed

* Parallelize part of neighbor search

* Check box size for LCPO with periodic boundary conditions

* Fixes for updating parameters in context

* GBSAOBCForce doesn't use first & last indices for updates, so no need for this optimization here

* Changes to neighbor checking and optimization

* Fixes and minor changes

* Add global surface tension parameter

* Only process half of the pairs in the neighbor list

* Remove unnecessary checks

* Initial version of common platform implementation

* Asynchronously download neighbor list size

* Debugging

* Do pair precomputation in copyPairsToNeighborList

* Recompute interactions instead of scanning neighbor list in inner loop

* Condense position array before computations

* Also make neighbor count download asynchronous on device

* Fixes for kernel launching

* Topology-based LCPO parameter assignment

* Fixes, and use test system for LCPO with nucleic acids

* Always raise instead of warn when LCPO parameters can't be assigned

* Use Amber convention for phosphates

* Minor optimization to PME

* Minor simplification

* Initial implementation of PythonForce

* Continuing implementation of PythonForce

* Tests for PythonForce

* Fix memory leaks

* Documentation for PythonForce

* Fixed incorrect return type

* Fix compilation error on Python older than 3.12

* Handle all dtypes

* Optimizations to PythonForce

* Optimized getPositions()

* Test all platforms

* Fix test failures

* Ensure that neighbor list is valid before solving for charges

* Add test with neighbor list that needs to be resized

* Try another approach to skip interactions for neighbor list generation only

* Increase CG error tolerance for test

* Replace SimTK-containing file headers

* Update file headers for new Tinker reader files added

* Initial implementation of C++ API

* Add kernel interface and information for API generation

* API updates for updating electrode parameters

* Add serialization proxy for ConstantPotentialForce

* Update file headers

* Add CG error tolerance and fix units on getCharges() return value

* Initial implementation of matrix solver

* Fixes and conjugate gradient solver

* Try to fix Linux and Windows builds

* Make sure charge constraint target is on total charge

* Restore handling of exceptions like NonbondedForce since they won't involve electrode atoms

* Ameliorate numerical instability in constrained conjugate gradient

* Fix uninitialized pointers, memory leak, and style

* Set CG tolerance units in Python API

* Test ConstantPotentialForce serialization

* Read/write ExceptionsUsePeriodicBoundaryConditions as bool

* Improve constrained conjugate gradient robustness to roundoff error accumulation

* Recompute matrix if electrode atoms move due to setPositions()

* Tolerance is now in gradient (potential) units again

* Add neutralizing background correction

* Add Python API tests

* Fixes for CG and nonbonded exceptions

* Add initial tests checking against existing NonbondedForce behavior

* Expand test suite and fix some implementation issues

* Add additional tests using larger reference system

* Add Gaussian test

* Finish test against reference computation

* CPU platform implementation

* Fixes for compilation on some platforms

* Fixes for constant potential with AVX/AVX2

* Test linking CPU PME library to constant potential test directly

* Older SWIG versions don't support Python set to C++ set conversion

* Add user guide entry

* Increase speed of reference test

* Conditional building constant potential CPU test is unreliable

* Debugging

* Miscellaneous fixes and improvements for CI

* Cache charges so solver will not run if system and coordinates have not changed

* Preconditioner flag, stability, and automatic detection improvements

* Add GPU platform-specific constant potential kernel classes

* PME and device-host I/O changes to support constant potential

* Initial common constant potential implementation

* Constant potential fixes:

* Fix preconditioner PME position/charge save/restore logic

* Fix reduction synchronization in constant potential solver kernels

* Add double-float accumulation for conjugate gradient solver when
  double unsupported by hardware

* Improve conditioning of a test system, and make sure particles are in or
out of cutoff for consistency and ease of comparing between platforms

* Reorder guess charges for CG when atom reordering changes positions

* Remove PME queue for now

* Trying to debug optimized direct space derivative kernel

* Remove extraneous debugging lines

* Style updates; just make CPU preconditioner double precision

* Debugging updated optimized direct derivatives kernel for all but OpenCL CPU

* OpenCL CPU implementation of direct space derivatives, and cleanup

* Try to make test even shorter to not time out on CI

* Temporary - Debugging

* Debugging

* Debugging

* Debugging

* Debugging

* Remove debugging code and fix reduction synchronization

* Fix other reductions

* Debugging - are tests hanging or just slow on CI?

* Debugging

* Debugging

* Fix macro for case when double precision is available on hardware

* Remove changes for debugging again

* Try to improve matrix solver cache locality by uploading transpose

* Fixes for atom ordering and periodic images

* Can't rely on reorder listener for cell offset updates

* Test reducing number of contexts and timing for CI

* Debugging

* Remove timing code and revert debugging changes

* Matrix solver and plasma term optimizations

* Reduce CG solver kernel calls and downloads

* Don't read back convergence flag from global memory

* Update PME due to refactoring in master branch

* Faster matrix solver (1st step)

* Faster matrix solver for CUDA

* Faster matrix solver compatibility with non-CUDA platforms

* Matrix solver fixes

* Use warp shuffle reductions when possible

* Attempt to work around intermittent compiler crash in Intel CPU OpenCL

* Optimize CG solver kernel 1

* Rework CG solver so some kernels can use more than 1 block

* Don't run out of shared memory

* Asynchronously download convergence flag while clearing buffers

---------
Co-authored-by: Evan Pretti <pretti@sh03-17n15.int>