GitLab

* 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

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

* 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

* 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

* 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

* 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

* 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

* Unified lots of parallel computation code between platforms

* Unified test code between platforms

* Eliminated duplicated timing code

CustomCPPForceImpl for writing forces in C++

    https://github.com/openmm/openmm/commit/9a0db72
    https://github.com/openmm/openmm/pull/4231

Virtual sites can depend on other virtual sites

    https://github.com/openmm/openmm/commit/71f4b3f

Use LF-Middle for LangevinIntegrator and VariableLangevinIntegrator

    https://github.com/openmm/openmm/commit/86988b9

Merged more code into common platform

    https://github.com/openmm/openmm/commit/5739788

    * Common implementation of BondedUtilities
    * Common implementation of UpdateStateDataKernel

Fixed periodic box changing from rectangular to triclinic

    https://github.com/openmm/openmm/commit/75d4f29

Co-authored-by: Emilio Gallicchio <emilio.gallicchio@gmail.com>

Use cuCtxPushCurrent() and cuCtxPopCurrent() for selecting CUDA context

    https://github.com/openmm/openmm/pull/3258

Fixed uninitialized memory access

    https://github.com/openmm/openmm/issues/3392
    https://github.com/openmm/openmm/pull/3399

Fixed potential invalid memory access

    See https://github.com/openmm/openmm/pull/3428

Improved temperature reporting for Drude particles

    https://github.com/openmm/openmm/pull/3486
    https://github.com/openmm/openmm/commit/a5e42f5

Fixed race condition with multiple GPUs

    https://github.com/openmm/openmm/commit/6fb1c8a41edff980862750bc086f6a204eb50941

Use blocking sync when creating events

    https://github.com/openmm/openmm/commit/fe21d5ee4f14673a4ea38b7244991772a64ceec2

Very minor optimizations

    https://github.com/openmm/openmm/commit/109f6b2535da4e0c0dd88007d6ca06b4add2ce81

Use PocketFFT

    https://github.com/openmm/openmm/commit/1dac981a63300a2a53a7925f570995914f7163ed

Improved logic for deciding when to reorder atoms

    https://github.com/openmm/openmm/commit/48664a1f1a4490a4dabc277757545ac070e7b898

Ensure valid atom order after loading a checkpoint

    https://github.com/openmm/openmm/commit/a056d5a3754e193105409afa12c9f0c9a2d972a2

Improve performance running on multiple GPUs

    https://github.com/openmm/openmm/commit/0c82c2647de98da5c6dab7bf7a7b8b19705aadc0

Fixed errors when running on multiple GPUs

    https://github.com/openmm/openmm/commit/ed9df876d43c037c08d4762721e73e5caae086d9

Optimized reducing energy

    https://github.com/openmm/openmm/commit/2975f44

* VkFFT-based 3D FFT;
* Caching of compiled VkFFT kernels;
* Extend FFT tests with more sizes.

* Unload all loaded modules in HipContext's destructor,
  HIP modules keep file desctriptors opened, but OpenMM never unloads
  modules leaking these file descriptors. This can cause crashinf of
  some scripts like test-openmm-platforms from openmmtools.
* ROCm 6.0 defines operator* for complex types (that are typedefs for
  float2 and double2), they conflict with operators defined for vectors.
  This is fixed in newer ROCm versions.
* Revert HIP_DYNAMIC_SHARED back to extern __shared__ (the macro is
  in the headers).
* Reduce the speed of the HIP platform if there are no HIP devices in
  the system.

* Compile kernels with max block size of 256 threads:
  The default hipcc behavior since ROCm 4.2 is to compile kernels
  with 1024 threads unless __launch_bounds__ is specified. This
  significantly increases register pressure especially in heavy kernels
  (double precision, for example), requiring register spilling;
* Optimize computeRange by using multiple blocks for reduction;
* Use blocks of 1024 threads for computeBucketPositions - it is executed
  as a single work group so larger block size is faster;
* Sort up-to lenghtNextPow2 instead of blockDim.x (faster for short
  buckets);
* Optimize sortShortList2;
* Optimize sortBuckets with bit instructions;
* Decrease bucket size for non-uniform sorting: too many buckets may
  have sizes too large to sort in shared memory;
* Add more sizes in tests.

* Remove setting of link libraries, include and link dirs and compile
  flags for each target, instead let Cmake deal with them by linking the
  main library to hip::host hiprtc::hiprtc hip::hipfft;
* Fix: custom command without ADD_CUSTOM_TARGET and ADD_DEPENDENCIES is
  executed for both static and shared targets;
* Remove IF(APPLE) parts.

Port changes in CUDA backend to HIP

Fix a warning about arithmetic operations on void* in HipArray::uploadSubArray

Fix "Error Initializing context ROCm 5.3.0"

    https://github.com/StreamHPC/openmm-hip/issues/3


    hipDeviceSetCacheConfig returns hipErrorNotSupported on 5.3
Co-authored-by: Nick Curtis <nicholas.curtis@amd.com>