GitLab

* 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>

* Big fix to QTB

* Fixed race condition

* Created OrientationRestraintForce

* Python API

* Documentation

* Serialization

* GPU implementation of OrientationRestraintForce

* Bug fix

* Remove unused code

* Reference implementation of RGForce

* GPU implementation of RGForce

* Serialization

* Documentation

* Fix compilation error

* Fixed error building API docs

* Variable displacements based on particle positions

* set variable displacements when adding particles

* update documentation

* address compilation error in OpenMMFortranWrapper

* update python API tests

* fix stray 'and'

* addParticle() without arguments add a particle that is not displaced

* pack displacement particles into int4

* put back default displacement removed in error

* ATMForce interface with coordinate transformation objects

* revise variable displacement API

* documentation, formatting, serialization

* Fixed C and Fortran wrappers

* Fixed Python wrappers

* Fixed factory

* Sort files to ensure classes are listed in the correct order

* Converted APIUnits test to new ATMForce API

* write class name

* skip the documentation for forward declarations

* undo 9e91d0b since it does not fix the doc build

* remove temporary doc files for nested classes

* Clean away tabs

---------
Co-authored-by: Peter Eastman <peastman@stanford.edu>

* Began implementing QTBIntegrator

* Adaptation and deconvolution

* Continuing reference implementation

* Continuing to implement QTBIntegrator

* Use common thread pool

* More tests, documentation, and threading

* Fix segfault

* Serialize adapted friction when creating a State

* Beginning of GPU implementation

* Added missing files

* Bug fixes

* Fixed inverse FFT

* Continuing GPU implementation

* Checkpointing

* Bug fixes

* Test cases run faster

* Changes needed for latest main branch

* Minor optimizations

* Documentation

* Fixed atom reordering

* Added parahydrogen test case

* Workaround for bug in Microsoft's compiler

* Added a Python test

* Normalize kernel in deconvolution

* Minor documentation improvements

* Reference implementation of SymmetrySite

* Common implementation of SymmetrySite

* Removed duplicated code

* Serialization for SymmetrySite

* Fixed compilation error building C wrapper

* Added SymmetrySite to user guide

* Bug fix

* Added P21 test case

* Unified storage of global parameters

* Fixes to CUDA and HIP

* Store global parameters as real instead of float

* Use fixed point spread charge on RDNA4 as it is faster

Even though RDNA4 (gfx12) has global_atomic_add_f32, micro-benchmarks and OpenMM benchmarks show
that it is very slow compared to global_atomic_add_u64.

* Add a workaround for fixed point gridSpreadCharge on RDNA4

Workaround for rare cases when few values of pmeGrid are very large and
incorrect. The cause is unknown. Why this workaround or other irrelevant
changes like printf help is also unknown.

* Add a workaround for infinite loop in computeNonbonded (HIP)

computeNonbonded hangs in some tests (without neighbor list).
Reproducible on ROCm 6.4 and 6.4.1 (maybe on older versions too) on various architectures (both CDNA and RDNA).
Affected tests: TestHipATMForce, TestHipMonteCarloBarostat, TestHipNonbondedForce, TestHipVirtualSites.

Disassembly shows that the compiler splits branches of `if (skipBase+tgx < NUM_TILES_WITH_EXCLUSIONS)` and does
`SHFL(skipTiles, TILE_SIZE-1) < pos` checks in them separately, even though `__builtin_amdgcn_ds_bpermute`
is a convergent function. Apparently in this case not all lanes participate in each call.

* Simplify includeTile check using ballot

* set box vectors of the inner contexts before atom reordering

* test for changing box vectors

* Optimized setPositions() and setVelocities()

* Fix test failures

* Fix GPU memory leak

* Undo CUDA change

* Use common API for kernels

* More code uses common interface

* Bug fixes

* Unified interface for sorting

* Simplified interface for FFT

* Use common event API for synchronization

* Minor changes to make code more consistent between platforms

* Common implementation of NonbondedForce

* Bug fixes

* Flag to enable list of single pairs

* CUDA and OpenCL use common implementation of NonbondedForce

* Fixed compilation error

* HIP uses common implementation of NonbondedForce

* Unified interface for queues

* Simplified stream handling in CudaFFT3D

* HIP implementation of ComputeQueue

* Added computeCurrentPressure() to MonteCarloBarostat

* Use instantaneous temperature to compute pressure

* Added computeCurrentPressure() to MonteCarloAnisotropicBarostat

* Added computeCurrentPressure() to MonteCarloMembraneBarostat

* Fixed compilation error

* Fixed error in typemap

* Added documentation on computing pressure

* Fixed CUDA compilation errors

* Made test case more robust

* Made a test case more robust

* Added computeCurrentPressure() to MonteCarloFlexibleBarostat

* Fixed compilation error

* More documentation on computing pressure

* Unified interface for FFTs

* AMOEBA uses unified interface for FFTs

* HIP implementation of common FFT interface

* Add correction for self energy of neutralizing plasma

* Fixed compilation errors

* Update total charge in copyParametersToContext()

* Bug fixes

* Fixed compilation errors in HIP

* Bug fix

* Created DPDIntegrator class

* Reference implementation of DPDIntegrator

* Build neighbor list for DPDIntegrator

* Minor fixes

* Documentation for DPDIntegrator

* Python API for DPDIntegrator

* Preliminary OpenCL implementation of DPDIntegrator

* Enable USE_PERIODIC

* Use updated positions in DPD thermostat

* Working on neighbor list for OpenCL DPDIntegrator

* ReorderListener for particle types

* Serialization for DPDIntegrator

* CUDA implementation of DPDIntegrator

* HIP implementation of DPDIntegrator

* Fixed compile error in Python wrapper

* Fixed compile error in wrappers

* Fixed uninitialized memory in reference neighbor list

* Added DPDIntegrator to C++ API docs

* Fixed incorrect launch size

* Fixed nan in DPD random number generator

* Minor optimizations

* Improved load balancing

* Fixed an indexing error

* Neighbor list uses the maximum cutoff of any force

* Fixed HIP compilation error

* Fixed access to invalid memory

* Added test case for diffusion coefficient

* Try to debug segfaults on CI

* Debugging

* Debugging

* Debugging

* Debugging

* Debugging

* Debugging

* Possible fix

* Debugging

* Debugging

* Debugging

* Use correct block size on CPU OpenCL

* Workaround for bug in Intel's OpenCL for CPUs

* Removed an unnecessary define

* Removed debugging code

* Include Dart

* More Intel workarounds

* Workaround for error in NVIDIA OpenCL

* Began splitting CommonKernels into multiple files

* Moved two more kernels into separate files

* Moved two integrators into separate files

* Fix compilation error on Windows

* reset overflowed state energies at the alchemical endpoints

* address formatting, complete clash test

* Fixed indentation

---------
Co-authored-by: Peter Eastman <peter.eastman@gmail.com>